在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带来的性能增益大于缩减操作占用的时间时,可以使用这种方法…