Commit 5ff7fd4e authored by Tomáš Oberhuber's avatar Tomáš Oberhuber

Added HIP tests for smart pointers.

parent 96e2ad03
/***************************************************************************
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"
/***************************************************************************
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"
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment