CUDA 共享内存的 Bank Conflict 实例分析与优化
【摘要】
引言
CUDA 将 shared memory 按照 4 字节或 8 字节（默认 4 字节，可以设置为 8 字节）被划分到 32 个 bank （楼）中，不同 bank 之间的内存能同时读写，但是同一个 bank 的不同地址（同一栋楼的不同层）的数据则只能串行读写（如果是同一个...
引言
CUDA 将 shared memory 按照 4 字节或 8 字节（默认 4 字节，可以设置为 8 字节）被划分到 32 个 bank （楼）中，不同 bank 之间的内存能同时读写，但是同一个 bank 的不同地址（同一栋楼的不同层）的数据则只能串行读写（如果是同一个 bank 的同一个地址则可以 broadcast，不会出现 bank conflict），因此当同一个 warp 的线程去访问 shared memory 数据时，如果有两个以上线程访问了同一个 bank 的不同地址的数据，就会产生多余的内存事务（transaction）请求（后面有具体实例图示）影响程序的性能.
假设在共享内存上申请了 1024 个 float 数据 —— __shared__ float data[32][32]，由于每个 float 正好是 4 字节，且 data 按行存储，则 data[0][0] 就位于第 0 个bank，data[0][1] 位于第 1 个 bank，以此类推 data[row][col] 就被划分在了第 col 个 bank 中，即 col 相同的数据划分至了同一个 bank 的不同地址上。如果一个 warp 的线程按 col 处理 data 那么就会造成 bank conflict.
文章来源: panda1234lee.blog.csdn.net，作者：panda1234lee，版权归原作者所有，如需转载，请联系作者。
原文链接：panda1234lee.blog.csdn.net/article/details/110185753
【版权声明】本文为华为云社区用户转载文章，如果您发现本社区中有涉嫌抄袭的内容，欢迎发送邮件进行举报，并提供相关证据，一经查实，本社区将立刻删除涉嫌侵权内容，举报邮箱：
cloudbbs@huaweicloud.com
- 点赞
- 收藏
- 关注作者
作者其他文章
评论（0）