CUDA中Bank conflict冲突

本文详细解释了CUDA中Bank Conflict的概念及其对性能的影响,并提供了几种避免Bank Conflict的方法,包括通过调整数据布局和利用广播特性来优化共享内存访问。

    其实这两天一直不知道什么叫bank conflict冲突,这两天因为要看那个矩阵转置优化的问题,里面有讲到这些问题,但是没办法,为了要看懂那个bank conflict冲突,我不得不去找资料,说句实话我现在不是完全弄明白,但是应该说有点眉目了,现在我就把网上找的整理一下,放在这边,等哪天完全弄明白了我就在修改里面的错误。

    Tesla 的每个 SM 拥有 16KB 共享存储器,用于同一个线程块内的线程间通信。为了使一
half-warp 内的线程能够在一个内核周期中并行访问,共享存储器被组织成 16 bank
每个 bank 拥有 32bit 的宽度,故每个 bank 可保存 256 个整形或单精度浮点数,或者说目前 bank 组织成了 256 16 列的矩阵。如果一个 half-warp 中有一部分线程访问属于同一bank 的数据,则会产生 bank conflict,降低访存效率,在冲突最严重的情况下,速度会比全局显存还慢,但是如果 half-warp 的线程访问同一地址的时候,会产生一次广播,其速度反而没有下降。在不发生 bank conflict 时,访问共享存储器的速度与寄存器相同。在不同的块之间,共享存储器是毫不相关的。 ------风辰的 CUDA 入门教程 

   里面说的很清楚就是每个bank有1KB的存储空间。

   Shared memory 是以 4 bytes 为单位分成 banks。因此,假设以下的数据:
     __shared__ int data[128];
    那么,data[0] bank 0data[1] bank 1data[2] bank 2data[15] bank 15,而 data[16] 又回到 bank 0。由于 warp 在执行时是以 half-warp 的方式执行,因此分属于不同的 half warp threads,不会造成 bank conflict
    因此,如果程序在存取 shared memory 的时候,使用以下的方式:
      int number = data[base + tid];
    那就不会有任何 bank conflict,可以达到最高的效率。但是,如果是以下的方式:
      int number = data[base + 4 * tid];
    那么,thread 0 thread 4 就会存取到同一个 bankthread 1 thread 5 也是同 样,这样就会造成 bank conflict。在这个例子中,一个 half warp 16 threads 会有四个 threads 存取同一个 bank,因此存取 share memory 的速度会变成原来的 1/4
    一个重要的例外是,当多个 thread 存取到同一个 shared memory 的地址时,shared memory 可以将这个地址的 32 bits 数据「广播」到所有读取的 threads,因此不会造成 bank conflict。例如:
      int number = data[3];
    这样不会造成 bank conflict,因为所有的 thread 都读取同一个地址的数据。
很多时候 shared memory bank conflict 可以透过修改数据存放的方式来解决。例如,以下的程序:
      data[tid] = global_data[tid];
     ...
      int number = data[16 * tid];

    会造成严重的 bank conflict,为了避免这个问题,可以把数据的排列方式稍加修改,把存取方式改成:
      int row = tid / 16;
      int column = tid % 16;
      data[row * 17 + column] = global_data[tid];
      ...
     int number = data[17 * tid];
   这样就不会造成 bank conflict 了。

    

    简单的说,矩阵中的数据是按照bank存储的,第i个数据存储在第i%16bank中。一个block要访问shared memory,只要能够保证以其中相邻的16个线程一组访问thread,每个线程与bank是一一对应就不会产生bank conflict。否则会产生bankconflict,访存时间成倍增加,增加的倍数由一个bank最多被多少个thread同时访问决定。有一种极端情况,就是所有的16thread同时访问同一bank时反而只需要一个访问周期,此时产生了一次广播。

    下面有一些小技巧可以避免bank conflict 或者提高global存储器的访问速度

       1. 尽量按行操作,需要按列操作时可以先对矩阵进行转置

       2. 划分子问题时,使每个block处理的问题宽度恰好为16的整数倍,使得访存可以按照 s_data[tid]=i_data[tid]的形式进行

       3. 使用对齐的数据格式,尽量使用nvidia定义的格式如float3,int2等,这些格式本身已经对齐。

       4. 当要处理的矩阵宽度不是16的整数倍时,将其补为16的整数倍,或者用malloctopitch而不是malloc

        5. 利用广播,例如s_odata[tid] = tid%16 < 8 ? s_idata[tid] : s_idata[15];会产生8路的块访问冲突而用:s_odata[tid]=s_idata[15];s_odata[tid]= tid%16 < 8 ? s_idata[tid] : s_data[tid]; 则不会产生块访问冲突

 

 

### CUDA Shared Memory Bank Conflict 和 Architecture #### 什么是CUDA共享内存? CUDA中的共享内存是一种高速片上存储器,位于GPU芯片内部。它由同一个线程块(block)内的所有线程共享,用于减少全局内存访问延迟并提高性能[^1]。 #### 共享内存的Bank结构 为了实现高效的并发访问,共享内存被划分为多个独立的存储单元,称为 **banks** 或者存储体。通常情况下,现代GPU拥有32个银行(banks)。这些银行允许同时进行无冲突的数据读写操作,从而最大化吞吐量[^1]。 #### Bank宽度与配置 根据GPU计算能力的不同,每个bank的宽度可能为32位(对于计算能力2.x及更早版本)或64位(计算能力3.x及以上版本)。这意味着连续的32-bit或64-bit数据会被分布到不同的banks中。开发者可以通过`cudaDeviceSetSharedMemConfig()`函数来调整bank大小,默认设置通常是每bank 4字节(32位),但在某些架构下也可以将其配置为8字节(64位),这对于处理双精度浮点数特别有用,能够显著降低因数据对齐不当而导致的bank冲突概率。 #### Tesla P100实例分析 以Tesla P100为例,在其默认配置下,如果我们将bank宽度设定为32比特,则理论上全部32 banks仅能覆盖128字节的空间。然而实际上P100提供了高达48 KB的共享内存容量。这是因为除了初始划分外,剩余部分继续按照相同的模式重复利用这32个banks来进行扩展管理。 #### 发生Bank Conflicts的情况 当两个或更多来自相同warp的线程试图在同一周期内访问属于同一bank的位置时就会引发bank冲突。这种现象降低了有效带宽利用率,并可能导致程序运行效率下降。具体来说: - 如果一个warp中有超过一条路径尝试存取同样的bank位置,则会产生串行化问题; - 数据未正确排列使得相邻元素映射到了相同的bank也会造成此类情况的发生[^2]。 #### 减少Bank Conflicts的方法 为了避免上述提到的各种类型的bank冲突,可以采取如下策略之一或多组合使用: - 调整数组索引方式使各线程间请求尽量均匀分布在各个banks之间; - 使用编译指令提示优化器重新安排加载顺序; - 增加padding填充空白区域改变原始布局以便更好地匹配目标硬件特性等等。 ```cpp // 示例代码展示如何通过增加pad避免简单的bank conflict __global__ void kernel(float* d_out, const float* __restrict__ d_in){ extern __shared__ float s_data[]; int idx = threadIdx.x + blockIdx.x * blockDim.x; // 添加额外空间作为垫补(padding),防止潜在conflict unsigned int paddedIdx = (idx / NUM_BANKS) * PADDED_SIZE + (idx % NUM_BANKS); s_data[paddedIdx] = d_in[idx]; __syncthreads(); } ```
评论 3
成就一亿技术人!
拼手气红包6.0元
还能输入1000个字符
 
红包 添加红包
表情包 插入表情
 条评论被折叠 查看
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

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

余额充值