summaryrefslogtreecommitdiff
path: root/expt_0503.cu
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2023-05-10 00:39:18 +0800
committerKunoiSayami <[email protected]>2023-05-10 00:39:18 +0800
commit901db91584ceb50a5c3f2a1c06620035a064cb01 (patch)
treec36f3631e2a30d8414b737d7ff9441e8f8210017 /expt_0503.cu
parenta2a1c2db484ef519d732209fadc6252b17a90fbc (diff)
feat(exp): Make length can change from arguments
Signed-off-by: KunoiSayami <[email protected]>
Diffstat (limited to 'expt_0503.cu')
-rw-r--r--expt_0503.cu79
1 files changed, 53 insertions, 26 deletions
diff --git a/expt_0503.cu b/expt_0503.cu
index 96cf44e..58f35aa 100644
--- a/expt_0503.cu
+++ b/expt_0503.cu
@@ -9,7 +9,7 @@
std::vector<unsigned long long> population_vector, sample_vector,
le_sample_vector;
-constexpr size_t SAMPLE_LENGTH = 1023;
+constexpr long DEFAULT_SAMPLE_LENGTH = 65535;
constexpr size_t TEST_LENGTH = 32'768;
// constexpr size_t TEST_LENGTH = 4096;
constexpr size_t RESERVED_BLOCK = 10'000;
@@ -28,6 +28,7 @@ inline void store_into_vector(unsigned long long value) {
}
// #define TEST_BOUNDS
+__device__ long cuda_sample_length;
__device__ key_type *cudaSampleItem, *cudaNormalSampleItem, *cudaPopulationItem;
__device__ bool *cdf_result, *cdf_normal_result;
#ifdef TEST_BOUNDS
@@ -51,14 +52,14 @@ __global__ void initSample(key_type *sample, key_type *normal_sample,
__global__ void kernel(unsigned long step, const double slice_size,
const unsigned long split_size) {
- auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(key_type) * 8);
+ auto custom_sort = CustomSort(cuda_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.sample_cdf_custom_version(
- cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH,
+ cudaSampleItem, cudaSampleItem + cuda_sample_length,
cudaPopulationItem[tid]) /
slice_size) -
1;
@@ -71,7 +72,7 @@ __global__ void kernel2(unsigned long step, const double slice_size) {
for (int i = 0; i < step; i++) {
auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
auto index =
- (int)(FactorySort::sample_cdf(cudaNormalSampleItem, SAMPLE_LENGTH,
+ (int)(FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length,
cudaPopulationItem[tid]) /
slice_size) -
1;
@@ -84,7 +85,7 @@ __global__ void print2() {
printf("%llu ", cudaSampleItem[i]);
}
printf("\n");*/
- printf("%lu\n", CustomSort::fast_log(SAMPLE_LENGTH));
+ printf("%lu\n", CustomSort::fast_log(cuda_sample_length));
}
__global__ void print_function() {
@@ -95,7 +96,7 @@ __global__ void print_function() {
}
__global__ void check_items() {
- for (int i = 0; i < SAMPLE_LENGTH; i++) {
+ for (int i = 0; i < cuda_sample_length; i++) {
printf("%lld ", cudaSampleItem[i]);
}
printf("\n");
@@ -122,13 +123,13 @@ __global__ void freeStorage() {
__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++) {
+ auto custom_sort = CustomSort(cuda_sample_length, 0);
+ cudaMalloc(&tmp, sizeof(key_type) * cuda_sample_length);
+ for (size_t i = 0; i < cuda_sample_length; i++) {
tmp[i] = cudaSampleItem[custom_sort.calculate_index(i) - 1];
}
// printf("copy\n");
- memcpy(cudaSampleItem, tmp, sizeof(key_type) * SAMPLE_LENGTH);
+ memcpy(cudaSampleItem, tmp, sizeof(key_type) * cuda_sample_length);
cudaFree(tmp);
tmp = nullptr;
// memset(cdf_result, 0, sizeof(key_type) * TEST_LENGTH);
@@ -172,47 +173,73 @@ 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++) {
+ x *= 2;
+ }
+ return x - 1;
+}
+
+__global__ void applyCudaSampleLength(long length) {
+ cuda_sample_length = length;
+}
+
int main(int argc, char const *argv[]) {
auto test_size = TEST_LENGTH;
+ auto sample_length = DEFAULT_SAMPLE_LENGTH;
if (argc >= 2) {
test_size = strtol(argv[1], nullptr, 10);
}
+ if (argc >= 3) {
+ sample_length = pow_for_sample(strtol(argv[2], nullptr, 10));
+ }
- FILE *file = fopen("normal_distribution.txt", "r");
- assert(file);
- for (long long i; fscanf(file, "%lld ", &i) != EOF; store_into_vector(i))
- ;
- fclose(file);
+ read_file("normal_distribution.txt", test_size);
- printf("population: %zu, test size: %zu\n", population_vector.size(),
- test_size);
+ printf("population: %zu, test size: %zu, sample length: %zu\n",
+ population_vector.size(), test_size, sample_length);
sample_vector = std::vector<key_type>(
- population_vector.begin(), population_vector.begin() + SAMPLE_LENGTH - 2);
+ 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;
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,
+ 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);
- cudaMemcpy(cudaPopulation, population_vector.data() + SAMPLE_LENGTH,
+ cudaMemcpy(cudaPopulation, population_vector.data() + sample_length,
sizeof(key_type) * test_size, cudaMemcpyHostToDevice);
key_type *cudaNormalSample = nullptr;
- cudaMalloc(&cudaNormalSample, sizeof(key_type) * SAMPLE_LENGTH);
+ cudaMalloc(&cudaNormalSample, sizeof(key_type) * sample_length);
cudaMemcpy(cudaNormalSample, le_sample_vector.data(),
- sizeof(key_type) * SAMPLE_LENGTH, cudaMemcpyHostToDevice);
+ sizeof(key_type) * sample_length, cudaMemcpyHostToDevice);
// auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
// custom_sort.testCalculation();
- testCustomCalculation<<<1, 1>>>(SAMPLE_LENGTH);
+ applyCudaSampleLength<<<1, 1>>>(sample_length);
+ testCustomCalculation<<<1, 1>>>(sample_length);
cudaDeviceSynchronize();
initSample<<<1, 1>>>(cudaSample, cudaNormalSample, cudaPopulation);
@@ -220,7 +247,7 @@ int main(int argc, char const *argv[]) {
initCustomSample<<<1, 1>>>();
cudaDeviceSynchronize();
- cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * SAMPLE_LENGTH,
+ cudaMemcpy(sample_vector.data(), cudaSample, sizeof(key_type) * sample_length,
cudaMemcpyDeviceToHost);
// check_items<<<1, 1>>>();