榕树贷款GPU 硬件架构

榕树贷款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 执行第二个分配资源