diff options
Diffstat (limited to 'db/skiplist_test.cu')
| -rw-r--r-- | db/skiplist_test.cu | 169 |
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. |