aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2021-12-03 20:22:07 +0800
committerKunoiSayami <[email protected]>2021-12-03 20:22:07 +0800
commita2a9e40d917e3a932e5f563743a0f908f38ac2c6 (patch)
tree25f27381453b3effd914937b7bb58736df5242bc
parent8316f42c232a7cbfa6c77de85d0001d144d61561 (diff)
fix(skiplist): Fix constructor
Signed-off-by: KunoiSayami <[email protected]>
-rw-r--r--db/skiplist.cuh12
-rw-r--r--db/skiplist_test.cu14
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();