brucefan1983 / CUDA-Programming

Sample codes for my CUDA programming book
GNU General Public License v3.0
1.55k stars 320 forks source link

bank冲突避免 #23

Closed li199603 closed 1 year ago

li199603 commented 1 year ago

你好,有一个问题需要请教 在8.3中,通过shared real S[TILE_DIM][TILE_DIM + 1]来解决bank冲突的问题。原先S[TILE_DIM][TILE_DIM]的数据刚好可以存进32个bank(具有32层)。现在数组变大了,那多出来的那部分数据被存到哪里? 有一种解释是,每个bank的层数并不是固定的32,是硬件决定的。所以多出来那部分数据在第33层?

如果这个说法正确的话,那么书上可能要明确写出bank层数不是32,及其层数计算方法。就我自己而言,看到图8.1的时候,就认为bank只会是32层。

希望采纳~

fever-Wong commented 1 year ago

谢谢,您发给我的邮件已经收到,我会尽快处理。Thank you,the email you sent me has been received and I will handle it as soon as possible.王景博fever wong

brucefan1983 commented 1 year ago

感谢您的提问。书中没有说共享内存的层数一定是32层,只不过图8.1显示的是32层(因为是以数组转置为例进行讲解)。共享内存只规定每一层有32个bank,至于有多少层,那是由共享内存数组的数据量决定的。如果您仔细看,会发现书中第87页的第一段所举的例子就只有4层。

brucefan1983 commented 1 year ago

再明确回答一下,如果定义了

#define TILE_DIM 32

// in the kernel:
__shared__ float S[TILE_DIM][TILE_DIM+1];

那么该共享内存数组 S 的数据需要使用33层,每层有32个4字节的bank。