【并行计算-CUDA开发】关于共享内存(shared memory)和存储体(bank)的事实和疑惑...
关于共享内存(shared memory)和存储体(bank)的事实和疑惑
主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑。对于这点疑惑,网上都没有相关描述,
不管是国内还是国外的网上资料。貌似大家都是当作一个事实,一个公理,而没有对其仔细研究。还是我自己才学疏浅,不知道某些知识。
比如下面这篇讲解bank conflict的文章。
http://cuda-programming.blogspot.com/2013/02/bank-conflicts-in-shared-memory-in-cuda.html
我这里重点不在bank conflict,而是主要讨论shared memory和 memory bank的对应关系。
文中有这么一段描述:
shared_a[threadIdx.x] = global_a[threadIdx.x];
|
First let’s assume your arrays are say for example of the type int (a 32-bit word). Your codesaves these ints into shared memory, acrossany half warp the Kth thread is saving to the Kth memory bank. Sofor example thread 0 of the first half warp will save to shared_a[0] which isin the first memory bank, thread 1 will save to shared_a[1], each half warp has16 threads these map to the 16 4byte banks. In the next half warp, the firstthread will now save its value into shared_a[16] which is in the first memory bankagain. So if you use a 4byte word such int, float etc, then this example willnot result in a bank conflict.
翻译过来的意思大概是这样子。
有一个数组,元素类型为整型,个数为256,开始这个数组存储在全局内存里面。现在我们一个线程块里有256个线程,我们想把这个数组拷贝到共享内存。因此每个线程负责拷贝一个元素。
- shared_a[threadIdx.x] = global_a[threadIdx.x];
想一下,这种访问是否会导致bank conflict呢?(看答案之前,先想想)
好的!
首先,我们假设你的数组元素是int类型的,占32位。你的代码把这些元素放进共享内存中,在任意一个half-warp,第k个线程刚好把元素放进第k个memory bank。
比如,第一个half warp中的线程0会放进shared_a[0],她刚好在第一个memory bank中,线程1把放进shared_a[1],每一个half warp有16个线程,刚好跟16个大小为4byte的bank对应。在下一个half warp中,第一个线程(线程0)会把值放进shared_a[16],她刚好也是在第一个memory bank中。所以在这个例子中,如果你使用4byte的字,比如int,float等,最后是不会产生bank confict的。
好了,回到我的讨论。
从上面描述,我们知道一些事实。
假如一个线程块有一块共享内存 int shared_a[256],该显卡设备的memory bank有16个。那么这块共享内存跟memory bank的对应关系是怎么样的?
例子说明一切,显然shared_a[0]在第1个bank中,shared_a[1]在第2个bank中,shared_a[15]在第16个bank中。
那么shared_a[16]呢?shared_a[17]呢?
根据文中的介绍,shared_a[16]在第1个bank中,shared_a[17]在第2个bank中。
规律是shared_a[index]在第(index%16+1)个bank中。
现在疑问来了,每一个bank的大小不是刚好为32位吗?(开普勒是64位)。
既然,shared_a[0]在第1个bank中,shared_a[0]已经是32位的了,那么shared_a[16]又是32位,放哪里?
shared_a[32]也是在第1个bank中,又放哪里?
一个bank怎么可以对应几个元素呢?
还是说bank只是缓存的地方,有其她地方存储,会自动切换的,类似缓存那样。
但是,貌似我没有找到任何资料有关这方面的解释。找了书,找了国内外的网上资料,都没有。
现在只好先记住这么一个事实了:shared_a[index]在第(index%16+1)个bank中。
本文作者:linger
本文链接:http://blog.csdn.NET/lingerlanlan/article/details/32712749
转载于:https://www.cnblogs.com/huty/p/8517830.html
【并行计算-CUDA开发】关于共享内存(shared memory)和存储体(bank)的事实和疑惑...相关推荐
- linux shared,从 0 开始学习 Linux 系列之「22.共享内存 Shared Memory」
共享内存 版权声明:本文为 cdeveloper 原创文章,可以随意转载,但必须在明确位置注明出处! 共享内存 Shared Memory 这次我们来学习在 Linux 中最快的一种 IPC 方式:共 ...
- Linux——详解共享内存shared memory
目录 一.共享内存介绍 (一).什么是共享内存 (二).共享内存优点 (三).共享内存缺点 二.共享内存使用 (一).创建-shmget ①key ②size ③shmflg ④返回值 (二).连接- ...
- 进程间通信之-共享内存Shared Memory--linux内核剖析(十一)
共享内存 共享内存是进程间通信中最简单的方式之中的一个. 共享内存是系统出于多个进程之间通讯的考虑,而预留的的一块内存区. 共享内存同意两个或很多其他进程訪问同一块内存,就如同 malloc() 函数 ...
- 【并行计算-CUDA开发】CUDA线程、线程块、线程束、流多处理器、流处理器、网格概念的深入理解
GPU的硬件结构,也不是具体的硬件结构,就是与CUDA相关的几个概念:thread,block,grid,warp,sp,sm. sp: 最基本的处理单元,streaming processor 最 ...
- Linux应用开发之共享内存
一.概述 共享内存是一种最为高效的进程间通信方式.因为进程可以直接读写内存, 不需要任何数据的拷贝.为了在多个进程间交换信息,内核专门留出了一块内存区.这段 内存区可以由需要访问的进程将其映射到自 ...
- 【并行计算-CUDA开发】GPU 的硬体架构
GPU 的硬体架构 这里我们会简单介绍,NVIDIA 目前支援CUDA 的GPU,其在执行CUDA 程式的部份(基本上就是其shader 单元)的架构.这里的资料是综合NVIDIA 所公布的资讯, ...
- 【并行计算-CUDA开发】GPU 的硬体架构 关于存储体的介绍比较好bank conflicts
GPU 的硬体架构 这里我们会简单介绍,NVIDIA 目前支援CUDA 的GPU,其在执行CUDA 程式的部份(基本上就是其shader 单元)的架构.这里的资料是综合NVIDIA 所公布的资讯, ...
- CUDA总结:共享内存
共享内存是片上内存(on-chip),所以速度比一般的显存快很多,如(全局内存.常亮内存.纹理内存).共享内存是gpu中,带宽仅次于寄存器的存储器. 共享内存是有限的,与L1 Cache公用一块on- ...
- 【并行计算-CUDA开发】从零开始学习OpenCL开发(一)架构
多谢大家关注 转载本文请注明:http://blog.csdn.net/leonwei/article/details/8880012 本文将作为我<从零开始做OpenCL开发>系列文章的 ...
最新文章
- 欧盟中止对中国数据卡产品“两反一保”调查
- 创建hadoop用户并配置ssh免密码登录
- 前端实现连连看小游戏(1)
- 数组中的forEach和map的区别
- python字符串三,删除空白符,对齐函数及判断函数
- java 反射 判断是否存在_如何判断Javascript对象是否存在
- 将超星PDG文件转换成PDF文件的方法
- Oracle之rollUp函数
- 4.2 制定项目章程
- 淘宝/京东/苏宁/拼多多/唯品会 返利消息批量转链思路
- WIFI手机使用正常电脑使用卡顿解决方案
- linux定制欢迎界面motd,linux 界面 /etc/motd
- [数学]二维对数正态分布的概率分布,期望,方差和相关系数
- jsp中给div加背景_html中给元素添加背景图片或者gif动图
- 教你如何使用WinCE CAB Manager制作PPC绿色软件
- NetSniper网络尖兵:宽带网络运营维护管理器
- .NET程序员的技能分析参考V1.0(有附件)
- 利用AppInventor实现登录功能(完整版)
- CSCAD AutoCAD 修改光标颜色
- 中文字符和英文字符判断
热门文章
- emoji隐藏表情_MacBook 7大隐藏功能曝光!强大到逆天!
- rabbitmq消息的序列化与反序列化
- MYSQL重置密码遇到ERROR 1045 (28000): Access denied for user ‘root‘@‘localhost‘ (using passwor:yes)问题
- 查看Linux磁盘文件占用大小
- js- 引用和复制(传值和传址)
- Lucene就是这么简单
- scrapy 伪装代理和fake_userAgent的使用
- CSU 1329: 一行盒子
- 关于ECLIPSE中JSP代码无提示
- 文件错误关于hibernate中报Duplicate class/entity mapping org.model.User错的问题