CUDA bank 及bank conflict

bank 是CUDA中一个重要概念，是内存的访问时一种划分方式，在CPU中，访问某个地址的内存时，为了减少读写内次次数，访问地址并不是随机的，而是一次性访问bank内的内存地址，类似于内存对齐一样，一次性获取到该bank内的所有地址内存，以提高内存带宽利用率，一般CPU认为如果一个程序要访问某个内存地址时，其附近的数据也有很大概率会在接下来会被访问到。

在CUDA中 在理解bank之前，需要了解共享内存。

shared memory

shared memory为CUDA中内存模型中的一中内存模式，为一个片上内存，比全局内存（global memory)要快很多，在同一个block内的所有线程都可以访问到该内存中的数据，与local 或者global内存相比具有高带宽、低延迟的作用。

Because it is on-chip, shared memory has much higher bandwidth and much lower latency than local or global memory.

为了提高share memory的访问速度 除了在硬件上采用片上内存的方式之外，还采用了很多其他技术。其中为了提高内存带宽，共享内存被划分为相同大小的内存模型，称之为bank,，这样就可以将n个地址读写合并成n个独立的bank，这样就有效提高了带宽。

To achieve high bandwidth, shared memory is divided into equally-sized memory modules, called banks, which can be accessed simultaneously. Any memory read or write request made of n addresses that fall in n distinct memory banks can therefore be serviced simultaneously, yielding an overall bandwidth that is n times as high as the bandwidth of a single module.

映射关系如下所图：

如上图共享内存映射为bank采用列映射方式，例如warp size = 32, banks = 16,（计算能力1.x的设备）数据映射关系如下

例如对于一个 32*32大小的float数组，

__shared__ float sData[32][32];

在一个warp size = 32,bank=32的GPU中 中bank的映射关系为：

上述例子中每一列为一个bank分布，同一个bank一次只能访问一次，不同bank可以同时访问。

Bank conflicts

如果在block内多个线程访问的地址落入到同一个bank内，那么就会访问同一个bank就会产生bank conflict，这些访问将是变成串行，在实际开发调式中非常主要bank conflict.

However, if two addresses of a memory request fall in the same memory bank, there is a bank conflict and the access has to be serialized. The hardware splits a memory request with bank conflicts into as many separate conflict-free requests as necessary, decreasing throughput by a factor equal to the number of separate memory requests. If the number of separate memory requests is n, the initial memory request is said to cause n-way bank conflicts.

时

上述图中part1、part2、part3都会访问到同一个bank，将会产生bank conflict ,造成程序串行化，会发现此时性能会产生严重下降。

在上述数组例子中，如果有多个线程同时访问一个列中的不同数组将会产生bank conflict

如果多个线程同时访问同一列中相同的数组元素 不会产生bank conflict，将会出发广播，这是CUDA中唯一的解决方案，在一个warp内访问到相同内存地址，将会将内存广播到其他线程中，同一个warp内访问同一个bank内的不同地址貌似还没看到解决方案。

不同的线程访问不同的bank，不会产生bank conflict

如下图所示：

图中左侧和右侧都没有发生bank conflict。而在中间村子bank conflict.

 如果warp中的线程经过最多两次冲突就能得到所要的数据则成为2-way bank conflict，如果同一个warp中的所有线程访问一个bank中的32个不同的地址，则需要分32此，则称为32-way bank conflict,

注意:只要同一个 warp 的不同线程会访问到同一个 bank 的不同地址就会发生 bank conflict，除此之外的都不会发生 bank conflict。

上述图中均未产生bank conflict。 

修改bank中bank size

注意共享内存中 每个SM的bank数量目前是无法修改的，但是可以修改bank中单个数组元素容纳的字节数，例如上个例子中

__shared__ float sData[32][32];，每个数组元素大小为4个字节，一般cuda中默认是按照4个字节进行组织被划分到bank中，CUDA提供可修改按照8个字节进行组织API：

__host__ ​cudaError_t cudaDeviceSetSharedMemConfig ( cudaSharedMemConfig config )

其中 cudaSharedMemConfi为一个枚举型：

cudaSharedMemBankSizeDefault = 0

cudaSharedMemBankSizeFourByte = 1

cudaSharedMemBankSizeEightByte = 2

 只支持在host端进行调用，不支持在device端调用。

CUDA API中还支持获取bank size大小：

__host__ ​ __device__ ​cudaError_t cudaDeviceGetSharedMemConfig ( cudaSharedMemConfig ** pConfig )

该接口在host和device端都可调用。 

参考资料

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

https://www.microway.com/hpc-tech-tips/gpu-shared-memory-performance-optimization/

https://blog.csdn.net/xysjj/article/details/103885803

https://blog.csdn.net/kebu12345678/article/details/82979934?utm_medium=distribute.pc_relevant_right.none-task-blog-BlogCommendFromMachineLearnPai2-4.nonecase&depth_1-utm_source=distribute.pc_relevant_right.none-task-blog-BlogCommendFromMachineLearnPai2-4.nonecase