581
社区成员
发帖
与我相关
我的任务
分享请发表友善的回复…
发表回复
zalyer 2012-09-03
- 打赏
- 举报
比喻成通道容易理解
[Quote=引用 10 楼 的回复:]
直观来看，16个bank就相当于16个并行存取通道，每个通道内存放着若干数据，可以并行存取，
如1楼所述，数组数据被依次分散存放在16个通道内，原文是这样说的，
Shared memory has 32 banks that are organized such that
successive 32-bit words are assigned to successive ba……
[/Quote]
[Quote=引用 10 楼 的回复:]
直观来看，16个bank就相当于16个并行存取通道，每个通道内存放着若干数据，可以并行存取，
如1楼所述，数组数据被依次分散存放在16个通道内，原文是这样说的，
Shared memory has 32 banks that are organized such that
successive 32-bit words are assigned to successive ba……
[/Quote]
hohaics 2012-08-03
- 打赏
- 举报
直观来看，16个bank就相当于16个并行存取通道，每个通道内存放着若干数据，可以并行存取，
如1楼所述，数组数据被依次分散存放在16个通道内，原文是这样说的，
Shared memory has 32 banks that are organized such that
successive 32-bit words are assigned to successive banks, i.e. interleaved.
接连的 32-bit words assigned 接连的 banks
如1楼所述，数组数据被依次分散存放在16个通道内，原文是这样说的，
Shared memory has 32 banks that are organized such that
successive 32-bit words are assigned to successive banks, i.e. interleaved.
接连的 32-bit words assigned 接连的 banks
ddweidong 2012-06-04
- 打赏
- 举报
[Quote=引用 1 楼 的回复:]
这类的文章挺多的。
Shared memory 是以 4 bytes 为单位分成 banks。因此，假设以下的数据：
__shared__ int data[128];
那么，data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是 bank 15，而 data[16] 又回到 bank 0。由于 wa……
[/Quote]讲的很好，“因此，如果程序在存取 shared memory 的时候，使用以下的方式： int number = data[base + tid];”这句解决了我的疑惑
这类的文章挺多的。
Shared memory 是以 4 bytes 为单位分成 banks。因此，假设以下的数据：
__shared__ int data[128];
那么，data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是 bank 15，而 data[16] 又回到 bank 0。由于 wa……
[/Quote]讲的很好，“因此，如果程序在存取 shared memory 的时候，使用以下的方式： int number = data[base + tid];”这句解决了我的疑惑
frank0712105003 2012-05-26
- 打赏
- 举报
恩，很清楚了。。
jmailhappy 2011-03-28
- 打赏
- 举报
[Quote=引用 5 楼 xhsh001 的回复:]
既然一个SM里只有16个bank，且每个bank只有4Byte，也就是只能保存16个整数，那int data[128]中这128个整数是如何保存在共享存储器里的？
[/Quote]
Tesla 的每个 SM 拥有 16KB 共享存储器，用于同一个线程块内的线程间通信。为了使一
个 half-warp 内的线程能够在一个内核周期中并行访问，共享存储器被组织成 16 个 bank，
每个 bank 拥有 32bit 的宽度，故每个 bank 可保存 256 个整形或单精度浮点数，或者说目前
的 bank 组织成了 256 行 16 列的矩阵。如果一个 half-warp 中有一部分线程访问属于同一
bank 的数据，则会产生 bank conflict，降低访存效率，在冲突最严重的情况下，速度会比
全局显存还慢，但是如果 half-warp 的线程访问同一地址的时候，会产生一次广播，其速度
反而没有下降。在不发生 bank conflict 时，访问共享存储器的速度与寄存器相同。在不同
的块之间，共享存储器是毫不相关的。 ------风辰的 CUDA 入门教程
既然一个SM里只有16个bank，且每个bank只有4Byte，也就是只能保存16个整数，那int data[128]中这128个整数是如何保存在共享存储器里的？
[/Quote]
Tesla 的每个 SM 拥有 16KB 共享存储器，用于同一个线程块内的线程间通信。为了使一
个 half-warp 内的线程能够在一个内核周期中并行访问，共享存储器被组织成 16 个 bank，
每个 bank 拥有 32bit 的宽度，故每个 bank 可保存 256 个整形或单精度浮点数，或者说目前
的 bank 组织成了 256 行 16 列的矩阵。如果一个 half-warp 中有一部分线程访问属于同一
bank 的数据，则会产生 bank conflict，降低访存效率，在冲突最严重的情况下，速度会比
全局显存还慢，但是如果 half-warp 的线程访问同一地址的时候，会产生一次广播，其速度
反而没有下降。在不发生 bank conflict 时，访问共享存储器的速度与寄存器相同。在不同
的块之间，共享存储器是毫不相关的。 ------风辰的 CUDA 入门教程
jmailhappy 2011-03-28
- 打赏
- 举报
同意楼上的，我也很疑惑
xhsh001 2011-02-21
- 打赏
- 举报
既然一个SM里只有16个bank，且每个bank只有4Byte，也就是只能保存16个整数，那int data[128]中这128个整数是如何保存在共享存储器里的？
xhsh001 2011-02-21
- 打赏
- 举报
[Quote=引用 1 楼 zhui22222 的回复:]
很多时候 shared memory 的 bank conflict 可以透过修改数据存放的方式来解决。例如，以下的程序：
data[tid] = global_data[tid];
...
int number = data[16 * tid];
会造成严重的 bank conflict，...
[/Quote]
我不大理解，这里怎么会造成严重的bank conflict呢？16*tid都是16的倍数，都对应着bank 0，这应当会执行一次广播啊。
很多时候 shared memory 的 bank conflict 可以透过修改数据存放的方式来解决。例如，以下的程序：
data[tid] = global_data[tid];
...
int number = data[16 * tid];
会造成严重的 bank conflict，...
[/Quote]
我不大理解，这里怎么会造成严重的bank conflict呢？16*tid都是16的倍数，都对应着bank 0，这应当会执行一次广播啊。
开飞机的贝塔塔 2010-07-22
- 打赏
- 举报
MARK..
penyunwudong 2010-07-19
- 打赏
- 举报
谢了，讲得很清楚，可以理解了～
zhui22222 2010-07-16
- 打赏
- 举报
这类的文章挺多的。
Shared memory 是以 4 bytes 为单位分成 banks。因此，假设以下的数据：
__shared__ int data[128];
那么，data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是 bank 15，而 data[16] 又回到 bank 0。由于 warp 在执行时是以 half-warp 的方式执行，因此分属于不同的 half warp 的 threads，不会造成 bank conflict。
因此，如果程序在存取 shared memory 的时候，使用以下的方式：
int number = data[base + tid];
那就不会有任何 bank conflict，可以达到最高的效率。但是，如果是以下的方式：
int number = data[base + 4 * tid];
那么，thread 0 和 thread 4 就会存取到同一个 bank，thread 1 和 thread 5 也是同样，这样就会造成 bank conflict。在这个例子中，一个 half warp 的 16 个 threads 会有四个 threads 存取同一个 bank，因此存取 share memory 的速度会变成原来的 1/4。
一个重要的例外是，当多个 thread 存取到同一个 shared memory 的地址时，shared memory 可以将这个地址的 32 bits 数据「广播」到所有读取的 threads，因此不会造成 bank conflict。例如：
int number = data[3];
这样不会造成 bank conflict，因为所有的 thread 都读取同一个地址的数据。
很多时候 shared memory 的 bank conflict 可以透过修改数据存放的方式来解决。例如，以下的程序：
data[tid] = global_data[tid];
...
int number = data[16 * tid];
会造成严重的 bank conflict，为了避免这个问题，可以把数据的排列方式稍加修改，把存取方式改成：
int row = tid / 16;
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid];
这样就不会造成 bank conflict 了。
Shared memory 是以 4 bytes 为单位分成 banks。因此，假设以下的数据：
__shared__ int data[128];
那么，data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是 bank 15，而 data[16] 又回到 bank 0。由于 warp 在执行时是以 half-warp 的方式执行，因此分属于不同的 half warp 的 threads，不会造成 bank conflict。
因此，如果程序在存取 shared memory 的时候，使用以下的方式：
int number = data[base + tid];
那就不会有任何 bank conflict，可以达到最高的效率。但是，如果是以下的方式：
int number = data[base + 4 * tid];
那么，thread 0 和 thread 4 就会存取到同一个 bank，thread 1 和 thread 5 也是同样，这样就会造成 bank conflict。在这个例子中，一个 half warp 的 16 个 threads 会有四个 threads 存取同一个 bank，因此存取 share memory 的速度会变成原来的 1/4。
一个重要的例外是，当多个 thread 存取到同一个 shared memory 的地址时，shared memory 可以将这个地址的 32 bits 数据「广播」到所有读取的 threads，因此不会造成 bank conflict。例如：
int number = data[3];
这样不会造成 bank conflict，因为所有的 thread 都读取同一个地址的数据。
很多时候 shared memory 的 bank conflict 可以透过修改数据存放的方式来解决。例如，以下的程序：
data[tid] = global_data[tid];
...
int number = data[16 * tid];
会造成严重的 bank conflict，为了避免这个问题，可以把数据的排列方式稍加修改，把存取方式改成：
int row = tid / 16;
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid];
这样就不会造成 bank conflict 了。
