Share memory中bank conflict问题

Share memory是片上资源,生命周期是整个block中,它的数据读写十分快,有1个cycle latency。在Share memory中,经常存在bank conflict问题,如果没有bank conflict问题,它的数据读写可以和片上的寄存器(Register)一样快。因此,我们需要尽量减少bank conflicts.
首先,什么是bank?我们以capability 1.x为例,Share memory被等分成同等尺寸大小的存储器模式,即banks,以下图为例:
这里写图片描述

每一个bank的带宽为32-bit,即4 bytes。连续的32-bit字节可以被分到连续的bank中去。我们以G80为例,它具有16个banks,所以每一个bank=4-byte address % 16,此外,bank的数量和半个warp的thread数量一致,所以share memory对一个warp的请求,分成了两个先后请求来做,不同的半个warp之间没有bank conflicts。如果同时有两个请求在同一个bank中,就会出现bank conflicts,如下图:
这里写图片描述这里写图片描述

对于share memory,有数据读写有快慢两种情况:

快:1)半个warp的所有threads在同一时刻读写share memory的不同bank。2)半个warp中的所有threads在同一时刻以广播的形式读写share memory中的同一个bank。

慢:半个warp的所有threads在同一时刻读写share memory的多个bank(不是全部),这时必须串行读写。

以G80为例,其share memory具有16个bank,在下面的图示中,只有S为奇数时,才不会存在bank conflicts。
这里写图片描述
这里写图片描述
当s=4时,
这里写图片描述

同样,在存储结构体数据时,也有可能会出现bank conflicts。

例如如下代码存在三路冲突:

struct vector { float x, y, z; };

__shared__ struct vector vectors[32];

struct vector v = vectors[baseIndex + threadIdx.x];

这里写图片描述
因此,对于share memory,最好的读写方式每个thread读取block中连续的元素。

以如下例程为例:

对于一个二维数组,半个warp的threads读取每一列,这存在了16路bank conflicts。
这里写图片描述

我们有两种方式可以解决:

1) 在每一行的后面加一个元素,例如:
这里写图片描述

2)在处理矩阵之前转置矩阵

我们对上一节global memory中的coalescing的例程,作进一步优化,分配额外的一列:
这里写图片描述

代码如下:

__global__ void transpose_exp(float *odata, float *idata, int width, int height){

__shared__ float block[BLOCK_DIM][BLOCK_DIM+1];

unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;

unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;

if( (xIndex < width)&&(yIndex < height) {

unsigned int index_in = xIndex + yIndex * width;

block[threadIdx.y][threadIdx.x] = idata[index_in];

}

__syncthreads();

xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;

yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;

if( (xIndex < height)&&(yIndex < width) ){

unsigned int index_out = yIndex * height + xIndex;

odata[index_out] = block[threadIdx.x][threadIdx.y];

}

}

实验结果为:

 128x128: 0.011ms vs. 0.022ms (2.0X speedup)

 512x512: 0.07ms vs. 0.33ms (4.5X speedup)

 1024x1024: 0.30ms vs. 1.92ms (6.4X speedup)

 1024x2048: 0.79ms vs. 6.6ms (8.4X speedup)

我们再来考虑矩阵相乘时,会不会出现bank conflicts问题,判断技巧在于固定k,考虑同一时刻thread的读些问题

1)当在读一个Ms时,在同一时刻,半个warp的16个threads以广播的形式读取同一个bank

for (int k = 0; k < 16; ++k) //ty is constant over a half-warp
Csub += Ms[ty][k] * Ns[k][tx];//here k is fixed over a half-warp

这里写图片描述

2)当在读一个Ns时,在同一时刻,半个warp的16个threads连续读取不同的bank,stride=1.

for (int k = 0; k < 16; ++k) //tx is constant over a half-warp
Csub += Ms[ty][k] * Ns[k][tx];

这里写图片描述

所以综上所述,矩阵相乘的例子并没有bank conflicts问题。

需要注意的是,我们上述所说的均为capability 1.x的情况,在2.x 3.x中,每个share memory有32个bank,具体还是要看技术手册的。

[转]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=http://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=http://img.bbs.csdn.net/upload/201407/16/1405523956_261689.jpg][/img]rnrn[img=http://img.bbs.csdn.net/upload/201407/16/1405523974_182781.jpg][/img]rnrn[img=http://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=http://img.bbs.csdn.net/upload/201407/16/1405524074_971678.jpg][/img]rnrn[img=http://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 论坛

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