summaryrefslogtreecommitdiff
path: root/expt_0406.cu
diff options
context:
space:
mode:
Diffstat (limited to 'expt_0406.cu')
-rw-r--r--expt_0406.cu221
1 files changed, 75 insertions, 146 deletions
diff --git a/expt_0406.cu b/expt_0406.cu
index fd7915d..96533c9 100644
--- a/expt_0406.cu
+++ b/expt_0406.cu
@@ -1,13 +1,21 @@
#include <algorithm>
#include <cassert>
#include <cstdio>
-#include <iostream>
#include <vector>
constexpr size_t length = 2097152;
+constexpr long TEST_LENGTH = 32'768;
std::vector<unsigned long long> population_vector, sample_vector,
sample_vector_into_cuda;
+#define CHECK_CUDA_ERR(x) \
+ { \
+ auto ret_ = (x); \
+ if (ret_ != 0) { \
+ printf("%d: CUDA ERR: %d\n", __LINE__, ret_); \
+ } \
+ }
+
unsigned long long max_value = 0, min_value = 0xfffffffff;
inline void store_into_vector(unsigned long long value) {
@@ -25,68 +33,20 @@ typedef unsigned long long *key_type_ptr;
constexpr long SAMPLE_LENGTH = 1024;
-constexpr long TEST_SIZE = 32'768;
-
-__device__ key_type *SampleItem, *PopulationItem;
+__device__ key_type *cudaSampleItem, *cudaPopulationItem;
__device__ bool *cdf_result;
__device__ unsigned insert_value;
-#ifdef TEST_BOUNDS
-__device__ unsigned int index_max, index_min;
-#endif
-
-__device__ const key_type *cudaBinarySearch(key_type *start, key_type *end,
- const key_type val) {
- auto begin = start;
- key_type *last_known_point = nullptr;
- assert(begin < end);
- while (begin < end) {
- auto mid = (end - begin) / 2;
- auto mid_val = *(begin + mid);
- if (val == mid_val) {
- return begin + mid;
- } else if (val > mid_val) {
- begin = begin + mid + 1;
- } else {
- end = end - mid - 1;
- }
- last_known_point = begin;
- }
- return last_known_point;
-}
-
-__device__ double sample_cdf(double x) {
- auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x);
- if (it == SampleItem + SAMPLE_LENGTH) {
- return 1;
- }
- if (it == SampleItem) {
- return 0;
- }
- auto it_prev = it - 1;
- return (double(it_prev - SampleItem) +
- (x - (double)*it_prev) / (double)(*it - *it_prev)) /
- double(SAMPLE_LENGTH - 1);
-}
__global__ void initStorage(unsigned scale, unsigned test_size) {
- cudaFree(cdf_result);
- cudaMalloc(&cdf_result, scale * test_size * sizeof(bool));
+ CHECK_CUDA_ERR(cudaFree(cdf_result));
+ CHECK_CUDA_ERR(cudaMalloc(&cdf_result, scale * test_size * sizeof(bool)));
memset(cdf_result, 0, scale * test_size * sizeof(bool));
insert_value = 0;
- // printf("initCuda storage\n");
-#ifdef TEST_BOUNDS
- index_max = 0;
- index_min = 0x7fffffff;
-#endif
}
__global__ void initCuda(key_type *sample_item, key_type *population_item) {
- SampleItem = sample_item;
- PopulationItem = population_item;
- /*for (int i = 0; i < SAMPLE_LENGTH; i++) {
- }*/
- // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long));
- // cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type));
+ cudaSampleItem = sample_item;
+ cudaPopulationItem = population_item;
cdf_result = nullptr;
}
@@ -95,126 +55,95 @@ __global__ void kernel(unsigned long step, const double slice_size) {
auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
// printf("%d\n", tid);
- auto index = (int)(sample_cdf(PopulationItem[tid]) / slice_size);
- /*if (cdf_result[index]) {
- while (cdf_result[++index]) {
- assert(index < split_size);
- }
- }*/
- cdf_result[index] = true;
- // atomicAdd(&insert_value, 1);
+ auto index = (int)(sample_cdf(cudaPopulationItem[tid]) / slice_size);
-#ifdef TEST_BOUNDS
- while (true) {
- unsigned tmp = index_min;
- if (tmp < tid) {
- break;
- }
- if (atomicCAS(&index_min, tmp, tid) == tmp) {
- break;
- }
- }
- while (true) {
- unsigned tmp = index_max;
- if (tmp > tid) {
- break;
- }
- if (atomicCAS(&index_max, tmp, tid) == tmp) {
- break;
- }
- }
-#endif
+ cdf_result[index] = true;
}
}
__global__ void print_function() {
+
// printf("%u\n", insert_value);
#ifdef TEST_BOUNDS
printf("%u %u\n", index_min, index_max);
#endif
}
-int main(int argc, char const *argv[]) {
- size_t test_size;
- if (argc == 1) {
- test_size = 1048576;
- } else {
- try {
- test_size = std::stol(argv[1]);
- } catch (...) {
- return 1;
+__global__ void freeStorage() {
+ cudaFree(cudaSampleItem);
+ cudaFree(cudaPopulationItem);
+ cudaFree(cdf_result);
+}
+
+typedef std::pair<int, int> dim_pair_type;
+const dim_pair_type DIM_PAIR[] = {dim_pair_type(32, 32)};
+
+void run_kernel(size_t test_size) {
+ for (auto &pair : DIM_PAIR) {
+ dim3 grid_dim = pair.first, block_dim = pair.second;
+ unsigned long step = test_size / (grid_dim.x * block_dim.x);
+ printf("current dim: %d %d %lu\n", pair.first, pair.second, step);
+ for (int scale = 2; scale <= 16; scale += 2) {
+ const unsigned long split_size = test_size * scale;
+ const auto slice_size = 1.0 / (double)split_size;
+ printf("scale: %i ", scale);
+ initStorage<<<1, 1>>>(scale, test_size);
+ cudaDeviceSynchronize();
+ cudaEvent_t start, stop;
+ cudaEventCreate(&start);
+ cudaEventCreate(&stop);
+ cudaEventRecord(start, nullptr);
+ kernel<<<grid_dim, block_dim>>>(step, slice_size);
+
+ cudaDeviceSynchronize();
+ cudaEventRecord(stop, nullptr);
+ cudaEventSynchronize(stop);
+ float time;
+ cudaEventElapsedTime(&time, start, stop);
+ cudaEventDestroy(start);
+ cudaEventDestroy(stop);
+ if (time == 0) {
+ // printf("last error: %u\n", cudaGetLastError());
+ }
+ printf("time: %lf\n", time);
+
+ print_function<<<1, 1>>>();
+ cudaDeviceSynchronize();
}
}
+}
+
+int main() {
+ auto test_size = TEST_LENGTH;
FILE *file = fopen("normal_distribution.txt", "r");
assert(file);
for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i))
;
fclose(file);
- assert(population_vector.size() == length);
-
- puts("Read success");
-
- sample_vector = std::vector<unsigned long long>(
+ printf("test size: %d\n", 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);
-
- printf("Sample vector length: %zu\n", sample_vector.size());
-
std::sort(sample_vector.begin(), sample_vector.end());
- sample_vector_into_cuda.resize(SAMPLE_LENGTH);
- assert(sample_vector.size() == SAMPLE_LENGTH);
- /*for (int i = 0; i < SAMPLE_LENGTH; i++) {
- sample_vector_into_cuda[i] = sample_vector[calculate_location(i) - 1];
- }*/
-
- key_type *cudaSample = nullptr, *cudaPopulation = nullptr;
- cudaMalloc(&cudaSample, sizeof(key_type) * (SAMPLE_LENGTH));
+ key_type *cudaSample = nullptr, *cudaPopulation;
+ cudaMalloc(&cudaSample, sizeof(key_type) * SAMPLE_LENGTH);
+ assert(SAMPLE_LENGTH == sample_vector.size());
cudaMemcpy(cudaSample, sample_vector.data(), sizeof(key_type) * SAMPLE_LENGTH,
cudaMemcpyHostToDevice);
-
- cudaMalloc(&cudaPopulation, sizeof(key_type) * TEST_SIZE);
+ cudaMalloc(&cudaPopulation, sizeof(key_type) * test_size);
cudaMemcpy(cudaPopulation, population_vector.data() + SAMPLE_LENGTH,
- sizeof(key_type) * TEST_SIZE, cudaMemcpyHostToDevice);
+ sizeof(key_type) * test_size, cudaMemcpyHostToDevice);
- // memcpy(sample_heap, sample_vector.cbegin(),sizeof(key_type) *(SAMPLE_LENGTH
- // + 2));
- puts("Copy successful");
+ // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
+ // custom_sort.testCalculation();
+ // testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH);
+ // cudaDeviceSynchronize();
+ run_kernel(test_size);
- initCuda<<<1, 1>>>(cudaSample, cudaPopulation);
+ freeStorage<<<1, 1>>>();
cudaDeviceSynchronize();
-
- dim3 grid_dim = 32, block_dim = 32;
-
- unsigned long step = test_size / (grid_dim.x * block_dim.x);
- printf("(old) step: %lu\n", step);
- assert(!(test_size % (grid_dim.x * block_dim.x)));
-
- for (int scale = 2; scale <= 8; scale += 2) {
- const unsigned long split_size = test_size * scale;
- const auto slice_size = 1.0 / (double)split_size;
- printf("scale: %i ", scale);
- initStorage<<<1, 1>>>(scale, test_size);
- cudaDeviceSynchronize();
- cudaEvent_t start, stop;
- cudaEventCreate(&start);
- cudaEventCreate(&stop);
- cudaEventRecord(start, nullptr);
- kernel<<<grid_dim, block_dim>>>(step, slice_size);
- cudaDeviceSynchronize();
- cudaEventRecord(stop, nullptr);
- cudaEventSynchronize(stop);
- float time;
- cudaEventElapsedTime(&time, start, stop);
- cudaEventDestroy(start);
- cudaEventDestroy(stop);
-
- printf("time: %lf\n", time);
- print_function<<<1, 1>>>();
- cudaDeviceSynchronize();
- }
- return 0;
-} \ No newline at end of file
+}
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 11:37:37 +0000