summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2023-05-28 23:20:31 +0800
committerKunoiSayami <[email protected]>2023-05-28 23:20:31 +0800
commit9c0ce2f5570dedd5ad68c7e8fc222928046c5908 (patch)
tree6ee48dfbd2bfb7607f87dcc3d96c332043327402
parent44a035ec3ebb094adb162ba0c0d374a3fe153094 (diff)
fix(exp): Fix search action
Signed-off-by: KunoiSayami <[email protected]>
-rw-r--r--expt_0525.cu63
-rw-r--r--sortlib.cuh8
2 files changed, 45 insertions, 26 deletions
diff --git a/expt_0525.cu b/expt_0525.cu
index b4e4da5..cff4f73 100644
--- a/expt_0525.cu
+++ b/expt_0525.cu
@@ -65,20 +65,21 @@ Conference on Parallel and Distributed Systems, December 2012.
// #include"cutil.h" // Comment this if cutil.h is not available
// #include "cuda_runtime.h"
#include "sortlib.cuh"
+#include <algorithm>
#include <cassert>
#include <cstdio>
#include <cstdlib>
#include <random>
#include <set>
-#define READ_NO_OUTPUT
-#include "read_helper.h"
-
typedef unsigned long long LL;
// #define MEASURE_TIME
// #define MEASURE_ACCESS
+#define READ_NO_OUTPUT
+#include "read_helper.h"
+
#if (defined(MEASURE_ACCESS) && defined(MEASURE_TIME))
#error "Shouldn't define MEASURE_TIME and MEASURE_ACCESS at the same time"
#endif
@@ -260,12 +261,14 @@ class LockFreeSkipList {
CustomSort customSort;
+ double scaleSize;
+
public:
Node *head = nullptr;
- Node *tail = nullptr;
- LockFreeSkipList(key_type *_sample, size_t sample_length)
+ Node *tail;
+ LockFreeSkipList(key_type *_sample, size_t sample_length, double scale_size)
: sampleLength(sample_length),
- customSort(sample_length, sizeof(key_type) * 8) {
+ customSort(sample_length, sizeof(key_type) * 8), scaleSize(scale_size) {
Node *h = new Node(0);
// size_ = 0;
Node *t = new Node(std::numeric_limits<key_type>::max() - 1);
@@ -309,6 +312,13 @@ public:
return bits;
}
+ __device__ unsigned calcIndex(key_type key) {
+ auto result = customSort.sample_cdf_custom_version(this->sample, key);
+ auto cdf_index = result / scaleSize;
+ auto index = trailing_zeroes((size_t)cdf_index);
+ return index;
+ }
+
#ifdef MEASURE_ACCESS
unsigned access_times = 0;
@@ -588,6 +598,7 @@ __global__ void kernel(const LL *items, size_t search_length, LL *result) {
unsigned long long start_time = clock64();
#endif
result[tid] = lockFreeSkipList->Search(item);
+ assert(result[tid]);
#ifdef MEASURE_TIME
unsigned long long end_time = clock64() - start_time;
if (lockFreeSkipList->spend_time[tid]) {
@@ -611,14 +622,6 @@ __global__ void kernelAdd(LL *item, size_t insertion_length) {
}
}
-/*LL Randomlevel() {
- LL v = 1;
- double p = 0.5;
- while (((rand() / (double)(RAND_MAX)) < p) && (v < MAX_LEVEL))
- v++;
- return v;
-}*/
-
// Generate the level of a newly created node
__global__ void print_function() {
@@ -631,6 +634,10 @@ __global__ void print_function() {
memcpy(spend_time, lockFreeSkipList->spend_time, sizeof(int) * NUM_ITEMS);
}*/
+inline double calcSliceSize(size_t insertion_size) {
+ return 1.0 / (double)insertion_size;
+}
+
inline size_t calcBlocks(size_t input) {
return (input % (NUM_THREADS * FACTOR) == 0)
? input / (NUM_THREADS * FACTOR)
@@ -639,27 +646,32 @@ inline size_t calcBlocks(size_t input) {
int main(int argc, char **argv) {
if (argc != 4) {
- printf("Usage %s [sample] [insertion] [search]\n", argv[0]);
+ printf("Usage %s [sample] [search] [insertion]\n", argv[0]);
exit(1);
}
auto sample_length = strtol(argv[1], nullptr, 10);
- auto insertion_length = strtol(argv[2], nullptr, 10);
- auto search_length = strtol(argv[3], nullptr, 10);
+ auto search_length = strtol(argv[2], nullptr, 10);
+ auto insertion_length = strtol(argv[3], nullptr, 10);
+
+ if (insertion_length < sample_length) {
+ printf("Search should smaller than insertion\n");
+ }
+
printf("Sample: %ld, Insertion: %ld, Search: %ld\n", sample_length,
insertion_length, search_length);
ReadHelper readHelper("normal_distribution.txt", sample_length,
- insertion_length + search_length);
+ insertion_length);
readHelper.readFile(nullptr);
std::vector<key_type> _sample, _population;
readHelper.split_into(_sample, _population);
/*printf("%lu, Create search vector: %ld\n", _population.size(),
_population.end() - (_population.begin() + insertion_length));*/
- rebuild(_sample);
- std::vector<key_type> _search(_population.begin() + insertion_length,
- _population.end());
- _population.resize(insertion_length);
+ rebuildSort(_sample);
+ std::vector<key_type> _search(_population.begin(),
+ _population.begin() + search_length);
+ //_population.resize(insertion_length);
// Allocate necessary arrays
// LL *op = new LL[NUM_ITEMS]; //(LL *)malloc(sizeof(LL) * NUM_ITEMS);
@@ -673,8 +685,8 @@ int main(int argc, char **argv) {
// LL *Cop;
LL *cudaResult;
- cudaMalloc(&cudaResult, sizeof(LL) * search_length);
- cudaMalloc(&cudaOperatorItems, sizeof(LL) * insertion_length);
+ cudaMalloc(&cudaResult, sizeof(key_type) * search_length);
+ cudaMalloc(&cudaOperatorItems, sizeof(key_type) * insertion_length);
// cudaMalloc(&Cop, sizeof(LL) * NUM_ITEMS);
// cudaMemcpy(Clevels, levels, sizeof(LL) * NUM_ITEMS,
// cudaMemcpyHostToDevice);
@@ -697,7 +709,8 @@ int main(int argc, char **argv) {
// Allocate the skip list
LockFreeSkipList *Clist;
- auto *list = new LockFreeSkipList(_sample.data(), _sample.size());
+ auto *list = new LockFreeSkipList(_sample.data(), _sample.size(),
+ calcSliceSize(insertion_length));
cudaMalloc(&Clist, sizeof(LockFreeSkipList));
cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
diff --git a/sortlib.cuh b/sortlib.cuh
index 6b19d3e..a7d3e2d 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -1,6 +1,7 @@
#ifndef LOCKFREE_SORTLIB_CUH
#define LOCKFREE_SORTLIB_CUH
+#include <algorithm>
#include <cassert>
#include <cstdio>
#include <vector>
@@ -182,7 +183,7 @@ public:
__host__ __device__ size_t length() const { return this->LENGTH; }
};
-template <typename T> void rebuild(std::vector<T> original) {
+template <typename T> void rebuild(std::vector<T> &original) {
auto sample_length = original.size();
auto sorter = CustomSort(sample_length, sizeof(T) * 8);
auto tmp = new T[sample_length];
@@ -195,6 +196,11 @@ template <typename T> void rebuild(std::vector<T> original) {
delete[] tmp;
}
+template <typename T> void rebuildSort(std::vector<T> &original) {
+ std::sort(original.begin(), original.end());
+ rebuild(original);
+}
+
__global__ void testCustomCalculation(size_t length) {
CustomSort(length, sizeof(long) * 8).testCalculation();
}