CUDA

您所在的位置：网站首页 › 什么是共享内存 › CUDA

CUDA

2024-07-14 13:34| 来源: 网络整理| 查看: 265

一、共享内存的结构 1)什么是共享内存？

共享内存是GPU的一种稀缺资源，它位于芯片上，所以共享内存空间要比本地和全局内存空间快得多。对于warp里的所有线程，只要线程之间没有任何存储体冲突(bank conflict)，访问共享内存就与访问寄存器一样快。

2)什么是存储体（bank）？

共享内存被划分为同样大小的、可以同时访问的内存块，名为存储体。在计算能力为1.x的设备上，存储体数为16，在2.0及以上的设备，存储体数为32。

存储体的存在可使共享内存获得高内存带宽，假设有n个存储体，此时访问n个分别位于不同bank的地址时是同时进行的，最后获得的有效带宽就是单个模块的带宽的N倍，因为只需要发射一次指令。

3）共享内存是如何映射到存储体的？

共享内存空间的存储体组织为：连续的32位（4个字节）分配到连续的存储体中，每个存储体的带宽为 32位/2个时钟周期

一维共享内存

如下的方式申请一个一维的共享内存 (32个存储体)

__shared__ float sData[64]; // sData[0]-sData[31]分别对应bank[0]-bank[31]; // sData[32]-sData[63]分别对应bank[0]-bank[31];

此时， sData[0]与sData[32]位于同一个存储体bank[0] sData[1]与sData[33]位于同一个存储体bank[1] 。。。。。。 sData[31]与sData[63]位于同一个存储体bank[31]

二维共享内存

二维共享内存其实可以展开成一维内存，其映射方式跟一维一样。

__shared__ float sData[32][32];

由于上面共享内存的每一行大小为32，刚好等于存储体个数，那么此时共享内存的每一列就是一个bank。

二、避免共享存储体冲突（bank conflict）

同一个warp里的线程访问同一个bank里不同的地址时，会出现bank conflict，如果访问的不同的地址的个数为n，那么此种情况称为n路存储器冲突（n-way bank conflict）。

1）同一个warp访问共享内存的同一个bank的不同地址，所产生的bank conflict。

a.通过改变数据在共享内存中的排列方式，使其映射到bank时，原本在同一个bank的不同地址的数据分开到不同的bank。

二维共享内存例子：

size_t ix=threadIdx.x+blockIdx.x*blockDim.x;//0-31 size_t iy=threadIdx.y+blockIdx.y*blockDim.y;//0-31 __shared__ float sData[32][32+1]; //此处的共享内存改变了数据的排列方式，通过映射到bank上，使得原本在同一个bank的地址偏移到了不同的bank上。 if(ix (int)log2(32);//0-1 size_t ix = tid & (32 - 1);//0-31 __shared__ float sData[64 + 2]; float data; if (tid

【本文地址】

CUDA

CUDA

今日新闻

推荐新闻