关于gpu:经验分享GPU-CUDA-使用-memory-padding-避免-bank-conflict

  欢送关注我的公众号 [极智视界]，回复 001 获取 Google 编程标准

  O_o   >_<   o_O   O_o   ~_~   o_O

  本文聊一下如何在 GPU CUDA 编程里应用 memory padding 来防止 bank conflict。

1、Shared memory

  Shared memory 是一块很小、低提早的 on-chip memory，比 global memory 快上百倍，能够把 shared memory 当作可编程的 cache，次要作用有：

• An intra-block thread communication channel：线程间的交换通道；
• A program-managed cache for global memory data：可编程的 cache；
• Scratch pad memory for transforming data to improve global memory access patterns：通过缓存数据缩小 glabal memory 访存次数。

   能够动静或者动态的调配 shared memory，其申明既能够在 kernel 外部也能够作为全局变量，能够通过以下关键字进行申明：

__shared__     /// 标识符

__shared__ float tile[_y][_x];     /// 动态的申明了一个 2D 浮点型数组
    
/// kernel 内申明
extern __shared__ int tile[];

kernel<<<grid, block, isize * sizeof(int)>>>(...);

  为了取得高带宽，shared memory 被分成 32 个等大小的内存卡，对应 warp 中的 thread，他们能够同时被拜访。

2、应用 memory padding 防止 bank conflict

  如果没有 bank 抵触的话，shared memory 跟 registers 一样快。

  疾速的状况：

• warp 内所有线程拜访不同 banks，没有抵触；
• warp 内所有线程读取同一地址，触发播送机制，没有抵触。

  慢速的状况：

• bank conflict：warp 内多个线程拜访同一个 bank；
• 访存必须串行化；
• 多个线程同时拜访同一个 bank 的线程数最大值。

  举个 bank conflict 的例子，如下是一块共享内存：

  没有 bank conflict 的状况：

int x_id = blockDim.x * blockIdx.x + threadIdx.x;     // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y;     // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if(x_id < col && y_id < row)
{sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.y][threadIdx.x];
}

   有 bank conflict 的状况：

int x_id = blockDim.x * blockIdx.x + threadIdx.x;     // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y;     // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if(x_id < col && y_id < row)
{sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}

   以上例子中从没有 bank conflict 到有 bank conflict 只是做了很小的改变，上面看看如何解决上述的 bank conflict。

  以下面的例子为例，能够简略的通过 memory padding 的形式来防止 bank conflict，如下图：

  从代码角度来看一下，是怎么样通过 memory padding 来把下面有 bank conflict 的代码进行性能改善的：

int x_id = blockDim.x * blockIdx.x + threadIdx.x;     // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y;     // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE + 1];     // memory padding

if(x_id < col && y_id < row)
{sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}

  以上分享了在 GPU CUDA 编程中应用 memory padding 来防止 bank conflict 的办法，心愿我的分享对你的学习有一点帮忙。

【公众号传送】
《【教训分享】GPU CUDA 应用 memory padding 防止 bank conflict》