CUDA 共享内存 bank conflict

原文链接为:http://blog.csdn.net/endlch/article/details/47043069

1. bank conflict

本文所有的实验针对 GTX980 显卡,Maxwell 架构,计算能力 5.2。

GPU 共享内存是基于存储体切换的架构(bank-switched-architecture)。在 Femi,Kepler,Maxwell 架构的设备上有 32 个存储体(也就是常说的共享内存分成 32 个bank),而在 G200 与 G80 的硬件上只有 16 个存储体。

每个存储体(bank)每个周期只能指向一次操作(一个 32bit 的整数或者一个单精度的浮点型数据),一次读或者一次写,也就是说每个存储体(bank)的带宽为 每周期 32bit。

如下图所示,在一个线程块中申请如下的共享内存:

__shared__ float sData[32][32];

这里写图片描述

也就是说在上述的 32 * 32 的二维数组共享内存中,每一列对应同一个 bank。

  1. 同常量内存一样,当一个 warp 中的所有线程访问同一地址的共享内存时,会触发一个广播(broadcast)机制到 warp 中所有线程,这是最高效的。
  2. 如果同一个 warp 中的线程访问同一个 bank 中的不同地址时将发生 bank conflict。
  3. 每个 bank 除了能广播(broadcast)还可以多播(mutilcast)(计算能力 >= 2.0),也就是说,如果一个 warp 中的多个线程访问同一个 bank 的同一个地址时(其他线程也没有访问同一个bank 的不同地址)不会发生 bank conflict。
  4. 即使同一个 warp 中的线程 随机的访问不同的 bank,只要没有访问同一个 bank 的不同地址就不会发生 bank conflict。

这里写图片描述

如上图所示,左侧和右侧的都没有发生 bank conflict。而中间的存在 bank conflcit,由于经过最多两次,该 warp 中的线程就都可以得到所要的数据,所有称为 2-way bank conflict,如果同一个 warp 中的所有线程访问一个 bank 中的 32 个不同地址,则需要分 32 次,称为 32-way bank conflict。

这里写图片描述

如上图所示,左中右均未发生 bank conflict。

依次我们可以总结:只要同一个 warp 的不同线程会访问到同一个 bank 的不同地址就会发生 bank conflict,除此之外的都不会发生 bank conflict。

既然广播是针对同一个 warp 而言的,那么如果不同的 warp 访问同一个 bank 中的同一个地址呢?由于 每个 SM 中有 4 个 warp scheduler (GTX980),可以很好的调度 warp,使其 warp 之间的访问冲突可以充分的隐藏,因此对效率的影响很小,远远小于 warp 内的 bank conflict。至于 warp scheduler 的调度机制,NVIDIA 没有说的特别清楚,可能也是想要开发者不要过于关注于此。

2. 实验 1

实现定义如下图所示的 32 * 32 线程块,共 1024 个线程,32 个 warp。

这里写图片描述

申请如 1 中所示的 32 * 32 的共享内存,共 32 个 bank,每个 bank 对应 32 个元素。

  • 实验 1.1
    该线程块中的每个 warp 读写不同的 bank,不同的 warp 不会访问一个地址,也就是一一对应的关系。图中的数字就表示上图中的线程标号。经分析可知,此时是没有 bank conflict 的。

这里写图片描述

代码如下:

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.y][threadIdx.x];
}
  • 实验 1.2
    该线程块中的每个 warp 读写相同的 bank 的不同地址,不同的 warp 访问不同,也就是一一对应的关系。图中的数字就表示上图中的线程标号。经分析可知,此时是存在很严重的 bank conflict 。

这里写图片描述

代码如下:

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}
  • 实验 1.3(避免 bank conflict 的技巧)
    针对实验 1.2 中出现的严重的 bank conflict,我们可以通过添加一个附加列来避免 bank conflict,如下图所示,左图为申请的共享内存矩阵形式,右图是表示成 bank 后的形式,通过这种方式,原来在一个 bank 中的同一个 warp 都正好偏移到了不同的 bank 中。

这里写图片描述

代码如下:

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE+1];

if (x_id < col && y_id < row)
{
    sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}

上述三个小实验的运行时间为:

实验 1.1 :0.052416 ms
实验 1.2 :0.131072 ms
实验 1.3 :0.053280 ms

除去公共代码后的时间为:

实验 1.1 :0.034816 ms
实验 1.2 :0.113472 ms
实验 1.3 :0.035680 ms

结论:

  1. 通过额外的一行,可以避免 bank conflict,运行时间与完全没有 bank conflict 的运行时间差距很小。
  2. 存在 bank conflict 的,运行时间几乎是没有 bank conflict 的运行时间的 4 倍。

其实只要添加的是奇数列就可以,只不过 1 列是最节省空间(共享内存太宝贵)的。

3. 实验 2

  • 实验 2.1
    同一个 block 中所有第 i 列的线程都计算第 i 行的元素的和,此时所有同一个warp 会访问同一个 bank 的不同地址。如下图所示,分别表示第 0 列访问 bank 0 中的第一个地址,第 1 列访问 bank 1 中的第 1 个地址,依次类推。

这里写图片描述

代码如下:

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;
    for (int j = 0; j < BLOCKSIZE; j++)
    {
        data += sData[threadIdx.x][j];
    }
    matrixTest[index] = data;
}
  • 实验 2.2
    同实验 1.3 类似,添加额外的一列,如下图所示:

这里写图片描述

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE+1];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;

    for (int j = 0; j < BLOCKSIZE; j++)
    {
        data += sData[threadIdx.x][j];
    }
    matrixTest[index] = data;
}

上述两个实验的运行时间如下所示:

实验 2.1 :0.458144 ms
实验 2.2 :0.090848 ms

从上图也可以看出,修改后的带宽相当于修改前的 32 倍。修改后的运行时间也明显得到改善。

4. 实验 3

  • 实验 3.1
    采用实验 1.1 的方式,同一个 warp 访问不同的 bank,不同的 warp 访问不同的地址。

代码如下:

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;
    for (int j = 0; j < 1000; j++)
    {
        data = sData[threadIdx.y][threadIdx.x];
    }
    matrixTest[index] = data;
}
  • 实验 3.2
    同一个 warp 访问不同的 bank,所有 warp 访问同一个地址,也就是说所有的行都会访问第 0 行。

代码如下:

int x_id = blockDim.x * blockIdx.x + threadIdx.x; // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y; // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if (x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();

    float data = 0.0f;
    for (int j = 0; j < 1000; j++)
    {
        data = sData[0][threadIdx.x];
    }
    matrixTest[index] = data;
}

上述两个实验的运行时间如下所示:

实验 2.1 :0.053800 ms
实验 2.2 :0.055328 ms

在实验 2.2 中存在明显的不同 warp 间的冲突,但是运行时间差距很小,也就是说 warp 间冲突的影响很小。

展开阅读全文

[转]CUDA bank conflict in shared memory

07-16

http://hi.baidu.com/pengkuny/item/c8070b388d75d481b611db7arnrn以前以为 shared memory 是一个万能的 L1 cache,速度很快,只要数据的 size 够小,能够放到 shared memory,剩下的事情我就不用操心啦。实际上不是这样,bank conflict 是一个绕不过去的问题,否则,性能会降得很低,很低,很低。。。rnrn---------------------------------------------------------------rnrn为什么 shared memory 存在 bank conflict,而 global memory 不存在?因为访问 global memory 的只能是 block,而访问 shared memory 的却是同一个 half-warp 中的任意线程。rnhttp://stackoverflow.com/questions/3843032/why-arent-there-bank-conflicts-in-global-memory-for-cuda-opencl rnrn----------------------------------------------------------------rnrn摘自这个要翻墙的网页:Introduction to GPU Programming (HPC Summer Institute at Rice University) http://davidmedinawiki.wordpress.com/2012/06/08/introduction-to-gpu-programming/rnrnNow that we know a little about shared memory, we need to see how this memory is accessed within the thread block. There are these memory managers called “memory banks” that are in charge of distributing the memory they manage.rnrnOk, that sentence sounded like it just stated the obvious but that is what memory banks do. The question now is, what memory are they in charge of?rnrnrn[img=https://img-bbs.csdn.net/upload/201407/16/1405523860_517620.jpg][/img]rnrnMemory Bank Architecture (From the HPC Session)rnrnAbove is a diagram that shows how a GPU with 8 memory banks would store shared memory. Using basic math we get the following equation:rnrn// mem is the memory locationrnbank = mem/8;rnrnSo why are these memory banks so important?rnrnWell, the memory banks distribute data stored in their bank of shared memory one call at a time. This means that a parallel code can easily be turned into serial code due to bank conflicts (when each thread accesses from the same bank at the same time). There is, however, one exception to bank conflicts” which is when threads access the same memory from the same memory bank.rnHere are some examples that show good and bad uses of bank memory (Images taken from the HPC Session):rnrn[img=https://img-bbs.csdn.net/upload/201407/16/1405523956_261689.jpg][/img]rnrn[img=https://img-bbs.csdn.net/upload/201407/16/1405523974_182781.jpg][/img]rnrn[img=https://img-bbs.csdn.net/upload/201407/16/1405523986_842165.jpg][/img]rnrn---------------------------------------------------------------rnrn下面的文字来自:http://hi.baidu.com/dwdxdy/item/e5d66f40168f852810ee1ef7 rn rn共享存储器被组织为16个bank,每个bank拥有32bit的宽度。rn一个warp中的线程对共享存储器的访问请求会被划分为两个half-warp的访问请求。rn无 bank conflict 时,一个half-warp内的线程可以在一个内核周期中并行访问rn对同一 bank 的同时访问导致 bank conflict 只能顺序处理 访存效率降低rn如果half-warp的线程访问同一地址时,会产生一次广播,不会产生 bank conflictrnrnrn[img=https://img-bbs.csdn.net/upload/201407/16/1405524074_971678.jpg][/img]rnrn[img=https://img-bbs.csdn.net/upload/201407/16/1405524086_721675.jpg][/img]rnrn__shared__ float shared[256];rnfloat foo = shared[threadIdx.x];rn没有访问冲突rnrnrnrn__shared__ float shared[256];rnfloat foo = shared[threadIdx.x * 2];rn产生2路访问冲突rnrn__shared__ float shared[256];rnfloat foo = shared[threadIdx.x*8];rn产生8路访问冲突rnrn---------------------------------------------------------------rnrnNumber of shared memory banksrn来源:http://en.wikipedia.org/wiki/CUDA rnGPU device 1.x : 16rnGPU device 2.x : 32rnrn---------------------------------------------------------------rnrn书上说:“每个 bank 的宽度固定为 32 bit,相邻的 32 bit 字被组织在相邻的 bank 中,每个 bank 在每个时钟周期可以提供 32 bit 的带宽。”rnrn由上面这句话可以看出来:将 shared memory 看成一个二维存储空间,每个 bank 就是一列,每一行就是 16(或 32)个 banks。要么,尽量让一个 half-warp(或 full warp)中的线程分散访问不同列(即访问不同的 bank,同行不同行没有关系);要么,让一个 half-warp(或 full warp)中的线程务必全部访问同一列且同一行(即访问同一个地址,仅对读操作有效)。rnrn对于计算能力 1.0 的设备,前个 half-warp 和 后个 half-warp 不存在 bank conflict;rn对于计算能力 2.0 的设备,前个 half-warp 和 后个 half-warp 可能存在 bank conflict,因为 shared memory 可以同时让 32 个 bank 响应请求;rnrn如果是写操作,一个 half-warp(或 full warp) 中所有线程访问同一地址的时候,此时会产生不确定的结果(也只需要一个 clock cycle,不确定哪个线程会胜出),发生这种情况时应使用原子操作——但是原子操作对性能影响太大。rnrn“Shared memory features a broadcast mechanism whereby a 32-bit word can be read and broadcast to several threadssimultaneously when servicing one memory read request. ”——从这个描述来看,只要是多个线程访问同一地址都可以产生一次广播,多个线程访问同一地址将有效减少 bank conflict 的数量。若 half-warp(或 full warp) 中所有线程都要访问同一地址,则完全没有 bank conflict。rnrn对于大于 32 bit 的 struct 来说,对它的访问将编译成多个独立的存储器访问。– “Share memory only supports 32 bit reads/writes”rnrn因此,shared memory 的写操作的 bank conflict 是一个很头疼的问题。rn 论坛

看不到bank conflict,求指教

04-01

#include rn#include rn#include rn#include rn#include rn#include rn#include rn#include rn#include rn//void checkCUDAError(const char *msg);rn#define TILE_DIM 16rnrn__global__ void transposeCoalesced(float *idata, float *odata,int width, int height)rnrn // __shared__ float tile[TILE_DIM][TILE_DIM+1];rn int xIndex = blockIdx.x * TILE_DIM + threadIdx.x;rn int yIndex = blockIdx.y * TILE_DIM + threadIdx.y;rn int index_in = xIndex + width * yIndex;rn int index_out = yIndex + height * xIndex;rn odata[index_out]=idata[index_in];rn rn //odata[index_out]=2;rn /* xIndex = blockIdx.y * TILE_DIM + threadIdx.x; rn yIndex = blockIdx.x * TILE_DIM + threadIdx.y; rn tile[threadIdx.y][threadIdx.x] = idata[index_in];rnrn __syncthreads();rnrn odata[index_out] = tile[threadIdx.x][threadIdx.y];rn */rnrnrn/**rn * Host function that prepares data array and passes it to the CUDA kernel.rn */rnint main(void) rnrn // pointer for host memoryrn float *odata, * h_a, * h_b;rn float *idata;rn int width;rn int height;rn printf("please enter the width of A\n");rn scanf("%d", &width);rn printf("please enter the height of A\n");rn scanf("%d", &height);rn h_a = (float *)malloc(sizeof(float)*width*height);rn h_b = (float *)malloc(sizeof(float)*height*width);rnrn for(int n=0;n>>( idata,odata,width,height );rnrn cudaEventRecord(stop, 0); rn cudaEventSynchronize(stop); rn float elapsedTime; rn cudaEventElapsedTime(&elapsedTime, start, stop);rn cudaThreadSynchronize();rn cudaMemcpy(h_b,odata,sizeof(float)*height*width,cudaMemcpyDeviceToHost);rn rn rn // Part 5 of 5: verify the data returned to the host is correctrn for(int i=0;i 论坛

没有更多推荐了,返回首页