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

CUDA bank 及bank conflict相关推荐

难理解的bank conflict
之前看Nvidia-OpenCL-SDK里有一个例子讲到过bank conflict,但没怎么明白,它选择的是用奇数来避免. #define BLOCK_DIM 16// This kernel is ...
内存颗粒位宽和容量_SDRAM的逻辑Bank与芯片容量表示方法
1.逻辑Bank与芯片位宽讲完SDRAM的外在形式,就该深入了解SDRAM的内部结构了.这里主要的概念就是逻辑Bank.简单地说,SDRAM的内部是一个存储阵列.因为如果是管道式存储(就如排队买票) ...
DDR从channel/rank/chip/bank/row/col/cell,DDR/GDDR/HBM
DDR/GDDR/HBM请参考文献 1.ddr的层级结构一个soc或者PC上的ddr都是有很多颗ddr single chip组成的.这么多颗ddr又组成了不同的层级.这些层级从大到小分为: cha ...
图解RAM结构与原理，系统内存的Channel、Chip与Bank
文章目录转载正文标题挥发性内存分2种,SRAM和DRAM 主内存子系统 channel 和 DIMM rank 和 chip bank.row.column 内存的读写方式越多越好,加速读写能力 ...
内存条的逻辑BANK和RANK（物理BANK）概念
在内存上有一个参数:2R X16,1R X16, 1R X8-.. 这里的R就是Rank 既物理BANK,X16,X8,是指芯片位宽何谓内存BANK: 内存的BANK其实分为两部分,逻辑BANK和物 ...
对一些内存名词术语的解释（bank ECC等）——转载
内存 : BANK Bank (内存库) 在内存行业里,Bank至少有三种意思,所以一定要注意. 1.在SDRAM内存模组上,"bank 数"表示该内存的物理存储体的数量. ...
Xilinx FPGA中HR、HD、HP bank说明
HR bank HP bank HD bank 全称 High Range High Performance High Desity 名称高范围bank 高性能bank 高密度bank 电压范围 ...
R12现金管理系统：银行对账单（Bank Statement）的标准处理流程
·银行对账单概念银行对账单,是由银行发出的一种单据,详细记录着某个银行账户的资金进出情况.它是银行和企业核对账务的联系单,证实企业业务往来的记录. R12的现金管理模块(Cash Managemen ...
01_设计一个Bank类，银行某账号的资金往来账管理
题目设计一个Bank类,实现银行某账号的资金往来账管理,包括建账号.存入.取出等.Bank类包括私有数据成员top(当前指针).date(日期).money(金额).rest(余额)和sum(累计余 ...

CUDA bank 及bank conflict

shared memory

Bank conflicts

修改bank中bank size

参考资料

CUDA bank 及bank conflict相关推荐

最新文章

热门文章