CUDA ---- Shared Memory

时间 2019-11-14 标签 cuda shared memory

CUDA SHARED MEMORY

shared memory在以前的博文有些介绍，这部分会专门讲解其内容。在global Memory部分，数据对齐和连续是很重要的话题，当使用L1的时候，对齐问题能够忽略，可是非连续的获取内存依然会下降性能。依赖于算法本质，某些状况下，非连续访问是不可避免的。使用shared memory是另外一种提升性能的方式。程序员

GPU上的memory有两种：算法

· On-board memory编程

· On-chip memory数组

global memory就是一块很大的on-board memory，而且有很高的latency。而shared memory正好相反，是一块很小，低延迟的on-chip memory，比global memory拥有高得多的带宽。咱们能够把他当作可编程的cache，其主要做用有：缓存

· An intra-block thread communication channel 线程间交流通道架构

· A program-managed cache for global memory data可编程cacheapp

· Scratch pad memory for transforming data to improve global memory access patterns函数

本文主要涉及两个例子做解释：reduction kernel，matrix transpose kernel。oop

shared memory（SMEM）是GPU的重要组成之一。物理上，每一个SM包含一个当前正在执行的block中全部thread共享的低延迟的内存池。SMEM使得同一个block中的thread可以相互合做，重用on-chip数据，而且可以显著减小kernel须要的global memory带宽。因为APP能够直接显式的操做SMEM的内容，因此又被称为可编程缓存。性能

因为shared memory和L1要比L2和global memory更接近SM，shared memory的延迟比global memory低20到30倍，带宽大约高10倍。

当一个block开始执行时，GPU会分配其必定数量的shared memory，这个shared memory的地址空间会由block中的全部thread 共享。shared memory是划分给SM中驻留的全部block的，也是GPU的稀缺资源。因此，使用越多的shared memory，可以并行的active就越少。

关于Program-Managed Cache：在C语言编程里，循环（loop transformation）通常都使用cache来优化。在循环遍历的时候使用从新排列的迭代顺序能够很好利用cache局部性。在算法层面上，咱们须要手动调节循环来达到使人满意的空间局部性，同时还要考虑cache size。cache对于程序员来讲是透明的，编译器会处理全部的数据移动，咱们没有能力控制cache的行为。shared memory则是一个可编程可操做的cache，程序员能够彻底控制其行为。

Shared Memory Allocation

咱们能够动态或者静态的分配shared Memory，其声明便可以在kernel内部也能够做为全局变量。

其标识符为：__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的组织方式，以便研究其对性能的影响。

Memory Banks

为了得到高带宽，shared Memory被分红32（对应warp中的thread）个相等大小的内存块，他们能够被同时访问。不一样的CC版本，shared memory以不一样的模式映射到不一样的块（稍后详解）。若是warp访问shared Memory，对于每一个bank只访问很少于一个内存地址，那么只须要一次内存传输就能够了，不然须要屡次传输，所以会下降内存带宽的使用。

Bank Conflict

当多个地址请求落在同一个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。这样的话就会致使带宽利用率低。

下图是最优状况的访问图示：

下图一种随机访问，一样没有conflict：

下图则是某些thread访问到同一个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

Access Mode

根据不一样的CC版本，bank的配置也不一样，具体为：

· 4 bytes for devices of CC 2.x

· 8 bytes for devices of CC3.x

对于Fermi，一个bank是4bytes。每一个bank的带宽是32bits每两个cycle。连续的32位字映射到连续的bank中，也就是说，bank的索引和shared memory地址的映射关系以下：

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

下图是Fermi的地址映射关系，注意到，bank中每一个地址相差32，相邻的word分到不一样的bank中以便使warp可以得到更多的并行获取内存操做（获取连续内存时，连续地址分配到了不一样bank中）。

当同一个warp的两个thread要获取同一个地址（注意是同一个地址仍是同一个bank）的时候并不发生bank conflict。对于读操做，会用一次transaction得到结果后广播给全部请求，当写操做时，只有一个thread会真正去写，可是哪一个thread执行了写是没法知道的（undefined）。

在8bytes模式中，同理4bytes，连续的64-bits字会映射到连续的bank。每一个bank带宽是64bite/1个clock。其映射关系公式：

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

这里，若是两个thread访问同一个64-bit中的任意一个两个相邻word（1byte）也不会致使bank conflict，由于一次64-bit（bank带宽64bit/cycle）的读就能够知足请求了。也就是说，同等状况下，64-bit模式通常比32-bit模式更少碰到bank conflict。

下图是64-bit的关系图。尽管word0和word32都在bank0中，同时读这两个word也不会致使bank conflict（64-bit/cycle）：

下图是64-bit模式下，conflict-free的状况，每一个thread获取不一样的bank：

下图是另外一种conflict-free状况，两个thread或获取同一个bank中的word：

下图红色箭头是bank conflict发生的状况：

Memory Padding

memory padding是一种避免bank conflict的方法，以下图所示，全部的thread分别访问了bank0的五个不一样的word，这时就会致使bank conflict，咱们采起的方法就是在每N（bank数目）个word后面加一个word，这样就以下面右图那样，本来bank0的每一个word转移到了不一样的bank中，从而避免了bank conflict。

增长的这写word不会用来存储数据，其惟一的做用就是移动原始bank中的word，使用memory padding会致使block可得到shared memory中有用的数量减小。还有就是，要从新计算数组索引来获取正确的数据元素。

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。

Synchronization

由于shared Memory能够被同一个block中的不一样的thread同时访问，当同一个地址的值被多个thread修改就致使了inter-thread conflict，因此咱们须要同步操做。CUDA提供了两类block内部的同步操做，即：

· Barriers

· Memory fences

对于barrier，全部thread会等待其余thread到达barrier point；对于Memory fence，全部thread会阻塞到全部修改Memory的操做对其余thread可见，下面解释下CUDA须要同步的主要缘由：weakly-ordered。

Weakly-Ordered Memory Model

现代内存架构有很是宽松的内存模式，也就是意味着，Memory的获取没必要按照程序中的顺序来执行。CUDA采用了一种叫作weakly-ordered Memory model来获取更激进的编译器优化。

GPU thread写数据到不一样的Memory的顺序（好比shared Memory，global Memory，page-locked host memory或者另外一个device上的Memory）一样不必跟程序里面顺序呢相同。一个thread的读操做的顺序对其余thread可见时也可能与实际上执行写操做的thread顺序不一致。

为了显式的强制程序以一个确切的顺序运行，就须要用到fence和barrier。他们也是惟一能保证kernel对Memory有正确的行为的操做。

Explicit Barrier

同步操做在咱们以前的文章中也提到过很多，好比下面这个：

void __syncthreads();

__syncthreads就是做为一个barrier point起做用，block中的thread必须等待全部thread都到达这个point后才能继续下一步。这也保证了全部在这个point以前获取global Memory和shared Memory的操做对同一个block中全部thread可见。__syncthreads被用来协做同一个block中的thread。当一些thread获取Memory相同的地址时，就会致使潜在的问题（读后写，写后读，写后写）从而引发未定义行为状态，此时就可使用__syncthreads来避免这种状况。

使用__syncthreads要至关当心，只有在全部thread都会到达这个point时才能够调用这个同步，显而易见，若是同一个block中的某些thread永远都到达该点，那么程序将一直等下去，下面代码就是一种错误的使用方式：

if (threadID % 2 == 0) {
    __syncthreads();
    } else {
        __syncthreads();
}

Memory Fence

这种方式保证了任何在fence以前的Memory写操做对fence以后thread均可见，也就是，fence以前写完了，fence以后其它thread就都知道这块Memory写后的值了。fence的设置范围比较广，分为：block，grid和system。

能够经过下面的API来设置fence：

void __threadfence_block();

看名字就知道，这个函数是对应的block范围，也就是保证同一个block中thread在fence以前写完的值对block中其它的thread可见，不一样于barrier，该function不须要全部的thread都执行。

下面是grid范围的API，做用同理block范围，把上面的block换成grid就是了：

void __threadfence();

下面是system的，其范围针对整个系统，包括device和host：

void __threadfence_system();

Volatile Oualifier

声明一个使用global Memory或者shared Memory的变量，用volatile修饰符来修饰该变量的话，会组织编译器作一个该变量的cache的优化，使用该修饰符后，编译器就会认为该变量可能在某一时刻被别的thread改变，若是使用cache优化的话，获得的值就缺少时效，所以使用volatile强制每次都到global 或者shared Memory中去读取其绝对有效值。

CHECKING THE DATA LAYOUT OF SHARED MEMORY

该部分会试验一些使用shared Memory的例子，包括如下几个方面：

· 方阵vs矩阵数组

· Row-major vs column-major access

· 静态vs动态shared Memory声明

· 全局vs局部shared Memory

· Memory padding vs no Memory padding

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

· Mapping data elements across Memory banks

· Mapping from thread index to shared Memory offset

搞明白这两点，就能够掌握shared Memory的使用了，从而构建出牛逼的代码。

Square Shared Memory

下图展现了一个每一维度有32个元素并以row-major存储在shared Memory，图的最上方是该矩阵实际的一维存储图示，下方的逻辑的二维shared Memory：

咱们可使用下面的语句静态声明一个2D的shared Memory变量：

__shared__ int tile[N][N];

可使用下面的方式来数据，相邻的thread获取相邻的word：

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。

Accessing 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) {
    // static 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] ;
}

观察代码可知，咱们有三个内存操做：

· 向shared Memory存数据

· 从shared Memor取数据

· 向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];
}

编译运行：

$ nvcc checkSmemSquare.cu –o smemSquare
$ nvprof ./smemSquare

在Tesla K40c（4-byte模式）上的结果以下，正如咱们所想的，row-major表现要出色：

./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*)

而后使用nvprof的下面的两个参数来衡量相应的bank-conflict：

shared_load_transactions_per_request

shared_store_transactions_per_request

结果以下（8 bytes模式，4 bytes应该是32），row-major只有一次transaction：

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

本节的kernel实现以row-major写shared Memory，以Column-major读shared Memory，下图指明了这两种操做的实现：

kernel代码：

__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];
}

查看nvprof结果：

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模式，可能须要屡次试验才能获得正确结果。

参考书《professional cuda c programming》