【发布时间】：2012-09-25 19:58:56
【问题描述】：
我有一个基于 warp 的并行减少的想法，因为根据定义，warp 的所有线程都是同步的。
所以想法是输入数据可以减少 64 倍（每个线程减少两个元素），而无需任何同步。
与 Mark Harris 的原始实现相同，减少应用于块级，数据位于共享内存上。 http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf
我创建了一个内核来测试他的版本和我的基于 warp 的版本。
内核本身完全相同地将 BLOCK_SIZE 元素存储在共享内存中，并在输出数组中的唯一块索引处输出其结果。
算法本身运行良好。用一个完整的数组进行测试以测试“计数”。
实现的函数体：
/**
* 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
}
1.他的版本的实现：
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];
}
资源：
使用了 4 个同步线程
12个if语句使用
11个读+加+写操作
1 次最终写入操作
5 寄存器用法
性能：
5 次测试运行平均：~ 19.54 毫秒
2。基于 Warp 的方法：（与上述相同的函数体）
/*
* 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];
}
资源：
使用了 1 个同步线程
7 if 语句
10个读写操作
1 次最终写入操作
5 寄存器用法
5 位移位
1 个添加
1个子
性能：
5 次测试运行平均：~ 20.82 毫秒
在具有 256 mb 浮点值的 Geforce 8800 GT 512 mb 上多次测试两个内核。 并以 每块 256 个线程（100 % 占用率）运行内核。
基于 warp 的版本慢了 ~ 1.28 毫秒。
如果未来的卡允许更大的块大小，基于 warp 的方法仍然不需要进一步的同步语句，因为最大值是 4096，它会减少到 64，而最终的 warp 会减少到 1
为什么不快？或者说这个想法，内核的缺陷在哪里？
从资源使用情况来看，warp 方法应该领先吗？
Edit1：更正内核只有一半线程处于活动状态，不会导致超出范围的读取，添加了新的性能数据
【问题讨论】：