详细源码见本人github
点这里
比较经典的树状加法示意图
具体的计算流程如下图所示:
继续进一步的kernel函数解析为:
其中引入两个变量offset与mask
offset用来决定偏移量的大小,mask用来判断当前位置的元素是否参与运算。
#include <iostream>
#include <device_launch_parameters.h>
#include <cuda_runtime.h>
#define N_size 256
using namespace std;
#define THREAD_NUM 16
#define BLOCK_NUM 1
int offset=1,mask=1;
while(tid<THREAD_NUM){
if((tid&MASK)==0)
{
shared[tid]+=sharerd[tid+offset];
}
offset+=offset;
mask+=offset;
__syncthreads();
}
另外一种归约加法
__shared__ sdata[128];
sdata[tid]=A[bid * 128 + tìd];
__syncthreads();
//N归约求和
for(int i = 64; i > 0 ; i/= 2) {
if( tid <i)
sdata[tid] =s_data[tid]+ s_data[tid + i];
__syncthreads();
}
if(tid =0)
B[bid]=sdata[0];
参考:
https://blog.csdn.net/orchestra56/article/details/80981767