From e98fc6bf301bb3bf38aab4ccce51edc16953024b Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Thu, 17 Feb 2022 13:28:17 +0800 Subject: --- db/skiplist.cuh | 2 +- db/skiplist_test.cu | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/db/skiplist.cuh b/db/skiplist.cuh index 3f7b7e9..2501af8 100644 --- a/db/skiplist.cuh +++ b/db/skiplist.cuh @@ -453,7 +453,7 @@ __device__ void SkipList::Insert(const Key& key) { // NoBarrier_SetNext() suffices since we will add a barrier when // we publish a pointer to "x" in prev[i]. while (true) { - assert(prev[i]->key < key); + //assert(prev[i]->key < key); if (prev[i]->SetNextSafe(i, x)) break; //printf("failed\n"); diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu index 06675cd..d73e433 100644 --- a/db/skiplist_test.cu +++ b/db/skiplist_test.cu @@ -397,14 +397,14 @@ TEST(SkipTest, TestLock) { cudaDeviceSynchronize(); } -constexpr size_t BLOCK_COUNT_X = 40; -constexpr size_t BLOCK_COUNT_Y = 256; +constexpr size_t BLOCK_COUNT_X = 8; +constexpr size_t BLOCK_COUNT_Y = 128; constexpr size_t TEST_STEP = 1; constexpr size_t SKIPLIST_TEST_SIZE = BLOCK_COUNT_X * BLOCK_COUNT_Y * TEST_STEP; __global__ void testParallel(SkipList * skipList, Key * keys, CudaSpinLock * lock) { unsigned int start = blockIdx.x * blockDim.x + threadIdx.x; - printf("%u\t%d\t%d\t%d\n", start, blockIdx.x, blockDim.x, threadIdx.x); + //printf("%u\t%d\t%d\t%d\n", start, blockIdx.x, blockDim.x, threadIdx.x); //printf("%u\n", start); //lock->lock(); //printf("start insert: %u\n", start); -- cgit v1.3.1