summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--expt_0502.cu78
-rw-r--r--expt_0503.cu32
2 files changed, 24 insertions, 86 deletions
diff --git a/expt_0502.cu b/expt_0502.cu
index ef0414d..004813d 100644
--- a/expt_0502.cu
+++ b/expt_0502.cu
@@ -9,6 +9,8 @@
#include <cstdio>
#include <vector>
+#include "read_helper.h"
+
std::vector<unsigned long long> population_vector, sample_vector,
le_sample_vector;
@@ -52,36 +54,6 @@ __global__ void initSample(key_type *sample, key_type *normal_sample,
#endif
}
-__global__ void kernel(unsigned long step, const double slice_size) {
- auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8);
- // printf("kernel1 step: %ld\n", step);
- for (int i = 0; i < step; i++) {
- auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
-
- // printf("%d\n", tid);
- auto index =
- (int)(custom_sort.cdf(cudaSampleItem, cudaPopulationItem[tid]) /
- slice_size) -
- 1;
-
- auto index2 = (int)(FactorySort::cdf(cudaNormalSampleItem, SAMPLE_LENGTH,
- cudaPopulationItem[tid]) /
- slice_size) -
- 1;
- printf("x => %llu %d %d\n", cudaPopulationItem[tid], index, index2);
- // printf("%llu %d\n", cudaPopulationItem[tid], index);
- /*if (cdf_result[index]) {
- while (cdf_result[++index]) {
- if (index >= split_size) {
- //printf("escape: %llu %lu\n", cudaPopulationItem[tid], split_size);
- }
- assert(index < (split_size + RESERVED_BLOCK));
- }
- }*/
- // cdf_result[index] = true;
- }
-}
-
__global__ void kernel2(unsigned long step, const double _slice_size) {
auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8);
// printf("kernel1 step: %ld\n", step);
@@ -165,24 +137,6 @@ __global__ void freeStorage() {
cudaFree(cdf_result);
}
-__global__ void initCustomSample() {
- // printf("init\n");
- key_type *tmp = nullptr;
- auto custom_sort = CustomSort(SAMPLE_LENGTH, 0);
- cudaMalloc(&tmp, sizeof(key_type) * SAMPLE_LENGTH);
- for (size_t i = 0; i < SAMPLE_LENGTH; i++) {
- tmp[i] = cudaSampleItem[custom_sort.calculate_index(i) - 1];
- }
- // printf("copy\n");
- memcpy(cudaSampleItem, tmp, sizeof(key_type) * SAMPLE_LENGTH);
- cudaFree(tmp);
- tmp = nullptr;
- // memset(cdf_result, 0, sizeof(key_type) * TEST_LENGTH);
- cudaFree(cdf_result);
- cdf_result = nullptr;
- // printf("finalize\n");
-}
-
const dim_pair_type DIM_PAIR[] = {dim_pair_type(32, 32)};
void run_kernel(size_t test_size) {
@@ -226,19 +180,17 @@ 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);
+ ReadHelper readHelper("normal_distribution.txt", SAMPLE_LENGTH, TEST_LENGTH,
+ 0);
+ readHelper.readFile();
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);
+
+ readHelper.split_into(sample_vector, population_vector);
+
std::sort(sample_vector.begin(), sample_vector.end());
le_sample_vector = sample_vector;
+ rebuild(sample_vector);
key_type *cudaSample = nullptr, *cudaPopulation;
cudaMalloc(&cudaSample, sizeof(key_type) * SAMPLE_LENGTH);
@@ -256,16 +208,16 @@ int main() {
// auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
// custom_sort.testCalculation();
- testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH);
- cudaDeviceSynchronize();
+ // testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH);
+ // cudaDeviceSynchronize();
initSample<<<1, 1>>>(cudaSample, cudaNormalSample, cudaPopulation);
cudaDeviceSynchronize();
- initCustomSample<<<1, 1>>>();
- cudaDeviceSynchronize();
- cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * SAMPLE_LENGTH,
- cudaMemcpyDeviceToHost);
+ // initCustomSample<<<1, 1>>>();
+ // cudaDeviceSynchronize();
+ // cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) *
+ // SAMPLE_LENGTH,cudaMemcpyDeviceToHost);
// check_items<<<1, 1>>>();
// cudaDeviceSynchronize();
diff --git a/expt_0503.cu b/expt_0503.cu
index c738e43..2785392 100644
--- a/expt_0503.cu
+++ b/expt_0503.cu
@@ -2,6 +2,7 @@
#define ENABLE_SORT_TEST
#include "sortlib.cuh"
+#include "read_helper.h"
#include <algorithm>
#include <cassert>
#include <cstdio>
@@ -51,8 +52,7 @@ __global__ void initSample(key_type *sample, key_type *normal_sample,
#endif
}
-__global__ void kernel(unsigned long step, const double slice_size,
- const unsigned long split_size) {
+__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);
for (int i = 0; i < step; i++) {
@@ -151,7 +151,7 @@ void run_kernel(size_t test_size, bool normal = true) {
cudaEventCreate(&stop);
cudaEventRecord(start, nullptr);
if (normal) {
- kernel<<<grid_dim, block_dim>>>(step, slice_size, split_size);
+ kernel<<<grid_dim, block_dim>>>(step, slice_size);
} else {
kernel2<<<grid_dim, block_dim>>>(step, slice_size);
}
@@ -173,19 +173,6 @@ void run_kernel(size_t test_size, bool normal = true) {
cudaDeviceSynchronize();
}
-void read_file(char const *filename, long sample_length,
- unsigned long max_number = TEST_LENGTH) {
- max_number += sample_length + 256;
- auto read_number = 0;
- 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++;
- fclose(file);
-}
-
long pow_for_sample(long n) {
auto x = 2;
for (int i = 0; i < n; i++) {
@@ -210,17 +197,16 @@ int main(int argc, char const *argv[]) {
sample_length = pow_for_sample(strtol(argv[2], nullptr, 10));
}
- read_file("normal_distribution.txt", test_size);
+ ReadHelper readHelper("normal_distribution.txt", sample_length, test_size, 0);
+
+ printf("test size: %zu, sample length: %zu\n", test_size, sample_length);
- printf("population: %zu, test size: %zu, sample length: %zu\n",
- population_vector.size(), test_size, sample_length);
+ readHelper.readFile();
+ readHelper.split_into(sample_vector, population_vector);
- 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;
+ rebuild(sample_vector);
key_type *cudaSample = nullptr, *cudaPopulation;
cudaMalloc(&cudaSample, sizeof(key_type) * sample_length);
generated by cgit v1.3.1 (git 2.54.0) at 2026-09-04 06:20:29 +0000