summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--expt_0510.cu36
-rw-r--r--expt_0516.cu158
-rw-r--r--expt_0520.cu5
-rw-r--r--expt_0525.cu21
-rw-r--r--expt_0528.cu12
-rw-r--r--expt_0529.cu2
-rw-r--r--read_helper.h36
-rw-r--r--sortlib.cuh3
8 files changed, 111 insertions, 162 deletions
diff --git a/expt_0510.cu b/expt_0510.cu
index 82f2624..eca1b92 100644
--- a/expt_0510.cu
+++ b/expt_0510.cu
@@ -1,6 +1,7 @@
// Experimental content: Test branch performance (optimized memory)
#include "sortlib.cuh"
+#include "read_helper.h"
#include <algorithm>
#include <cassert>
#include <cstdio>
@@ -167,34 +168,6 @@ unsigned long randomRow(unsigned long max_value_) {
return dst(mt19937);
}
-void readFile(char const *filename, long sample_length,
- unsigned long &total_row,
- unsigned long max_number = TEST_LENGTH) {
- auto read_number = 0UL;
- FILE *file = fopen(filename, "r");
- assert(file);
- for (long long i;
- read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
- store_into_vector(i))
- read_number++;
-
- if (total_row > 0) {
- read_number = randomRow(total_row - sample_length - max_number - 256);
- total_row = read_number;
- // fprintf(stderr, "Skip %lu\n", read_number);
- for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;)
- read_number--;
- }
-
- read_number = 0;
-
- auto remain = sample_length + 256;
- for (long long i; read_number < remain && fscanf(file, "%lld ", &i) != EOF;
- store_into_vector(i))
- read_number++;
- fclose(file);
-}
-
long pow_for_sample(long n) {
auto x = 2;
for (int i = 0; i < n; i++) {
@@ -225,10 +198,13 @@ int main(int argc, char const *argv[]) {
// CustomSort::testSelf(sample_length);
- readFile("normal_distribution.txt", sample_length, total_row, test_size);
+ ReadHelper readHelper("normal_distribution.txt", sample_length, test_size,
+ total_row);
+ readHelper.readFile();
printf("population: %zu, skip: %lu, test size: %zu, sample length: %zu | ",
- population_vector.size(), total_row, test_size, sample_length);
+ population_vector.size(), readHelper.random_number, test_size,
+ sample_length);
sample_vector = std::vector<key_type>(
population_vector.begin(), population_vector.begin() + sample_length - 2);
diff --git a/expt_0516.cu b/expt_0516.cu
index 82755ac..351a8c4 100644
--- a/expt_0516.cu
+++ b/expt_0516.cu
@@ -1,18 +1,13 @@
// Experimental content: Test branch performance (optimized memory and output)
+// #define PRINT_READ_PROCESS
+#include "read_helper.h"
#include "sortlib.cuh"
#include <algorithm>
#include <cassert>
#include <cstdio>
-#include <random>
#include <vector>
-#ifndef NDEBUG
-#define P_ERR(...) fprintf(stderr, __VA_ARGS__)
-#else
-#define P_ERR(...)
-#endif
-
std::vector<unsigned long long> population_vector, sample_vector,
le_sample_vector;
@@ -21,27 +16,10 @@ constexpr size_t TEST_LENGTH = 32'768;
// constexpr size_t TEST_LENGTH = 4096;
// constexpr size_t RESERVED_BLOCK = 100;
-unsigned long long max_value = 0, min_value = 0xfffffffff;
typedef std::pair<int, int> dim_pair_type;
-inline void store_into_vector(unsigned long long value) {
- if (max_value < value) {
- max_value = value;
- }
- if (min_value > value) {
- min_value = value;
- }
- population_vector.push_back(value);
-}
-// #define TEST_BOUNDS
-
-__device__ long cuda_sample_length;
+__device__ unsigned long cuda_sample_length;
__device__ key_type *cudaSampleItem, *cudaNormalSampleItem, *cudaPopulationItem;
-//__device__ bool *cdf_result;
-#ifdef TEST_BOUNDS
-__device__ unsigned insert_value;
-__device__ unsigned int index_max, index_min;
-#endif
__global__ void initSample(key_type *sample, key_type *population) {
cudaSampleItem = sample;
@@ -55,6 +33,27 @@ __global__ void initNormalSample(key_type *normal_sample) {
cudaNormalSampleItem = normal_sample;
}
+/*
+#define CudaCheckError() cudaCheckError(__FILE__, __LINE__)
+inline void cudaCheckError(const char *file, const int line) {
+ cudaError err = cudaGetLastError();
+ if (cudaSuccess != err) {
+ fprintf(stderr, "cudaCheckError() failed at %s:%i : %s\n", file, line,
+ cudaGetErrorString(err));
+ exit(-1);
+ }
+
+ // More careful checking. However, this will affect performance.
+ // Comment away if needed.
+ err = cudaDeviceSynchronize();
+ if (cudaSuccess != err) {
+ fprintf(stderr, "cudaCheckError() with sync failed at %s:%i : %s\n", file,
+ line, cudaGetErrorString(err));
+ exit(-1);
+ }
+}
+*/
+
__global__ void kernel(unsigned long step, const double slice_size) {
auto custom_sort = CustomSort(cuda_sample_length, sizeof(key_type) * 8);
// printf("kernel1 step: %ld\n", step);
@@ -100,25 +99,8 @@ __global__ void checkItems() {
printf("\n");
}
-void initCustomSampleCPU() {
- auto sample_length = sample_vector.size();
- auto tmp = new key_type[sample_length];
-
- auto custom_sort = CustomSort(sample_length, 0);
-
- for (size_t i = 0; i < sample_length; i++) {
- tmp[i] = sample_vector[custom_sort.calculate_index(i) - 1];
- }
-
- memcpy(sample_vector.data(), tmp, sizeof(key_type) * sample_length);
-
- delete[] tmp;
-}
-
constexpr dim3 grid_dim = 32, block_dim = 32;
-bool checkTestSize() {}
-
void run_kernel(size_t test_size, bool custom = true) {
// need description
unsigned long step = test_size / (grid_dim.x * block_dim.x);
@@ -163,45 +145,6 @@ __global__ void freeStorageStage2() {
// cudaFree(cdf_result);
}
-unsigned long randomRow(unsigned long max_value_) {
- std::random_device randomDevice;
- std::mt19937 mt19937(randomDevice());
- std::uniform_int_distribution<std::mt19937::result_type> dst(0, max_value_);
- return dst(mt19937);
-}
-
-void readFile(char const *filename, long sample_length,
- unsigned long &total_row,
- unsigned long max_number = TEST_LENGTH) {
- auto read_number = 0UL;
- FILE *file = fopen(filename, "r");
- assert(file);
- P_ERR("Reading sample");
- for (long long i;
- read_number < max_number && fscanf(file, "%lld ", &i) != EOF;
- store_into_vector(i))
- read_number++;
-
- if (total_row > 0) {
- P_ERR("\rReading skip");
- read_number = randomRow(total_row - sample_length - max_number - 256);
- total_row = read_number;
- // fprintf(stderr, "Skip %lu\n", read_number);
- for (long long i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;)
- read_number--;
- }
-
- P_ERR("\rReading population");
- read_number = 0;
-
- auto remain = sample_length + 256;
- for (long long i; read_number < remain && fscanf(file, "%lld ", &i) != EOF;
- store_into_vector(i))
- read_number++;
- fclose(file);
- P_ERR("\r");
-}
-
long pow_for_sample(long n) {
auto x = 2;
for (int i = 1; i < n; i++) {
@@ -210,7 +153,7 @@ long pow_for_sample(long n) {
return x - 1;
}
-__global__ void applyCudaSampleLength(long length) {
+__global__ void applyCudaSampleLength(unsigned long length) {
cuda_sample_length = length;
}
@@ -230,36 +173,46 @@ int main(int argc, char const *argv[]) {
total_row = strtol(argv[3], nullptr, 10);
}
- readFile("normal_distribution.txt", sample_length, total_row, test_size);
+ // printf("%ld %ld %ld\n", test_size, sample_length, total_row);
+
+ ReadHelper readHelper("normal_distribution.txt", sample_length, test_size,
+ total_row);
- printf("population: %zu, skip: %lu, test size: %zu, sample length: %zu\n",
- population_vector.size(), total_row, test_size, sample_length);
+ printf("population: %zu, skip: %lu, sample length: %zu ", test_size,
+ readHelper.random_number, sample_length);
+ readHelper.readFile();
+
+ readHelper.split_into(sample_vector, population_vector);
+ assert(sample_vector.size() == sample_length);
+ assert(population_vector.size() == test_size);
- sample_vector = std::vector<key_type>(
- population_vector.begin(), population_vector.begin() + sample_length - 2);
- sample_vector.push_back(min_value);
- sample_vector.push_back(max_value);
std::sort(sample_vector.begin(), sample_vector.end());
le_sample_vector = sample_vector;
- initCustomSampleCPU();
+ rebuild(sample_vector);
+
+ /*printf("population: %zu sample length: %zu\n", population_vector.size(),
+ sample_vector.size());*/
key_type *cudaSample = nullptr, *cudaPopulation;
P_ERR("Coping sample");
- cudaMalloc(&cudaSample, sizeof(key_type) * sample_length);
+ cudaMalloc(&cudaSample, sizeof(key_type) * sample_vector.size());
assert(sample_length == sample_vector.size());
- cudaMemcpy(cudaSample, sample_vector.data(), sizeof(key_type) * sample_length,
- cudaMemcpyHostToDevice);
- cudaMalloc(&cudaPopulation, sizeof(key_type) * test_size);
+ cudaMemcpy(cudaSample, sample_vector.data(),
+ sizeof(key_type) * sample_vector.size(), cudaMemcpyHostToDevice);
+ cudaMalloc(&cudaPopulation, sizeof(key_type) * population_vector.size());
+
+ assert(population_vector.size() == test_size);
P_ERR("\rCoping population");
- cudaMemcpy(cudaPopulation, population_vector.data() + sample_length,
- sizeof(key_type) * test_size, cudaMemcpyHostToDevice);
+ cudaMemcpy(cudaPopulation, population_vector.data(),
+ sizeof(key_type) * population_vector.size(),
+ cudaMemcpyHostToDevice);
// auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
// custom_sort.testCalculation();
P_ERR("\rApply custom length and test calculation");
- applyCudaSampleLength<<<1, 1>>>(sample_length);
+ applyCudaSampleLength<<<1, 1>>>(sample_vector.size());
// testCustomCalculation<<<1, 1>>>(sample_length);
cudaDeviceSynchronize();
P_ERR("\rApply custom length and test calculation completed");
@@ -276,15 +229,18 @@ int main(int argc, char const *argv[]) {
"kernel\n");
run_kernel(test_size);
- freeStorageStage1<<<1, 1>>>();
- cudaDeviceSynchronize();
+ // freeStorageStage1<<<1, 1>>>();
+ // cudaDeviceSynchronize();
+
+ cudaFree(cudaSample);
key_type *cudaNormalSample = nullptr;
P_ERR("\nCoping new sample");
- cudaMalloc(&cudaNormalSample, sizeof(key_type) * sample_length);
+ cudaMalloc(&cudaNormalSample, sizeof(key_type) * le_sample_vector.size());
cudaMemcpy(cudaNormalSample, le_sample_vector.data(),
- sizeof(key_type) * sample_length, cudaMemcpyHostToDevice);
+ sizeof(key_type) * le_sample_vector.size(),
+ cudaMemcpyHostToDevice);
/*cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) *
sample_length, cudaMemcpyDeviceToHost);*/
diff --git a/expt_0520.cu b/expt_0520.cu
index d31dc78..bd1439f 100644
--- a/expt_0520.cu
+++ b/expt_0520.cu
@@ -208,8 +208,9 @@ int main(int argc, char **argv) {
printf("insert_length %ld, sample_length: %ld\n", insertion_length,
sample_length);
- ReadHelper reader("normal_distribution.txt", sample_length, insertion_length);
- reader.readFile(total_row);
+ ReadHelper reader("normal_distribution.txt", sample_length, insertion_length,
+ total_row);
+ reader.readFile();
key_type *cudaPopulation;
diff --git a/expt_0525.cu b/expt_0525.cu
index c2ee041..8867cee 100644
--- a/expt_0525.cu
+++ b/expt_0525.cu
@@ -76,7 +76,6 @@ 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))
@@ -108,12 +107,6 @@ constexpr size_t FACTOR = 1;
typedef LL key_type;
-#ifdef RANDOM_TARGET
-constexpr const char *TARGET_STRING = "RANDOM";
-#else
-// constexpr const char *TARGET_STRING = "PERFECT";
-#endif
-
#define CUDA_ERROR_CHECK
#define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__)
@@ -153,7 +146,13 @@ inline void __cudaCheckError(const char *file, const int line) {
// Definition of generic node class
-class __attribute__((aligned(16))) Node {
+class
+#ifndef _MSC_VER
+ __attribute__((aligned(16)))
+#else
+ __declspec(align(16))
+#endif
+ Node {
public:
int topLevel; // Level of the node
LL key; // Key value
@@ -636,7 +635,7 @@ inline double calcSliceSize(size_t insertion_size) {
return 1.0 / (double)insertion_size;
}
-inline size_t calcBlocks(size_t input) {
+inline auto calcBlocks(size_t input) {
return (input % (NUM_THREADS * FACTOR) == 0)
? input / (NUM_THREADS * FACTOR)
: (input / (NUM_THREADS * FACTOR)) + 1;
@@ -660,8 +659,8 @@ int main(int argc, char **argv) {
insertion_length, search_length);
ReadHelper readHelper("normal_distribution.txt", sample_length,
- insertion_length);
- readHelper.readFile(nullptr);
+ insertion_length, 0);
+ readHelper.readFile();
std::vector<key_type> _sample, _population;
readHelper.split_into(_sample, _population);
/*printf("%lu, Create search vector: %ld\n", _population.size(),
diff --git a/expt_0528.cu b/expt_0528.cu
index 9ac7059..d74c93a 100644
--- a/expt_0528.cu
+++ b/expt_0528.cu
@@ -153,7 +153,13 @@ inline void __cudaCheckError(const char *file, const int line) {
// Definition of generic node class
-class __attribute__((aligned(16))) Node {
+class
+#ifndef _MSC_VER
+ __attribute__((aligned(16)))
+#else
+ __declspec(align(16))
+#endif
+ Node {
public:
int topLevel; // Level of the node
LL key; // Key value
@@ -652,8 +658,8 @@ int main(int argc, char **argv) {
printf("Insertion: %ld, Search: %ld\n", insertion_length, search_length);
- ReadHelper readHelper("normal_distribution.txt", 0, insertion_length);
- readHelper.readFile(nullptr);
+ ReadHelper readHelper("normal_distribution.txt", 0, insertion_length, 0);
+ readHelper.readFile();
std::vector<key_type> _search(readHelper.population_vector.begin(),
readHelper.population_vector.begin() +
search_length);
diff --git a/expt_0529.cu b/expt_0529.cu
index 985648f..bd64ad0 100644
--- a/expt_0529.cu
+++ b/expt_0529.cu
@@ -46,7 +46,7 @@ int main(int argc, char const *argv[]) {
population_length);
ReadHelper readHelper("normal_distribution.txt", sample_length, 0);
- readHelper.readFile(nullptr);
+ readHelper.readFile();
std::vector<key_type> sample, population;
std::vector<unsigned int> result(RESULT_LENGTH);
diff --git a/read_helper.h b/read_helper.h
index d19834a..f8396f2 100644
--- a/read_helper.h
+++ b/read_helper.h
@@ -2,6 +2,7 @@
#ifndef LOCKFREE_READ_HELPER_H
#define LOCKFREE_READ_HELPER_H
+#include <cassert>
#include <cstdio>
#include <random>
#include <vector>
@@ -34,12 +35,22 @@ class ReadHelper {
char const *filename;
static constexpr size_t REVERSED_BLOCK = 256;
+ auto genRandomRow(size_t total_row) const {
+ if (total_row != 0) {
+ assert(total_row > (population_length + sample_length + REVERSED_BLOCK));
+ return randomRow(total_row - population_length - sample_length -
+ REVERSED_BLOCK);
+ }
+ return 0UL;
+ }
+
public:
- const size_t sample_length, needed_read_length;
+ const size_t sample_length, population_length, random_number;
ReadHelper(char const *filename, size_t sample_length,
- size_t insertion_length)
+ size_t population_length, size_t total_row = 0)
: filename(filename), sample_length(sample_length),
- needed_read_length(insertion_length + sample_length) {
+ population_length(population_length),
+ random_number(genRandomRow(total_row)) {
// printf("%zu %zu\n", sample_length, needed_read_length);
}
key_type maxValue() const { return max_value; }
@@ -54,10 +65,11 @@ public:
return dst(mt19937);
}
- void readFile(unsigned long *total_row) {
+ void readFile() {
this->population_vector.clear();
auto read_number = 0UL;
FILE *file = fopen(filename, "r");
+
if (file == nullptr) {
fprintf(stderr, "Unable to open file %s\n", filename);
exit(1);
@@ -69,12 +81,10 @@ public:
store_into_vector(i))
read_number++;
- if (total_row != nullptr) {
+ if (random_number > 0) {
P_ERR("\rReading skip");
- read_number = randomRow(*total_row - sample_length - needed_read_length -
- REVERSED_BLOCK);
- *total_row = read_number;
// fprintf(stderr, "Skip %lu\n", read_number);
+ read_number = random_number;
for (key_type i; read_number > 0 && fscanf(file, "%lld ", &i) != EOF;)
read_number--;
}
@@ -82,7 +92,7 @@ public:
P_ERR("\rReading population");
read_number = 0;
- auto remain = needed_read_length - sample_length + REVERSED_BLOCK;
+ auto remain = population_length + REVERSED_BLOCK;
for (key_type i; read_number < remain && fscanf(file, "%lld ", &i) != EOF;
store_into_vector(i))
read_number++;
@@ -92,25 +102,25 @@ public:
void split_into(std::vector<key_type> &sample, std::vector<key_type> &p) {
sample.resize(sample_length - 2);
- p.resize(needed_read_length - sample_length);
+ p.resize(population_length);
// printf("%zu\n", needed_read_length - sample_length);
memcpy(sample.data(), population_vector.data(),
sizeof(key_type) * (sample_length - 2));
sample.push_back(this->max_value);
sample.push_back(this->min_value);
memcpy(p.data(), population_vector.data() + sample_length,
- sizeof(key_type) * (needed_read_length - sample_length));
+ sizeof(key_type) * population_length);
}
void split_into(key_type *&sample, key_type *&p) {
sample = new key_type[sample_length];
- p = new key_type[needed_read_length - sample_length];
+ p = new key_type[population_length];
memcpy(sample, population_vector.data(),
sizeof(key_type) * (sample_length - 2));
sample[sample_length - 2] = this->max_value;
sample[sample_length - 1] = this->min_value;
memcpy(p, population_vector.data() + sample_length,
- sizeof(key_type) * (needed_read_length - sample_length));
+ sizeof(key_type) * (population_length));
}
size_t size() const { return this->population_vector.size(); }
diff --git a/sortlib.cuh b/sortlib.cuh
index 64f355b..8d6788e 100644
--- a/sortlib.cuh
+++ b/sortlib.cuh
@@ -239,7 +239,8 @@ public:
return last_known_point;
}
- __device__ __host__ double static sample_cdf(key_type *begin, long length,
+ __device__ __host__ double static sample_cdf(key_type *begin,
+ unsigned long length,
key_type x) {
// printf("%f\n", x);
auto end = begin + length;
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 09:46:26 +0000