3 files changed, 104 insertions, 5 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 06e2eec..baf37e7 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -215,3 +215,12 @@ set_target_properties(expt_0516_2 PROPERTIES CUDA_ARCHITECTURES "75")
 set_target_properties(expt_0516_2 PROPERTIES LINKER_LANGUAGE CUDA)
 
 
+add_executable(expt_0517 expt_0517.cu)
+target_link_libraries(expt_0517 m stdc++)
+
+set_target_properties(expt_0517 PROPERTIES
+        CUDA_SEPARABLE_COMPILATION ON)
+set_target_properties(expt_0517 PROPERTIES CUDA_ARCHITECTURES "75")
+set_target_properties(expt_0517 PROPERTIES LINKER_LANGUAGE CUDA)
+
+
diff --git a/expt_0510.cu b/expt_0510.cu
index 3efad47..82f2624 100644
--- a/expt_0510.cu
+++ b/expt_0510.cu
@@ -15,7 +15,7 @@ constexpr size_t TEST_LENGTH = 32'768;
 // constexpr size_t TEST_LENGTH = 4096;
 // constexpr size_t RESERVED_BLOCK = 100;
 
-unsigned long long max_value = 0, min_value = 0xfffffffff;
+unsigned long long max_value = 0, min_value = 0xffffffffffff;
 typedef std::pair<int, int> dim_pair_type;
 
 inline void store_into_vector(unsigned long long value) {
@@ -61,8 +61,7 @@ __global__ void kernel(unsigned long step, const double slice_size) {
 
     // printf("%d\n", tid);
     auto _index = (int)(custom_sort.sample_cdf_custom_version(
-                            cudaSampleItem, cudaSampleItem + cuda_sample_length,
-                            cudaPopulationItem[tid]) /
+                            cudaSampleItem, cudaPopulationItem[tid]) /
                         slice_size) -
                   1;
     // cdf_result[index] = true;
@@ -224,7 +223,7 @@ int main(int argc, char const *argv[]) {
     total_row = strtol(argv[3], nullptr, 10);
   }
 
-  CustomSort::testSelf(sample_length);
+  // CustomSort::testSelf(sample_length);
 
   readFile("normal_distribution.txt", sample_length, total_row, test_size);
 
@@ -251,7 +250,7 @@ int main(int argc, char const *argv[]) {
   // auto custom_sort = CustomSort(SAMPLE_LENGTH, sizeof(long) * 8);
   // custom_sort.testCalculation();
   applyCudaSampleLength<<<1, 1>>>(sample_length);
-  testCustomCalculation<<<1, 1>>>(sample_length);
+  // testCustomCalculation<<<1, 1>>>(sample_length);
   cudaDeviceSynchronize();
 
   initSample<<<1, 1>>>(cudaSample, cudaPopulation);
diff --git a/expt_0517.cu b/expt_0517.cu
new file mode 100644
index 0000000..ea73b08
--- /dev/null
+++ b/expt_0517.cu
@@ -0,0 +1,91 @@
+// Experimental content: Test log performance
+#include "sortlib.cuh"
+#include <algorithm>
+#include <cstdio>
+
+#ifndef LOCKFREE_SORTLIB_CUH
+#include <cassert>
+class CustomSort {
+public:
+  __device__ __host__ static size_t fast_log(size_t a) {
+#ifdef __CUDA_ARCH__
+    return (size_t)log2((double)a);
+#else
+    return (size_t)std::log2(a);
+#endif
+  }
+};
+
+#endif
+
+class OriginalFastLog {
+public:
+  __device__ __host__ static size_t fast_log(size_t a) {
+    float t = a;
+    return (((*(int *)&t) >> 23) + 1) & 127;
+  }
+};
+
+constexpr int TEST_SIZE = 262144;
+
+__global__ void init(int *result1, int *result2) {
+  memset(result1, 0, sizeof(int) * 32);
+  memset(result2, 0, sizeof(int) * 32);
+}
+
+__global__ void testNew(int *result) {
+  for (int i = 2; i < TEST_SIZE; i++) {
+    result[CustomSort::fast_log(i)]++;
+  }
+}
+
+__global__ void testOld(int *result) {
+  for (int i = 2; i < TEST_SIZE; i++) {
+    result[OriginalFastLog::fast_log(i)]++;
+  }
+}
+
+__global__ void testResult(const int *result1, const int *result2) {
+  for (int i = 0; i < 32; i++) {
+    assert(result1[i] == result2[i]);
+  }
+}
+
+void run_kernel(int *dst, bool custom = false) {
+
+  cudaEvent_t start, stop;
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+  cudaEventRecord(start, nullptr);
+  if (custom) {
+    testNew<<<1, 1>>>(dst);
+  } else {
+    testOld<<<1, 1>>>(dst);
+  }
+
+  cudaDeviceSynchronize();
+  cudaEventRecord(stop, nullptr);
+  cudaEventSynchronize(stop);
+  float time;
+  cudaEventElapsedTime(&time, start, stop);
+  cudaEventDestroy(start);
+  cudaEventDestroy(stop);
+  printf("%stime: %lf ", custom ? "custom " : "", time);
+  cudaDeviceSynchronize();
+}
+
+int main() {
+  int *result1, *result2;
+  cudaMalloc(&result1, sizeof(int) * 32);
+  cudaMalloc(&result2, sizeof(int) * 32);
+
+  init<<<1, 1>>>(result1, result2);
+  cudaDeviceSynchronize();
+
+  run_kernel(result1);
+  run_kernel(result2);
+
+  cudaFree(result1);
+  cudaFree(result2);
+  puts("");
+}