CUDA并行计算的高效策略

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;
    }
  }
}