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.

映射关系如下所图:

Diagram of NVIDIA Kepler Shared Memory Banks Parallel Accesses

如上图共享内存映射为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.

Diagram of NVIDIA Kepler Shared Memory Banks Serialized Accesses

上述图中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

  • 5
    点赞
  • 37
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

Huo的藏经阁

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值