2 files changed, 53 insertions, 25 deletions

diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index a285b8e..b12305f 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -170,7 +170,7 @@ struct SkipList<Key, Comparator>::Node {
     // DO NOTHING HERE
   }
 
-  // Accessors/mutators for links.  Wrapped in methods so we can
+  // Accessors/mutators for links. Wrapped in methods, so we can
   // add the appropriate barriers as necessary.
   __device__ Node* Next(int n) {
     assert(n >= 0);
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index a9f7e08..b8fcb5e 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -294,25 +294,26 @@ __device__ void ASSERT_EQ_dev(T a, U b) {
   assert(a == b);
 }
 
-__global__ void insert_and_lookup(SkipList<Key, Comparator> * skipList) {
+__global__ void insert_and_lookup(SkipList<Key, Comparator> * list) {
 
+  printf("Lookup success\n");
   const int N = 2000;
   const int R = 5000;
   Random rnd(1000);
 
   MemorySet<Key> keys;
-  Arena arena;
-  Comparator cmp;
-  SkipList<Key, Comparator> list(cmp, &arena);
+  //auto * arena = new Arena();
+  //Comparator cmp;
+  //SkipList<Key, Comparator> list(cmp, &*arena);
   for (int i = 0; i < N; i++) {
     Key key = rnd.Next() % R;
     if (keys.insert(key)) {
-      list.Insert(key);
+      list->Insert(key);
     }
   }
 
   for (int i = 0; i < R; i++) {
-    if (list.Contains(i)) {
+    if (list->Contains(i)) {
       ASSERT_EQ_dev(keys.count(i), 1);
     } else {
       ASSERT_EQ_dev(keys.count(i), 0);
@@ -321,9 +322,10 @@ __global__ void insert_and_lookup(SkipList<Key, Comparator> * skipList) {
 
   Node * cur = keys.get_first();
   while (cur != nullptr) {
-    assert(list.Contains(cur->num));
+    assert(list->Contains(cur->num));
     cur = cur->next;
   }
+  printf("Lookup success\n");
 /*
 
   // Forward iteration test
@@ -391,44 +393,76 @@ class TestClass {
   cuda::atomic<size_t> atomic;
 };
 
-__global__ void init(Arena ** pArena, SkipList<Key, Comparator> ** pSkipList) {
+__global__ void init(Arena * pArena, SkipList<Key, Comparator> * pSkipList) {
   Comparator cmp;
   printf("init\n");
-  //cuda::atomic<size_t> a;
-  *pArena = new Arena();
+  pArena = new Arena();
+  //new TestClass();
   printf("init arena\n");
-  //*pSkipList = new SkipList<Key, Comparator>(cmp, reinterpret_cast<Arena*>(&*pArena));
+  pSkipList = new SkipList<Key, Comparator>(cmp, pArena);
   printf("init2\n");
+  pSkipList->Insert(16807);
+  printf("skiplist: %p\n", pSkipList);
+  printf("init3\n");
 }
 
+__global__ void testParallel(SkipList<Key, Comparator> * skipList, Key * keys) {
+  unsigned int start = threadIdx.x;
+  printf("skiplist: %p\n", skipList);
+  printf("start: %u\n", start);
+  for (unsigned i = start; i < start + 100 ; i++) {
+    printf("key: %lu\n", keys[i]);
+    skipList->Insert(keys[i]);
+  }
+  printf("done: %u\n", start);
+}
 
 
-__host__ void host_insert_test() {
+void host_insert_test() {
   //Key * keys;
   //SkipList<Key, Comparator> list(cmp, &arena);
   /*
   for (int i = 0; i < 1000; i++) {
     keys[i] = .Next();
   }*/
+  Key * keys = new Key[1000];
   Arena * pArena;
-  cudaMalloc((void**)&pArena, sizeof(*pArena));
   SkipList<Key, Comparator> * skipList;
+  std::set<Key> k;
+
+  cudaMallocManaged((void**)&pArena, sizeof(Arena));
+  cudaMallocManaged((void**)&skipList, sizeof(SkipList<Key, Comparator>));
   auto * device_rnd = new Random(test::RandomSeed());
+  Key * device_keys = nullptr;
+  cudaMallocManaged((void**)&device_keys, sizeof(Key) * 1000 );
+
+  for (int i = 0; i< 1000; i++) {
+    Key tmp;
+    size_t current = k.size();
+    do {
+      tmp = device_rnd->Next();
+      k.insert(tmp);
+    } while (k.size() == current);
+    keys[i] = tmp;
+    //printf("%ld\n", tmp);
+  }
 
-  dim3 blockSize(1, 1);
+  cudaMemcpy(device_keys, keys, 1000, cudaMemcpyHostToDevice);
 
   dim3 gridSize(1, 1);
+  dim3 blockSize(10, 1);
 
-  init<<<gridSize, blockSize>>>(&pArena, &skipList);
+  init<<<1, 1>>>(pArena, skipList);
+  //sleep(5);
   cudaDeviceSynchronize();
   //insert_skiplist<<<gridSize, blockSize>>>(skipList, device_rnd);
 
-
+  testParallel<<<gridSize, 1>>>(skipList, device_keys);
   //cudaDeviceSynchronize();
 
-  insert_and_lookup<<<gridSize, blockSize>>>(skipList);
+  //insert_and_lookup<<<gridSize, blockSize>>>(skipList);
   cudaDeviceSynchronize();
-  std::cout << "test";
+  std::cout << "test\n";
 }
 
 
@@ -436,9 +470,6 @@ TEST(SkipTest, TestCudaInsert) {
   host_insert_test();
 }
 
-#define cudaAssert(condition) \
-  if (!(condition)){ printf("Assertion %s failed!\n", #condition); asm("trap;"); }
-
 __global__ void test_empty() {
   Arena arena;
   Comparator cmp;
@@ -459,9 +490,6 @@ __global__ void test_fail() {
   assert(0);
 }
 
-CUDA_TEST(SkipTest, Empty) {
-
-}
 
 /*
 // Needed when building in C++11 mode.