fix: Fix expt_0729

Signed-off-by: KunoiSayami <[email protected]>

1 files changed, 28 insertions, 19 deletions

diff --git a/expt_0729.cu b/expt_0729.cu
index 4e28281..31aac7f 100644
--- a/expt_0729.cu
+++ b/expt_0729.cu
@@ -55,7 +55,7 @@ __device__ const key_type *cudaBinarySearch(key_type *start, key_type *end,
 }
 
 __device__ double sample_cdf(double x) {
-  auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH, x);
+  auto it = cudaBinarySearch(SampleItem, SampleItem + SAMPLE_LENGTH + 2, x);
   if (it == SampleItem + SAMPLE_LENGTH) {
     return 1;
   }
@@ -73,7 +73,7 @@ __global__ void initStorage(unsigned scale, unsigned test_size) {
   cudaMalloc(&cdf_result, scale * test_size * sizeof(bool));
   memset(cdf_result, 0, scale * test_size * sizeof(bool));
   insert_value = 0;
-  printf("init storage\n");
+  // printf("init storage\n");
 #ifdef TEST_BOUNDS
   index_max = 0;
   index_min = 0x7fffffff;
@@ -90,19 +90,19 @@ __global__ void init(key_type *sample_item, key_type *population_item) {
 
 __global__ void kernel(unsigned long step, const double slice_size,
                        const long split_size) {
-  for (int i = 1; i <= step; i++) {
-    auto tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x +
-               threadIdx.x - 1024;
+  for (int i = 0; i < step; i++) {
+    auto tid = step * gridDim.x * blockDim.x + blockIdx.x * blockDim.x +
+               threadIdx.x + i;
 
     // printf("%d\n", tid);
-    auto index = (int)(sample_cdf(PopulationItem[tid - 1024]) / slice_size);
+    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);
+    // atomicAdd(&insert_value, 1);
 
 #ifdef TEST_BOUNDS
     while (true) {
@@ -128,21 +128,16 @@ __global__ void kernel(unsigned long step, const double slice_size,
 }
 
 __global__ void print_function() {
-  printf("%u\n", insert_value);
+  // printf("%u\n", insert_value);
 #ifdef TEST_BOUNDS
   printf("%u %u\n", index_min, index_max);
 #endif
 }
 
-__device__ inline long double safe_ceil(double value) {
-  auto c = ceil(value);
-  return c == 0 ? 1 : c;
-}
-
 int main(int argc, char const *argv[]) {
   size_t test_size;
   if (argc == 1) {
-    test_size = 16384;
+    test_size = 1048576;
   } else {
     try {
       test_size = std::stol(argv[1]);
@@ -179,18 +174,32 @@ int main(int argc, char const *argv[]) {
   init<<<1, 1>>>(cudaSample, cudaPopulation);
   cudaDeviceSynchronize();
 
-  unsigned long step = test_size / 1024;
+  dim3 grid_dim = 2, block_dim = 512;
+
+  unsigned long step = test_size / (grid_dim.x * block_dim.x);
   printf("step: %lu\n", step);
-  assert(!(test_size % 1024));
+  assert(!(test_size % (grid_dim.x * block_dim.x)));
 
-  for (int scale = 2; scale <= 16; scale++) {
+  for (int scale = 2; scale <= 8; scale++) {
     const long split_size = test_size * scale;
     const auto slice_size = 1.0 / (double)split_size;
-    printf("scale: %i\n", scale);
+    printf("scale: %i ", scale);
     initStorage<<<1, 1>>>(scale, test_size);
     cudaDeviceSynchronize();
-    kernel<<<2, 512>>>(step, slice_size, split_size);
+    cudaEvent_t start, stop;
+    cudaEventCreate(&start);
+    cudaEventCreate(&stop);
+    cudaEventRecord(start, nullptr);
+    kernel<<<grid_dim, block_dim>>>(step, slice_size, split_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();
   }