aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2021-12-19 15:25:11 +0800
committerKunoiSayami <[email protected]>2021-12-19 15:25:11 +0800
commit9406aaf2566cc68e75cc01ba02dfa83b2179d9f6 (patch)
treeb80c521318002cde4848c947a9f9314cf2021b32
parent4e02041bb14d92d8e724bd81430f8f48658c37ca (diff)
test(skiplist): Make cuda code work in test
Signed-off-by: KunoiSayami <[email protected]>
-rw-r--r--db/skiplist_test.cu49
-rw-r--r--util/cuda_gtest_plugin.h147
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_ */