我有关于基于扭曲的并行减少的想法，因为根据定义，扭曲的所有线程都是同步的。
因此，我们的想法是输入数据可以减少 64 倍（每个线程减少两个元素），而无需任何同步。
与 Mark Harris 的原始实现相同，减少应用于块级，数据位于共享内存上。http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf
我创建了一个内核来测试他的版本和我的基于扭曲的版本。
内核本身完全相同地将 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 寄存器的使用
表现：
五次测试运行平均值：~ 19.54 ms
2. 基于扭曲的方法：（与上面相同的函数体）
/*
* 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 add
1 sub
表现：
五次测试运行平均值：~ 20.82 ms
在 a 上多次测试两个内核Geforce 8800 GT 512 MB with 256mb 浮点值。
并运行内核每块 256 个线程（100% 入住率）。
基于扭曲的版本是〜1.28慢几毫秒。
如果未来的卡允许更大的块大小，基于扭曲的方法仍然不需要进一步的同步语句，因为最大值为 4096，它减少到 64，最终扭曲减少到 1
为什么它不更快？或者这个想法、内核的缺陷在哪里？
从资源使用来看，扭曲方法应该领先吗？
Edit1：更正了内核，仅一半线程处于活动状态，不会导致越界读取，添加了新的性能数据
