深海的小鱼儿

  博客园 :: 首页 :: 博问 :: 闪存 :: 新随笔 :: 联系 :: 订阅 订阅 :: 管理 ::

下面简单介绍一些cuda中的共享存储器和全局存储器 

共享存储器,shared memory,可以被同一块中的所有线程访问的可读写存储器,生存期是块的生命期。

Tesla的每个SM拥有16KB共享存储器。

在编程过程中,有静态的shared memory 动态的shared memory

静态的shared memory 在程序中定义   __shared__ type shared[SIZE];

动态的shared memory 通过内核函数的每三个参数设置大小 extern __shared__ type shared[];

共享存储器被组织为16个bank,每个bank拥有32bit的宽度。

无bank conflict时,一个half-warp内的线程可以在一个内核周期中并行访问

对同一bank的同时访问导致bank conflict   只能顺序处理 访存效率降低

如果half-warp的线程访问同一地址时,会产生一次广播,不会产生bank conflict

__shared__ float shared[256];

float foo = shared[threadIdx.x];

没有访问冲突

__shared__ float shared[256];

float foo = shared[threadIdx.x * 2];

产生2路访问冲突

__shared__ float shared[256];

float foo = shared[threadIdx.x*8];

产生8路访问冲突

posted on 2015-05-16 19:33  深海的小鱼儿  阅读(395)  评论(0编辑  收藏  举报