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