5.4 本地内存

本地内存包括CUDA内核中每一个线程的栈，在实现以下的目的中使用：

·实现应用程序二进制接口（ABI）——这是一个调用惯例。
·容纳寄存器溢出的数据。
·保存编译器不能解析其索引的数组。

在早期的CUDA硬件中，任何使用本地内存的操作都是一次“死亡之吻”。它使机器的速度大幅减慢，所以提倡开发者动用一切手段避免使用本地内存。随着费米架构引入一级缓存，只要本地内存传输仅限于在一级缓存中进行[1]，这一性能问题就变得不那么严重。

为了让编译器报告给定内核需要的本地内存数量，开发者可以使用nvcc选项：-Xptxas-v, abi=no。在运行时中，内核使用的本地内存数量可以使用函数查询：

cuFuncGetAttribute(CU FUNC_ATTRIBUTE_LOCAL_SIZEBytes)

英伟达的Paulius Micikevicius介绍过一个关于确定本地内存是否对性能产生了影响和解决这个问题的方法 [2]。寄存器溢出会导致

两项开销：指令数增加和内存传输数量增加。

一级和二级缓存性能计数器可以用来确定内存传输是否影响到了性能。这里有一些在这种情况下提升性能的策略。

·在编译时，对-maxregcount指定一个较高的限制。通过增加线程中可用的寄存器数量，指令数和内存传输数量均会减少。当内核使用PTXAS在线编译时，_launch_bounds_指令可以用来调整这一参数。

·为全局内存使用无缓存加载，例如nvcc-Xptxas-dlcm=cg。

· 增加一级缓存到48KB。（调用函数CUDAFuncSetCacheConfig()或 CUDADeviceSetCache config）。）

当启动一个内核，如果使用超过了默认的本地内存分配数量，那么在内核可以启动前，CUDA驱动程序必须分配一块新的本地内存缓冲区。因此，内核启动可能需要额外的时间；可能会导致预想不到的GPU/CPU同步；并且，如果驱动程序不能为本地内存分配缓冲区，内核启动会以失败告终[3]。默认的，CUDA驱动程序会在内核启动后释放这些分配的本地内存。这一行为可以在函数cuCtxCreate()中指定标志CUCTX_RESIZE_LMEM_TO_MAX予以禁止，或者同样可以使用标志CUDADeviceLmemResizeToMax调用函数CUDASetDeviceFlags()达到同样目的。

建立一个函数模板描述寄存器溢出时发生的“性能突然退化”并不困难。代码清单5-10中的内核模板GlobalCopy实现了一个简单的内存复制程序，使用了本地数组temp来中转全局内存引用。模板参数n指定了temp中的元素数量，这也决定了内存复制的内循环中执行加载/存储的数量。

如果我们快速的浏览由编译器提交的SASS微码，我们会证实，编译器可以在寄存器中保存temp直到n变得非常大。

代码清单5-10 GlobalCopy内核

template<class T, const int n> global void GlobalCopy(T *out, const T *in, size_t N) {
    T temp[n];
    size_t i;
    for (i = n*blockIdx.x*blockDim.x+threadIdx.x; i < N-n*blockDim.x*gridDim.x; i += n*blockDim.x*gridDim.x) {
        for (int j = 0; j < n; j++) {
            size_t index = i+j*blockDim.x;
            temp[j] = in[index];
        }
    for (int j = 0; j < n; j++) {
        size_t index = i+j*blockDim.x;
        out[index] = temp[j];
    }
} // to avoid the (index<N) conditional in the inner loop,
// we left off some work at the end
for (int j = 0; j < n; j++) {
    for (int j = 0; j < n; j++) {
        size_t index = i+j*blockDim.x;
        if (index<N) temp[j] = in(index);
    }
    for (int j = 0; j < n; j++) {
        size_t index = i+j*blockDim.x;
        if (index<N) out[index] = temp[j];
    }
}

代码清单5-11展示了在GK104 GPU上来自globalCopy.cu的输出的一部分摘要：64位操作数的复制性能。由寄存器溢出导致的性能退化在循环展开到12时变得尤为明显，这时，传输带宽由117GB/s锐减到不足90GB/s，并在循环展开到16时已退化到30GB/s以下。

代码清单5-11 globalCopy.cu输出（64位）

表5-9总结了循环展开中内核的寄存器和本地内存使用情况。复制的性能退化与本地内存的使用情况有关。在这种情况下，每一个线程总是在内层循环中溢出；推测来看，当其中的一些线程在溢出时，性能不会退化的太多（例如执行一个分支的代码路径）。

表5-9 globalCopy寄存器和本地内存使用情况

[1] 一级缓存在每个SM上独立存在，与共享内存有着同样的物理内存实现。
[2] http://bit.ly/ZAeHc5。
[3] 在运行时中，由于大多数的资源是预分配的，无法分配本地内存是少数几个可以使内核启动失败的情况。

5.4_本地内存

5.4 本地内存