榕树贷款L2为所有SM都能访问到,速度比全局内存块,所以为了提高速度有些小的数据可以缓存到L2上面;L1为SM内的数据,SM内的运算单元能够共享,但跨SM之间的L1不能相互访问。
L2 可以被显式的使用(cuda 11 ),去优化性能
榕树贷款Nvida示例: Cuda 11 L2 示例
榕树贷款共享内存(shared memory)
共享内存是片内内存,被 SM 独享,SM 内的块所共享。
共享内存是片内存储, 和 L1的速度相当
榕树贷款共享内存与L1的位置、速度极其类似,区别在于共享内存的控制与生命周期管理与L1不同,共享内存的使用受用户控制,L1受系统控制,shared memory更利于block之间数据交互。
Kernel <<< * , 线程, 共享内存>>>>,kernel 执行第三个分配资源
__global__ void staticReverse(int *d, int n)
__shared__ int s[64]; // 分配的共享内存
int t = threadIdx.x;
int tr = n-t-1;
s[t] = d[t];
__syncthreads();
d[t] = s[tr];
局部内存(local memory)
榕树贷款局部内存(local memory) 是线程独享的内存资源,线程之间不可以相互访问,硬件位置是off chip状态,所以访问速度跟全局内存一样。局部内存主要是用来解决当寄存器不足时的场景,即在线程申请的变量超过可用的寄存器大小时,会将变量存储在局部内存中。
寄存器(register)
榕树贷款寄存器(register)是线程能独立访问的资源,它所在的位置与局部内存不一样,是在片上(on chip)的存储,用来存储一些线程的暂存数据。寄存器的速度是访问中最快的,但是它的容量较小。以目前最新的Ampere架构的GA102为例,每个SM上的寄存器总量256KB,使用时被均分为了4块,且该寄存器块的64KB空间需要被warp中线程平均分配,所以在线程多的情况下,每个线程拿到的寄存器空间相当小。寄存器的分配对SM的占用率(occupancy)存在影响,可以通过CUDA Occupancy Calculator 计算比较,举例:如图当registers从32增加到128时,occupancy从100%降低到了33.0。 一般 Occupancy 越高,kernel 优化的越好。
Kernel <<< * , 线程>>>>,kernel 执行第二个分配资源