铜川市建设集团网站,免费建站建站,汇赢网站建设,最新军事动态CUDA SHARED MEMORY
shared memory在之前的博文有些介绍#xff0c;这部分会专门讲解其内容。在global Memory部分#xff0c;数据对齐和连续是很重要的话题#xff0c;当使用L1的时候#xff0c;对齐问题可以忽略#xff0c;但是非连续的获取内存依然会降低性能。依赖于…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可编程cache
· Scratch pad memory for transforming data to improve global memory access patterns
本文主要涉及两个例子作解释reduction kernelmatrix transpose kernel。
shared memorySMEM是GPU的重要组成之一。物理上每个SM包含一个当前正在执行的block中所有thread共享的低延迟的内存池。SMEM使得同一个block中的thread能够相互合作重用on-chip数据并且能够显著减少kernel需要的global memory带宽。由于APP可以直接显式的操作SMEM的内容所以又被称为可编程缓存。
由于shared memory和L1要比L2和global memory更接近SMshared 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也就是最开始提及的第三个参数
kernelgrid, 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中的任意一个两个相邻word1byte也不会导致bank conflict因为一次64-bitbank带宽64bit/cycle的读就可以满足请求了。也就是说同等情况下64-bit模式一般比32-bit模式更少碰到bank conflict。
下图是64-bit的关系图。尽管word0和word32都在bank0中同时读这两个word也不会导致bank conflict64-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我们采取的方法就是在每Nbank数目个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 Memoryglobal Memorypage-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的设置范围比较广分为blockgrid和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的block32,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 operationtile[threadIdx.y][threadIdx.x] idx;// 这里同步是为了使下面shared Memory的获取以row-major执行//若有的线程未完成而其他线程已经在读shared Memory。。。__syncthreads();// shared memory load operationout[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 indexunsigned int idx threadIdx.y * blockDim.x threadIdx.x;// shared memory store operationtile[threadIdx.x][threadIdx.y] idx;// wait for all threads to complete__syncthreads();// shared memory load operationout[idx] tile[threadIdx.x][threadIdx.y];
} 编译运行
$ nvcc checkSmemSquare.cu –o smemSquare
$ nvprof ./smemSquare
在Tesla K40c4-byte模式上的结果如下正如我们所想的row-major表现要出色
./smemSquare at device 0 of Tesla K40c with Bank Mode:4-bytegrid (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应该是32row-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 indexunsigned int idx threadIdx.y * blockDim.x threadIdx.x;// shared memory store operationtile[threadIdx.y][threadIdx.x] idx;// wait for all threads to complete__syncthreads();// shared memory load operationout[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_idx1D row-major 内存的偏移
· col_idx1D column-major内存偏移
kernel代码 __global__ void setRowReadColDyn(int *out) {// dynamic shared memoryextern __shared__ int tile[];// mapping from thread index to global memory indexunsigned int row_idx threadIdx.y * blockDim.x threadIdx.x;unsigned int col_idx threadIdx.x * blockDim.y threadIdx.y;// shared memory store operationtile[row_idx] row_idx;// wait for all threads to complete__syncthreads();// shared memory load operationout[row_idx] tile[col_idx];
} kernel调用时配置的shared Memory
setRowReadColDyngrid, 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][BDIMXIPAD];// mapping from thread index to global memory offsetunsigned int idx threadIdx.y * blockDim.x threadIdx.x;// shared memory store operationtile[threadIdx.y][threadIdx.x] idx;// wait for all threads to complete__syncthreads();// shared memory load operationout[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模式可能需要多次试验才能得到正确结果。
