test(skiplist): Make cuda code work in test

Signed-off-by: KunoiSayami <[email protected]>

2 files changed, 188 insertions, 8 deletions

diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index 0022c84..a9f7e08 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -17,7 +17,7 @@
 #include "util/random.cuh"
 #include "util/testutil.h"
 
-#include "gtest/gtest.h"
+#include "util/cuda_gtest_plugin.h"
 
 namespace leveldb {
 
@@ -364,10 +364,41 @@ __global__ void insert_and_lookup(SkipList<Key, Comparator> * skipList) {
 
 }
 
+
+__device__ void new_cuda_atomic() {
+}
+
+class TestClass {
+ public:
+  explicit __device__ TestClass(): atomic(0), alloc_ptr_(nullptr), alloc_bytes_remaining_(0),
+  head_(nullptr), blocks_(nullptr) {
+
+  }
+
+  TestClass(const TestClass&) = delete;
+  TestClass& operator=(const TestClass&) = delete;
+
+    char* alloc_ptr_;
+    size_t alloc_bytes_remaining_;
+
+    // Array of new[] allocated memory blocks
+    //thrust::host_vector<char *> blocks_;
+    //std::vector<char*> blocks_;
+
+    void * head_;
+    void * blocks_;
+
+  cuda::atomic<size_t> atomic;
+};
+
 __global__ void init(Arena ** pArena, SkipList<Key, Comparator> ** pSkipList) {
   Comparator cmp;
+  printf("init\n");
+  //cuda::atomic<size_t> a;
   *pArena = new Arena();
-  *pSkipList = new SkipList<Key, Comparator>(cmp, reinterpret_cast<Arena*>(&*pArena));
+  printf("init arena\n");
+  //*pSkipList = new SkipList<Key, Comparator>(cmp, reinterpret_cast<Arena*>(&*pArena));
+  printf("init2\n");
 }
 
 
@@ -380,12 +411,13 @@ __host__ void host_insert_test() {
     keys[i] = .Next();
   }*/
   Arena * pArena;
+  cudaMalloc((void**)&pArena, sizeof(*pArena));
   SkipList<Key, Comparator> * skipList;
   auto * device_rnd = new Random(test::RandomSeed());
 
   dim3 blockSize(1, 1);
 
-  dim3 gridSize(32, 32);
+  dim3 gridSize(1, 1);
 
   init<<<gridSize, blockSize>>>(&pArena, &skipList);
   cudaDeviceSynchronize();
@@ -396,7 +428,7 @@ __host__ void host_insert_test() {
 
   insert_and_lookup<<<gridSize, blockSize>>>(skipList);
   cudaDeviceSynchronize();
-
+  std::cout << "test";
 }
 
 
@@ -427,9 +459,8 @@ __global__ void test_fail() {
   assert(0);
 }
 
-TEST(SkipTest, Empty) {
-  test_fail<<<1,1>>>();
-  cudaDeviceSynchronize();
+CUDA_TEST(SkipTest, Empty) {
+
 }
 
 /*
@@ -522,5 +553,7 @@ TEST(SkipTest, Concurrent5) { RunConcurrent(5); }
 
 int main(int argc, char** argv) {
   testing::InitGoogleTest(&argc, argv);
-  return RUN_ALL_TESTS();
+  int i = RUN_ALL_TESTS();
+  sleep(1);
+  return i;
 }
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_ */