Write comments on algorithms (#1158)

Co-authored-by: Stephen Nicholas Swatman <[email protected]>

Author: beomki-yeo

device/cuda/src/ambiguity_resolution/greedy_ambiguity_resolution_algorithm.cu

@@ -31,6 +31,8 @@ __global__ void add_block_offset(device::add_block_offset_payload payload) {
         return;
     }
 
+    // Add the scanned block offsets to block-wise prefix sums of the number of
+    // updated tracks.
     prefix_sums[globalIndex] += block_offsets[blockIdx.x - 1];
 }
 

device/cuda/src/ambiguity_resolution/kernels/add_block_offset.cu

@@ -22,6 +22,7 @@ __global__ void block_inclusive_scan(
         return;
     }
 
+    // temporary buffer where the block-wise prefix sum will be calculated
     extern __shared__ int shared_temp[];
 
     vecmem::device_vector<const unsigned int> sorted_ids(
@@ -37,13 +38,16 @@ __global__ void block_inclusive_scan(
 
     if (globalIndex >= n_accepted) {
         shared_temp[threadIndex] = 0;
-    } else {
+    }
+    // Start with boolean number depending on whether track id corresponding to
+    // the current thread is updated during the iteration
+    else {
         shared_temp[threadIndex] = is_updated[sorted_ids[globalIndex]];
     }
 
     __syncthreads();
 
-    // inclusive scan in shared memory
+    // Inclusive scan the boolean numbers to calculate the block-wise prefix sum
     for (int stride = 1; stride < blockDim.x; stride *= 2) {
         int val = 0;
         if (threadIndex >= stride) {
@@ -56,12 +60,15 @@ __global__ void block_inclusive_scan(
         __syncthreads();
     }
 
+    // Move the block-wise prefix_sums to global memory
     if (globalIndex < n_accepted) {
         prefix_sums[globalIndex] = shared_temp[threadIndex];
     }
 
     __syncthreads();
 
+    // Block offset, the last element of block-wise prefix sums, is also
+    // recorded to calculate full prefix sums later
     if (threadIndex == blockDim.x - 1) {
         block_offsets[blockIdx.x] = shared_temp[threadIndex];
     }

device/cuda/src/ambiguity_resolution/kernels/block_inclusive_scan.cu

@@ -31,6 +31,8 @@ __global__ void fill_inverted_ids(device::fill_inverted_ids_payload payload) {
         return;
     }
 
+    // Fill the inverted_ids vector which converts a track id to the index of
+    // sorted ids
     inverted_ids[sorted_ids[globalIndex]] = globalIndex;
 }
 

device/cuda/src/ambiguity_resolution/kernels/fill_inverted_ids.cu

@@ -96,6 +96,9 @@ __launch_bounds__(1024) __global__
         int ini_idx = stride * (threadIdx.x % nThreads_per_track);
         int fin_idx = std::min(ini_idx + stride, static_cast<int>(N));
 
+        // If it is an updated track, find new sorted index by using the binary
+        // search. The index is also shifted by using the bitonic sort result
+        // from sort_updated_tracks and prefix sums
         if (is_updated[tid]) {
 
             if (gid > 0) {
@@ -188,7 +191,10 @@ __launch_bounds__(1024) __global__
             if (offset != 0) {
                 atomicAdd(&shifted_idx, offset);
             }
-        } else {
+        }
+        // If it is not an updated track, it is enough to count the number of
+        // updated tracks which need to come earlier.
+        else {
 
             for (int i = ini_idx; i < fin_idx; i++) {
 
@@ -209,6 +215,7 @@ __launch_bounds__(1024) __global__
 
     __syncthreads();
 
+    // Save the result of new indices into a temporary buffer
     if (is_valid_thread && (threadIdx.x % nThreads_per_track) == 0) {
         temp_sorted_ids.at(shifted_idx) = tid;
     }

device/cuda/src/ambiguity_resolution/kernels/rearrange_tracks.cu

@@ -39,7 +39,8 @@ __global__ void scan_block_offsets(device::scan_block_offsets_payload payload) {
     }
     __syncthreads();
 
-    // 2. Inclusive scan (Hillis-Steele style)
+    // 2. Inclusive scan to caculated the scanned block offset which is the
+    // prefix sum of block offsets
     for (int offset = 1; offset < n_blocks_prev; offset *= 2) {
         int temp = 0;
         if (threadIndex >= offset) {

device/cuda/src/ambiguity_resolution/kernels/scan_block_offsets.cu

@@ -32,7 +32,7 @@ __launch_bounds__(512) __global__
 
     const unsigned int tid = threadIdx.x;
 
-    // Load to shared memory
+    // Load updated track indices into shared memory (for sorting)
     shared_mem_tracks[tid] = std::numeric_limits<unsigned int>::max();
 
     if (tid < *(payload.n_updated_tracks)) {
@@ -90,6 +90,7 @@ __launch_bounds__(512) __global__
         }
     }
 
+    // Write back the sorted result from shared memory to global memory
     if (tid < *(payload.n_updated_tracks)) {
         updated_tracks[tid] = shared_mem_tracks[tid];
     }