1 files changed, 24 insertions, 5 deletions

diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index a48bd06..8d924f4 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -404,6 +404,25 @@ class CudaSpinLock {
   }
 };
 
+__global__ void testCudaAtomic() {
+  cuda::atomic<void *> node;
+  void * test_point = reinterpret_cast<void*>(0xdeadbeef);
+  void * rep_point = reinterpret_cast<void*>(0xdead);
+  node.store(rep_point);
+  assert(node.compare_exchange_weak(rep_point, test_point, cuda::memory_order_acquire));
+  void * null_ptr = nullptr;
+  node.store(nullptr);
+  assert(node.compare_exchange_weak(null_ptr, test_point, cuda::memory_order_acquire));
+}
+
+TEST(SkipTest, TestCudaAtomic) {
+  testCudaAtomic<<<1, 1>>>();
+  cudaDeviceSynchronize();
+  int last_error = cudaGetLastError();
+  if (last_error != cudaSuccess) {
+    printf("Got error: %d\n", last_error);
+  }
+}
 
 __global__ void testLock() {
   CudaSpinLock lock;
@@ -421,16 +440,16 @@ TEST(SkipTest, TestLock) {
 
 __global__ void testParallel(SkipList<Key, Comparator> * skipList, Key * keys, CudaSpinLock * lock) {
   unsigned int start = threadIdx.x;
-  //printf("start: %u\n", start);
-  lock->lock();
+  printf("start: %u\n", start);
+  //lock->lock();
   //printf("start insert: %u\n", start);
   for (unsigned i = start * TEST_STEP; i < (start + 1) * TEST_STEP; i++) {
-    //printf("%u %02u %lu\n", start, i, keys[i]);
+    printf("%u %02u %lu\n", start, i, keys[i]);
     //printf("key: %lu\n", keys[i]);
     skipList->Insert(keys[i]);
   }
-  lock->unlock();
-  //printf("done: %u\n", start);
+  //lock->unlock();
+  printf("done: %u\n", start);
 }
 
 __global__ void testSingle(SkipList<Key, Comparator>* skipList, Key * keys) {