feat(exp): Add expt_0525.cu

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

1 files changed, 28 insertions, 30 deletions

diff --git a/expt_0520.cu b/expt_0520.cu
index 88815a7..d31dc78 100644
--- a/expt_0520.cu
+++ b/expt_0520.cu
@@ -111,18 +111,12 @@ constexpr size_t FACTOR = 1;
 
 // should change this to dynamic next time
 // constexpr size_t KEY_INDEX_SIZE = 32;
-constexpr size_t SAMPLE_SIZE = 1024;
+// constexpr size_t SAMPLE_SIZE = 1023;
 
 constexpr int block_size = STEP_SIZE;
 
 typedef LL key_type;
 
-#ifdef RANDOM_TARGET
-constexpr const char *TARGET_STRING = "RANDOM";
-#else
-constexpr const char *TARGET_STRING = "PERFECT";
-#endif
-
 #define CUDA_ERROR_CHECK
 
 #define CudaSafeCall(err) __cudaSafeCall(err, __FILE__, __LINE__)
@@ -158,30 +152,38 @@ inline void __cudaCheckError(const char *file, const int line) {
 #endif
 }
 
-__device__ key_type SampleStorage[SAMPLE_SIZE];
+//__device__ key_type SampleStorage[SAMPLE_SIZE];
+__device__ LockFreeSkipList *cudaGlobalSkipList;
 
 // Kernel for initializing device memory
 
-__global__ void init(Node **n) { nodes = n; }
+__global__ void init(LockFreeSkipList *list, Node **n) {
+  cudaGlobalSkipList = list;
+  nodes = n;
+}
+
+__global__ void testFunctions() { cudaGlobalSkipList->testSample(); }
 
 // The main kernel
 
-__global__ void kernel(LockFreeSkipList *skipList, const key_type *population,
-                       size_t insertion_length) {
+__global__ void kernel(const key_type *population, size_t insertion_length) {
   // The array items holds the sequence of keys
   // The array op holds the sequence of operations
   // The array result, at the end, will hold the outcome of the operations
 
   for (int i = 0; i < FACTOR;
        i++) { // FACTOR is the number of operations per thread
-    auto tid =
-        i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x;
+    // auto tid = i * gridDim.x * blockDim.x + blockIdx.x * blockDim.x +
+    // threadIdx.x;
+    auto tid = FACTOR * (blockIdx.x * blockDim.x + threadIdx.x) + i;
+    // printf("%lu\n", tid);
     if (tid >= insertion_length)
       return;
 
     // Grab the operation and the associated key and execute
     key_type item = population[tid];
-    skipList->Add(item);
+    // printf("%llu\n", population[tid]);
+    cudaGlobalSkipList->Add(item);
   }
 }
 
@@ -207,20 +209,16 @@ int main(int argc, char **argv) {
          sample_length);
 
   ReadHelper reader("normal_distribution.txt", sample_length, insertion_length);
-  printf("%d\n", __LINE__);
   reader.readFile(total_row);
 
-  printf("%d\n", __LINE__);
   key_type *cudaPopulation;
 
-  printf("%d\n", __LINE__);
   cudaMalloc(&cudaPopulation, sizeof(key_type) * insertion_length);
   std::vector<key_type> _sample, _population;
   reader.split_into(_sample, _population);
   cudaMemcpy(cudaPopulation, _population.data(),
              sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice);
 
-  printf("%d\n", __LINE__);
   // Allocate device memory
 
   // cudaMalloc((void **)&Clevels, sizeof(LL) * NUM_ITEMS);
@@ -230,8 +228,6 @@ int main(int argc, char **argv) {
       (Node **)new LL[insertion_length]; // malloc(sizeof(LL) * adds);
   Node **cudaNodePointers;
 
-  printf("%d\n", __LINE__);
-
   // Allocate the pool of free nodes
 
   for (int i = 0; i < insertion_length; i++) {
@@ -242,14 +238,15 @@ int main(int argc, char **argv) {
              cudaMemcpyHostToDevice);
   CudaCheckError();
 
-  printf("%d\n", __LINE__);
   // Allocate the skip list
 
-  LockFreeSkipList *Clist;
-  auto *list = new LockFreeSkipList(_sample.data(), SAMPLE_SIZE);
+  // LockFreeSkipList *cudaLockFreeSkipList;
+  auto *list = new LockFreeSkipList(_sample.data(), sample_length);
 
-  cudaMalloc((void **)&Clist, sizeof(LockFreeSkipList));
-  cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
+  LockFreeSkipList *cudaSkipList = nullptr;
+  cudaMalloc(&cudaSkipList, sizeof(LockFreeSkipList));
+  cudaMemcpy(cudaSkipList, list, sizeof(LockFreeSkipList),
+             cudaMemcpyHostToDevice);
   CudaCheckError();
   // Calculate the number of thread blocks
   // NUM_ITEMS = total number of operations to execute
@@ -267,7 +264,10 @@ int main(int argc, char **argv) {
   }
 
   // Initialize the device memory
-  init<<<1, 32>>>(cudaNodePointers);
+  init<<<1, 32>>>(cudaSkipList, cudaNodePointers);
+  cudaDeviceSynchronize();
+
+  testFunctions<<<1, 1>>>();
   cudaDeviceSynchronize();
 
   // Launch main kernel
@@ -277,7 +277,7 @@ int main(int argc, char **argv) {
   cudaEventCreate(&stop);
   cudaEventRecord(start, nullptr);
 
-  kernel<<<blocks, NUM_THREADS>>>(Clist, cudaPopulation, insertion_length);
+  kernel<<<blocks, NUM_THREADS>>>(cudaPopulation, insertion_length);
   CudaCheckError();
   cudaDeviceSynchronize();
   cudaEventRecord(stop, nullptr);
@@ -289,8 +289,6 @@ int main(int argc, char **argv) {
 
   // Print kernel execution time in milliseconds
 
-  printf("%s %d ", TARGET_STRING, block_size);
-
   printf("%lu: %lf", NUM_ITEMS, time);
 
 #if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS))
@@ -322,7 +320,7 @@ int main(int argc, char **argv) {
   //   printf("%d\n", element);
 #endif
 #endif
-  cudaFree(Clist);
+  // cudaFree(cudaLockFreeSkipList);
   for (int i = 0; i < insertion_length; i++) {
     cudaFree(pointers[i]);
   }