aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2022-01-09 23:44:59 +0800
committerKunoiSayami <[email protected]>2022-01-09 23:44:59 +0800
commitf7d12c744d3bb264ff0f6611c752d3cd117891cb (patch)
tree4430220730caffeb97ca36d542b5095628ec3149
parentcca8b238147951036574c49d911cea20ea124cac (diff)
test(skiplist): Make sure atomic start with unlock
Signed-off-by: KunoiSayami <[email protected]>
-rw-r--r--db/skiplist_test.cu9
1 files changed, 6 insertions, 3 deletions
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index b7823d0..5233ecd 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -367,8 +367,8 @@ class TestClass {
constexpr size_t SKIPLIST_TEST_SIZE = 10000;
constexpr size_t TEST_STEP = SKIPLIST_TEST_SIZE / 10;
-constexpr unsigned UNLOCKED = 0;
+// source: https://stackoverflow.com/a/22598599
class CudaSpinLock {
static constexpr int UNLOCKED = 0;
static constexpr int LOCKED = 1;
@@ -377,14 +377,15 @@ class CudaSpinLock {
public:
- __device__ __host__ explicit CudaSpinLock(): m_value(0) {}
+ __device__ __host__ explicit CudaSpinLock(): m_value(UNLOCKED) {}
__device__ void lock()
{
while (true)
{
int expected = UNLOCKED;
- if (this->m_value.compare_exchange_strong(expected, LOCKED))
+ //this->m_value.wait(LOCKED);
+ if (this->m_value.compare_exchange_weak(expected, LOCKED))
break;
}
}
@@ -395,12 +396,14 @@ class CudaSpinLock {
}
__device__ bool isLock() {
+ //printf("%d\n", this->m_value.load());
return this->m_value.load() == LOCKED;
}
};
__global__ void testLock() {
CudaSpinLock lock;
+ assert(!lock.isLock());
lock.lock();
assert(lock.isLock());
lock.unlock();