JeremyJ · 2021年11月22日

【经验分享】GPU CUDA 使用 memory padding 避免 bank conflict

欢迎关注我的公众号 [极智视界],回复001获取Google编程规范

本文聊一下如何在 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

推荐阅读
关注数
5
内容数
25
目录
极术微信服务号
关注极术微信号
实时接收点赞提醒和评论通知
安谋科技学堂公众号
关注安谋科技学堂
实时获取安谋科技及 Arm 教学资源
安谋科技招聘公众号
关注安谋科技招聘
实时获取安谋科技中国职位信息