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

Added HIP tests for smart pointers.

parent 3cb394a0

src/UnitTests/Pointers/CMakeLists.txt

+12 −0
Original line number Diff line number Diff line
@@ -18,3 +18,15 @@ if( BUILD_CUDA ) 
      add_test( ${target} ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}/${target}${CMAKE_EXECUTABLE_SUFFIX} )

   endforeach()

endif()



if( BUILD_HIP )

   HIP_ADD_EXECUTABLE( SharedPointerHipTest SharedPointerHipTest.cpp )

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

   TARGET_LINK_LIBRARIES( SharedPointerHipTest ${GTEST_BOTH_LIBRARIES} )

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



   HIP_ADD_EXECUTABLE( DevicePointerHipTest DevicePointerHipTest.cpp )

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

   TARGET_LINK_LIBRARIES( DevicePointerHipTest ${GTEST_BOTH_LIBRARIES} )

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

endif( BUILD_HIP )

src/UnitTests/Pointers/DevicePointerHipTest.cpp

0 → 100644
+157 −0
Original line number Diff line number Diff line 
/***************************************************************************

                          DevicePointerHipTest.cpp  -  description

                             -------------------

    begin                : Jun 26, 2020

    copyright            : (C) 2020 by Tomas Oberhuber

    email                : tomas.oberhuber@fjfi.cvut.cz

 ***************************************************************************/



/* See Copyright Notice in tnl/Copyright */



#include <cstdlib>

#include <TNL/Devices/Host.h>

#include <TNL/Pointers/DevicePointer.h>

#include <TNL/Containers/StaticArray.h>

#include <TNL/Containers/Array.h>



#ifdef HAVE_GTEST

#include <gtest/gtest.h>

#endif



#ifdef HAVE_HIP

#include "hip/hip_runtime.h"

#endif



#include <TNL/Devices/Hip.h>



using namespace TNL;



#ifdef HAVE_GTEST

TEST( DevicePointerHipTest, ConstructorTest )

{

#ifdef HAVE_HIP

   using TestType =  TNL::Containers::StaticArray< 2, int  >;

   TestType obj1;

   Pointers::DevicePointer< TestType, Devices::Hip > ptr1( obj1 );



   ptr1->x() = 0;

   ptr1->y() = 0;

   ASSERT_EQ( ptr1->x(), 0 );

   ASSERT_EQ( ptr1->y(), 0 );



   TestType obj2( 1,2 );

   Pointers::DevicePointer< TestType, Devices::Hip > ptr2( obj2 );

   ASSERT_EQ( ptr2->x(), 1 );

   ASSERT_EQ( ptr2->y(), 2 );



   ptr1 = ptr2;

   ASSERT_EQ( ptr1->x(), 1 );

   ASSERT_EQ( ptr1->y(), 2 );

#endif

};



TEST( DevicePointerHipTest, getDataTest )

{

#ifdef HAVE_HIP

   using TestType = TNL::Containers::StaticArray< 2, int  >;

   TestType obj1( 1, 2 );

   Pointers::DevicePointer< TestType, Devices::Hip > ptr1( obj1 );



   Pointers::synchronizeSmartPointersOnDevice< Devices::Hip >();



   TestType aux;



   hipMemcpy( ( void*) &aux, &ptr1.getData< Devices::Hip >(), sizeof( TestType ), hipMemcpyDeviceToHost );



   ASSERT_EQ( aux[ 0 ], 1 );

   ASSERT_EQ( aux[ 1 ], 2 );

#endif  // HAVE_HIP

};



#ifdef HAVE_HIP

__global__ void copyArrayKernel( const TNL::Containers::Array< int, Devices::Hip >* inArray,

                                 int* outArray )

{

   if( threadIdx.x < 2 )

   {

      outArray[ threadIdx.x ] = ( *inArray )[ threadIdx.x ];

   }

}



__global__ void copyArrayKernel2( const Pointers::DevicePointer< TNL::Containers::Array< int, Devices::Hip > > inArray,

                                  int* outArray )

{

   if( threadIdx.x < 2 )

   {

      outArray[ threadIdx.x ] = ( *inArray )[ threadIdx.x ];

   }

}

#endif



TEST( DevicePointerHipTest, getDataArrayTest )

{

#ifdef HAVE_HIP

   using TestType = TNL::Containers::Array< int, Devices::Hip  >;

   TestType obj;

   Pointers::DevicePointer< TestType > ptr( obj );



   ptr->setSize( 2 );

   ptr->setElement( 0, 1 );

   ptr->setElement( 1, 2 );



   Pointers::synchronizeSmartPointersOnDevice< Devices::Hip >();



   int *testArray_device, *testArray_host;

   hipMalloc( ( void** ) &testArray_device, 2 * sizeof( int ) );

   hipLaunchKernelGGL(copyArrayKernel, dim3(1), dim3(2 ), 0, 0,  &ptr.getData< Devices::Hip >(), testArray_device );

   testArray_host = new int [ 2 ];

   hipMemcpy( testArray_host, testArray_device, 2 * sizeof( int ), hipMemcpyDeviceToHost );



   ASSERT_EQ( testArray_host[ 0 ], 1 );

   ASSERT_EQ( testArray_host[ 1 ], 2 );



   hipLaunchKernelGGL(copyArrayKernel2, dim3(1), dim3(2 ), 0, 0,  ptr, testArray_device );

   hipMemcpy( testArray_host, testArray_device, 2 * sizeof( int ), hipMemcpyDeviceToHost );



   ASSERT_EQ( testArray_host[ 0 ], 1 );

   ASSERT_EQ( testArray_host[ 1 ], 2 );



   delete[] testArray_host;

   hipFree( testArray_device );



#endif

};



TEST( DevicePointerHipTest, nullptrAssignement )

{

#ifdef HAVE_HIP

   using TestType = Pointers::DevicePointer< double, Devices::Hip >;

   double o1 = 5;

   TestType p1( o1 ), p2( nullptr );



   // This should not crash

   p1 = p2;



   ASSERT_FALSE( p1 );

   ASSERT_FALSE( p2 );

#endif

}



TEST( DevicePointerHipTest, swap )

{

#ifdef HAVE_HIP

   using TestType = Pointers::DevicePointer< double, Devices::Hip >;

   double o1( 1 ), o2( 2 );

   TestType p1( o1 ), p2( o2 );



   p1.swap( p2 );



   ASSERT_EQ( *p1, 2 );

   ASSERT_EQ( *p2, 1 );

#endif

}



#endif





#include "../main.h"

src/UnitTests/Pointers/SharedPointerHipTest.cpp

0 → 100644
+151 −0
Original line number Diff line number Diff line 
/***************************************************************************

                          SharedPointerHipTest.cpp  -  description

                             -------------------

    begin                : Aug 22, 2018

    copyright            : (C) 2018 by Tomas Oberhuber

    email                : tomas.oberhuber@fjfi.cvut.cz

 ***************************************************************************/



/* See Copyright Notice in tnl/Copyright */



#include <cstdlib>

#include <TNL/Devices/Host.h>

#include <TNL/Pointers/SharedPointer.h>

#include <TNL/Containers/StaticArray.h>

#include <TNL/Containers/Array.h>



#ifdef HAVE_GTEST

#include <gtest/gtest.h>

#endif



#ifdef HAVE_HIP

#include "hip/hip_runtime.h"

#endif



#include <TNL/Devices/Hip.h>



using namespace TNL;



#ifdef HAVE_GTEST

TEST( SharedPointerHipTest, ConstructorTest )

{

#ifdef HAVE_HIP

   typedef TNL::Containers::StaticArray< 2, int  > TestType;

   Pointers::SharedPointer< TestType, Devices::Hip > ptr1;



   ptr1->x() = 0;

   ptr1->y() = 0;

   ASSERT_EQ( ptr1->x(), 0 );

   ASSERT_EQ( ptr1->y(), 0 );



   Pointers::SharedPointer< TestType, Devices::Hip > ptr2( 1, 2 );

   ASSERT_EQ( ptr2->x(), 1 );

   ASSERT_EQ( ptr2->y(), 2 );



   ptr1 = ptr2;

   ASSERT_EQ( ptr1->x(), 1 );

   ASSERT_EQ( ptr1->y(), 2 );

#endif

};



TEST( SharedPointerHipTest, getDataTest )

{

#ifdef HAVE_HIP

   typedef TNL::Containers::StaticArray< 2, int  > TestType;

   Pointers::SharedPointer< TestType, Devices::Hip > ptr1( 1, 2 );



   Pointers::synchronizeSmartPointersOnDevice< Devices::Hip >();



   TestType aux;



   hipMemcpy( ( void*) &aux, &ptr1.getData< Devices::Hip >(), sizeof( TestType ), hipMemcpyDeviceToHost );



   ASSERT_EQ( aux[ 0 ], 1 );

   ASSERT_EQ( aux[ 1 ], 2 );

#endif  // HAVE_HIP

};



#ifdef HAVE_HIP

__global__ void copyArrayKernel( const TNL::Containers::Array< int, Devices::Hip >* inArray,

                                 int* outArray )

{

   if( threadIdx.x < 2 )

   {

      outArray[ threadIdx.x ] = ( *inArray )[ threadIdx.x ];

   }

}



__global__ void copyArrayKernel2( const Pointers::SharedPointer< TNL::Containers::Array< int, Devices::Hip > > inArray,

                                  int* outArray )

{

   if( threadIdx.x < 2 )

   {

      outArray[ threadIdx.x ] = ( *inArray )[ threadIdx.x ];

   }

}

#endif



TEST( SharedPointerHipTest, getDataArrayTest )

{

#ifdef HAVE_HIP

   typedef TNL::Containers::Array< int, Devices::Hip  > TestType;

   Pointers::SharedPointer< TestType > ptr;



   ptr->setSize( 2 );

   ptr->setElement( 0, 1 );

   ptr->setElement( 1, 2 );



   Pointers::synchronizeSmartPointersOnDevice< Devices::Hip >();



   int *testArray_device, *testArray_host;

   hipMalloc( ( void** ) &testArray_device, 2 * sizeof( int ) );

   hipLaunchKernelGGL(copyArrayKernel, dim3(1), dim3(2 ), 0, 0,  &ptr.getData< Devices::Hip >(), testArray_device );

   testArray_host = new int [ 2 ];

   hipMemcpy( testArray_host, testArray_device, 2 * sizeof( int ), hipMemcpyDeviceToHost );



   ASSERT_EQ( testArray_host[ 0 ], 1 );

   ASSERT_EQ( testArray_host[ 1 ], 2 );



   hipLaunchKernelGGL(copyArrayKernel2, dim3(1), dim3(2 ), 0, 0,  ptr, testArray_device );

   hipMemcpy( testArray_host, testArray_device, 2 * sizeof( int ), hipMemcpyDeviceToHost );



   ASSERT_EQ( testArray_host[ 0 ], 1 );

   ASSERT_EQ( testArray_host[ 1 ], 2 );



   delete[] testArray_host;

   hipFree( testArray_device );



#endif

};



TEST( SharedPointerHipTest, nullptrAssignement )

{

#ifdef HAVE_HIP

   using TestType = Pointers::SharedPointer< double, Devices::Hip >;

   TestType p1( 5 ), p2( nullptr );



   // This should not crash

   p1 = p2;



   ASSERT_FALSE( p1 );

   ASSERT_FALSE( p2 );

#endif

}



TEST( SharedPointerHipTest, swap )

{

#ifdef HAVE_HIP

   using TestType = Pointers::SharedPointer< double, Devices::Hip >;

   TestType p1( 1 ), p2( 2 );



   p1.swap( p2 );



   ASSERT_EQ( *p1, 2 );

   ASSERT_EQ( *p2, 1 );

#endif

}



#endif





#include "../main.h"