refs #27: Add CUDA-side tests for ShiftedInt

* It is working as expected.
ANLAB-KAIST · Feb 23, 2016 · e9e8770 · e9e8770
1 parent eb8bcef
commit e9e8770
Show file tree

Hide file tree

Showing 3 changed files with 70 additions and 5 deletions.
diff --git a/include/nba/engines/cuda/test.hh b/include/nba/engines/cuda/test.hh
@@ -1,5 +1,7 @@
 namespace nba {
 extern void* get_test_kernel_noop();
+extern void* get_test_kernel_shiftedint_size_check();
+extern void* get_test_kernel_shiftedint_value_check();
 }
 
 // vim: ts=8 sts=4 sw=4 et
diff --git a/src/engines/cuda/test.cu b/src/engines/cuda/test.cu
@@ -1,4 +1,5 @@
 #include <nba/engines/cuda/test.hh>
+#include <nba/core/shiftedint.hh>
 
 using namespace std;
 using namespace nba;
@@ -9,8 +10,23 @@ __global__ void noop()
 }
 
 void *nba::get_test_kernel_noop()
+{ return reinterpret_cast<void *> (noop); }
+
+__global__ void shiftedint_size_check(size_t *sz_measured_in_device)
+{
+    *sz_measured_in_device = sizeof(nba::dev_offset_t);
+}
+
+__global__ void shiftedint_value_check
+(nba::dev_offset_t *v, uint64_t *raw_v)
 {
-    return reinterpret_cast<void *> (noop);
+    *raw_v = v->as_value<uint64_t>();
 }
 
+void *nba::get_test_kernel_shiftedint_size_check()
+{ return reinterpret_cast<void *> (shiftedint_size_check); }
+
+void *nba::get_test_kernel_shiftedint_value_check()
+{ return reinterpret_cast<void *> (shiftedint_value_check); }
+
 // vim: ts=8 sts=4 sw=4 et
diff --git a/tests/test_cuda.cc b/tests/test_cuda.cc
@@ -14,16 +14,63 @@ using namespace nba;
 #ifdef USE_CUDA
 
 TEST(CUDADeviceTest, Initialization) {
-    EXPECT_EQ(cudaSuccess, cudaSetDevice(0));
-    EXPECT_EQ(cudaSuccess, cudaDeviceReset());
+    ASSERT_EQ(cudaSuccess, cudaSetDevice(0));
+    ASSERT_EQ(cudaSuccess, cudaDeviceReset());
 }
 
 TEST(CUDADeviceTest, NoopKernel) {
-    EXPECT_EQ(cudaSuccess, cudaSetDevice(0));
+    ASSERT_EQ(cudaSuccess, cudaSetDevice(0));
     void *k = get_test_kernel_noop();
     EXPECT_EQ(cudaSuccess, cudaLaunchKernel(k, dim3(1), dim3(1), nullptr, 0, 0));
     EXPECT_EQ(cudaSuccess, cudaDeviceSynchronize());
-    EXPECT_EQ(cudaSuccess, cudaDeviceReset());
+    ASSERT_EQ(cudaSuccess, cudaDeviceReset());
+}
+
+class CUDAStructTest : public ::testing::Test {
+protected:
+    virtual void SetUp() {
+        cudaSetDevice(0);
+    }
+
+    virtual void TearDown() {
+        cudaDeviceReset();
+    }
+};
+
+TEST(CUDAStructTest, ShfitedIntSizeCheck) {
+    void *k = get_test_kernel_shiftedint_size_check();
+    void *output_d;
+    ASSERT_EQ(cudaSuccess, cudaMalloc(&output_d, sizeof(size_t)));
+    ASSERT_NE(nullptr, output_d);
+    size_t output_h = 0;
+    void *raw_args[1] = { &output_d };
+    EXPECT_EQ(cudaSuccess, cudaLaunchKernel(k, dim3(1), dim3(1), raw_args, 0, 0));
+    EXPECT_EQ(cudaSuccess, cudaMemcpy(&output_h, output_d, sizeof(size_t), cudaMemcpyDeviceToHost));
+    EXPECT_EQ(sizeof(nba::dev_offset_t), 2);
+    EXPECT_EQ(sizeof(nba::dev_offset_t), output_h);
+    EXPECT_EQ(cudaSuccess, cudaFree(output_d));
+    EXPECT_EQ(cudaSuccess, cudaDeviceSynchronize());
+}
+
+TEST(CUDAStructTest, ShfitedIntValueCheck) {
+    void *k = get_test_kernel_shiftedint_value_check();
+    void *input_d;
+    void *output_d;
+    ASSERT_EQ(cudaSuccess, cudaMalloc(&input_d, sizeof(nba::dev_offset_t)));
+    ASSERT_NE(nullptr, input_d);
+    ASSERT_EQ(cudaSuccess, cudaMalloc(&output_d, sizeof(uint64_t)));
+    ASSERT_NE(nullptr, output_d);
+    nba::dev_offset_t input_h = 165321;
+    EXPECT_EQ(165320, input_h.as_value<uint64_t>());
+    size_t output_h = 0;
+    void *raw_args[2] = { &input_d, &output_d };
+    EXPECT_EQ(cudaSuccess, cudaMemcpy(input_d, &input_h, sizeof(nba::dev_offset_t), cudaMemcpyHostToDevice));
+    EXPECT_EQ(cudaSuccess, cudaLaunchKernel(k, dim3(1), dim3(1), raw_args, 0, 0));
+    EXPECT_EQ(cudaSuccess, cudaMemcpy(&output_h, output_d, sizeof(uint64_t), cudaMemcpyDeviceToHost));
+    EXPECT_EQ(165320, output_h);
+    EXPECT_EQ(cudaSuccess, cudaFree(input_d));
+    EXPECT_EQ(cudaSuccess, cudaFree(output_d));
+    EXPECT_EQ(cudaSuccess, cudaDeviceSynchronize());
 }
 
 #else