doc: Add doc

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

1 files changed, 59 insertions, 133 deletions

diff --git a/publish_0630.cu b/publish_0630.cu
index 1befc23..107b024 100644
--- a/publish_0630.cu
+++ b/publish_0630.cu
@@ -66,16 +66,17 @@ Conference on Parallel and Distributed Systems, December 2012.
 
 // #include"cutil.h"		// Comment this if cutil.h is not available
 // #include "cuda_runtime.h"
+#include <algorithm>
 #include <cassert>
 #include <cstdio>
 #include <cstdlib>
+#include <fstream>
 #include <random>
 #include <set>
+#include <vector>
 
 typedef unsigned long long key_type;
 
-#include "read_helper.h"
-
 // Maximum level of a node in the skip list
 constexpr size_t MAX_LEVEL = 32;
 
@@ -175,7 +176,6 @@ public:
       return false;
     }
 
-    P_ERR("Reading sample");
     for (key_type i; read_number < sample_length && !fin.eof();
          store_into_vector(i)) {
       fin >> i;
@@ -183,13 +183,11 @@ public:
     }
 
     if (random_number > 0) {
-      P_ERR("\rReading skip");
       read_number = random_number;
       for (key_type i; read_number > 0 && !fin.eof(); fin >> i)
         read_number--;
     }
 
-    P_ERR("\rReading population");
     read_number = 0;
 
     auto remain = population_length + REVERSED_BLOCK;
@@ -199,7 +197,6 @@ public:
     }
 
     fin.close();
-    P_ERR("\r");
     return true;
   }
 
@@ -231,11 +228,6 @@ public:
 
 typedef ReadHelper_<unsigned long long> ReadHelper;
 
-#include <algorithm>
-#include <cassert>
-#include <cstdio>
-#include <vector>
-
 typedef unsigned long long key_type;
 
 class CustomSort {
@@ -312,7 +304,7 @@ public:
 };
 
 template <typename T> void rebuild(std::vector<T> &original) {
-  auto sample_length = original.size();
+  const auto sample_length = original.size();
   auto sorter = CustomSort(sample_length, sizeof(T) * 8);
   auto tmp = new T[sample_length];
 
@@ -366,7 +358,7 @@ public:
     }
     auto it_prev = it - 1;
     return (double(it_prev - begin) +
-            (x - (double)*it_prev) / (double)(*it - *it_prev)) /
+            ((double)x - (double)*it_prev) / (double)(*it - *it_prev)) /
            double(length - 1);
   }
 };
@@ -499,34 +491,10 @@ public:
     auto index = trailing_zeroes((size_t)cdf_index);
     return index;
   }
-
-#ifdef MEASURE_ACCESS
-  unsigned access_times = 0;
-
-  __device__ unsigned getAccessCount() const { return this->access_times; }
-  __device__ void increaseAccessCount(unsigned count = 1) {
-    atomicAdd(&this->access_times, count);
-  }
-#else
-  __device__ void increaseAccessCount(unsigned = 1) {}
-#endif
-
-#ifdef MEASURE_TIME
-  unsigned round = 0;
-  __device__ void increaseRoundCount(unsigned count = 1) {
-    atomicAdd(&this->round, count);
-  }
-  int spend_time[NUM_ITEMS]{0};
-  unsigned long long total_time = 0;
-  __device__ unsigned getRoundCount() const { return this->round; }
-#endif
 };
 
 __device__ Node **nodes;                  // Pool of pre-allocated nodes
 __device__ unsigned int pointerIndex = 0; // Index into pool of free nodes
-//__device__ key_type *randoms; // Array storing the levels of the nodes in the
-// free
-// pool
 __device__ unsigned int NODE_LIMIT;
 
 // Function for creating a new node when requested by an add operation
@@ -546,10 +514,6 @@ __device__ Node *GetNewNode(key_type key, size_t topLevel) {
 
 __device__ LockFreeSkipList *lockFreeSkipList; // The lock-free skip list
 
-//__device__ key_type KeyIndex[KEY_INDEX_SIZE];
-
-//__device__ key_type SampleStorage[SAMPLE_SIZE];
-
 // Kernel for initializing device memory
 
 __global__ void init(LockFreeSkipList *l1, Node **n,
@@ -620,20 +584,11 @@ __device__ bool LockFreeSkipList::Search(key_type key) {
   int level;
   for (level = MAX_LEVEL; level >= bottomLevel; level--) {
     curr = pred->GetReference(level);
-#ifdef MEASURE_ACCESS
-    this->increaseAccessCount();
-#endif
     while (true) {
       succ = curr->Get(level, &marked);
-#ifdef MEASURE_ACCESS
-      this->increaseAccessCount();
-#endif
       while (marked) {
         curr = curr->GetReference(level);
         succ = curr->Get(level, &marked);
-#ifdef MEASURE_ACCESS
-        this->increaseAccessCount(2);
-#endif
       }
       if (curr->key < key) {
         pred = curr;
@@ -752,18 +707,8 @@ __global__ void kernel(const key_type *items, size_t search_length,
 
     // Grab the operation and the associated key and execute
     key_type item = items[tid];
-#ifdef MEASURE_TIME
-    unsigned long long start_time = clock64();
-#endif
     result[tid] = lockFreeSkipList->Search(item);
     assert(result[tid]);
-#ifdef MEASURE_TIME
-    unsigned long long end_time = clock64() - start_time;
-    if (lockFreeSkipList->spend_time[tid]) {
-      printf("conflict: %d\n", tid);
-    }
-    lockFreeSkipList->spend_time[tid] = (int)end_time;
-#endif
   }
 }
 
@@ -780,22 +725,12 @@ __global__ void kernelAdd(key_type *item, size_t insertion_length) {
   }
 }
 
-// Generate the level of a newly created node
-
-__global__ void print_function() {
-#ifdef MEASURE_ACCESS
-  printf("count: %u\n", l->getAccessCount());
-#endif
-}
-
-/*__global__ void copy_function(int *spend_time) {
-  memcpy(spend_time, lockFreeSkipList->spend_time, sizeof(int) * NUM_ITEMS);
-}*/
-
 inline double calcSliceSize(size_t insertion_size) {
   return 1.0 / (double)insertion_size;
 }
 
+/// 計算 blocks 的大小，該大小和 NUM_THREADS
+/// 相乘應能正好大於等於需要插入的大小
 inline auto calcBlocks(size_t input) {
   return (input % (NUM_THREADS * FACTOR) == 0)
              ? input / (NUM_THREADS * FACTOR)
@@ -808,121 +743,145 @@ int main(int argc, char **argv) {
     exit(1);
   }
 
+  /// 從命令行中讀取樣例的大小
   auto sample_length = strtol(argv[1], nullptr, 10);
+  /// 從命令行中讀取需要搜尋的大小
   auto search_length = strtol(argv[2], nullptr, 10);
+  /// 從命令行中讀取需要插入節點的大小
   auto insertion_length = strtol(argv[3], nullptr, 10);
+
+  /// 定義一個數用來考慮是否需要隨機化插入的元素
   auto total_row = 0L;
 
+  /// 因為該參數可選，所以需要進行判斷輸入的參數個數是否大於 4
   if (argc > 4) {
+    /// 從命令行中讀取讀入資料的總大小
     total_row = strtol(argv[4], nullptr, 10);
   }
 
+  /// 參數檢查
+  /// 因為搜尋是從插入的資料中得來，如果搜尋數大於插入數，則沒有其他的意義
   if (insertion_length < search_length) {
     printf("Search should smaller than insertion\n");
   }
 
+  /// 列印出程式的參數
   printf("Sample: %ld, Insertion: %ld, Search: %ld ", sample_length,
          insertion_length, search_length);
+  /// 刷新輸出流
   fflush(stdout);
 
+  /// 宣告一個讀入 unsigned long long 的幫助類
+  /// 將文件名和上面讀入的數值作為參數
   ReadHelper readHelper("normal_distribution.txt", sample_length,
                         insertion_length, total_row);
 
+  /// 如果有指定需要隨機化讀入，則輸出跳過了多少的數才開始正式地讀取
   if (total_row) {
     printf("Skip: %ld ", readHelper.random_number);
     fflush(stdout);
   }
 
+  /// 讀入文件
   readHelper.readFile();
+
+  /// 宣告兩個 unsigned long long 的 vector ，用來存放樣例和插入的數值
   std::vector<key_type> _sample, _population;
+  /// 調用讀入幫助類，將兩個數值分別賦值到上述宣告的 vector 中
   readHelper.split_into(_sample, _population);
-  /*printf("%lu, Create search vector: %ld\n", _population.size(),
-         _population.end() - (_population.begin() + insertion_length));*/
+  /// 重新排序
   rebuildSort(_sample);
   std::vector<key_type> _search(_population.begin(),
                                 _population.begin() + search_length);
-  //_population.resize(insertion_length);
-
-  // Allocate necessary arrays
-  // key_type *op = new key_type[NUM_ITEMS];     //(key_type
-  // *)malloc(sizeof(key_type) * NUM_ITEMS); key_type *levels = new
-  // key_type[NUM_ITEMS]; //(key_type *)malloc(sizeof(key_type) * NUM_ITEMS);
-  // key_type *items = new key_type[NUM_ITEMS];  //(key_type
-  // *)malloc(sizeof(key_type) * NUM_ITEMS);
-  auto *result =
-      new key_type[search_length]; //(key_type *)malloc(sizeof(key_type) *
-                                   // NUM_ITEMS);
 
   // Allocate device memory
 
+  /// 宣告存儲需要操作的陣列指標
   key_type *cudaOperatorItems;
   // key_type *Cop;
+  /// 宣告存儲結果陣列的指標
   key_type *cudaResult;
 
+  /// 分配存儲查詢結果的記憶體
   cudaMalloc(&cudaResult, sizeof(key_type) * search_length);
+  /// 分配用於查詢的記憶體
   cudaMalloc(&cudaOperatorItems, sizeof(key_type) * insertion_length);
-  // cudaMalloc(&Cop, sizeof(key_type) * NUM_ITEMS);
-  // cudaMemcpy(Clevels, levels, sizeof(key_type) * NUM_ITEMS,
-  // cudaMemcpyHostToDevice);
+  /// 將需要查詢的物件複製到記憶體中
   cudaMemcpy(cudaOperatorItems, _population.data(),
              sizeof(key_type) * insertion_length, cudaMemcpyHostToDevice);
-  // cudaMemcpy(Cop, op, sizeof(key_type) * NUM_ITEMS, cudaMemcpyHostToDevice);
-  Node **pointers =
-      (Node *
-           *)new key_type[insertion_length]; // malloc(sizeof(key_type) * adds);
+  /// 分配插入操作所需要的節點陣列
+  Node **pointers = (Node **)new key_type[insertion_length];
   Node **Cpointers;
 
   // Allocate the pool of free nodes
 
+  /// 給每個指標分配節點的記憶體
   for (int i = 0; i < insertion_length; i++) {
     cudaMalloc(&pointers[i], sizeof(Node));
   }
+  /// 分配 GPU 處的節點記憶體
   cudaMalloc(&Cpointers, sizeof(Node *) * insertion_length);
+  /// 將存儲節點的陣列複製到記憶體中
   cudaMemcpy(Cpointers, pointers, sizeof(Node *) * insertion_length,
              cudaMemcpyHostToDevice);
 
   // Allocate the skip list
 
+  /// 宣告一個指向 GPU 記憶體的 Skip lists 的指標
   LockFreeSkipList *Clist;
+  /// 在 CPU 中先把 Skip Lists 創建出來
   auto *list = new LockFreeSkipList(_sample.data(), _sample.size(),
                                     calcSliceSize(insertion_length));
 
+  /// 分配 GPU 記憶體給 Skip Lists
   cudaMalloc(&Clist, sizeof(LockFreeSkipList));
+  /// 将在 CPU 处建立的 Skip Lists 複製到 GPU 記憶體中
   cudaMemcpy(Clist, list, sizeof(LockFreeSkipList), cudaMemcpyHostToDevice);
   // Calculate the number of thread blocks
   // NUM_ITEMS = total number of operations to execute
   // NUM_THREADS = number of threads per block
   // FACTOR = number of operations per thread
 
+  /// 計算 blocks 的大小，該大小和 NUM_THREADS
+  /// 相乘應能正好大於等於需要插入的大小
   size_t blocks = calcBlocks(insertion_length);
 
-  CudaCheckError();
-
   // Initialize the device memory
+  /// 將需要的值賦值到 GPU 中的變數中
   init<<<1, 32>>>(Clist, Cpointers, insertion_length);
   cudaDeviceSynchronize();
 
   // Insertion to skiplist
+  /// 該 GPU 語句啓動的 GPU 執行緒會將所需要的數據插入到 Skip lists 中
   kernelAdd<<<blocks, NUM_THREADS>>>(cudaOperatorItems, insertion_length);
   cudaDeviceSynchronize();
 
   // Re-allocate memory for search
+  /// 將原來的記憶體釋放
   cudaFree(cudaOperatorItems);
+  /// 分配新的記憶體大小給用於搜尋的數組
   cudaMalloc(&cudaOperatorItems, sizeof(key_type) * search_length);
+  /// 將需要搜尋的數值複製到 GPU 的記憶體中
   cudaMemcpy(cudaOperatorItems, _search.data(),
              sizeof(key_type) * search_length, cudaMemcpyHostToDevice);
 
-  // Launch main kernel
+  /// 計算 blocks 的大小，該大小和 NUM_THREADS
+  /// 相乘應能正好大於等於需要搜尋的大小
   blocks = calcBlocks(search_length);
 
+  /// 創建用來計算時間的變數
   cudaEvent_t start, stop;
   cudaEventCreate(&start);
   cudaEventCreate(&stop);
   cudaEventRecord(start, nullptr);
 
+  /// 執行 Kernel
   kernel<<<blocks, NUM_THREADS>>>(cudaOperatorItems, search_length, cudaResult);
+  /// 檢測執行緒是否有錯誤
   CudaCheckError();
+  /// 同步執行緒等待其完成
   cudaDeviceSynchronize();
+
   cudaEventRecord(stop, nullptr);
   cudaEventSynchronize(stop);
   float time;
@@ -930,50 +889,17 @@ int main(int argc, char **argv) {
   cudaEventDestroy(start);
   cudaEventDestroy(stop);
 
-  // Print kernel execution time in milliseconds
-
+  /// 列印出搜尋的長度以及所耗費的時間（單位為毫秒）
   printf("%lu: %lf\n", search_length, time);
-  // Check for errors
 
-  // Move results back to host memory
+  /// 分配用來存儲結果的記憶體
+  auto *result = new key_type[search_length];
 
+  /// 將存儲結果的陣列複製回 CPU 的記憶體中
   cudaMemcpy(result, cudaResult, sizeof(key_type) * search_length,
              cudaMemcpyDeviceToHost);
 
-  // Uncomment the following for debugging
-  // print<<<1,32>>>();
-  cudaDeviceSynchronize();
-
-#if (defined(MEASURE_TIME) || defined(MEASURE_ACCESS))
-
-  print_function<<<1, 1>>>();
-
-  cudaDeviceSynchronize();
-#ifdef MEASURE_TIME
-  {
-    int *cuda_tmp = nullptr;
-    cudaMalloc(&cuda_tmp, sizeof(int) * NUM_ITEMS);
-    copy_function<<<1, 1>>>(cuda_tmp);
-    cudaDeviceSynchronize();
-    CudaCheckError();
-    FILE *file = fopen("spend_time.txt", "w");
-    int *tmp = new int[NUM_ITEMS];
-    cudaMemcpy(tmp, cuda_tmp, sizeof(int) * NUM_ITEMS, cudaMemcpyDeviceToHost);
-    // memcpy(tmp, SpendTime, sizeof(int) * NUM_ITEMS);
-    for (int i = 0; i < NUM_ITEMS; i++) {
-      if (tmp[i] == 0)
-        break;
-      fprintf(file, "%d\n", tmp[i]);
-    }
-    // printf("%d\n", i);
-    delete[] tmp;
-    fclose(file);
-  }
-  // for (auto element : SpendTimeVec)
-  //   printf("%d\n", element);
-#endif
-#endif
-
+  /// 釋放記憶體
   cudaFree(Clist);
   cudaFree(cudaResult);
   cudaFree(cudaOperatorItems);