Commit ad4e179b authored by Tomáš Oberhuber's avatar Tomáš Oberhuber Committed by Jakub Klinkovský
Browse files

Added HIP tests for vectors.

parent a85da5ed

src/UnitTests/Containers/ArrayViewTest.h

+49 −1
Original line number Diff line number Diff line
@@ -93,6 +93,24 @@ using ViewTypes = ::testing::Types< 
   ,ArrayView< double, Devices::Cuda, long >

   ,ArrayView< MyData, Devices::Cuda, long >

#endif

#ifdef HAVE_HIP

  ,ArrayView< int,    Devices::Hip, short >,

   ArrayView< long,   Devices::Hip, short >,

   ArrayView< float,  Devices::Hip, short >,

   ArrayView< double, Devices::Hip, short >,

   ArrayView< MyData, Devices::Hip, short >,

   ArrayView< int,    Devices::Hip, int >,

   ArrayView< long,   Devices::Hip, int >,

   ArrayView< float,  Devices::Hip, int >,

   ArrayView< double, Devices::Hip, int >,

   ArrayView< MyData, Devices::Hip, int >,

   ArrayView< int,    Devices::Hip, long >,

   ArrayView< long,   Devices::Hip, long >,

   ArrayView< float,  Devices::Hip, long >,

   ArrayView< double, Devices::Hip, long >,

   ArrayView< MyData, Devices::Hip, long >

#endif





   // all ArrayView tests should also work with VectorView

   // (but we can't test all types because the argument list would be too long...)
@@ -106,6 +124,11 @@ using ViewTypes = ::testing::Types< 
   VectorView< float,  Devices::Cuda, long >,

   VectorView< double, Devices::Cuda, long >

#endif

#ifdef HAVE_HIP

   ,

   VectorView< float,  Devices::Hip, long >,

   VectorView< double, Devices::Hip, long >

#endif

>;



TYPED_TEST_SUITE( ArrayViewTest, ViewTypes );
@@ -257,6 +280,31 @@ void testArrayViewElementwiseAccess( Array< Value, Devices::Cuda, Index >&& u ) 
#endif

}



#ifdef HAVE_HIP

template< typename ValueType, typename IndexType >

__global__ void testSetGetElementKernel( ArrayView< ValueType, Devices::Hip, IndexType > v )

{

   if( threadIdx.x < v.getSize() )

      v[ threadIdx.x ] = threadIdx.x;

}

#endif /* HAVE_HIP */



template< typename Value, typename Index >

void testArrayViewElementwiseAccess( Array< Value, Devices::Hip, Index >&& u )

{

#ifdef HAVE_HIP

   u.setSize( 10 );

   using ArrayType = Array< Value, Devices::Hip, Index >;

   using ViewType = ArrayView< Value, Devices::Hip, Index >;

   ViewType v( u );

   testSetGetElementKernel<<< 1, 16 >>>( v );

   TNL_CHECK_HIP_DEVICE;

   for( int i = 0; i < 10; i++ ) {

      EXPECT_EQ( u.getElement( i ), i );

   }

#endif

}



TYPED_TEST( ArrayViewTest, elementwiseAccess )

{

   using ArrayType = typename TestFixture::ArrayType;

src/UnitTests/Containers/CMakeLists.txt

+29 −0
Original line number Diff line number Diff line
@@ -48,6 +48,35 @@ IF( BUILD_HIP ) 
   HIP_ADD_EXECUTABLE( ArrayTestHip ArrayTest.cpp )

   TARGET_COMPILE_OPTIONS( ArrayTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( ArrayTestHip ${GTEST_BOTH_LIBRARIES} )



   HIP_ADD_EXECUTABLE( ArrayViewTestHip ArrayViewTest.cpp )

   TARGET_COMPILE_OPTIONS( ArrayViewTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( ArrayViewTestHip ${GTEST_BOTH_LIBRARIES} )



   HIP_ADD_EXECUTABLE( VectorTestHip VectorTest.cpp )

   TARGET_COMPILE_OPTIONS( VectorTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( VectorTestHip ${GTEST_BOTH_LIBRARIES} )



   HIP_ADD_EXECUTABLE( VectorPrefixSumTestHip VectorPrefixSumTest.cpp )

   TARGET_COMPILE_OPTIONS( VectorPrefixSumTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( VectorPrefixSumTestHip ${GTEST_BOTH_LIBRARIES} )



   HIP_ADD_EXECUTABLE( VectorEvaluateAndReduceTestHip VectorEvaluateAndReduceTest.cpp )

   TARGET_COMPILE_OPTIONS( VectorEvaluateAndReduceTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( VectorEvaluateAndReduceTestHip ${GTEST_BOTH_LIBRARIES} )



   HIP_ADD_EXECUTABLE( VectorBinaryOperationsTestHip VectorBinaryOperationsTest.cpp )

   TARGET_COMPILE_OPTIONS( VectorBinaryOperationsTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( VectorBinaryOperationsTestHip ${GTEST_BOTH_LIBRARIES} )



   HIP_ADD_EXECUTABLE( VectorUnaryOperationsTestHip VectorUnaryOperationsTest.cpp )

   TARGET_COMPILE_OPTIONS( VectorUnaryOperationsTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( VectorUnaryOperationsTestHip ${GTEST_BOTH_LIBRARIES} )



   HIP_ADD_EXECUTABLE( VectorVerticalOperationsTestHip VectorVerticalOperationsTest.cpp )

   TARGET_COMPILE_OPTIONS( VectorVerticalOperationsTestHip PRIVATE ${CMAKE_HIPCXX_FLAGS} ${CXX_TESTS_FLAGS} )

   TARGET_LINK_LIBRARIES( VectorVerticalOperationsTestHip ${GTEST_BOTH_LIBRARIES} )



   ADD_TEST( ArrayTestHip ${EXECUTABLE_OUTPUT_PATH}/ArrayTestHip${CMAKE_EXECUTABLE_SUFFIX} )

ENDIF( BUILD_HIP )

src/UnitTests/Containers/VectorBinaryOperationsTest.h

+41 −6
Original line number Diff line number Diff line
@@ -156,7 +156,7 @@ protected: 
// types for which VectorBinaryOperationsTest is instantiated

#if defined(DISTRIBUTED_VECTOR)

   using VectorPairs = ::testing::Types<

   #ifndef HAVE_CUDA

   #if !defined HAVE_CUDA && !defined HAVE_HIP

      Pair< DistributedVector<     int,   Devices::Host, int >,

            DistributedVector<     short, Devices::Host, int > >,

      Pair< DistributedVector<     int,   Devices::Host, int >,
@@ -165,7 +165,8 @@ protected: 
            DistributedVector<     short, Devices::Host, int > >,

      Pair< DistributedVectorView< int,   Devices::Host, int >,

            DistributedVectorView< short, Devices::Host, int > >

   #else

   #endif

   #ifdef HAVE_CUDA

      Pair< DistributedVector<     int,   Devices::Cuda, int >,

            DistributedVector<     short, Devices::Cuda, int > >,

      Pair< DistributedVector<     int,   Devices::Cuda, int >,
@@ -175,6 +176,17 @@ protected: 
      Pair< DistributedVectorView< int,   Devices::Cuda, int >,

            DistributedVectorView< short, Devices::Cuda, int > >

   #endif

   // TODO HIP: Add tests for HIP distributed vectors when it is implemented

   /*#ifdef HAVE_HIP

      Pair< DistributedVector<     int,   Devices::Hip, int >,

            DistributedVector<     short, Devices::Hip, int > >,

      Pair< DistributedVector<     int,   Devices::Hip, int >,

            DistributedVectorView< short, Devices::Hip, int > >,

      Pair< DistributedVectorView< int,   Devices::Hip, int >,

            DistributedVector<     short, Devices::Hip, int > >,

      Pair< DistributedVectorView< int,   Devices::Hip, int >,

            DistributedVectorView< short, Devices::Hip, int > >,

   #endif*/

   >;

#elif defined(STATIC_VECTOR)

   #ifdef VECTOR_OF_STATIC_VECTORS
@@ -202,21 +214,28 @@ protected: 
#else

   #ifdef VECTOR_OF_STATIC_VECTORS

      using VectorPairs = ::testing::Types<

      #ifndef HAVE_CUDA

      #if !defined HAVE_CUDA && !defined HAVE_HIP

         Pair< Vector<     StaticVector< 3, int >, Devices::Host >, Vector<     StaticVector< 3, short >, Devices::Host > >,

         Pair< VectorView< StaticVector< 3, int >, Devices::Host >, Vector<     StaticVector< 3, short >, Devices::Host > >,

         Pair< Vector<     StaticVector< 3, int >, Devices::Host >, VectorView< StaticVector< 3, short >, Devices::Host > >,

         Pair< VectorView< StaticVector< 3, int >, Devices::Host >, VectorView< StaticVector< 3, short >, Devices::Host > >

      #else

      #endif

      #ifdef HAVE_CUDA

         Pair< Vector<     StaticVector< 3, int >, Devices::Cuda >, Vector<     StaticVector< 3, short >, Devices::Cuda > >,

         Pair< VectorView< StaticVector< 3, int >, Devices::Cuda >, Vector<     StaticVector< 3, short >, Devices::Cuda > >,

         Pair< Vector<     StaticVector< 3, int >, Devices::Cuda >, VectorView< StaticVector< 3, short >, Devices::Cuda > >,

         Pair< VectorView< StaticVector< 3, int >, Devices::Cuda >, VectorView< StaticVector< 3, short >, Devices::Cuda > >

      #endif

      #ifdef HAVE_HIP

         Pair< Vector<     StaticVector< 3, int >, Devices::Hip >, Vector<     StaticVector< 3, short >, Devices::Hip > >,

         Pair< VectorView< StaticVector< 3, int >, Devices::Hip >, Vector<     StaticVector< 3, short >, Devices::Hip > >,

         Pair< Vector<     StaticVector< 3, int >, Devices::Hip >, VectorView< StaticVector< 3, short >, Devices::Hip > >,

         Pair< VectorView< StaticVector< 3, int >, Devices::Hip >, VectorView< StaticVector< 3, short >, Devices::Hip > >

      #endif

      >;

   #else

      using VectorPairs = ::testing::Types<

      #ifndef HAVE_CUDA

      #if ! defined HAVE_CUDA && ! defined HAVE_HIP

         Pair< Vector<     int,       Devices::Host >, Vector<     int,       Devices::Host > >,

         Pair< VectorView< int,       Devices::Host >, Vector<     int,       Devices::Host > >,

         Pair< VectorView< const int, Devices::Host >, Vector<     int,       Devices::Host > >,
@@ -230,7 +249,8 @@ protected: 
         Pair< VectorView< double,    Devices::Host >, Vector<     double,    Devices::Host > >,

         Pair< Vector<     double,    Devices::Host >, VectorView< double,    Devices::Host > >,

         Pair< VectorView< double,    Devices::Host >, VectorView< double,    Devices::Host > >

      #else

      #endif

      #ifdef HAVE_CUDA

         Pair< Vector<     int,       Devices::Cuda >, Vector<     int,       Devices::Cuda > >,

         Pair< VectorView< int,       Devices::Cuda >, Vector<     int,       Devices::Cuda > >,

         Pair< VectorView< const int, Devices::Cuda >, Vector<     int,       Devices::Cuda > >,
@@ -245,6 +265,21 @@ protected: 
         Pair< Vector<     double,    Devices::Cuda >, VectorView< double,    Devices::Cuda > >,

         Pair< VectorView< double,    Devices::Cuda >, VectorView< double,    Devices::Cuda > >

      #endif

      #ifdef HAVE_HIP

         Pair< Vector<     int,       Devices::Hip >, Vector<     int,       Devices::Hip > >,

         Pair< VectorView< int,       Devices::Hip >, Vector<     int,       Devices::Hip > >,

         Pair< VectorView< const int, Devices::Hip >, Vector<     int,       Devices::Hip > >,

         Pair< Vector<     int,       Devices::Hip >, VectorView< int,       Devices::Hip > >,

         Pair< Vector<     int,       Devices::Hip >, VectorView< const int, Devices::Hip > >,

         Pair< VectorView< int,       Devices::Hip >, VectorView< int,       Devices::Hip > >,

         Pair< VectorView< const int, Devices::Hip >, VectorView< int,       Devices::Hip > >,

         Pair< VectorView< const int, Devices::Hip >, VectorView< const int, Devices::Hip > >,

         Pair< VectorView< int,       Devices::Hip >, VectorView< const int, Devices::Hip > >,

         Pair< Vector<     double,    Devices::Hip >, Vector<     double,    Devices::Hip > >,

         Pair< VectorView< double,    Devices::Hip >, Vector<     double,    Devices::Hip > >,

         Pair< Vector<     double,    Devices::Hip >, VectorView< double,    Devices::Hip > >,

         Pair< VectorView< double,    Devices::Hip >, VectorView< double,    Devices::Hip > >

      #endif

      >;

   #endif

#endif

src/UnitTests/Containers/VectorPrefixSumTest.h

+120 −2
Original line number Diff line number Diff line
@@ -135,6 +135,66 @@ TYPED_TEST( VectorTest, scan ) 
         EXPECT_EQ( v_host[ i ], (i * (i + 1)) / 2 ) << "i = " << i;



      Algorithms::detail::CudaScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::resetMaxGridSize();

#endif

   }



   ////

   // With HIP, perform tests with multiple HIP grids.

   if( std::is_same< DeviceType, Devices::Hip >::value )

   {

#ifdef HAVE_HIP

      Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::maxGridSize() = 3;



      setConstantSequence( v, 0 );

      v_host = -1;

      v.scan();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::gridsCount() ), 1  );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], 0 ) << "i = " << i;



      setConstantSequence( v, 1 );

      v_host = -1;

      v.scan();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::gridsCount() ), 1  );

      v_host = v_view;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], i + 1 ) << "i = " << i;



      setLinearSequence( v );

      v_host = -1;

      v.scan();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::gridsCount() ), 1  );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], (i * (i + 1)) / 2 ) << "i = " << i;



      // test views

      setConstantSequence( v, 0 );

      v_host = -1;

      v_view.scan();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::gridsCount() ), 1  );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], 0 ) << "i = " << i;



      setConstantSequence( v, 1 );

      v_host = -1;

      v_view.scan();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::gridsCount() ), 1  );

      v_host = v_view;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], i + 1 ) << "i = " << i;



      setLinearSequence( v );

      v_host = -1;

      v_view.scan();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::gridsCount() ), 1  );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], (i * (i + 1)) / 2 ) << "i = " << i;



      Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Inclusive, RealType, IndexType >::resetMaxGridSize();

#endif

   }

}
@@ -260,11 +320,69 @@ TYPED_TEST( VectorTest, exclusiveScan ) 
      Algorithms::detail::CudaScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::resetMaxGridSize();

#endif

   }

}



// TODO: test scan with custom begin and end parameters

   ////

   // With HIP, perform tests with multiple HIP grids.

   if( std::is_same< DeviceType, Devices::Hip >::value )

   {

#ifdef HAVE_HIP

      Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::maxGridSize() = 3;



      setConstantSequence( v, 0 );

      v_host = -1;

      v.template scan< Algorithms::ScanType::Exclusive >();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::gridsCount() ), 1 );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], 0 ) << "i = " << i;



      setConstantSequence( v, 1 );

      v_host = -1;

      v.template scan< Algorithms::ScanType::Exclusive >();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::gridsCount() ), 1 );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], i ) << "i = " << i;



      setLinearSequence( v );

      v_host = -1;

      v.template scan< Algorithms::ScanType::Exclusive >();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::gridsCount() ), 1 );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], (i * (i - 1)) / 2 ) << "i = " << i;



      // test views

      setConstantSequence( v, 0 );

      v_host = -1;

      v_view.template scan< Algorithms::ScanType::Exclusive >();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::gridsCount() ), 1 );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], 0 ) << "i = " << i;



      setConstantSequence( v, 1 );

      v_host = -1;

      v_view.template scan< Algorithms::ScanType::Exclusive >();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::gridsCount() ), 1 );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], i ) << "i = " << i;



      setLinearSequence( v );

      v_host = -1;

      v_view.template scan< Algorithms::ScanType::Exclusive >();

      EXPECT_GT( ( Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::gridsCount() ), 1 );

      v_host = v;

      for( int i = 0; i < size; i++ )

         EXPECT_EQ( v_host[ i ], (i * (i - 1)) / 2 ) << "i = " << i;



      Algorithms::HipScanKernelLauncher< Algorithms::ScanType::Exclusive, RealType, IndexType >::resetMaxGridSize();

#endif

   }

}



// TODO: test scan with custom begin and end parameters

template< typename FlagsView >

void setupFlags( FlagsView& f )

{

src/UnitTests/Containers/VectorTestSetup.h

+21 −1
Original line number Diff line number Diff line
@@ -36,7 +36,7 @@ protected: 
// types for which VectorTest is instantiated

// TODO: Quad must be fixed

using VectorTypes = ::testing::Types<

#ifndef HAVE_CUDA

#if !defined HAVE_CUDA && ! defined HAVE_HIP

   Vector< int,            Devices::Host, short >,

   Vector< long,           Devices::Host, short >,

   Vector< float,          Devices::Host, short >,
@@ -76,6 +76,26 @@ using VectorTypes = ::testing::Types< 
   //Vector< Quad< float >,  Devices::Cuda, long >,

   //Vector< Quad< double >, Devices::Cuda, long >

#endif

#ifdef HAVE_HIP

   Vector< int,            Devices::Hip, short >,

   Vector< long,           Devices::Hip, short >,

   Vector< float,          Devices::Hip, short >,

   Vector< double,         Devices::Hip, short >,

   //Vector< Quad< float >,  Devices::Hip, short >,

   //Vector< Quad< double >, Devices::Hip, short >,

   Vector< int,            Devices::Hip, int >,

   Vector< long,           Devices::Hip, int >,

   Vector< float,          Devices::Hip, int >,

   Vector< double,         Devices::Hip, int >,

   //Vector< Quad< float >,  Devices::Hip, int >,

   //Vector< Quad< double >, Devices::Hip, int >,

   Vector< int,            Devices::Hip, long >,

   Vector< long,           Devices::Hip, long >,

   Vector< float,          Devices::Hip, long >,

   Vector< double,         Devices::Hip, long >

   //Vector< Quad< float >,  Devices::Hip, long >,

   //Vector< Quad< double >, Devices::Hip, long >

#endif

>;



TYPED_TEST_SUITE( VectorTest, VectorTypes );