summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--CMakeLists.txt6
-rw-r--r--expt_0525.cu2
-rw-r--r--expt_0529.cu71
-rw-r--r--read_helper.h23
-rw-r--r--sortlib.cuh10
5 files changed, 97 insertions, 15 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 18b9530..6fb4891 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -1,4 +1,4 @@
-cmake_minimum_required(VERSION 3.22)
+cmake_minimum_required(VERSION 3.24)
project(lockfree CUDA CXX)
set(CMAKE_CUDA_STANDARD 14)
@@ -174,4 +174,8 @@ set_cuda_target_base(expt_0525)
add_executable(expt_0528 expt_0528.cu read_helper.h)
set_cuda_target_base(expt_0528)
+add_executable(expt_0529 expt_0529.cu read_helper.h)
+set_cuda_target_base(expt_0529)
+
+
diff --git a/expt_0525.cu b/expt_0525.cu
index 36c8df6..8e803e0 100644
--- a/expt_0525.cu
+++ b/expt_0525.cu
@@ -264,7 +264,7 @@ class LockFreeSkipList {
public:
Node *head = nullptr;
- Node *tail;
+ Node *tail = nullptr;
LockFreeSkipList(key_type *_sample, size_t sample_length, double scale_size)
: sampleLength(sample_length),
customSort(sample_length, sizeof(key_type) * 8), scaleSize(scale_size) {
diff --git a/expt_0529.cu b/expt_0529.cu
new file mode 100644
index 0000000..11524a7
--- /dev/null
+++ b/expt_0529.cu
@@ -0,0 +1,71 @@
+#include "sortlib.cuh"
+#define READ_NO_OUTPUT
+#include "read_helper.h"
+
+long pow_for_sample(long n) {
+ auto x = 2;
+ for (int i = 1; i < n; i++) {
+ x *= 2;
+ }
+ return x - 1;
+}
+
+static unsigned trailing_zeroes(size_t index) {
+ constexpr auto block_size = 2;
+ unsigned bits = 0;
+ auto x = index / block_size;
+
+ if (x) {
+ while (x % block_size == 0) {
+ ++bits;
+ x /= block_size;
+ }
+ }
+ return bits;
+}
+
+inline double calcSliceSize(size_t insertion_size) {
+ return 1.0 / (double)insertion_size;
+}
+
+constexpr auto RESULT_LENGTH = 35;
+
+int main(int argc, char const *argv[]) {
+
+ if (argc != 3) {
+ printf("Usage %s [sample(pow)] [population]\n", argv[0]);
+ return 1;
+ }
+
+ auto sample_length = pow_for_sample(strtol(argv[1], nullptr, 10));
+ auto population_length = strtol(argv[2], nullptr, 10);
+
+ printf("sample length: %ld, population length: %ld\n", sample_length,
+ population_length);
+
+ ReadHelper readHelper("normal_distribution.txt", sample_length, 0);
+ readHelper.readFile(nullptr);
+
+ std::vector<key_type> sample, population;
+ std::vector<unsigned int> result(RESULT_LENGTH);
+ readHelper.split_into(sample, population);
+ rebuildSort(sample);
+
+ auto scale_size = calcSliceSize(population_length);
+
+ auto sort = CustomSort(sample_length, sizeof(key_type) * 8);
+
+ for (auto element : readHelper.population_vector) {
+ auto ret = sort.sample_cdf_custom_version(sample.data(), element);
+ auto cdf_index = ret / scale_size;
+ auto index = trailing_zeroes((size_t)cdf_index);
+ result[index]++;
+ }
+
+ for (int i = 0; i < 32; i++) {
+ if (!result[i]) {
+ continue;
+ }
+ printf("%d: %d\n", i, result[i]);
+ }
+}
diff --git a/read_helper.h b/read_helper.h
index eb7d2e3..2e87b0b 100644
--- a/read_helper.h
+++ b/read_helper.h
@@ -5,8 +5,9 @@
#include <cstdio>
#include <random>
#include <vector>
+
#ifndef P_ERR
-#if !(defined(NDEBUG) || defined(READ_NO_OUTPUT))
+#ifdef PRINT_READ_PROCESS
#define P_ERR(...) fprintf(stderr, __VA_ARGS__)
#else
#define P_ERR(...)
@@ -20,6 +21,16 @@ class ReadHelper {
key_type max_value = std::numeric_limits<key_type>::min(),
min_value = std::numeric_limits<key_type>::max();
+ inline void store_into_vector(key_type value) {
+ if (max_value < value) {
+ max_value = value;
+ }
+ if (min_value > value) {
+ min_value = value;
+ }
+ population_vector.push_back(value);
+ }
+
char const *filename;
static constexpr size_t REVERSED_BLOCK = 256;
@@ -35,16 +46,6 @@ public:
key_type minValue() const { return min_value; }
std::vector<key_type> population_vector;
- inline void store_into_vector(key_type value) {
- if (max_value < value) {
- max_value = value;
- }
- if (min_value > value) {
- min_value = value;
- }
- population_vector.push_back(value);
- }
-
void finish_read() {
this->population_vector.push_back(this->max_value);
this->population_vector.push_back(this->min_value);
diff --git a/sortlib.cuh b/sortlib.cuh
index a7d3e2d..64f355b 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -37,7 +37,7 @@ public:
size_t low_bit = index & (-index);
return ((LENGTH + index) / low_bit) >> 1;
}
-
+#ifdef ENABLE_SORT_TEST
__device__ __host__ __attribute__((unused)) void testCalculation() const {
for (size_t i = 0; i < LENGTH; i++) {
auto left = calculate_index(i);
@@ -50,6 +50,7 @@ public:
assert(i + 1 == right);
}
}
+#endif
__device__ __host__ const key_type *binary_search(key_type *const start,
// const key_type *end,
@@ -127,7 +128,7 @@ public:
(double)(this->LENGTH - 1);
// printf("tmp: %llu %lf %lu\n", *it - *it_prev, tmp, prev_real_location);
}
-
+#ifdef SORT_FIRST_VERSION
/// First version
__device__ __host__ key_type *original_binary_search(key_type *start,
const key_type *end,
@@ -172,13 +173,16 @@ public:
(double)(x - *it_prev) / (double)(*it - *it_prev)) /
(double)((end - start) - 1);
}
+#endif
+#ifdef ENABLE_SORT_TEST
static void testSelf(size_t length) {
CustomSort(length, sizeof(long) * 8).testCalculation();
#ifndef DISABLE_TEST_WARNING
puts("You can remove this function if passed already");
#endif
}
+#endif
__host__ __device__ size_t length() const { return this->LENGTH; }
};
@@ -201,6 +205,7 @@ template <typename T> void rebuildSort(std::vector<T> &original) {
rebuild(original);
}
+#ifdef ENABLE_SORT_TEST
__global__ void testCustomCalculation(size_t length) {
CustomSort(length, sizeof(long) * 8).testCalculation();
}
@@ -209,6 +214,7 @@ __global__ void testCustomCalculationWithPrint(size_t length) {
CustomSort(length, sizeof(long) * 8).testCalculation();
printf("%lu pass\n", length);
}
+#endif
class FactorySort {
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 16:17:26 +0000