kernel内的缩减操作范例

在kernel内完成缩减操作,个人研究的一点小成果,分享…
例如,下列程序是在一个 < < <256,1,1>>>的block里求512个元素的和,存在__shared__ c中(实际上如果一个kernel内只有缩减操作时,更小的block更有效率,此处只是范例)
效率不是特别高,不过写起来容易那么一点…

#define ADD temp1[tid] = temp[(tid>>1) < <1] + temp[((tid>>1) < <1) + 1]; __syncthreads();\
temp[tid] = temp1[(tid>>2) < <2] + temp1[((tid>>2) < <2) + 2]; __syncthreads();\
temp1[tid] = temp[(tid>>3) < <3] + temp[((tid>>3) < <3) + 4]; __syncthreads();\
temp[tid] = temp1[(tid>>4) < <4] + temp1[((tid>>4) < <4) + 8]; __syncthreads();\
temp1[tid] = temp[(tid>>5) < <5] + temp[((tid>>5) < <5) + 16]; __syncthreads();\
temp[tid] = temp1[(tid>>6) < <6] + temp1[((tid>>6) < <6) + 32]; __syncthreads();\
temp1[tid] = temp[(tid>>7) < <7] + temp[((tid>>7) < <7) + 64]; __syncthreads();

global static void bjrot(float *result, float *data)
{
shared float temp[256];
shared float temp1[256];

__syncthreads();
temp[tid] = value[tid];
temp[tid] += value[256 + tid];
__syncthreads();
ADD;
__syncthreads();
result = temp1[0] + temp1[128];
}

以上程序可以修改成求最大值最小值,连乘等的程序
例如,求最大值时可以这样修改:
temp1[tid] = temp[(tid>>1) < <1] 〉 temp[((tid>>1) < <1) + 1]?temp[(tid>>1) < <1] : temp[((tid>>1) < <1) + 1];

目前在cuda上做缩减的效率还不是很高,上述方法中大多数的计算是多余的,不过比for循环有效率一些,比手写方便,在增加block中thread带来的性能增益大于缩减操作占用的时间时,可以使用这种方法…

上述方法是用于在kernel内主要是进行并行数据不相关计算,但是需要一些kernel内缩减操作,同时你又舍不得把数据进行一次昂贵的global - shared交换时的情况…
纯粹的cuda上缩减操作目前看到一种方法,在host上循环,每次都将调用的block数目减半(见 http://forums.nvidia.com/lofiversion/index.php?t63637.html ),但我觉得这样会进行太多的数据交换。

例如
我的一个应用中进行80万点求和的方法是500个block,每block 32个线程。
进行一下操作:

const int offset = bid * 1600;
shared float temp[32];
shared float temp1[32];

temp[tid] += data[offset + tid + 32];
temp[tid] += data[offset + tid + 64];
temp[tid] += data[offset + tid + 96];


temp[tid] += data[offset + tid + 1536];
temp[tid] += data[offset + tid + 1568];

temp1[tid] = temp[(tid>>1) < <1] + temp[((tid>>1) < <1) + 1]; __syncthreads();\
temp[tid] = temp1[(tid>>2) < <2] + temp1[((tid>>2) < <2) + 2]; __syncthreads();\
temp1[tid] = temp[(tid>>3) < <3] + temp[((tid>>3) < <3) + 4]; __syncthreads();\
temp[tid] = temp1[(tid>>4) < <4] + temp1[((tid>>4) < <4) + 8]; __syncthreads();\
result = temp[(tid>>5) < <5] + temp[((tid>>5) < <5) + 16]; __syncthreads();\

当然,真正需要纯粹缩减加,最大,最小操作的时候用cublas里面的命令就行了,多快好省。可惜没有连乘,不过真正要连乘的时候估计单精度也很难装得下就是了...