fix(skiplist): Fix constructor

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

2 files changed, 16 insertions, 10 deletions

diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index 2520a13..4fcd220 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -150,22 +150,26 @@ struct SkipList<Key, Comparator>::Node {
 
   Key const key;
 
-  explicit __device__ Node(const Key& k): key(k) {}
+  explicit __device__ __host__ Node(const Key& k): key(k) {}
 
-  __device__ void* operator new(size_t bytes) {
+  __device__ __host__ void* operator new(size_t bytes) {
     Node * ptr = nullptr;
     cudaMalloc((void**)&ptr, bytes);
     return ptr;
   }
 
-  __device__ void* operator new(size_t bytes, void * const ptr) {
+  __device__ __host__ void* operator new(size_t _bytes, void * const ptr) {
     return ptr;
   }
 
-  __device__ void operator delete(void *ptr) {
+  __device__ __host__ void operator delete(void *ptr) {
       cudaFree(ptr);
   }
 
+  __device__ __host__ void operator delete(void *_ptr, void *_ptr2) {
+    // DO NOTHING HERE
+  }
+
   // Accessors/mutators for links.  Wrapped in methods so we can
   // add the appropriate barriers as necessary.
   __device__ Node* Next(int n) {
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index e83af79..d18f11f 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -287,11 +287,11 @@ __device__ void update_list(SkipList<Key, Comparator> * l, Key key) {
   l->Insert(key);
 }
 
-__global__ void insert_skiplist(int random_seed) {
+__global__ void insert_skiplist(SkipList<Key, Comparator> * skipList, Random *device_rnd) {
   //unsigned row = threadIdx.x + blockIdx.x * blockDim.x;
-  /*for (int i = 0; i < 1000; i++ ) {
-    update_list(&skipList, device_rnd->Next());
-  }*/
+  for (int i = 0; i < 1000; i++) {
+    update_list(skipList, device_rnd->Next());
+  }
 }
 
 __global__ void init(Arena ** pArena, SkipList<Key, Comparator> ** pSkipList) {
@@ -315,10 +315,12 @@ __host__ void host_insert_test() {
 
   dim3 blockSize(1, 1);
 
-  dim3 gridSize(1, 1);
+  dim3 gridSize(32, 32);
 
   init<<<gridSize, blockSize>>>(&pArena, &skipList);
-  insert_skiplist<<<gridSize, blockSize>>>(test::RandomSeed());
+  cudaDeviceSynchronize();
+  insert_skiplist<<<gridSize, blockSize>>>(skipList, device_rnd);
+
 
   cudaDeviceSynchronize();