标签：cuda 共享内存 memory bank bank conflict deep learning
主要是在研究访问共享内存会产生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];
|
有一个数组，元素类型为整型,个数为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
