CUDA 共享内存 bank conflict

最新推荐文章于 2024-09-08 14:19:55 发布
最新推荐文章于 2024-09-08 14:19:55 发布
阅读量1.4k 收藏 5
点赞数 3
分类专栏： 并行计算&GPU&CUDA&OpenMP

1. bank conflict

本文所有的实验针对 GTX980 显卡，Maxwell 架构，计算能力 5.2。

GPU 共享内存是基于存储体切换的架构（bank-switched-architecture）。在 Femi，Kepler，Maxwell 架构的设备上有 32 个存储体（也就是常说的共享内存分成 32 个bank），而在 G200 与 G80 的硬件上只有 16 个存储体。

每个存储体（bank）每个周期只能指向一次操作（一个 32bit 的整数或者一个单精度的浮点型数据），一次读或者一次写，也就是说每个存储体（bank）的带宽为 每周期 32bit。

如下图所示，在一个线程块中申请如下的共享内存：

__shared__ float sData[32][32];
  • 1

这里写图片描述

也就是说在上述的 32 * 32 的二维数组共享内存中，每一列对应同一个 bank。

  1. 同常量内存一样，当一个 warp 中的所有线程访问同一地址的共享内存时，会触发一个广播（broadcast）机制到 warp 中所有线程，这是最高效的。
  2. 如果同一个 warp 中的线程访问同一个 bank 中的不同地址时将发生 bank conflict。
  3. 每个 bank 除了能广播（broadcast）还可以多播（mutilcast）（计算能力 >= 2.0），也就是说，如果一个 warp 中的多个线程访问同一个 bank 的同一个地址时（其他线程也没有访问同一个bank 的不同地址）不会发生 bank conflict。
  4. 即使同一个 warp 中的线程 随机的访问不同的 bank，只要没有访问同一个 bank 的不同地址就不会发生 bank conflict。

这里写图片描述

如上图所示，左侧和右侧的都没有发生 bank conflict。而中间的存在 bank conflcit，由于经过最多两次，该 warp 中的线程就都可以得到所要的数据，所有称为 2-way bank conflict，如果同一个 warp 中的所有线程访问一个 bank 中的 32 个不同地址，则需要分 32 次，称为 32-way bank conflict。

这里写图片描述

如上图所示，左中右均未发生 bank conflict。

依次我们可以总结：只要同一个 warp 的不同线程会访问到同一个 bank 的不同地址就会发生 bank conflict，除此之外的都不会发生 bank conflict。

既然广播是针对同一个 warp 而言的，那么如果不同的 warp 访问同一个 bank 中的同一个地址呢？由于 每个 SM 中有 4 个 warp scheduler （GTX980），可以很好的调度 warp，使其 warp 之间的访问冲突可以充分的隐藏，因此对效率的影响很小，远远小于 warp 内的 bank conflict。至于 warp scheduler 的调度机制，NVIDIA 没有说的特别清楚，可能也是想要开发者不要过于关注于此。

2. 实验 1

实现定义如下图所示的 32 * 32 线程块，共 1024 个线程，32 个 warp。

这里写图片描述

申请如 1 中所示的 32 * 32 的共享内存，共 32 个 bank，每个 bank 对应 32 个元素。

  • 实验 1.1
    该线程块中的每个 warp 读写不同的 bank，不同的 warp 不会访问一个地址，也就是一一对应的关系。图中的数字就表示上图中的线程标号。经分析可知，此时是没有 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];
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 实验 1.2
    该线程块中的每个 warp 读写相同的 bank 的不同地址，不同的 warp 访问不同，也就是一一对应的关系。图中的数字就表示上图中的线程标号。经分析可知，此时是存在很严重的 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];
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 实验 1.3（避免 bank conflict 的技巧）
    针对实验 1.2 中出现的严重的 bank conflict，我们可以通过添加一个附加列来避免 bank conflict，如下图所示，左图为申请的共享内存矩阵形式，右图是表示成 bank 后的形式，通过这种方式，原来在一个 bank 中的同一个 warp 都正好偏移到了不同的 bank 中。

这里写图片描述

代码如下：

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];

if (x_id < col && y_id < row)
{
    sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12

上述三个小实验的运行时间为：

实验 1.1 ：0.052416 ms
实验 1.2 ：0.131072 ms
实验 1.3 ：0.053280 ms
  • 1
  • 2
  • 3

除去公共代码后的时间为：

实验 1.1 ：0.034816 ms
实验 1.2 ：0.113472 ms
实验 1.3 ：0.035680 ms
  • 1
  • 2
  • 3

结论：

  1. 通过额外的一行，可以避免 bank conflict，运行时间与完全没有 bank conflict 的运行时间差距很小。
  2. 存在 bank conflict 的，运行时间几乎是没有 bank conflict 的运行时间的 4 倍。

其实只要添加的是奇数列就可以，只不过 1 列是最节省空间（共享内存太宝贵）的。

3. 实验 2

  • 实验 2.1
    同一个 block 中所有第 i 列的线程都计算第 i 行的元素的和，此时所有同一个warp 会访问同一个 bank 的不同地址。如下图所示，分别表示第 0 列访问 bank 0 中的第一个地址，第 1 列访问 bank 1 中的第 1 个地址，依次类推。

这里写图片描述

代码如下：

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();

    float data = 0.0f;
    for (int j = 0; j < BLOCKSIZE; j++)
    {
        data += sData[threadIdx.x][j];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 实验 2.2
    同实验 1.3 类似，添加额外的一列，如下图所示：

这里写图片描述

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];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;

    for (int j = 0; j < BLOCKSIZE; j++)
    {
        data += sData[threadIdx.x][j];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 19

上述两个实验的运行时间如下所示：

实验 2.1 ：0.458144 ms
实验 2.2 ：0.090848 ms
  • 1
  • 2

从上图也可以看出，修改后的带宽相当于修改前的 32 倍。修改后的运行时间也明显得到改善。

4. 实验 3

  • 实验 3.1
    采用实验 1.1 的方式，同一个 warp 访问不同的 bank，不同的 warp 访问不同的地址。

代码如下：

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();

    float data = 0.0f;
    for (int j = 0; j < 1000; j++)
    {
        data = sData[threadIdx.y][threadIdx.x];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18
  • 实验 3.2
    同一个 warp 访问不同的 bank，所有 warp 访问同一个地址，也就是说所有的行都会访问第 0 行。

代码如下：

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();

    float data = 0.0f;
    for (int j = 0; j < 1000; j++)
    {
        data = sData[0][threadIdx.x];
    }
    matrixTest[index] = data;
}
  • 1
  • 2
  • 3
  • 4
  • 5
  • 6
  • 7
  • 8
  • 9
  • 10
  • 11
  • 12
  • 13
  • 14
  • 15
  • 16
  • 17
  • 18

上述两个实验的运行时间如下所示：

实验 2.1 ：0.053800 ms
实验 2.2 ：0.055328 ms
  • 1
  • 2

在实验 2.2 中存在明显的不同 warp 间的冲突，但是运行时间差距很小，也就是说 warp 间冲突的影响很小。

5. visual profiler

通过 visual profiler 可以判断程序中是否存在 bank conflict，在运行 visual profiler 前需要添加 -lineinfo选项，在 visual studio 中可以设置，如下所示：

这里写图片描述

在 visual profiler 中分析实验 1.2，结果如下所示，可以直接定位到出现 bank conflict 的行。

这里写图片描述

6. 完整代码

我的GitHub

7. 参考

  1. 《CUDA_C_Programming_Guide》7.0 Appendix G. COMPUTE CAPABILITIES / 4.1 / 4.2
  2. 《CUDA_C_Best_Practices_Guide》7.0
  3. 《CUDA 并行程序设计：GPU编程指南》6.4
  4. 《GPU 高性能运算之 CUDA》4.7.1.3/4.4.3
  5. 《Performance modeling of atomic additions on
    GPU scratchpad memory》
  6. 《stackoverflow》

--------------------- 作者：木子超同学 来源：CSDN 原文：https://blog.csdn.net/endlch/article/details/47043069?utm_source=copy 版权声明：本文为博主原创文章，转载请附上博文链接！

Tiger-Li
关注
专栏目录
Shared memory bank conflicts
qq_62704693的博客
08-30 686
这就像电影院的座位一样：一列的座位就相当于一个bank，所以每行有32个座位，在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个 char 型的数据，2个 short型的数据， 1 个 Uint32 数据)；所有的线程都访问了同一个bank，貌似产生了32路的bank冲突，但是由于广播(broadcast)机制， 当一个warp中的所有线程访问一个bank中的。当每个线程访问一个32-bits大小的数据类型的数据(如int，float)时，不会发生bank冲突。
CUDA 共享内存的 Bank Conflict 实例分析与优化
ShaderJoy 的兴趣技术杂货铺
11-26 2008
引言 CUDA将 shared memory 按照 4 字节或 8 字节（默认 4 字节，可以设置为 8 字节）被划分到 32 个 bank （楼）中，不同bank 之间的内存能同时读写，但是同一个 bank 的不同地址（同一栋楼的不同层）的数据则只能串行读写（如果是同一个 bank 的同一个地址则可以broadcast，不会出现 bank conflict），因此当同一个 warp 的线程去访问 shared memory 数据时，如果有两个以上线程访问了同一个 bank 的不同地址...
GPU高性能计算CUDA编程：使用共享内存作为缓冲区
最新发布
m0_51165837的博客
09-08 437
GPU高性能计算CUDA编程：使用共享内存作为缓冲区
[Z]CUDA中Bank conflict冲突
weixin_33975951的博客
02-19 428
    其实这两天一直不知道什么叫bank conflict冲突，这两天因为要看那个矩阵转置优化的问题，里面有讲到这些问题，但是没办法，为了要看懂那个bank conflict冲突，我不得不去找资料，说句实话我现在不是完全弄明白，但是应该说有点眉目了，现在我就把网上找的整理一下，放在这边，等哪天完全弄明白了我就在修改里面的错误。     Tesla 的每个 SM 拥有 16KB 共享...
CUDA ---- Shared Memory
weixin_30822451的博客
06-28 453
CUDA SHARED MEMORY shared memory在之前的博文有些介绍，这部分会专门讲解其内容。在global Memory部分，数据对齐和连续是很重要的话题，当使用L1的时候，对齐问题可以忽略，但是非连续的获取内存依然会降低性能。依赖于算法本质，某些情况下，非连续访问是不可避免的。使用shared memory是另一种提高性能的方式。 GPU上的memory有两种： · On...
CUDA bank 及bank conflict
weixin_42730667的博客
05-17 4804
bank 是CUDA中一个重要概念，是内存的访问时一种划分方式，在CPU中，访问某个地址的内存时，为了减少读写内次次数，访问地址并不是随机的，而是一次性访问bank内的内存地址，类似于内存对齐一样，一次性获取到该bank内的所有地址内存，以提高内存带宽利用率，一般CPU认为如果一个程序要访问某个内存地址时，其附近的数据也有很大概率会在接下来会被访问到。 在CUDA中 在理解bank之前，需要了解共享内存。 shared memory shared memory为CUDA中内存模型中的一中内存模式，为一
Cuda_bank-conflict
weixin_30340617的博客
08-07 253
一直不太理解GPU共享缓存中所谓的bank conflict是什么意思，知道今天对共享缓存操作时，发现简简单单的一句话：temp[tdx] = cach4[7];对程序的运行时间竟然有0.05ms的影响！！ 是时候彻底弄清楚到底何谓bank conflict了。 对于有八个memory-bank 的GPU，共享内存的存储方式如表一中所述。目前计算能力大于2.0的GPU中，大多有16个memor...
cuda 共享内存bank conflict详解
xysjj的博客
01-13 1万+
在cuda并行计算中，共享内存在GPU速度优化上扮演着重要作用，但是如果共享内存使用不当，也会导致速度不快反降或者提速效果不佳，如发生bank conflict; bank的中文翻译为存储体，GPU 共享内存是基于存储体切换的架构（bank-switched-architecture），一般现在的GPU都包含32个存储体,即共享内存被分成了32个bank；根据GPU计算能力的不同（Compute...
CUDA_共享内存
weixin_48151686的博客
11-09 547
typora-copy-images-to: images 共享内存简介 共享内存时受用户控制的一级缓存，共享存储器为片内高速存储器，是一块可以被同一block中的所有线程访问的可读写存储器。访问共享存储器的速度几乎和访问寄存器一样快(相对而言，不是十分严谨的说法，真实情况是，共享内存的延时极低，大约1.5T/s的带宽，远高于全局内存的190G/s,此速度是寄存器的1/10)，是实现线程间通信的延迟最小的方法。共享存储器可以用于实现多种功能，如果用于保存共用的计数器或者block的公用结果。 计算能力1..
CUDA进阶：内存模型
Shilong的博客
09-19 1260
CUDA内存模型
极智开发 | GPU CUDA 使用 memory padding 避免 bank conflict
极智视界
11-22 1001
本文聊一下如何在 GPU CUDA 编程里使用 memory padding 来避免 bank conflict。
DeepRoute Lab | CUDA算子优化：转置篇
DeepRoute_Lab的博客
02-21 1642
转置是深度模型中应用最广泛的算子之一。各种深度学习框架崛起的今天，它被“隔离”到了框架的底层，不再显式的活跃于用户面前。谈及转置，就不得不先聊聊卷积的优化算法以及它们所需要的图像数据格式。
cuda编程：关于共享内存（shared memory）和存储体(bank)的事实和疑惑
linger(心怀梦想，活在当下，知乎ID：linger liu)
06-20 6349
主要是在研究访问共享内存会产生bank conflict时，自己产生的疑惑。对于这
CUDA（15）之CUDA bank conflict in Shared Memory
林微的博客
09-30 997
摘要 本文主要讲述如何避免CUDA中的bank conflict。   1.  为什么会导致bank conflict 总而言之，目的就是避免出现：同一个bank中两个或者两个以上不同的元素被half warp中的并行的threads同时访问的状况，这里同一个bank两个不同的元素被同时访问性能减至1/2，三个不同元素被同时访问，性能减至1/3，以此类推。   2. 怎么解...
CUDA BANK冲突
dataowner的专栏
03-16 2778
为了获得较高的内存带宽，共享存储器被划分为多个大小相等的存储器模块，称为bank，可以被同时访问。因此任何跨越b个不同的内存bank的对n个地址进行读取和写入的操作可以被同时进行，这样就大大提高了整体带宽 ——可达到单独一个bank带宽的b倍。但是很多情况下，我们无法充分发挥bank的功能，以致于shared memory的带宽非常的小，这可能是因为我们遇到了bank冲突。 bank冲突 当一个warp中的不同线程访问一个bank中的不同的字地址时，就会发生bank冲突。 如果没有bank冲突的话，共享
cuda编程之共享内存的bank冲突
benben044的专栏
08-10 1427
bank冲突说明
CUDA中Bank conflict冲突
热门推荐
smsmn的专栏
04-20 1万+
<br />    其实这两天一直不知道什么叫bank conflict冲突，这两天因为要看那个矩阵转置优化的问题，里面有讲到这些问题，但是没办法，为了要看懂那个bank conflict冲突，我不得不去找资料，说句实话我现在不是完全弄明白，但是应该说有点眉目了，现在我就把网上找的整理一下，放在这边，等哪天完全弄明白了我就在修改里面的错误。<br />    Tesla 的每个 SM 拥有 16KB 共享存储器，用于同一个线程块内的线程间通信。为了使一<br />个 half-warp 内的线程能够在一个内
cuda合并访问的要求_CUDA 共享内存的bank conflict  以及合并存储器访问（Coalesced memory accesses）...
weixin_39823200的博客
12-20 698
上面两个概念不要搞混了，两个不同的概念大部分转自http://www.cnblogs.com/waytofall/archive/2013/02/19/2916996.htmlhttp://www.cnblogs.com/dwdxdy/p/3215187.html个人感觉sharedmemory可能是CUDA优化编程中最经常考虑的东西了。在编程过程中，有静态的shared memory 动态的s...
【CUDA学习】共享存储器
weixin_33933118的博客
07-25 154
下面简单介绍一些cuda中的共享存储器和全局存储器　 共享存储器，shared memory，可以被同一块中的所有线程访问的可读写存储器，生存期是块的生命期。 Tesla的每个SM拥有16KB共享存储器。 在编程过程中，有静态的shared memory 动态的shared memory 静态的shared memory 在程序中定义　　　__shared__ type shared[SI...
GPU内存架构优化：合并与bank conflict在AMD OpenCL课程中的应用
bank冲突（bank conflict）指的是当多个线程同时试图访问同一bank中的不同位置时，会阻碍并行执行，降低性能。为了避免bank冲突，程序员需要设计kernel代码，确保线程之间的内存访问尽可能分散在不同的bank上，从而...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值