test(skiplist): Make sure atomic start with unlock

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

1 files changed, 6 insertions, 3 deletions

diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index b7823d0..5233ecd 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -367,8 +367,8 @@ class TestClass {
 
 constexpr size_t SKIPLIST_TEST_SIZE = 10000;
 constexpr size_t TEST_STEP = SKIPLIST_TEST_SIZE / 10;
-constexpr unsigned UNLOCKED = 0;
 
+// source: https://stackoverflow.com/a/22598599
 class CudaSpinLock {
   static constexpr int UNLOCKED = 0;
   static constexpr int LOCKED = 1;
@@ -377,14 +377,15 @@ class CudaSpinLock {
 
  public:
 
-  __device__ __host__ explicit CudaSpinLock(): m_value(0) {}
+  __device__ __host__ explicit CudaSpinLock(): m_value(UNLOCKED) {}
 
   __device__ void lock()
   {
     while (true)
     {
       int expected = UNLOCKED;
-      if (this->m_value.compare_exchange_strong(expected, LOCKED))
+      //this->m_value.wait(LOCKED);
+      if (this->m_value.compare_exchange_weak(expected, LOCKED))
         break;
     }
   }
@@ -395,12 +396,14 @@ class CudaSpinLock {
   }
 
   __device__ bool isLock() {
+    //printf("%d\n", this->m_value.load());
     return this->m_value.load() == LOCKED;
   }
 };
 
 __global__ void testLock() {
   CudaSpinLock lock;
+  assert(!lock.isLock());
   lock.lock();
   assert(lock.isLock());
   lock.unlock();