From bcfe8c96899cfda8df98e540d7d9147b6d8db2a0 Mon Sep 17 00:00:00 2001
From: KunoiSayami <i@leanhe.dev>
Date: Sat, 20 May 2023 00:49:40 +0800
Subject: feat(exp): Add expt_0517

Signed-off-by: KunoiSayami <i@leanhe.dev>
---
 expt_0517.cu | 91 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 91 insertions(+)
 create mode 100644 expt_0517.cu

(limited to 'expt_0517.cu')

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("");
+}
-- 
cgit v1.3.1