我有一个关于基于扭曲的并行缩减的想法,因为根据定义,扭曲的所有线程都是同步的。
因此,想法是输入数据可以减少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如果使用语句
11个读+添加+写操作
1次最终写入操作
5寄存器使用
性能:
五次测试平均运行时间:〜19.54毫秒
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添加
1个子
性能:
五次测试平均运行时间:〜20.82毫秒
在 256 浮点值的 Geforce 8800 GT 512 mb 上多次测试两个内核。
并以每块256个线程(100%的占用率)运行内核。
基于warp的版本要慢〜 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;
您可以检查this article了解更多想法。
编辑:
这是基于扭曲的替代块减少方法。
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];
}
}
这应该更快一些,因为它使用了在第一阶段启动的所有warp,并且在最后阶段没有分支,但代价是在新的中间阶段需要额外的分支,共享的加载/存储和
__syncthreads()
。我尚未测试此代码。如果您运行它,请让我知道它的性能。如果您在原始代码中为blockDim使用模板,它可能会再次更快,但是我认为这段代码更加简洁。请注意,由于Fermi和更高版本的体系结构使用纯加载/存储体系结构,因此使用了临时变量
t
,因此从共享内存到共享内存的+=
会导致额外的负载(因为sdata
指针必须是 volatile 的)。显式加载到临时目录中可以避免这种情况。在G80上,它不会影响性能。关于CUDA-为什么基于扭曲的并行缩减速度较慢?,我们在Stack Overflow上找到一个类似的问题:https://stackoverflow.com/questions/12733084/