PROSAGA码农传奇-CF计算编织-CUDA-为什么基于扭曲的并行缩减速度较慢？

我有一个关于基于扭曲的并行缩减的想法，因为根据定义，扭曲的所有线程都是同步的。因此，我们的想法是输入数据可以减少64倍（每个线程减少两个元素），而无需任何同步。与马克·哈里斯（Mark Harris）的原始实现相同，减少量应用于块级，数据位于共享内存中。 <a href="http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf">http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf</a>我创建了一个内核来测试他的版本和基于warp的版本。 内核本身完全相同地将BLOCK_SIZE元素存储在共享内存中，并在输出数组的唯一块索引处输出其结果。该算法本身可以正常工作。经过全面测试，以测试“计数”。实现的功能主体：
<pre><code>/**
 * Performs a parallel reduction with operator add 
 * on the given array and writes the result with the thread 0
 * to the given target value
 *
 * @param inValues T* Input float array, length must be a multiple of 2 and equal to blockDim.x
 * @param targetValue float 
 */
__device__ void reductionAddBlockThread_f(float* inValues,
 float &outTargetVar)
{
 // code of the below functions
}
</code></pre>1.执行他的版本：
<pre><code> if (blockDim.x >= 1024 && threadIdx.x < 512)
 inValues[threadIdx.x] += inValues[threadIdx.x + 512];
__syncthreads();
if (blockDim.x >= 512 && threadIdx.x < 256)
 inValues[threadIdx.x] += inValues[threadIdx.x + 256];
__syncthreads();
if (blockDim.x >= 256 && threadIdx.x < 128)
 inValues[threadIdx.x] += inValues[threadIdx.x + 128];
__syncthreads();
if (blockDim.x >= 128 && threadIdx.x < 64)
 inValues[threadIdx.x] += inValues[threadIdx.x + 64];
__syncthreads();

//unroll last warp no sync needed
if (threadIdx.x < 32)
{
 if (blockDim.x >= 64) inValues[threadIdx.x] += inValues[threadIdx.x + 32];
 if (blockDim.x >= 32) inValues[threadIdx.x] += inValues[threadIdx.x + 16];
 if (blockDim.x >= 16) inValues[threadIdx.x] += inValues[threadIdx.x + 8];
 if (blockDim.x >= 8) inValues[threadIdx.x] += inValues[threadIdx.x + 4];
 if (blockDim.x >= 4) inValues[threadIdx.x] += inValues[threadIdx.x + 2];
 if (blockDim.x >= 2) inValues[threadIdx.x] += inValues[threadIdx.x + 1];

//set final value
 if (threadIdx.x == 0)
 outTargetVar = inValues[0];
}
</code></pre>资源：使用4个同步线程 12个if语句使用 11个读取+添加+写操作 1个最终写操作 5个寄存器使用性能：五次测试平均运行时间：〜19.54毫秒2.基于扭曲的方法：（与上面的功能相同）
<pre><code>/*
 * Perform first warp based reduction by factor of 64
 *
 * 32 Threads per Warp -> LOG2(32) = 5
 *
 * 1024 Threads / 32 Threads per Warp = 32 warps
 * 2 elements compared per thread -> 32 * 2 = 64 elements per warp
 *
 * 1024 Threads/elements divided by 64 = 16
 * 
 * Only half the warps/threads are active
 */
if (threadIdx.x < blockDim.x >> 1)
{
 const unsigned int warpId = threadIdx.x >> 5;
 // alternative threadIdx.x & 31
 const unsigned int threadWarpId = threadIdx.x - (warpId << 5);
 const unsigned int threadWarpOffset = (warpId << 6) + threadWarpId;

inValues[threadWarpOffset] += inValues[threadWarpOffset + 32];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 16];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 8];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 4];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 2];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 1];
}

// synchronize all warps - the local warp result is stored
// at the index of the warp equals the first thread of the warp
__syncthreads();

// use first warp to reduce the 16 warp results to the final one
if (threadIdx.x < 8)
{
 // get first element of a warp
 const unsigned int warpIdx = threadIdx.x << 6;

if (blockDim.x >= 1024) inValues[warpIdx] += inValues[warpIdx + 512];
    if (blockDim.x >= 512) inValues[warpIdx] += inValues[warpIdx + 256];
    if (blockDim.x >= 256) inValues[warpIdx] += inValues[warpIdx + 128];
    if (blockDim.x >= 128) inValues[warpIdx] += inValues[warpIdx + 64];

//set final value
 if (threadIdx.x == 0)
 outTargetVar = inValues[0];
}
</code></pre>资源：使用1个同步线程 7个if语句 10个读添加写操作 1个最终写操作 5个寄存器使用5位移位 1加 1子性能：五次测试运行平均值：〜20.82毫秒在带有256 mb浮点值的Geforce 8800 GT 512 mb上多次测试两个内核。并以每块256个线程（100％的占用率）运行内核。基于经线的版本要慢〜1.28毫秒。如果将来的卡允许更大的块大小，则基于扭曲的方法仍然不需要进一步的同步语句，因为最大值为4096，该数量减少为64，最终扭曲为1为什么它不快？或者内核的缺点在哪里？从资源使用情况来看，翘曲方法应该领先吗？Edit1：更正了只有一半线程处于活动状态而不导致读取超出范围的内核，添加了新的性能数据

我认为您的代码比我的代码慢的原因是，在我的代码中，在第一阶段中，对于每个ADD来说，活动的翘曲只有一半。在您的代码中，所有第一阶段的所有扭曲都处于活动状态。因此，总体而言，您的代码将执行更多的扭曲指令。在CUDA中，重要的是要考虑执行的“ warp指令”总数，而不仅仅是一个warp所执行的指令数。

另外，仅使用一半的扭曲是没有意义的。发起经线的开销只是让它们评估两个分支并退出。

另一个想法是，使用unsigned char和short实际上可能会降低您的性能。我不确定，但是肯定不会保存您的寄存器，因为它们没有打包到单个32位变量中。

同样，在我的原始代码中，我用模板参数BLOCKDIM替换了blockDim.x，这意味着它仅使用了5个运行时if语句（第二阶段的ifs被编译器消除了）。

顺便说一句，一种更便宜的计算您的方法threadWarpId是

```
const int threadWarpId = threadIdx.x & 31;
```

您可以查看本文以获取更多想法。

编辑： 这是一种替代的基于扭曲的块减少。

```
template <typename T, int level>
__device__
void sumReduceWarp(volatile T *sdata, const unsigned int tid)
{
 T t = sdata[tid];
 if (level > 5) sdata[tid] = t = t + sdata[tid + 32];
 if (level > 4) sdata[tid] = t = t + sdata[tid + 16];
 if (level > 3) sdata[tid] = t = t + sdata[tid + 8];
 if (level > 2) sdata[tid] = t = t + sdata[tid + 4];
 if (level > 1) sdata[tid] = t = t + sdata[tid + 2];
 if (level > 0) sdata[tid] = t = t + sdata[tid + 1];
}

template <typename T>
__device__
void sumReduceBlock(T *output, volatile T *sdata)
{
 // sdata is a shared array of length 2 * blockDim.x

const unsigned int warp = threadIdx.x >> 5;
 const unsigned int lane = threadIdx.x & 31;
 const unsigned int tid = (warp << 6) + lane;

sumReduceWarp<T, 5>(sdata, tid);
 __syncthreads();

// lane 0 of each warp now contains the sum of two warp's values
  if (lane == 0) sdata[warp] = sdata[tid];

__syncthreads();

if (warp == 0) {
 sumReduceWarp<T, 4>(sdata, threadIdx.x);
 if (lane == 0) *output = sdata[0];
 }
}
```

这应该更快一些，因为它使用了第一阶段中启动的所有扭曲，并且在最后阶段中没有分支，这是以额外分支，共享负载/存储以及__syncthreads()在新的中间阶段为代价的。我尚未测试此代码。如果您运行它，请让我知道它的性能。如果您在原始代码中为blockDim使用模板，则它可能会再次更快，但是我认为这段代码更为简洁。

请注意，使用临时变量t是因为Fermi和更高版本的体系结构使用纯加载/存储体系结构，因此+=从共享内存到共享内存会导致额外的负载（因为sdata指针必须是易失的）。显式加载到临时目录中可以避免这种情况。在G80上，它不会影响性能。