Shared Memory

1、引言

在global memory部分,数据对齐和连续是提升性能的很重要的因素,当使用L1 cache的时候,对齐问题不再是问题,但是非连续的获取内存依然会降低性能。依赖于算法本质,某些情况下,非连续访问是不可避免的。使用shared memory是另一种提高性能的方式。

2、Introduction CUDA Shared Memory

GPU上的memory有两种:

1、On-board memory

2、On-chip memory

global memory就是一块很大的on-board memory,并且有很高的latency;而shared memory正好相反,是一块很小、低延迟的on-chip memory,比global memory拥有高得多的带宽。我们可以把他当做可编程的cache,其主要作用有:

1、An intra-block thread communication channel   一个block中线程间交流通道

2、A program-managed cache for global memory data  可编程的cache

3、Scratch pad memory for transforming data to improve global memory access patterns 

shared memory(SMEM)是GPU的重要组成之一。物理上,每个SM包含一个当前正在执行的block中所有thread共享的低延迟的内存池。SMEM使得同一个block中的thread能够相互合作,重用on-chip数据,并且能够显著减少kernel需要的global memory带宽。由于CUDA可以直接显式的操作SMEM的内容,所以又被称为可编程缓存。
由于shared memory和L1要比L2和global memory更接近SM,shared memory的延迟比global memory低20到30倍,带宽大约高10倍。
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
当一个block开始执行时,GPU会分配其一定数量的shared memory,这个shared memory的地址空间会由block中的所有thread 共享。shared memory是划分给SM中驻留的所有block的,也是GPU的稀缺资源。所以,使用越多的shared memory,能够并行的active thread blocks就越少。
关于Program-Managed Cache:在C语言编程里,循环(loop transformation)一般都使用cache来优化。在循环遍历的时候使用重新排列的迭代顺序可以很好利用cache局部性。在算法层面上,我们需要手动调节循环来达到令人满意的空间局部性,同时还要考虑cache size。cache对于程序员来说是透明的,编译器会处理所有的数据移动,我们没有能力控制cache的行为。shared memory则是一个可编程可操作的cache,程序员可以完全控制其行为。

Shared Memory Allocation

我们可以动态或者静态的分配shared memory,其声明即可以在kernel内部也可以作为全局变量。CUDA支持声明1D、2D和3D的shared memory数组。其标识符为:__shared__
静态声明2D浮点型数组:

__shared__ float tile[size_y][size_x];

如果在kernel中声明的话,其作用域就是kernel内,否则是对所有kernel有效。如果shared memory的大小在编译期未知的话,可以使用extern关键字修饰,例如下面声明一个未知大小的1D数组:

extern __shared__ int tile[];

由于其大小在编译期未知,我们需要在每个kernel调用时,动态的分配其shared memory,也就是最开始提及的第三个参数:

kernel<<<grid, block, isize * sizeof(int)>>>(...)

注意:只有1D数组才能这样动态使用。

Shared Memory Banks and Access Mode

当优化内存性能时,有两个重要的因素来量化:latency和bandwidth。shared memory能够用来隐藏由于latency和bandwidth对性能的影响。下面将解释shared memory的组织方式,以便研究其对性能的影响。
(1) Memory Banks
为了获得高带宽,shared memory被分成32(计算能力1.x的device划分为16个banks)个相等大小的内存块,每块大小32-bit(4 bytes),他们可以被同时访问。不同的计算能力的device,shared memory以不同的模式映射到不同的块(稍后详解)。如果warp访问shared memory,对于每个bank只访问不多于一个内存地址,那么只需要一次内存传输就可以了,否则需要多次传输,因此会降低内存带宽的使用。
(2) Banks Conflict
当一个warp中多个地址请求落在同一个bank中就会发生bank conflict,从而导致请求多次执行。硬件会把这类请求分散到尽可能多的没有conflict的那些传输操作里面,降低有效带宽的因素是被分散到的传输操作个数。 warp有三种典型的获取shared memory的模式:

· Parallel access:多个地址分散在多个bank。

· Serial access:多个地址落在同一个bank。

· Broadcast access:一个地址读操作落在一个bank。

Parallel access是最通常的模式,这个模式表示,一些(也可能是全部)地址请求能够被一次传输解决。理想情况是,获取无conflict的shared memory的时,每个地址都在落在不同的bank中。
Serial access是最坏的模式,如果warp中的32个thread都访问了同一个bank中的不同位置,那就是32次单独的请求,而不是同时访问了。
Broadcast access也是只执行一次传输,然后传输结果会广播给所有发出请求的thread。这样的话就会导致带宽利用率低。
下图是最优情况的访问图示,每个线程访问一个32-bit的数据,不存在bank conflict:
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
 下图是不规律的随机访问模式,因为每个thread访问不同的bank,因此也没有冲突:
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
 下图是bank冲突的情况,几个threads访问同一个bank,会产生下列两种行为:

· Conflict-free broadcast access if threads access the same address within a bank

· Bank conflict access if threads access different addresses within a bank

CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
 (3) Access Mode
根据device不同的计算能力版本,bank的大小配置也不同,具体为:

· 4 bytes for devices of compute capability 2.x

· 8 bytes for devices of compute capability 3.x

以Fermi的GPU为例,它有32个banks,每个bank 32-bit,即4 bytes,每个bank的带宽是32bits每两个cycle。连续的32位数据映射到连续的bank中,也就是说,bank的索引和shared memory地址的映射关系如下:

 bank index = (byte address ÷ 4 bytes/bank) % 32 banks

下图是Fermi的地址映射关系,注意到,bank中每个地址相差32,相邻的word分到不同的bank中以便使warp能够获得更多的并行获取内存操作(获取连续内存时,连续地址分配到了不同bank中)。
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
当同一个warp的两个thread要获取同一个地址(注意是同一个地址,同一个bank会造成冲突)的时候并不发生bank conflict。对于读操作,会用一次transaction获得结果后广播给所有请求,当写操作时,只有一个thread会真正去写,但是哪个thread执行是无法确定的。
对于Kepler设备来说,shared memory有两种地址模式的32个banks:

· 64-bit mode

· 32-bit mode

在64-bit模式中,连续的64-bits字会映射到连续的bank。每个bank带宽是64bite/1个clock。其映射关系公式:

 bank index = (byte address ÷ 8 bytes/bank) % 32 banks

这里,如果两个thread访问同一个64-bit中的任意一个sub-word(1byte)也不会导致bank conflict,因为一次64-bit(bank带宽64bit/cycle)的读操作就可以满足请求了。也就是说,同等情况下,64-bit模式一般比32-bit模式更少碰到bank conflict。
下图表示了32-bit模式下从字节地址到bank索引的映射关系图。上面表示用字节地址和4-byte word索引标签的共享内存,下面表示从4-byte word索引到bank索引的映射关系。尽管word 0和word 32都在bank0中,同时读这两个word也不会导致bank conflict(64-bit/cycle):
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
下图是64-bit模式下,conflict-free的情况,每个thread获取不同的bank:
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
下图是另一种conflict-free情况,两个thread或获取同一个bank中的word:
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
下图红色箭头是三路bank conflict发生的情况:
 CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
 (3) Memory Padding
memory padding是一种避免bank conflict的方法,如下图所示,所有的thread分别访问了bank0的5个不同的word,这时就会导致bank conflict,我们采取的方法就是在每N(bank数目)个word后面加一个word,这样就如下面右图那样,原本bank0的每个word转移到了不同的bank中,从而避免了bank conflict。
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
增加的这些word不会用来存储数据,其唯一的作用就是移动原始bank中的word,避免冲突。使用memory padding会导致block可获得shared memory中有用的数量减少。内存填充后,需要重新计算数组索引以确保可以访问到正确的数据。注意Fermi和Kepler的bank宽度不同,所以针对于Fermi的填充模型用于Kepler可能导致bank conflict。
 (3) Access Mode Configuration
对Kepler来说,默认情况是4-byte模式,可以用下面的API来查看:

 cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);

返回结果放在pConfig中,其结果可以是下面两种:

cudaSharedMemBankSizeFourByte

cudaSharedMemBankSizeEightByte

可以使用下面的API来设置bank的大小:

 cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);

bank的配置参数如下三种:

cudaSharedMemBankSizeDefault

cudaSharedMemBankSizeFourByte

cudaSharedMemBankSizeEightByte

在启动不同的kernel之间修改bank配置会有一个隐式的device同步。修改shared memory的bank大小不会增加shared memory的利用率或者影响kernel的Occupancy,但是对性能是一个主要的影响因素。一个大的bank会产生较高的带宽,但是鉴于不同的access pattern,可能导致更多的bank conflict。

Configuring the Amount of Shared Memory

每个SM拥有64 KB的片上内存,shared memory和L1 cache共享这块内存。CUDA提供了两种方式配置它们各自的大小,参考CUDA学习----Memory Model,里面提到的是Per-kernel configuration的情况。

Per-device configuration

Per-kernel configuration

cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);

对于Per-device configuration的情况用以上API,两种情况类似,只是作用范围不同。配置方式孰优孰劣取决于kernel用的shared memory的多少。

Synchronization

因为shared memory可以被同一个block中的不同的thread同时访问,当同一个地址的值被多个thread修改就导致了inter-thread conflict,所以我们需要同步操作。CUDA提供了两类block内部的同步操作,即:

· Barriers

· Memory fences

对于barriers,所有thread会等待其他threads到达barrier point;对于Memory fence,所有threads会被阻塞直到所有修改memory的操作对其他threads可见。下面解释下CUDA需要同步的主要原因:weakly-ordered。
(1) Weakly-Ordered Memory Model
现代内存架构有非常宽松的内存模式,也就是意味着,memory的获取不必按照程序中的顺序来执行。CUDA采用了一种叫做weakly-ordered内存模型来获取更激进的编译器优化。
GPU thread写数据到不同的memory的顺序(比如shared memory,global memory,page-locked host memory或者另一个device上的memory)同样没必要跟程序里面顺序相同。一个thread的读操作的顺序对其他thread可见时也可能与实际执行写操作的thread顺序不一致。
为了显式的强制程序以一个确切的顺序运行,就需要用到fence和barrier。它们也是唯一的方式能保证kernel与其它线程分享资源时对memory有正确行为。
(2) Explicit Barrier
可以在kernel中设置一个barrier point通过调用下列函数:

void __syncthreads();

__syncthreads就是作为一个barrier point起作用,block中的threads必须等待所有thread都到达这个point后才能继续下一步。这也保证了所有在这个point之前获取global memory和shared memory的操作对同一个block中所有thread可见。 __syncthreads被用来协作同一个block中的thread通信。当一些thread获取memory相同的地址时,就会导致潜在的问题(read-after-write,write-after-read和write-after-write)从而引起未定义行为状态,此时就可以使用__syncthreads来避免这种情况。
使用__syncthreads要相当小心,只有在所有thread都会到达这个point时才可以调用这个同步,显而易见,如果同一个block中的某些thread永远都不能到达该point,那么程序将一直等下去,下面代码就是一种错误的使用方式:
if (threadID % 2 == 0)
__syncthreads();
else
__syncthreads();
如果在block之间不同步的话,thread blocks可能以任意顺序,并行或者串行,在任意的SM上被执行。如果一个CUDA kernel需要全局同步,可以通过在同步点分割kernel和启动多个kernel来达到这种期望的行为。
(3) Memory Fence
这种方式保证了任何在fence之前的memory写操作对fence之后thread都可见,也就是,fence之前写完了,fence之后其它thread就都知道这块memory写后的值了。fence的设置范围比较广,分为:block,grid和system。 可以通过下面的API来设置fence:

void __threadfence_block();

void __threadfence();

void __threadfence_system();

其中,第一个函数是对应的block范围的,也就是保证同一个block中thread在fence之前写完的值对block中其它的thread可见,不同于barrier,该函数不需要所有的thread都执行;第二个函数是对应grid范围的;第三个对用system的,其范围针对整个系统,包括device和host。
(4) Volatile Qualifier
声明一个使用global memory或者shared memory的变量,用volatile修饰符来修饰该变量的话,会组织编译器做一个该变量的cache优化。使用该修饰符后,编译器就会认为该变量可能在某一时刻被别的thread改变,如果使用cache优化的话,得到的值就缺乏时效,因此使用volatile强制每次都到global 或者shared memory中去读取其绝对有效值。

3、Checking the Data Layout of Shared Memory

我们在设计使用shared Memory的时候应该关注下面的信息:

· Mapping data elements across Memory banks

· Mapping from thread index to shared Memory offset

搞明白这两点,就可以掌握shared memory的使用了,从而构建出高性能的代码。

Square Shared Memory

我们可以以一种直接的方式用shared memory缓存全局内存中的方阵。下图展示了一个每一维度有32个元素并以row-major存储在shared memory的数组,图的最上方是该矩阵实际的一维存储图示,下方是通过映射4-byte数据和banks关系的逻辑二维shared memory:
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
 我们可以使用下面的语句静态声明一个2D的shared memory变量:

__shared__ int tile[N][N];

因为是方阵,可以从2D线程块中以相邻的thread获取相邻的元素的方式访问数据:

 tile[threadIdx.y][threadIdx.x]

 tile[threadIdx.x][threadIdx.y]

上面两种方式哪个性能更好呢?这就需要注意thread和bank的映射关系了,我们最希望看到的是同一个warp中的thread获取的是不同的bank。同一个warp中的thread可以使用连续的threadIdx.x来确定。不同bank中的元素同样是连续存储的,以word大小作为偏移。因此,最好是让连续的thread(由连续的threadIdx.x确定)获取shared memory中连续的地址,由此得知,tile[threadIdx.y][threadIdx.x]应该展现出更好的性能以及更少的bank conflict。
(1) Access Row-Major versus Column-Major
假设我们的grid有2D的block(32,32),定义如下:
#define BDIMX 32
#define BDIMY 32
dim3 block(BDIMX,BDIMY);
dim3 grid(
1,1);
我们对这个kernel有如下两个操作:

· 将thread索引以row-major写到2D的shared memory数组中;

· 从shared memory中读取这些值并写入到global memory中。

kernel代码如下:
__global__ void setRowReadRow(int *out) {
// declare static 2D shared memory
__shared__ int tile[BDIMY][BDIMX];
// 因为block只有一个
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
// shared memory store operation
tile[threadIdx.y][threadIdx.x] = idx;
// 这里同步是为了使下面shared memory的获取以row-major执行
// 避免若有的线程未完成,而其他线程已经在读shared memory的情况
__syncthreads();
// shared memory load operation
out[idx] = tile[threadIdx.y][threadIdx.x] ;
}
此段有三个内存操作:(1)向shared Memory存数据;(2)从shared Memor取数据;(3)向global Memory存数据。
因为在同一个warp中的thread使用连续的threadIdx.x来检索title,该kernel是没有bank conflict的。如果交换上述代码threadIdx.y和threadIdx.x的位置,就变成了column-major的顺序。每个shared memory的读写都会导致Fermi上32-way的bank conflict或者在Kepler上16-way的bank conflict。
__global__ void setColReadCol(int *out) {
// static shared memor
__shared__ int tile[BDIMX][BDIMY];
// mapping from thread index to global memory index
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
// shared memory store operation
tile[threadIdx.x][threadIdx.y] = idx;
// wait for all threads to complete
__syncthreads();
// shared memory load operation
out[idx] = tile[threadIdx.x][threadIdx.y];
}
编译运行结果如下(在K40上以4-byte模式运行):
$ nvcc checkSmemSquare.cu o smemSquare
$ nvprof .
/smemSquare
./smemSquare at device 0 of Tesla K40c with Bank Mode:4-byte
<<< grid (1,1) block (32,32)>>
Time(
%) Time Calls Avg Min Max Name
13.25% 2.6880us 1 2.6880us 2.6880us 2.6880us setColReadCol(int*)
11.36% 2.3040us 1 2.3040us 2.3040us 2.3040us setRowReadRow(int*)
从结果可以看出,row-major的kernel表现更出色。
然后使用nvprof命令的下面的两个参数来衡量相应的bank-conflict:

shared_load_transactions_per_request shared_store_transactions_per_request

运行结果如下(K40,8-byte模式下),row-major只有一次transaction,而column-major需要16次,如果4-byte模式下,可能需要32次:
Kernel:setColReadCol (int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 16.000000
Kernel:setRowReadRow(
int*)
1 shared_load_transactions_per_request 1.000000
1 shared_store_transactions_per_request 1.000000
Writing Row
-Major and Reading Column-Major
(2) Writing Row-Major and Reading Column-Major
下面代码实现以row-major写shared memory,以column-major读shared memory:
__global__ void setRowReadCol(int *out) {
// static shared memory
__shared__ int tile[BDIMY][BDIMX];
// mapping from thread index to global memory index
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
// shared memory store operation
tile[threadIdx.y][threadIdx.x] = idx;
// wait for all threads to complete
__syncthreads();
// shared memory load operation
out[idx] = tile[threadIdx.x][threadIdx.y];
}
下图展示了用简单的5路bank shared memory实现两种内存操作:
CUDA性能优化----Shared Memory - 樂不思蜀 - 樂不思蜀
 用nvprof命令查看相关bank conflict情况:
Kernel:setRowReadCol (int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 1.000000
从结果可以看出:写操作是没有conflict的,读操作则引起了一个16次的transaction。

Dynamic Shared Memory

正如前文所说,我们可以全局范围的动态声明shared Memory,也可以在kernel内部动态声明一个局部范围的shared Memory。注意,动态声明必须是未确定大小一维数组,因此,我们就需要重新计算索引。因为我们将要以row-major写,以colu-major读,所以就需要保持下面两个索引值:

· row_idx:1D row-major 内存的偏移

· col_idx:1D column-major内存偏移

kernel代码:

复制代码
__global__ void setRowReadColDyn(int *out) {
    // dynamic shared memory
    extern __shared__ int tile[];
    // mapping from thread index to global memory index
    unsigned int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned int col_idx = threadIdx.x * blockDim.y + threadIdx.y;
    // shared memory store operation
    tile[row_idx] = row_idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[row_idx] = tile[col_idx];
}            
复制代码

kernel调用时配置的shared Memory:

setRowReadColDyn<<<grid, block, BDIMX * BDIMY * sizeof(int)>>>(d_C);

查看transaction:

Kernel: setRowReadColDyn(int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 1.000000

该结果和之前的例子相同,不过这里使用的是动态声明。

Padding Statically Declared Shared Memory

直接看kernel代码:

复制代码
__global__ void setRowReadColPad(int *out) {
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX+IPAD];
    // mapping from thread index to global memory offset
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}                            
复制代码

改代码是setRowReadCol的翻版,查看结果:

Kernel: setRowReadColPad(int*)
1 shared_load_transactions_per_request 1.000000
1 shared_store_transactions_per_request 1.000000

正如期望的那样,load的bank_conflict已经消失。在Fermi上,只需要加上一列就可以解决bank-conflict,但是在Kepler上却不一定,这取决于2D shared Memory的大小,因此对于8-byte模式,可能需要多次试验才能得到正确结果。


原文链接:http://blog.163.com/wujiaxing009@126/blog/static/71988399201712735436357/

                https://www.cnblogs.com/1024incn/p/4605502.html

阅读更多
想对作者说点什么? 我来说一句

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