summaryrefslogtreecommitdiff
path: root/expt_0516.cu
diff options
context:
space:
mode:
Diffstat (limited to 'expt_0516.cu')
-rw-r--r--expt_0516.cu158
1 files changed, 57 insertions, 101 deletions
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);*/
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 13:27:21 +0000