update prefixSum section

DsoTsin · DsoTsin · commit bbd44428b5a1 · 2024-04-22T21:02:56.000+08:00
diff --git a/8.misc/gvdb_notes.md b/8.misc/gvdb_notes.md
@@ -70,6 +70,8 @@ GPU的构建逻辑：
     * Voxel点云排序
     * CUDA::FUNC_FIND_UNIQUE
     * CUDA::AUX_UNIQUE_CNT
+    * CUDA::FUNC_PREFIXSUM
+    * CUDA::FUNC_PREFIXFIXUP
     * CUDA::FUNC_COMPACT_UNIQUE
     ![](images/gvdb_activate_bricks.png)
     ![](images/gvdb_build_alloc_resources.png)
@@ -116,4 +118,58 @@ GPU的构建逻辑：
         }
         ```
 
-* 核心CUDA代码位于 `gvdb_library/src/cuda_gvdb_nodes.cuh`
+* 核心CUDA代码位于 `gvdb_library/src/cuda_gvdb_nodes.cuh`
+
+
+### PrefixSum
+
+```cuda
+#define SCAN_BLOCKSIZE		512
+
+extern "C" __global__ void prefixFixup ( uint *input, uint *aux, int len) 
+{
+    unsigned int t = threadIdx.x;
+	unsigned int start = t + 2 * blockIdx.x * SCAN_BLOCKSIZE; 	
+	if (start < len)					input[start] += aux[blockIdx.x] ;
+	if (start + SCAN_BLOCKSIZE < len)   input[start + SCAN_BLOCKSIZE] += aux[blockIdx.x];
+}
+
+extern "C" __global__ void prefixSum ( uint* input, uint* output, uint* aux, int len, int zeroff )
+{
+    __shared__ uint scan_array[SCAN_BLOCKSIZE << 1];    
+	unsigned int t1 = threadIdx.x + 2 * blockIdx.x * SCAN_BLOCKSIZE;
+	unsigned int t2 = t1 + SCAN_BLOCKSIZE;
+    
+	// Pre-load into shared memory
+    scan_array[threadIdx.x] = (t1<len) ? input[t1] : 0.0f;
+	scan_array[threadIdx.x + SCAN_BLOCKSIZE] = (t2<len) ? input[t2] : 0.0f;
+    __syncthreads();
+
+    // Reduction
+    int stride;
+    for (stride = 1; stride <= SCAN_BLOCKSIZE; stride <<= 1) {
+       int index = (threadIdx.x + 1) * stride * 2 - 1;
+       if (index < 2 * SCAN_BLOCKSIZE)
+          scan_array[index] += scan_array[index - stride];
+       __syncthreads();
+    }
+
+    // Post reduction
+    for (stride = SCAN_BLOCKSIZE >> 1; stride > 0; stride >>= 1) {
+       int index = (threadIdx.x + 1) * stride * 2 - 1;
+       if (index + stride < 2 * SCAN_BLOCKSIZE)
+          scan_array[index + stride] += scan_array[index];
+       __syncthreads();
+    }
+	__syncthreads();
+	
+	// Output values & aux
+	if (t1+zeroff < len)	output[t1+zeroff] = scan_array[threadIdx.x];
+	if (t2+zeroff < len)	output[t2+zeroff] = (threadIdx.x==SCAN_BLOCKSIZE-1 && zeroff) ? 0 : scan_array[threadIdx.x + SCAN_BLOCKSIZE];	
+	if ( threadIdx.x == 0 ) {
+		if ( zeroff ) output[0] = 0;
+		if (aux) aux[blockIdx.x] = scan_array[2 * SCAN_BLOCKSIZE - 1];				
+	}    	
+}
+```
+
