feat(exp): Remove useless code in expt_0516

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

1 files changed, 28 insertions, 44 deletions

diff --git a/expt_0516.cu b/expt_0516.cu
index 351a8c4..1954380 100644
--- a/expt_0516.cu
+++ b/expt_0516.cu
@@ -8,6 +8,9 @@
 #include <cstdio>
 #include <vector>
 
+constexpr auto NUM_THREADS = 512;
+constexpr auto FACTOR = 1;
+
 std::vector<unsigned long long> population_vector, sample_vector,
     le_sample_vector;
 
@@ -54,70 +57,57 @@ inline void cudaCheckError(const char *file, const int line) {
 }
 */
 
-__global__ void kernel(unsigned long step, const double slice_size) {
+__global__ void kernel(const size_t test_length) {
   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;
+  for (int i = 0; i < FACTOR; i++) {
+    auto tid =
+        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
 
+    if (tid >= test_length)
+      return;
     // printf("%d\n", tid);
-    auto _index = (int)(custom_sort.sample_cdf_custom_version(
-                            cudaSampleItem, cudaPopulationItem[tid]) /
-                        slice_size) -
-                  1;
+    custom_sort.sample_cdf_custom_version(cudaSampleItem,
+                                          cudaPopulationItem[tid]);
     // cdf_result[index] = true;
   }
 }
 
-__global__ void kernel2(unsigned long step, const double slice_size) {
+__global__ void kernel2(const size_t test_length) {
 
-  for (int i = 0; i < step; i++) {
-    auto tid = step * (blockIdx.x * blockDim.x + threadIdx.x) + i;
-    auto _index =
-        (int)(FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length,
-                                      cudaPopulationItem[tid]) /
-              slice_size) -
-        1;
-    // cdf_result[index] = true;
-  }
-}
-
-__global__ void print2() {
-  /*for (int i = 0; i< SAMPLE_LENGTH; i++) {
-    printf("%llu ", cudaSampleItem[i]);
-  }
-  printf("\n");*/
-  printf("%lu\n", CustomSort::fast_log(cuda_sample_length));
-}
+  for (int i = 0; i < FACTOR; i++) {
+    auto tid =
+        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
 
-__global__ void printFunction() {}
+    if (tid >= test_length)
+      return;
 
-__global__ void checkItems() {
-  for (int i = 0; i < cuda_sample_length; i++) {
-    printf("%lld ", cudaSampleItem[i]);
+    FactorySort::sample_cdf(cudaNormalSampleItem, cuda_sample_length,
+                            cudaPopulationItem[tid]);
+    // cdf_result[index] = true;
   }
-  printf("\n");
 }
 
-constexpr dim3 grid_dim = 32, block_dim = 32;
+inline size_t calcBlocks(size_t input) {
+  return (input % (NUM_THREADS * FACTOR) == 0)
+             ? input / (NUM_THREADS * FACTOR)
+             : (input / (NUM_THREADS * FACTOR)) + 1;
+}
 
 void run_kernel(size_t test_size, bool custom = true) {
   // need description
-  unsigned long step = test_size / (grid_dim.x * block_dim.x);
 
   // Only scale mode will use
-  auto scale = 2;
-  const unsigned long split_size = test_size * scale;
-  const auto slice_size = 1.0 / (double)split_size;
+  auto blocks = calcBlocks(test_size);
 
   cudaEvent_t start, stop;
   cudaEventCreate(&start);
   cudaEventCreate(&stop);
   cudaEventRecord(start, nullptr);
   if (custom) {
-    kernel<<<grid_dim, block_dim>>>(step, slice_size);
+    kernel<<<blocks, NUM_THREADS>>>(test_size);
   } else {
-    kernel2<<<grid_dim, block_dim>>>(step, slice_size);
+    kernel2<<<blocks, NUM_THREADS>>>(test_size);
   }
 
   cudaDeviceSynchronize();
@@ -132,9 +122,6 @@ void run_kernel(size_t test_size, bool custom = true) {
   }
   printf("%stime: %lf ", custom ? "custom " : "", time);
   cudaDeviceSynchronize();
-
-  printFunction<<<1, 1>>>();
-  cudaDeviceSynchronize();
 }
 
 __global__ void freeStorageStage1() { cudaFree(cudaSampleItem); }
@@ -229,9 +216,6 @@ int main(int argc, char const *argv[]) {
         "kernel\n");
   run_kernel(test_size);
 
-  // freeStorageStage1<<<1, 1>>>();
-  // cudaDeviceSynchronize();
-
   cudaFree(cudaSample);
 
   key_type *cudaNormalSample = nullptr;