diff options
Diffstat (limited to 'util')
| -rw-r--r-- | util/cuda_gtest_plugin.h | 147 |
1 files changed, 147 insertions, 0 deletions
diff --git a/util/cuda_gtest_plugin.h b/util/cuda_gtest_plugin.h new file mode 100644 index 0000000..0341f29 --- /dev/null +++ b/util/cuda_gtest_plugin.h @@ -0,0 +1,147 @@ +/** + * Created on: Feb 24, 2014 + * Author: vogt + */ + +#ifndef CUDA_GTEST_PLUGIN_H_ +#define CUDA_GTEST_PLUGIN_H_ + +#include "gtest/gtest.h" +#ifdef __CUDACC__ + +#ifndef CUDA_LAST_ERROR +#define CUDA_LAST_ERROR( msg )\ + {cudaDeviceSynchronize();\ + cudaError_t error = cudaGetLastError();\ + if(error!=cudaSuccess) {\ + fprintf(stderr,"ERROR: %s: %s\n", msg, cudaGetErrorString(error) );\ + exit(-1);\ + }} +#endif + +struct TestTransporter +{ + float2 tfloat[10]; + int2 tint[10]; + + int evaluateInt; + int evaluateFloat; + + __host__ __device__ TestTransporter(): evaluateInt(0), evaluateFloat(0) + { + }; +}; + + template<typename T> static __host__ __device__ void setTestTransporterValue( TestTransporter* transporter, T expected, T actual ); + + template<> __host__ __device__ void setTestTransporterValue( TestTransporter* transporter, float expected, float actual ) + { + transporter->tfloat[transporter->evaluateFloat].x = expected; + transporter->tfloat[transporter->evaluateFloat].y = actual; + transporter->evaluateFloat++; + } + + template<> __host__ __device__ void setTestTransporterValue( TestTransporter* transporter, int expected, int actual ) + { + transporter->tint[transporter->evaluateInt].x = expected; + transporter->tint[transporter->evaluateInt].y = actual; + transporter->evaluateInt++; + } + + template<> __host__ __device__ void setTestTransporterValue( TestTransporter* transporter, bool expected, bool actual ) + { + transporter->tint[transporter->evaluateInt].x = (int)expected; + transporter->tint[transporter->evaluateInt].y = (int)actual; + transporter->evaluateInt++; + } + +#define CUDA_TEST_CLASS_NAME_(test_case_name, test_name)\ + kernel_test_case_name##_##test_name##_Test + +#ifdef __CUDA_ARCH__ +#undef TEST +#define CUDA_DEAD_FUNCTION_NAME_(test_case_name, test_name)\ + MAKE_UNIQUE( dead_function_test_case_name##_##test_name##_Test ) +#define TEST(test_case_name, test_name) void CUDA_DEAD_FUNCTION_NAME_(test_case_name, test_name)( TestTransporter* testTransporter )//GTEST_TEST(test_case_name, test_name) +#define TESTTRANSPORTERDEFINITIONWITHCOMMA , TestTransporter* testTransporter +#define TESTTRANSPORTERDEFANDINSTANCE +#define TESTTRANSPORTERDEFINITION TestTransporter* testTransporter +#define TESTCALLHOST +#define TESTCALLDEVICE test( testTransporter ) +#else +#define CUDA_DEAD_FUNCTION_NAME_(test_case_name, test_name) +#define TESTTRANSPORTERDEFANDINSTANCE TestTransporter* testTransporter = new TestTransporter; +#define TESTTRANSPORTERDEFINITIONWITHCOMMA +#define TESTTRANSPORTERDEFINITION +#define TESTCALLHOST test() +#define TESTCALLDEVICE +#endif + +#define TESTKERNELCALL(test_case_name, test_name) CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name) test;CUDA_TEST_CLASS_NAME_(test_case_name, test_name)<<<1,1>>>(test,dTestTransporter) + +#define CUDA_ASSERT_EQ(expected,actual)\ + setTestTransporterValue( testTransporter, expected, actual ); + + +#ifdef __CUDA_ARCH__ +#undef ASSERT_EQ +#define ASSERT_EQ(val1, val2) CUDA_ASSERT_EQ(val1, val2) +#endif + +#ifdef __CUDA_ARCH__ +#undef ASSERT_FLOAT_EQ +#define ASSERT_FLOAT_EQ(val1, val2) CUDA_ASSERT_EQ(val1, val2) +#endif + +#define CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name)\ + test_function_test_case_name##_##test_name##_Test +#define TEST_NAME_CUDA( test_name )\ + test_name##_CUDA + +#define CONCATENATE_DETAIL(x, y) x##y +#define CONCATENATE(x, y) CONCATENATE_DETAIL(x, y) +#define MAKE_UNIQUE(x) CONCATENATE(x, __COUNTER__) + + +#define CUDA_TEST(test_case_name, test_name)\ + struct CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name)\ + {\ + __host__ __device__ void operator()( TestTransporter* testTransporter );\ + };\ + __global__ void CUDA_TEST_CLASS_NAME_(test_case_name, test_name)(CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name) test, TestTransporter* testTransporter);\ + GTEST_TEST(test_case_name, test_name)\ + {\ + CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name) test;\ + TestTransporter* testTransporter = new TestTransporter;\ + test( testTransporter );\ + };\ + TEST(test_case_name, test_name##_CUDA )\ + {\ + TestTransporter* dTestTransporter;\ + cudaMalloc( (void**)(&dTestTransporter), sizeof( TestTransporter ) ); \ + CUDA_LAST_ERROR( "malloc" ); \ + TESTTRANSPORTERDEFANDINSTANCE\ + cudaMemcpy( dTestTransporter, testTransporter, sizeof(TestTransporter), cudaMemcpyHostToDevice );\ + CUDA_LAST_ERROR( "memcopyhosttodevice" );\ + CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name) test;\ + CUDA_TEST_CLASS_NAME_(test_case_name, test_name)<<<1,1>>>(test,dTestTransporter);\ + CUDA_LAST_ERROR( "kernel call" );\ + cudaMemcpy( testTransporter, dTestTransporter, sizeof(TestTransporter), cudaMemcpyDeviceToHost );\ + CUDA_LAST_ERROR( "memcopydevicetohost" );\ + for( int i = 0; i < testTransporter->evaluateFloat; i++ )\ + ASSERT_FLOAT_EQ( testTransporter->tfloat[i].x, testTransporter->tfloat[i].y );\ + for( int i = 0; i < testTransporter->evaluateInt; i++ ) \ + GTEST_ASSERT_EQ( testTransporter->tint[i].x, testTransporter->tint[i].y );\ + };\ + __global__ void CUDA_TEST_CLASS_NAME_(test_case_name, test_name)( \ + CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name) test, TestTransporter* testTransporter)\ + {\ + test( testTransporter );\ + }\ + __host__ __device__ void CUDA_TEST_FUNCTION_NAME_(test_case_name, test_name)::operator()( TestTransporter* testTransporter ) +#else +#warning "To enable CUDA tests compile with nvcc" +#define CUDA_TEST(test_case_name, test_name) TEST(test_case_name, test_name) +#endif + +#endif /* CUDA_GTEST_PLUGIN_H_ */ |