summaryrefslogtreecommitdiff
path: root/expt_0729.cu
diff options
context:
space:
mode:
authorKunoiSayami <[email protected]>2023-05-01 19:50:15 +0800
committerKunoiSayami <[email protected]>2023-05-01 19:50:15 +0800
commit6f6359c6d9f1a05b509b0509762b77df127ebe37 (patch)
tree8a28f56b7e0b84fdc0e8cbe7ee88e0d6364b471f /expt_0729.cu
parent4c6a7764a97c6daf836197022ea154fdc387c015 (diff)
2023-05-01 19:50
Signed-off-by: KunoiSayami <[email protected]>
Diffstat (limited to 'expt_0729.cu')
-rw-r--r--expt_0729.cu20
1 files changed, 11 insertions, 9 deletions
diff --git a/expt_0729.cu b/expt_0729.cu
index 2efa575..5660c6c 100644
--- a/expt_0729.cu
+++ b/expt_0729.cu
@@ -29,7 +29,7 @@ constexpr long SAMPLE_LENGTH = 1024;
constexpr long TEST_SIZE = 1024;
-__device__ key_type *SampleItem, *PopulationItem;
+__device__ key_type *cudaSampleItem, *cudaPopulationItem;
__device__ bool *cdf_result;
__device__ unsigned insert_value;
#ifdef TEST_BOUNDS
@@ -56,15 +56,16 @@ __device__ const key_type *cudaBinarySearch(key_type *start, key_type *end,
}
__device__ double sample_cdf(double x) {
- auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x);
- if (it == SampleItem + SAMPLE_LENGTH) {
+ auto it =
+ cudaBinarySearch(cudaSampleItem, cudaSampleItem + SAMPLE_LENGTH + 2, x);
+ if (it == cudaSampleItem + SAMPLE_LENGTH) {
return 1;
}
- if (it == SampleItem) {
+ if (it == cudaSampleItem) {
return 0;
}
auto it_prev = it - 1;
- return (double(it_prev - SampleItem) +
+ return (double(it_prev - cudaSampleItem) +
(x - (double)*it_prev) / (double)(*it - *it_prev)) /
double(SAMPLE_LENGTH - 1);
}
@@ -82,9 +83,10 @@ __global__ void initStorage(unsigned scale, unsigned test_size) {
}
__global__ void init(key_type *sample_item, key_type *population_item) {
- SampleItem = sample_item;
- PopulationItem = population_item;
- // cudaMalloc(&SampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long long));
+ cudaSampleItem = sample_item;
+ cudaPopulationItem = population_item;
+ // cudaMalloc(&cudaSampleItem, (SAMPLE_LENGTH + 2) * sizeof(unsigned long
+ // long));
// cudaMalloc(&Storage, TEST_SIZE * sizeof(key_type));
cdf_result = nullptr;
}
@@ -96,7 +98,7 @@ __global__ void kernel(unsigned long step, const double slice_size,
threadIdx.x + i;
// printf("%d\n", tid);
- auto index = (int)(sample_cdf(PopulationItem[tid]) / slice_size);
+ auto index = (int)(sample_cdf(cudaPopulationItem[tid]) / slice_size);
if (cdf_result[index]) {
while (cdf_result[++index]) {
assert(index < split_size);