码迷,mamicode.com
首页 > 其他好文 > 详细

cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑

时间:2014-06-22 00:39:19      阅读:477      评论:0      收藏:0      [点我收藏+]

标签:cuda   共享内存   memory bank   bank conflict   deep learning   


关于共享内存(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的对应关系。


文中有这么一段描述:


Example
Scenario
Let’ssay we have an array of size 256 of integer type in global memory and we have256 threads in a single Block, and we want to copy the array to shared memory.Therefore every thread copies one element.

shared_a[threadIdx.x] = global_a[threadIdx.x];

So, what u think, does it trap into bank conflict? (Before readinganswer, think first)

Ok Ok!!
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



cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑,布布扣,bubuko.com

cuda编程:关于共享内存(shared memory)和存储体(bank)的事实和疑惑

标签:cuda   共享内存   memory bank   bank conflict   deep learning   

原文地址:http://blog.csdn.net/lingerlanlan/article/details/32712749

(0)
(0)
   
举报
评论 一句话评论(0
登录后才能评论!
© 2014 mamicode.com 版权所有  联系我们:gaon5@hotmail.com
迷上了代码!