aboutsummaryrefslogtreecommitdiff
path: root/db/skiplist_test.cu
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2022-01-09 23:34:26 +0800
committerKunoiSayami <[email protected]>2022-01-09 23:34:26 +0800
commitcca8b238147951036574c49d911cea20ea124cac (patch)
tree55067271662e32be09bbef02e9ae266281eb765e /db/skiplist_test.cu
parent8110ac13bd7d31c462c1f2607d29253bb5a34a92 (diff)
test(skiplist): Performing parallel insertion test
Signed-off-by: KunoiSayami <[email protected]>
Diffstat (limited to 'db/skiplist_test.cu')
-rw-r--r--db/skiplist_test.cu169
1 files changed, 126 insertions, 43 deletions
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index b5511f2..b7823d0 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -35,21 +35,7 @@ struct Comparator {
}
};
-/*TEST(SkipTest, Empty) {
- Arena arena;
- Comparator cmp;
- SkipList<Key, Comparator> list(cmp, &arena);
- ASSERT_TRUE(!list.Contains(10));
-
- SkipList<Key, Comparator>::Iterator iter(&list);
- ASSERT_TRUE(!iter.Valid());
- iter.SeekToFirst();
- ASSERT_TRUE(!iter.Valid());
- iter.Seek(100);
- ASSERT_TRUE(!iter.Valid());
- iter.SeekToLast();
- ASSERT_TRUE(!iter.Valid());
-}
+/*
TEST(SkipTest, InsertAndLookup) {
}
@@ -211,17 +197,6 @@ class ConcurrentTest {
};*/
-__device__ void update_list(SkipList<Key, Comparator> * l, Key key) {
- l->Insert(key);
-}
-
-__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());
- }
-}
-
struct Node {
Key num;
Node* next;
@@ -367,9 +342,6 @@ __global__ void insert_and_lookup(SkipList<Key, Comparator> * list) {
}
-__device__ void new_cuda_atomic() {
-}
-
class TestClass {
public:
explicit __device__ TestClass(): atomic(0), alloc_ptr_(nullptr), alloc_bytes_remaining_(0),
@@ -395,16 +367,63 @@ class TestClass {
constexpr size_t SKIPLIST_TEST_SIZE = 10000;
constexpr size_t TEST_STEP = SKIPLIST_TEST_SIZE / 10;
+constexpr unsigned UNLOCKED = 0;
+
+class CudaSpinLock {
+ static constexpr int UNLOCKED = 0;
+ static constexpr int LOCKED = 1;
+
+ cuda::atomic<int> m_value;
+
+ public:
+
+ __device__ __host__ explicit CudaSpinLock(): m_value(0) {}
-__global__ void testParallel(SkipList<Key, Comparator> * skipList, Key * keys) {
+ __device__ void lock()
+ {
+ while (true)
+ {
+ int expected = UNLOCKED;
+ if (this->m_value.compare_exchange_strong(expected, LOCKED))
+ break;
+ }
+ }
+
+ __device__ void unlock()
+ {
+ m_value.store(UNLOCKED);
+ }
+
+ __device__ bool isLock() {
+ return this->m_value.load() == LOCKED;
+ }
+};
+
+__global__ void testLock() {
+ CudaSpinLock lock;
+ lock.lock();
+ assert(lock.isLock());
+ lock.unlock();
+ assert(!lock.isLock());
+}
+
+TEST(SkipTest, TestLock) {
+ testLock<<<1, 1>>>();
+ cudaDeviceSynchronize();
+}
+
+__global__ void testParallel(SkipList<Key, Comparator> * skipList, Key * keys, CudaSpinLock * lock) {
unsigned int start = threadIdx.x;
- printf("start: %u\n", start);
+ //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("key: %lu\n", keys[i]);
skipList->Insert(keys[i]);
}
- printf("done: %u\n", start);
+ lock->unlock();
+ //printf("done: %u\n", start);
}
__global__ void testSingle(SkipList<Key, Comparator>* skipList, Key * keys) {
@@ -420,6 +439,7 @@ __global__ void testKeysIsEqualLists(SkipList<Key, Comparator> * skiplist, const
iter.SeekToFirst();
for (unsigned i = 0; i < SKIPLIST_TEST_SIZE ; i++ ) {
+ assert(iter.Valid());
assert(iter.key() == sorted_keys[i]);
iter.Next();
}
@@ -430,9 +450,6 @@ __global__ void initSkipList(Arena ** pArena, SkipList<Key, Comparator> ** pSkip
Comparator cmp;
*pArena = new Arena();
*pSkipList = new SkipList<Key, Comparator>(cmp, *pArena);
-
- printf("%p %p \n", *pArena, *pSkipList);
-
}
__global__ void freeSkipList(Arena *** pArena, SkipList<Key, Comparator> *** pSkipList) {
@@ -460,7 +477,12 @@ TEST(SkipTest, TestInitSkiplist) {
cudaFree(pSkipList);
}
-TEST(SkipTest, TestCudaInsert) {
+__global__ void resetLock(CudaSpinLock * lock) {
+ assert(!lock->isLock());
+ lock->unlock();
+}
+
+TEST(SkipTest, TestSingleCudaInsert) {
//Key * keys;
//SkipList<Key, Comparator> list(cmp, &arena);
/*
@@ -519,12 +541,78 @@ TEST(SkipTest, TestCudaInsert) {
testKeysIsEqualLists<<<1, 1>>>(*skipList, device_keys);
//insert_and_lookup<<<gridSize, blockSize>>>(skipList);
cudaDeviceSynchronize();
+ cudaFree(*skipList);
+ cudaFree(*pArena);
cudaFree(device_keys);
cudaFree(skipList);
cudaFree(pArena);
+ delete [] sorted_keys;
+ delete [] keys;
+}
+
+TEST(SkipTest, TestMultiThreadInsert) {
+ Key * keys = new Key[SKIPLIST_TEST_SIZE], * sorted_keys = new Key[SKIPLIST_TEST_SIZE];
+ Arena ** pArena;
+ SkipList<Key, Comparator> ** pSkipList;
+ std::set<Key> k;
+ Key * device_keys = nullptr;
+
+ cudaMallocManaged((void**)&pArena, sizeof(void*));
+ cudaMallocManaged((void**)&pSkipList, sizeof(void*));
+ auto * device_rnd = new Random(test::RandomSeed());
+ cudaMallocManaged((void**)&device_keys, sizeof(Key) * SKIPLIST_TEST_SIZE );
+
+
+ //cuda::atomic<unsigned int> lock;
+ CudaSpinLock lock;
+ CudaSpinLock * device_lock = nullptr;
+
+ cudaMallocManaged((void**)&device_lock, sizeof(CudaSpinLock));
+ cudaMemcpy(device_lock, &lock, sizeof(CudaSpinLock), cudaMemcpyHostToDevice);
+
+ resetLock<<<1, 1>>>(device_lock);
+ cudaDeviceSynchronize();
+
+ for (int i = 0; i < SKIPLIST_TEST_SIZE; i++) {
+ Key tmp;
+ size_t current = k.size();
+ do {
+ tmp = device_rnd->Next();
+ k.insert(tmp);
+ } while (k.size() == current);
+
+ keys[i] = tmp;
+ }
+ memcpy(sorted_keys, keys, SKIPLIST_TEST_SIZE * sizeof(Key));
+ cudaMemcpy(device_keys, keys, SKIPLIST_TEST_SIZE * sizeof(Key), cudaMemcpyHostToDevice);
+ dim3 gridSize(1, 1);
+ dim3 blockSize(10, 1);
+
+ initSkipList<<<1, 1>>>(pArena, pSkipList);
+ //sleep(5);
+ cudaDeviceSynchronize();
+ //insert_skiplist<<<gridSize, blockSize>>>(skipList, device_rnd);
+
+ //testParallel<<<gridSize, blockSize>>>(*skipList, device_keys);
+ testParallel<<<1, blockSize>>>(*pSkipList, device_keys, device_lock);
+ cudaDeviceSynchronize();
+
+ std::sort(sorted_keys, sorted_keys + SKIPLIST_TEST_SIZE);
+ cudaMemcpy(device_keys, sorted_keys, SKIPLIST_TEST_SIZE * sizeof(Key), cudaMemcpyHostToDevice);
+ testKeysIsEqualLists<<<1, 1>>>(*pSkipList, device_keys);
+ //insert_and_lookup<<<gridSize, blockSize>>>(skipList);
+ cudaDeviceSynchronize();
+ cudaFree(*pSkipList);
+ cudaFree(*pArena);
+ cudaFree(device_keys);
+ cudaFree(pSkipList);
+ cudaFree(pArena);
+ cudaFree(device_lock);
+ delete [] sorted_keys;
+ delete [] keys;
}
-__global__ void test_empty() {
+__global__ void globalTestEmpty() {
Arena arena;
Comparator cmp;
SkipList<Key, Comparator> list(cmp, &arena);
@@ -540,15 +628,10 @@ __global__ void test_empty() {
}
TEST(SkipTest, TestCudaEmpty) {
- test_empty<<<1, 1>>>();
+ globalTestEmpty<<<1, 1>>>();
cudaDeviceSynchronize();
}
-__global__ void test_fail() {
- std::printf("show some message");
- assert(0);
-}
-
/*
// Needed when building in C++11 mode.