用户
 找回密码
 立即注册
洛嘟嘟 该用户已被删除
发表于 2009-11-18 20:16:36
58832
在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带来的性能增益大于缩减操作占用的时间时,可以使用这种方法...
使用道具 举报 回复
发表于 2009-11-18 20:24:27
上述方法是用于在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();\
使用道具 举报 回复 支持 反对
发表于 2009-11-18 20:25:27
当然,真正需要纯粹缩减加,最大,最小操作的时候用cublas里面的命令就行了,多快好省。可惜没有连乘,不过真正要连乘的时候估计单精度也很难装得下就是了...
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册