gre*_*man 3 cuda shared-memory
我对来自NVidia的2份文件感到困惑."CUDA最佳实践"描述了共享存储器在存储体中组织,并且通常在32位模式下,每4个字节是一个存储体(这就是我所理解的).但是,"使用CUDA进行并行前缀求和(扫描)"(此处提供:http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html)详细说明了由于存储库冲突,应如何在扫描算法中添加填充.
对我来说问题是,这个算法的基本类型是float,它的大小是4个字节.因此,每个浮动都是银行,没有银行冲突.
所以我的理解是正确的 - 即如果你使用4*N字节类型,你不必担心银行冲突,因为根据定义,没有？如果不是,我应该如何理解它(何时使用填充)？
您可能会感兴趣的这个研讨会由NVIDIA CUDA研讨会页面 共享存储器,包括也被描述从幻灯片35-45银行本次网络研讨会.
通常,共享内存库冲突可能发生在两个不同的线程试图访问(来自相同的内核指令)共享内存中的较低4(pre-cc2.0设备)或5位(cc2.0和更高版本)的位置设备)的地址是相同的.当发生银行冲突时,共享内存系统将访问序列化到同一组中的位置,从而降低性能.Padding试图为某些访问模式避免这种情况.请注意,对于cc2.0和更高版本,如果所有位都相同(即相同位置),则不会导致存储体冲突.
从图像上看,我们可以这样看:
__shared__ int A[2048];
int my;
my = A[0]; // A[0] is in bank 0
my = A[1]; // A[1] is in bank 1
my = A[2]; // A[2] is in bank 2
...
my = A[31]; // A[31] is in bank 31 (cc2.0 or newer device)
my = A[32]; // A[32] is in bank 0
my = A[33]; // A[33] is in bank 1
现在,如果我们跨warp中的线程访问共享内存,我们可能会遇到银行冲突:
my = A[threadIdx.x]; // no bank conflicts or serialization - handled in one trans.
my = A[threadIdx.x*2]; // 2-way bank conflicts - will cause 2 level serialization
my = A[threadIdx.x*32]; // 32-way bank conflicts - will cause 32 level serialization
让我们仔细看看上面的双向银行冲突.由于我们是乘以threadIdx.x 由2,线程0的访问位置0中存储体0但螺纹16访问位置32,其也于存储体0,因此产生存储体冲突.对于上面的32路示例,所有地址都对应于bank 0.因此,共享内存的32个事务必须满足此请求,因为它们都是序列化的.
所以要回答这个问题,如果我知道我的访问模式就像这样:
my = A[threadIdx.x*32];
然后我可能想要填充我的数据存储,因此这A[32]是一个虚拟/填充位置A[64],A[96] 等等.然后我可以获取相同的数据,如下所示:
my = A[threadIdx.x*33];
并获得我的数据没有银行冲突.
希望这可以帮助.
你的理解是错误的。当来自同一经线的线程访问驻留在同一组中的不同值时，就会发生组冲突。
来自 CUDA C 编程指南：
为了实现高带宽，共享内存被分成大小相等的内存模块，称为存储体，可以同时访问。因此，任何由位于 n 个不同存储器组中的 n 个地址构成的存储器读取或写入请求都可以同时得到服务，从而产生比单个模块带宽高 n 倍的总带宽。
但是，如果一个内存请求的两个地址落在同一个内存bank中，就会发生bank冲突，访问就必须串行化。硬件根据需要将具有存储体冲突的内存请求拆分为多个单独的无冲突请求，从而将吞吐量降低与单独内存请求数量相等的因子。如果单独的内存请求数为 n，则称初始内存请求会导致 n 路 bank 冲突。
填充用于避免库冲突。当您知道共享内存访问模式时，您就可以确定如何填充共享内存阵列以避免存储区冲突。
例如，如果假设您有__shared__ float x[32][32];并且每个线程索引为 tid 的线程都像这样访问 x somevariable = x[tid][0];。这将导致 32 路 bank 冲突，因为所有线程都从同一个 bank 访问不同的值。
为避免冲突，您可以在第一个维度中使用另一个元素填充数组：__shared__ float x[32][33];。这将完全消除 bank 冲突，因为现在每一行都有一个 bank 位置，该位置与前一行偏移一个。