aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2021-12-26 00:25:53 +0800
committerKunoiSayami <[email protected]>2021-12-26 00:25:53 +0800
commit0d6e6d17e1c3f01386254613a059b5138fb01e52 (patch)
tree4e4246782b64dae7ef20efab98757d7428e2dd52
parent9406aaf2566cc68e75cc01ba02dfa83b2179d9f6 (diff)
test: Bump version
Signed-off-by: KunoiSayami <[email protected]>
-rw-r--r--db/skiplist.cuh2
-rw-r--r--db/skiplist_test.cu76
2 files changed, 53 insertions, 25 deletions
diff --git a/db/skiplist.cuh b/db/skiplist.cuh
index a285b8e..b12305f 100644
--- a/db/skiplist.cuh
+++ b/db/skiplist.cuh
@@ -170,7 +170,7 @@ struct SkipList<Key, Comparator>::Node {
// DO NOTHING HERE
}
- // Accessors/mutators for links. Wrapped in methods so we can
+ // Accessors/mutators for links. Wrapped in methods, so we can
// add the appropriate barriers as necessary.
__device__ Node* Next(int n) {
assert(n >= 0);
diff --git a/db/skiplist_test.cu b/db/skiplist_test.cu
index a9f7e08..b8fcb5e 100644
--- a/db/skiplist_test.cu
+++ b/db/skiplist_test.cu
@@ -294,25 +294,26 @@ __device__ void ASSERT_EQ_dev(T a, U b) {
assert(a == b);
}
-__global__ void insert_and_lookup(SkipList<Key, Comparator> * skipList) {
+__global__ void insert_and_lookup(SkipList<Key, Comparator> * list) {
+ printf("Lookup success\n");
const int N = 2000;
const int R = 5000;
Random rnd(1000);
MemorySet<Key> keys;
- Arena arena;
- Comparator cmp;
- SkipList<Key, Comparator> list(cmp, &arena);
+ //auto * arena = new Arena();
+ //Comparator cmp;
+ //SkipList<Key, Comparator> list(cmp, &*arena);
for (int i = 0; i < N; i++) {
Key key = rnd.Next() % R;
if (keys.insert(key)) {
- list.Insert(key);
+ list->Insert(key);
}
}
for (int i = 0; i < R; i++) {
- if (list.Contains(i)) {
+ if (list->Contains(i)) {
ASSERT_EQ_dev(keys.count(i), 1);
} else {
ASSERT_EQ_dev(keys.count(i), 0);
@@ -321,9 +322,10 @@ __global__ void insert_and_lookup(SkipList<Key, Comparator> * skipList) {
Node * cur = keys.get_first();
while (cur != nullptr) {
- assert(list.Contains(cur->num));
+ assert(list->Contains(cur->num));
cur = cur->next;
}
+ printf("Lookup success\n");
/*
// Forward iteration test
@@ -391,44 +393,76 @@ class TestClass {
cuda::atomic<size_t> atomic;
};
-__global__ void init(Arena ** pArena, SkipList<Key, Comparator> ** pSkipList) {
+__global__ void init(Arena * pArena, SkipList<Key, Comparator> * pSkipList) {
Comparator cmp;
printf("init\n");
- //cuda::atomic<size_t> a;
- *pArena = new Arena();
+ pArena = new Arena();
+ //new TestClass();
printf("init arena\n");
- //*pSkipList = new SkipList<Key, Comparator>(cmp, reinterpret_cast<Arena*>(&*pArena));
+ pSkipList = new SkipList<Key, Comparator>(cmp, pArena);
printf("init2\n");
+ pSkipList->Insert(16807);
+ printf("skiplist: %p\n", pSkipList);
+ printf("init3\n");
}
+__global__ void testParallel(SkipList<Key, Comparator> * skipList, Key * keys) {
+ unsigned int start = threadIdx.x;
+ printf("skiplist: %p\n", skipList);
+ printf("start: %u\n", start);
+ for (unsigned i = start; i < start + 100 ; i++) {
+ printf("key: %lu\n", keys[i]);
+ skipList->Insert(keys[i]);
+ }
+ printf("done: %u\n", start);
+}
-__host__ void host_insert_test() {
+void host_insert_test() {
//Key * keys;
//SkipList<Key, Comparator> list(cmp, &arena);
/*
for (int i = 0; i < 1000; i++) {
keys[i] = .Next();
}*/
+ Key * keys = new Key[1000];
Arena * pArena;
- cudaMalloc((void**)&pArena, sizeof(*pArena));
SkipList<Key, Comparator> * skipList;
+ std::set<Key> k;
+
+ cudaMallocManaged((void**)&pArena, sizeof(Arena));
+ cudaMallocManaged((void**)&skipList, sizeof(SkipList<Key, Comparator>));
auto * device_rnd = new Random(test::RandomSeed());
+ Key * device_keys = nullptr;
+ cudaMallocManaged((void**)&device_keys, sizeof(Key) * 1000 );
+
+ for (int i = 0; i< 1000; i++) {
+ Key tmp;
+ size_t current = k.size();
+ do {
+ tmp = device_rnd->Next();
+ k.insert(tmp);
+ } while (k.size() == current);
+ keys[i] = tmp;
+ //printf("%ld\n", tmp);
+ }
- dim3 blockSize(1, 1);
+ cudaMemcpy(device_keys, keys, 1000, cudaMemcpyHostToDevice);
dim3 gridSize(1, 1);
+ dim3 blockSize(10, 1);
- init<<<gridSize, blockSize>>>(&pArena, &skipList);
+ init<<<1, 1>>>(pArena, skipList);
+ //sleep(5);
cudaDeviceSynchronize();
//insert_skiplist<<<gridSize, blockSize>>>(skipList, device_rnd);
-
+ testParallel<<<gridSize, 1>>>(skipList, device_keys);
//cudaDeviceSynchronize();
- insert_and_lookup<<<gridSize, blockSize>>>(skipList);
+ //insert_and_lookup<<<gridSize, blockSize>>>(skipList);
cudaDeviceSynchronize();
- std::cout << "test";
+ std::cout << "test\n";
}
@@ -436,9 +470,6 @@ TEST(SkipTest, TestCudaInsert) {
host_insert_test();
}
-#define cudaAssert(condition) \
- if (!(condition)){ printf("Assertion %s failed!\n", #condition); asm("trap;"); }
-
__global__ void test_empty() {
Arena arena;
Comparator cmp;
@@ -459,9 +490,6 @@ __global__ void test_fail() {
assert(0);
}
-CUDA_TEST(SkipTest, Empty) {
-
-}
/*
// Needed when building in C++11 mode.