bank conflict

存储体冲突(bank conflict)：当被访问的存储体没有恢复时又出现对该存储体新访问的现象。

简介目前 CUDA 装置中,每个 multiprocessor 有 16KB 的 shared memory。Shared memory 分成 16 个 bank。如果同时每个 thread 是存取不同的 bank,就不会产生任何问题,存取shared memory 的速度和存取寄存器相同。不过,如果同时有两个(或更多个) threads 存取同一个 bank 的数据,就会发生 bank conflict,这些 threads 就必须照顺序去存取,而无法同时存取 shared memory 了。例子Shared memory 是以 4 bytes 为单位分成 banks。因此,假设以下的数据:　__shared__ int data[128];
那么,data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、...、data[15] 是 bank15,而 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 存在 bank  conflict，而 global memory 不存在？因为访问 global memory 的只能是 block，而访问 shared memory 的却是同一个 half-warp 中的任意线程。

参考链接

我的理解是存储器划分各个bank是为了减少地址译码器的负担，否则一个32位的地址译码器输出有2^32（参考3-8译码器），所以现在将32位选择高几位（bank片选）作为bank选择器，减少单个译码器的负担。

而对于每个bank，都有一个自己的地址译码器。如果连续访问单个bank，该译码器一直处于被占用中，多线程就无法并行，导致线程并行度不高。由于每个bank都有自己的行缓冲器，有效利用多bank可以大幅度提高性能。

bank conflict相关推荐

Share memory中bank conflict问题
Share memory是片上资源,生命周期是整个block中,它的数据读写十分快,有1个cycle latency.在Share memory中,经常存在bank conflict问题,如果没有ba ...
CUDA bank 及bank conflict
bank 是CUDA中一个重要概念,是内存的访问时一种划分方式,在CPU中,访问某个地址的内存时,为了减少读写内次次数,访问地址并不是随机的,而是一次性访问bank内的内存地址,类似于内存对齐一样,一 ...
share memory的bank conflict分析
背景在做高性能分析的时候,经常会出现一个什么bank conflict的名词,不仅是GPU的share memory会出现bank confict, 甚至连寄存器也会出现bank conflict, ...
难理解的bank conflict
之前看Nvidia-OpenCL-SDK里有一个例子讲到过bank conflict,但没怎么明白,它选择的是用奇数来避免. #define BLOCK_DIM 16// This kernel is ...
Warp divergence Bank conflict
文章目录 Warp divergence & Bank conflict 1.warp.bank概念 2.warp divergence 3.bank conflict Warp diverg ...
【并行计算-CUDA开发】关于共享内存（shared memory）和存储体(bank)的事实和疑惑...
关于共享内存(shared memory)和存储体(bank)的事实和疑惑主要是在研究访问共享内存会产生bank conflict时,自己产生的疑惑.对于这点疑惑,网上都没有相关描述, 不管是国内还 ...
【并行计算-CUDA开发】GPU 的硬体架构关于存储体的介绍比较好bank conflicts
GPU 的硬体架构这里我们会简单介绍,NVIDIA 目前支援CUDA 的GPU,其在执行CUDA 程式的部份(基本上就是其shader 单元)的架构.这里的资料是综合NVIDIA 所公布的资讯, ...
GPU指令集技术分析
GPU指令集技术分析本文将两篇文章整理了一下. 参考文章链接如下: https://zhuanlan.zhihu.com/p/391238629 https://zhuanlan.zhihu.com ...
CUDA之单thread单block多thread单block多thread多block
用简单的立方和归约来举例: //单thread单block #include <stdio.h> #include <stdlib.h> #include <cuda_r ...

bank conflict

bank conflict相关推荐

最新文章

热门文章