diff options
| -rw-r--r-- | db/skiplist_test.cu | 12 | ||||
| -rw-r--r-- | util/arena.cu | 30 | ||||
| -rw-r--r-- | util/arena.cuh | 6 |
3 files changed, 34 insertions, 14 deletions
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu index fc295d7..06675cd 100644 --- a/db/skiplist_test.cu +++ b/db/skiplist_test.cu @@ -404,7 +404,8 @@ constexpr size_t SKIPLIST_TEST_SIZE = BLOCK_COUNT_X * BLOCK_COUNT_Y * TEST_STEP; __global__ void testParallel(SkipList<Key, Comparator> * skipList, Key * keys, CudaSpinLock * lock) { unsigned int start = blockIdx.x * blockDim.x + threadIdx.x; - //printf("start: %u %d %d %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); /*for (unsigned i = start * TEST_STEP; i < (start + 1) * TEST_STEP; i++) { @@ -430,10 +431,17 @@ __global__ void testKeysIsEqualLists(SkipList<Key, Comparator> * skiplist, const SkipList<Key, Comparator>::Iterator iter(skiplist); iter.SeekToFirst(); + bool need_break = false; + int break_after = 20; for (unsigned i = 0; i < SKIPLIST_TEST_SIZE ; i++ ) { assert(iter.Valid()); //printf("%d %lu %lu\n", i, iter.key(), sorted_keys[i]); + if (iter.key() != sorted_keys[i]) { + need_break = true; + } + if (need_break && !--break_after) + break; assert(iter.key() == sorted_keys[i]); iter.Next(); } @@ -541,7 +549,7 @@ TEST(SkipTest, TestSingleCudaInsert) { #endif testSingle<<<1, 1>>>(*skipList, device_keys); cudaDeviceSynchronize(); - printf("%d\n", cudaGetLastError()); + //printf("%d\n", cudaGetLastError() == cudaSuccess); showTimeSpan(start_time); std::sort(sorted_keys, sorted_keys + SKIPLIST_TEST_SIZE); diff --git a/util/arena.cu b/util/arena.cu index 82dbfe6..6d1b76b 100644 --- a/util/arena.cu +++ b/util/arena.cu @@ -62,17 +62,29 @@ __device__ char* Arena::AllocateAligned(size_t bytes) { __device__ char* Arena::AllocateNewBlock(size_t block_bytes) { char* result = nullptr; + cuda::atomic<ArenaNode*>* block_alloc = nullptr; cudaMalloc((void **)&result, sizeof(char) * block_bytes); - if (this->blocks_ == nullptr) { - cudaMalloc((void**)&this->blocks_, sizeof(ArenaNode)); - // First alloc - this->head_ = this->blocks_; - } else { - cudaMalloc((void**)&this->blocks_->next, sizeof(ArenaNode)); - this->blocks_ = this->blocks_->next; + cudaMalloc((void **)&block_alloc, sizeof(cuda::atomic<ArenaNode*>)); + while (true) { + ArenaNode * current_ = this->blocks_.load(cuda::memory_order_acquire); + if (current_ == nullptr) { + // cudaMalloc((void**)&this->blocks_, sizeof(ArenaNode)); + ArenaNode * current_end = this->blocks_.load(); + if (!this->blocks_.compare_exchange_weak(current_end, reinterpret_cast<ArenaNode*>(block_alloc))) + continue; + if (!this->head_.compare_exchange_weak(current_end, reinterpret_cast<ArenaNode*>(block_alloc))) + continue; + break; + } + ArenaNode * except_next = current_->next; + if (except_next != nullptr) + continue; + if (!this->blocks_.compare_exchange_weak(current_, reinterpret_cast<ArenaNode*>(block_alloc))) + continue; + current_->block = result; + current_->next = nullptr; + break; } - this->blocks_->block = result; - this->blocks_->next = nullptr; memory_usage_.fetch_add(block_bytes + sizeof(char*), cuda::memory_order_relaxed); return result; diff --git a/util/arena.cuh b/util/arena.cuh index f9ccb1e..b70dcb9 100644 --- a/util/arena.cuh +++ b/util/arena.cuh @@ -42,7 +42,7 @@ class Arena { struct ArenaNode { char * block; - ArenaNode * next; + cuda::atomic<ArenaNode*> next; }; // Allocation state @@ -53,8 +53,8 @@ class Arena { //thrust::host_vector<char *> blocks_; //std::vector<char*> blocks_; - ArenaNode * head_; - ArenaNode * blocks_; + cuda::atomic<ArenaNode *> head_; + cuda::atomic<ArenaNode *> blocks_; // Total memory usage of the arena. |