test(skiplist): Add option for CudaSpinLock

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

2 files changed, 38 insertions, 58 deletions

diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index b12305f..a6f7bb7 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -27,7 +27,6 @@
 //
 // ... prev vs. next pointer ordering ...
 
-#include <atomic>
 #include <cassert>
 #include <cstdlib>
 #include <cuda/atomic>
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index 5233ecd..a48bd06 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -197,6 +197,27 @@ class ConcurrentTest {
 };*/
 
 
+__global__ void globalTestEmpty() {
+  Arena arena;
+  Comparator cmp;
+  SkipList<Key, Comparator> list(cmp, &arena);
+  assert(!list.Contains(10));
+  SkipList<Key, Comparator>::Iterator iter(&list);
+  assert(!iter.Valid());
+  iter.SeekToFirst();
+  assert(!iter.Valid());
+  iter.Seek(100);
+  assert(!iter.Valid());
+  iter.SeekToLast();
+  assert(!iter.Valid());
+}
+
+TEST(SkipTest, TestLinkListEmpty) {
+  globalTestEmpty<<<1, 1>>>();
+  cudaDeviceSynchronize();
+}
+
+
 struct Node {
   Key num;
   Node* next;
@@ -342,29 +363,6 @@ __global__ void insert_and_lookup(SkipList<Key, Comparator> * list) {
 }
 
 
-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;
-};
-
 constexpr size_t SKIPLIST_TEST_SIZE = 10000;
 constexpr size_t TEST_STEP = SKIPLIST_TEST_SIZE / 10;
 
@@ -374,33 +372,39 @@ class CudaSpinLock {
   static constexpr int LOCKED = 1;
 
   cuda::atomic<int> m_value;
+  bool isFake;
 
  public:
 
-  __device__ __host__ explicit CudaSpinLock(): m_value(UNLOCKED) {}
+  __device__ __host__ explicit CudaSpinLock(): m_value(UNLOCKED), isFake(false) {}
+
+  __device__ __host__ explicit CudaSpinLock(bool fake): m_value(UNLOCKED), isFake(fake) {}
 
   __device__ void lock()
   {
-    while (true)
-    {
-      int expected = UNLOCKED;
-      //this->m_value.wait(LOCKED);
-      if (this->m_value.compare_exchange_weak(expected, LOCKED))
-        break;
+    if (!isFake) {
+      while (true) {
+        int expected = UNLOCKED;
+        // this->m_value.wait(LOCKED);
+        if (this->m_value.compare_exchange_weak(expected, LOCKED)) break;
+      }
     }
   }
 
   __device__ void unlock()
   {
-    m_value.store(UNLOCKED);
+    if (!isFake) {
+      m_value.store(UNLOCKED);
+    }
   }
 
   __device__ bool isLock() {
     //printf("%d\n", this->m_value.load());
-    return this->m_value.load() == LOCKED;
+    return !isFake && this->m_value.load() == LOCKED;
   }
 };
 
+
 __global__ void testLock() {
   CudaSpinLock lock;
   assert(!lock.isLock());
@@ -527,8 +531,6 @@ TEST(SkipTest, TestSingleCudaInsert) {
   memcpy(sorted_keys, keys, SKIPLIST_TEST_SIZE * sizeof(Key));
   cudaMemcpy(device_keys, keys, SKIPLIST_TEST_SIZE * sizeof(Key), cudaMemcpyHostToDevice);
 
-  dim3 gridSize(1, 1);
-  dim3 blockSize(1, 1);
 
   initSkipList<<<1, 1>>>(pArena, skipList);
   //sleep(5);
@@ -567,7 +569,7 @@ TEST(SkipTest, TestMultiThreadInsert) {
 
 
     //cuda::atomic<unsigned int> lock;
-  CudaSpinLock lock;
+  CudaSpinLock lock(true);
   CudaSpinLock * device_lock = nullptr;
 
   cudaMallocManaged((void**)&device_lock, sizeof(CudaSpinLock));
@@ -597,7 +599,7 @@ TEST(SkipTest, TestMultiThreadInsert) {
   //insert_skiplist<<<gridSize, blockSize>>>(skipList, device_rnd);
 
   //testParallel<<<gridSize, blockSize>>>(*skipList, device_keys);
-  testParallel<<<1, blockSize>>>(*pSkipList, device_keys, device_lock);
+  testParallel<<<gridSize, blockSize>>>(*pSkipList, device_keys, device_lock);
   cudaDeviceSynchronize();
 
   std::sort(sorted_keys, sorted_keys + SKIPLIST_TEST_SIZE);
@@ -615,27 +617,6 @@ TEST(SkipTest, TestMultiThreadInsert) {
   delete [] keys;
 }
 
-__global__ void globalTestEmpty() {
-  Arena arena;
-  Comparator cmp;
-  SkipList<Key, Comparator> list(cmp, &arena);
-  assert(!list.Contains(10));
-  SkipList<Key, Comparator>::Iterator iter(&list);
-  assert(!iter.Valid());
-  iter.SeekToFirst();
-  assert(!iter.Valid());
-  iter.Seek(100);
-  assert(!iter.Valid());
-  iter.SeekToLast();
-  assert(!iter.Valid());
-}
-
-TEST(SkipTest, TestCudaEmpty) {
-  globalTestEmpty<<<1, 1>>>();
-    cudaDeviceSynchronize();
-}
-
-
 /*
 // Needed when building in C++11 mode.
 constexpr uint32_t ConcurrentTest::K;