aboutsummaryrefslogtreecommitdiff
path: root/db/skiplist_test.cu
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2022-01-19 15:47:20 +0800
committerKunoiSayami <[email protected]>2022-01-19 15:47:20 +0800
commitb0d2dff85a05a5ced0f910e220daef108358289e (patch)
tree3f6fe9ef2b3ee7aaf26086cceb0fb6d7b5962d52 /db/skiplist_test.cu
parentdfa8c5c43d1d473c2106885bbde86cd1e66a78f9 (diff)
Diffstat (limited to 'db/skiplist_test.cu')
-rw-r--r--db/skiplist_test.cu29
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) {