Shared memory bank conflicts

2024-08-31 06:12
文章标签 memory shared bank conflicts

本文主要是介绍Shared memory bank conflicts,希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!

共享内存和bank:

在CUDA架构中,共享内存是一个非常快速的内存类型,它位于每个线程块内部并为该线程块内的所有线程提供服务。为了实现高吞吐量的访问,共享内存被划分为多个独立的存储区域,称为“banks”。每个bank可以在单个时钟周期内独立地服务一个线程。

Shared memory 共享内存

「CUDA ON ARM」如何避免共享内存 Bank conflict - 知乎 (zhihu.com)icon-default.png?t=N7T8https://zhuanlan.zhihu.com/p/538335829在同一个线程块(thread block)中的线程共享一块 Shared memory。Shared memory 被分割为 32 个逻辑块(banks),不同的逻辑块可以被多个线程同时访问。连续的 32-bit 访存被分配到连续的逻辑块(bank)。

例如,声明共享内存 __shared__ float sData[32][32],那么 sData[0][0]sData[1][0]...sData[31][0] 位于 Bank[0]sData[31][0]sData[31][1]...sData[31][31] 位于 Bank[31]

Bank conflict 初探

以下两种情况不会发生 Bank conflict:

  • half-warp/warp 内所有线程访问不同 banks;
  • half-warp/warp 内所有线程读取同一地址(multicast)。

因此,我们的设计原则应当是使得同一个 warp 中的不同线程访问互不相同的 bank 中的数据,使得数据的访问并行执行,而不是串行执行。

如果同一个 warp 中的不同线程将不可避免地访问同一个 bank 中的数据,我们可以使用 Memory Padding 优化 bank 的分割,使得同一个 warp 中的线程访问不同 bank 中的数据。

 

warp:

为提高运行效率,内存块(thread block)中的线程将会按照线程 ID,以 32 个为一组,分割为若干个 warp,每个 warp 将被分配到 32 个 core 上运行。half-warp 用于指代一个 warp 的前半段或者后半段。

共享内存的地址映射方式

GPU shared local memory bank 冲突 - 知乎 (zhihu.com)icon-default.png?t=N7T8https://zhuanlan.zhihu.com/p/668474624在共享内存(SLM)中,连续的 4-bytes 被分配到连续的 32个bank中(每一个 bank 存放一个 32-bits 的数据),这就像电影院的座位一样:一列的座位就相当于一个bank,所以每行有32个座位,在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据,如4个 char 型的数据,2个 short型的数据, 1 个 Uint32 数据);

正常情况下,我们是按照先坐完一行再坐下一行的顺序来坐座位的,在shared memory中地址映射的方式也是这样的。下图中内存地址是按照箭头的方向依次映射的:

上图中蓝色块 0~31 为 bank 编号。如果你申请一个 int类型 共享内存数组 ,你的每个元素所对应的 bank 编号就是地址偏移量 (也就是数组下标) 对32取余所得的结果,比如大小为1024的一维数组myShMem:

  • myShMem[4]: 对应的bank id为#4 (相应的行偏移量为0)
  • myShMem[31]: 对应的bank id为#31 (相应的行偏移量为0)
  • myShMem[50]: 对应的bank id为#18 (相应的行偏移量为1)
  • myShMem[128]: 对应的bank id为#0 (相应的行偏移量为4)
  • myShMem[178]: 对应的bank id为#18 (相应的行偏移量为5)

Bank id = x % 32 行偏移: x / 32

同时产生 Bank conflict 主要有三种情况: 1)线程访问 bank 的方式产生的冲突,这个比较常见,2)数据类型产生的 bank 冲突,3)访问步长与bank冲突

1. 线程访问 bank 的方式产生的冲突

几种典型的 bank 访问的形式。

1)访问步长(stride)为1,线性访问方式,将每个warp中的线程ID与每个bank的ID一一对应,因此不会产生bank冲突。

2) 交叉的访问,每个线程并没有与bank一一对应,但每个线程都会对应一个唯一的bank,所以也不会产生bank冲突。 

3)访问步长(stride)为2,线性访问方式,造成了线程0与线程16都访问到了bank 0,线程1与线程17都访问到了bank 2...,于是就造成了2路的bank冲突。

 4)8路的bank冲突

5) GPU 广播机制

所有的线程都访问了同一个bank,貌似产生了32路的bank冲突,但是由于广播(broadcast)机制, 当一个warp中的所有线程访问一个bank中的同一个字(word)地址时,就会向所有的线程广播这个字(word)),这种情况并不会发生bank冲突。

6) GPU 多播机制

多播机制(multicast)——当一个warp中的几个线程访问同一个bank中的相同字地址时,会将该字广播给这些线程。这个特性得去查询当前的 GPU 是否支持这个特性。

2. 数据类型产生的 bank 冲突

当每个线程访问一个32-bits大小的数据类型的数据(如int,float)时,不会发生bank冲突。

extern __shared__ int shrd[];
foo = shrd[baseIndex + threadIdx.x]

但是如果每个线程访问一个字节(8-bits)的数据时,会不会发生bank冲突呢?

很明显这种情况会发生bank冲突的,因为四个线程访问了同一个bank,造成了四路bank冲突。同理,如果是short类型(16-bits)也会发生bank冲突,会产生两路的bank冲突,下面是这种情况的两个例子:

1)四路bank冲突

2)二路bank冲突

 

3. 访问步长与bank冲突

通常这样来访问数组:每个线程根据线程编号 tid 与 s 的乘积来访问数组的32-bits字(word):

extern __shared__ float shared[];
float data = shared[baseIndex + s * tid];

按照上面的方式, s 是访问的步长(offset),tid 为 wrap 中的线程号。

1) 那么当 s*tid 是bank的数量 (即32) 的整数倍时 ,(baseIndex + s * tid )% 32 = baseIndex 产生 Bank conflict。

2) 仔细思考你会发现,只有warp的大小(即32)小于等于 32/d 时,才不会有bank冲突,而只有当d等于1时才能满足这个条件。要想让32和s的最大公约数d为1,s必须为奇数。于是,这里有一个显而易见的结论:当访问步长s为奇数时,就不会发生bank冲突。

 NOTE: 不同warp中的线程之间不存在什么bank冲突。--> 原因是,不同 wrap 中线程的 shared local memory 不是同一个 。

CUDA:共享内存总结 - 知乎 (zhihu.com)icon-default.png?t=N7T8https://zhuanlan.zhihu.com/p/388823838

这篇关于Shared memory bank conflicts的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



http://www.chinasem.cn/article/1123108

相关文章

error while loading shared libraries: libnuma.so.1: cannot open shared object file:

腾讯云CentOS,安装Mysql时: 1.yum remove libnuma.so.1 2.yum install numactl.x86_64

【0324】Postgres内核 Shared Buffer Access Rules (共享缓冲区访问规则)说明

0. 章节内容 1. 共享磁盘缓冲区访问机制 (shared disk buffers) 共享磁盘缓冲区有两套独立的访问控制机制:引用计数(a/k/a pin 计数)和缓冲区内容锁。(实际上,还有第三级访问控制:在访问任何属于某个关系表的页面之前,必须持有该关系表的适当类型的锁。这里不讨论关系级锁。) Pins 在对缓冲区做任何操作之前,必须“对缓冲区pin”(即增加其引用计数, re

#error: Building MFC application with /MD[d] (CRT dll version) requires MFC shared dll version

昨天编译文件时出现了Building MFC application with /MD[d] (CRT dll version)requires MFC shared dll version~~~~的错误。   在网上很容易找到了解决的方案,公布如下:   对着你的项目点击右键,依次选择:属性、配置属性、常规,然后右边有个“项目默认值”,下面有个MFC的使用,选择“在共享 DLL 中使

Learning Memory-guided Normality for Anomaly Detection——学习记忆引导的常态异常检测

又是一篇在自编码器框架中研究使用记忆模块的论文,可以看做19年的iccv的论文的衍生,在我的博客中对19年iccv这篇论文也做了简单介绍。韩国人写的,应该是吧,这名字听起来就像。 摘要abstract 我们解决异常检测的问题,即检测视频序列中的异常事件。基于卷积神经网络的异常检测方法通常利用代理任务(如重建输入视频帧)来学习描述正常情况的模型,而在训练时看不到异常样本,并在测试时使用重建误

Windows用户取消共享文件夹密码方法(Method for Windows Users to Cancel Shared Folder Password)

Windows用户取消访问共享文件夹密码方法 💝💝💝欢迎来到我的博客,很高兴能够在这里和您见面!希望您在这里可以感受到一份轻松愉快的氛围,不仅可以获得有趣的内容和知识,也可以畅所欲言、分享您的想法和见解。 推荐:Linux运维老纪的首页,持续学习,不断总结,共同进步,活到老学到老 导航剑指大厂系列:全面总结 运维核心技术:系统基础、数据库、网路技术、系统安全、自动化运维、容器技术、监

【论文分享】GPU Memory Exploitation for Fun and Profit 24‘USENIX

目录 AbstractIntroductionResponsible disclosure BackgroundGPU BasicsGPU architectureGPU virtual memory management GPU Programming and ExecutionGPU programming modelGPU kernelDevice function NVIDIA

sqlplus: error while loading shared libraries: libnsl.so.1: cannot open shared object file: No such

在Zabbix Server服务器上安装oracle-instantclient11.2后,结果使用sqlplus命令时遇到“sqlplus: error while loading shared libraries: libnsl.so.1: cannot open shared object file: No such file or directory“错误,下面总结一下解决过程。

DDR的Controller、Channel、Chip、Rank、Bank、Row、Column、Sided、Dimm

目录 概览 1.概览             先从半导体生产开始,生产出来还没切割的叫晶圆(wafer)。切割出来还没封装的叫裸die(bare die)。封装好的叫颗粒(component)。做成内存条后叫模组(module)。下文我们也会按这样的称呼去区分。 2.Controller(内存控制器)         一开始内存控制器在主板上有独立的芯片;在英特尔微处理器

关于std::shared_ptr和enable_share_from_this的一个隐蔽的问题

在使用共享指针时,遇到了一个如下问题: #include <iostream>class traversal;class observer {std::shared_ptr<traversal> m_tra;public:observer(std::shared_ptr<traversal> t):m_tra(t) {};~observer() { std::cout << "releas

FUSEE: A Fully Memory-Disaggregated Key-Value Store——论文阅读

FAST 2023 Paper 论文阅读笔记整理 问题 分布式内存键值(KV)存储正在采用分离式内存(DM)体系结构以提高资源利用率。然而,现有的DM上的KV存储采用半分离式设计,在DM上存储KV对,但在单个元数据服务器上管理元数据,因此仍然在元数据服务器上遭受低资源效率的问题。 如图1a,Clover[60]采用半分离式设计,在计算节点(CN)上部署客户端,在内存节点(MN)上存储KV对,