尝试理解前缀和执行

Question

我正在尝试理解《CUDA 手册》一书中提到的扫描实现 scan-then-fan。

1. 有人可以解释一下该设备的功能吗

   scanWarp

   ？为什么指数为负？你能举一个数值例子吗？
2. 我对线路也有同样的问题

   warpPartials[16+warpid] = sum

   。任务进展如何？
3. 这行的贡献是哪

   if ( warpid==0 ) {scanWarp<T,bZeroPadded>( 16+warpPartials+tid ); }

4. 请有人解释一下吗

   sum += warpPartials[16+warpid-1];

   ？数值示例将受到高度赞赏。
5. 最后，一个更面向 C++ 的问题，我们如何知道

   *sPartials = sum;

   中使用的索引来存储

   sPartials

   中的值？

PS： 演示整个执行过程的数值示例将非常有帮助。

template < class T, bool bZeroPadded > 
inline __device__ T
scanBlock( volatile T *sPartials ){


   extern __shared__ T warpPartials[];
   const int tid = threadIdx.x;
   const int lane = tid & 31;
   const int warpid = tid >> 5;

   //
   // Compute this thread's partial sum
   //
   T sum = scanWarp<T,bZeroPadded>( sPartials );
   __syncthreads();

   //
   // Write each warp's reduction to shared memory
   // 
   if ( lane == 31 ) {
       warpPartials[16+warpid] = sum;
   }
   __syncthreads();

   //
   // Have one warp scan reductions
   //
   if ( warpid==0 ) {
       scanWarp<T,bZeroPadded>( 16+warpPartials+tid );
   }
   __syncthreads();

   //
   // Fan out the exclusive scan element (obtained
   // by the conditional and the decrement by 1)
   // to this warp's pending output
   //
   if ( warpid > 0 ) {
       sum += warpPartials[16+warpid-1];
   }
   __syncthreads();

   //
   // Write this thread's scan output
   //
   *sPartials = sum;
   __syncthreads();

   //
   // The return value will only be used by caller if it
   // contains the spine value (i.e. the reduction
   // of the array we just scanned).
   //
   return sum;
}


template < class T >
inline __device__ T 
scanWarp( volatile T *sPartials ){

   const int tid = threadIdx.x;
   const int lane = tid & 31;

   if ( lane >=  1 ) sPartials[0] += sPartials[- 1];
   if ( lane >=  2 ) sPartials[0] += sPartials[- 2];
   if ( lane >=  4 ) sPartials[0] += sPartials[- 4];
   if ( lane >=  8 ) sPartials[0] += sPartials[- 8];
   if ( lane >= 16 ) sPartials[0] += sPartials[-16];

   return sPartials[0];
}

Answer 1

先扫描后扇动策略在两个层面上应用。对于网格级扫描（在全局内存上操作），部分数据被写入主机代码中分配的临时全局内存缓冲区，通过递归调用主机函数进行扫描，然后通过单独的内核调用添加到最终输出。对于块级扫描（在共享内存上操作），部分数据被写入共享内存的基址 (

warpPartials[]

)，通过一个 warp 进行扫描，然后添加到块级扫描的最终输出中。您询问的代码正在执行块级扫描。

1. 您引用的

   scanWarp

   的实现是使用已添加

   threadIdx.x

   的共享内存指针调用的，因此每个线程的

   sPartials

   版本都指向不同的共享内存元素。在

   sPartials

   上使用固定索引会导致相邻线程对相邻共享内存元素进行操作。负索引是可以的，只要它们不会导致数组索引越界。该实现借鉴了用零填充共享内存的优化版本，因此每个线程都可以无条件地使用固定的负索引，并且低于某个索引的线程只读取零。 （清单 13.14）它可以轻松地在经纱中的最低线程上进行谓词执行并使用正索引。

2. 每个 32 线程扭曲的第 31 个线程包含该扭曲的部分和，必须将其存储在某处以便扫描然后添加到输出中。

   warpPartials[]

   为第一个元素的共享内存起别名，因此可用于保存每个扭曲的部分和。您可以使用共享内存的任何部分来进行此计算，因为每个线程已经在寄存器中拥有自己的扫描值（赋值

   T sum = scanWarp...

   ）。

3. 某些扭曲（可以是任何扭曲，因此也可能是扭曲0）必须扫描写入到

   warpPartials[]

   的部分。最多需要 1 个 warp，因为硬件限制为每块 1024 个线程 = 1024/32 或 32 个 warp。因此，这段代码利用了这样的巧合：每个块的最大线程数除以扭曲计数，不大于每个扭曲的最大线程数。

4. 此代码将扫描的每扭曲部分添加到每个输出元素。第一个扭曲已经具有正确的值，因此仅由第二个和后续扭曲完成加法。另一种看待这个问题的方法是，它将扭曲部分的独占扫描添加到输出中。

5. scanBlock

   是一个设备函数 - 地址算术由其调用者完成，

   scanAndWritePartials

   ：

   volatile T *myShared = sPartials+tid;

Answer 2

（现在重写答案，我有更多时间）

这是一个示例（基于我用 C++ AMP 编写的实现，而不是 CUDA）。为了使图表更小，每个扭曲有 4 个元素宽，一个块有 16 个元素。

以下论文也非常有用GPU 的高效并行扫描算法。正如流架构的并行扫描。