经验分享|【经验分享】GPU CUDA 使用 memory padding 避免 bank conflict nvidia|cuda|编程语言|c语言|gpu

? 欢迎关注我的公众号 [极智视界]，回复001获取Google编程规范
? O_o ? >_< ? o_O ? O_o ? ~_~ ? o_O
? 本文聊一下如何在 GPU CUDA 编程里使用 memory padding 来避免 bank conflict。

文章目录

- 1、Shared memory
- 2、使用 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<<>>(...);

? 为了获得高带宽，shared memory 被分成 32 个等大小的内存卡，对应 warp 中的 thread，他们可以同时被访问。

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

文章图片

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 paddingif(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 的方法，希望我的分享对你的学习有一点帮助。

?【公众号传送】