CUDA并行计算的高效策略
1.最大化计算强度(高效公式)
包括最大化计算量和最小化每个线程的内存读取速度。
2.合并全局内存
其中连续合并>间隔合并>随机合并
3.应该避免线程发散(同一线程块中的线程执行不同内容的代码)
其中条件语句和循环中,因为线程操作长短不一导致线程发散
并行化高效策略实例
1.reduce(归约)
算法思路
使用全局内存
__global__ void global_reduce_kernel(float * d_out, float * d_in) //全局内存
{
int myId = threadIdx.x + blockDim.x * blockIdx.x; //总线程里对应的线程号(全局变量的结果)
int tid = threadIdx.x; //当前线程的线程号(判断当前是否满足在前一半的线程号)
// do reduction in global mem
for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) //循环满足小于一半,做相加
{
if (tid < s)
{
d_in[myId] += d_in[myId + s];
}
__syncthreads(); // make sure all adds at one stage are done!(线程同步)
}
// only thread 0 writes result for this block back to global mem
if (tid == 0)
{
d_out[blockIdx.x] = d_in[myId]; //第0个线程的求和结果
}
}
把全局内存分配给共享内存
__global__ void shmem_reduce_kernel(float * d_out, const float * d_in)
{
// sdata is allocated in the kernel call: 3rd arg to <<<b, t, shmem>>>
extern __shared__ float sdata[]; //定义共享内存数组
int myId = threadIdx.x + blockDim.x * blockIdx.x;
int tid = threadIdx.x;
// load shared mem from global mem
sdata[tid] = d_in[myId]; //全局内存d_in对应线程块的值赋值给共享内存
__syncthreads(); // 线程同步
// do reduction in shared mem
for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1)
{
if (tid < s)
{
sdata[tid] += sdata[tid + s];
}
__syncthreads(); // 线程同步
}
// only thread 0 writes result for this block back to global mem
if (tid == 0)
{
d_out[blockIdx.x] = sdata[0];
}
}
2.扫描(scan)
算法思路
使用全局内存
__global__ void global_scan(float* d_out,float* d_in){
int idx = threadIdx.x; //取当前的线程号
float out = 0.00f;
d_out[idx] = d_in[idx];
__syncthreads(); //先做一次同步复制已经到位
for(int interpre=1;interpre<sizeof(d_in);interpre*=2){ //循环以间隔为标准0 1 2 4 8
if(idx-interpre>=0){ //当前线程号位置减去间隔
out = d_out[idx]+d_out[idx-interpre]; //例如6=5+1(4-2>=0)
}
__syncthreads();
if(idx-interpre>=0){
d_out[idx] = out;
out = 0.00f;
}
}
}