第六章 CUDA的内存组织

        本章将对于CUDA中的内存组织进行学习和了解。

6.1 CUDA的内存组织介绍

        现代计算机中的内存往往存在一种组织结构。在这种结构中，有多种类型的内存，每种内存均具有不同的容量和延迟。一般来说，延迟低(速度高)的内存容量小，延迟高(速度低)的内存容量大。当前被处理的数据一般处于低延迟、低容量的内存中；当前没有被处理但之后要被处理的大量数据一般存放在高延迟，高容量的内存中。相对于不分级内存，这种内存结构可以降低延迟，提高计算效率。
        GPU和CPU都有内存分级的设计。相对于CPU的编程来说，CUDA编程模型向程序员提供更多的控制权。从而对于CUDA编程而言，熟悉内存分级组织更为重要。
        表6.1列出了CUDA中的几种内存和他们的主要特征，这些特征包括物理位置、设备的访问权限、可见范围和变量的生命周期：

表6.1 CUDA中设备内存的分类与特征

内存类型	物理位置	访问权限	可见范围	生命周期
全局内存	在芯片外	可读可写	所有线程和主机端	由主机分配和释放
常量内存	在芯片外	仅可读	所有线程和主机端	由主机分配和释放
纹理和表面内存	在芯片外	一般仅可读	所有线程和主机端	由主机分配和释放
寄存器内存	在芯片内	可读可写	单个线程	所在线程
局部内存	在芯片外	可读可写	单个线程	所在线程
共享内存	在芯片内	可读可写	单个线程块	所在线程块

CUDA中的内存组织示意图如图所示：

        

6.2 CUDA中不同类型的内存

6.2.1 全局内存

         这里的全局内存(global memory)的含义是核函数中的所有线程都能够访问其中的数据，和C++中的全局变量不是一个东西。已经使用过这种内存，在数组相加的例子中，使用cudaMalloc()分配的d_x和d_y等都是指向的全局内存。全局内存由于没有存放在GPU的芯片上，具有较高的延迟和较低的访问速度，但是其容量是设备内存中是最大的，其容量基本上就是显存容量。常见GPU的显存容量如下所示
        

 

         全局内存的主要角色是为核函数提供数据，并在主机与设备之间传递数据。首先用cudaMalloc()函数为全局内存变量分配设备内存。然后可以直接在核函数中访问分配的内存，改变其中的数据。要尽量减少主机和设备之间的数据传输，但有时是不可避免的。可以使用cudaMemcpy()函数将主机数据复制到全局内存，或者反过来，例:
        cudaMemcpy(d_x,h_x,m,cudaMemcpyHostToDevice);
将M字节的数据从主机复制到设备中去，而语句：
        cudaMemcpy(h_x,d_x,m,cudaMemcpyDeviceToHost);
将m个字节数据从设备复制到主机内存中去。
        cudaMemcpy(d_x,d_y,m,cudaMemcpyDeviceToDevice);
将m个字节从d_y的全局内存复制到d_x的全局内存中去。注意这里必须将数据传输的方向定义为cudaMemcpyDeviceToDevice或cudaMemcpyDefault.
        全局内存可读可写，在数组相加的例子中，语句：
        d_z[n]=d_x[n]+d_y[n].
同时体现了全局内存的可读性和可写性。对于线程n而言，该语句将变量d_x和d_y所指全局缓存中的第n个元素读出，相加后写入变量d_z中所指的全局缓存中第n个元素。
        全局内存对整个网格的所有线程可见。也就是说一个网格中的所有线程均可以访问(读或写)传入核函数的设备指针所指向的全局内存中的全部数据。在上面的语句中，第n个线程刚好访问全局缓存区的第n个元素，也可以访问第一个元素。
        全局内存的生命周期不由核函数决定，由主机端决定，生命周期开始于cudaMalloc()而结束与cudaFree(),在此期间，可以由任意相同或不同的核函数多次访问这些全局内存中的数据。
        在处理逻辑上的两维和三维问题上时可以用cudaMallocPitch()和cudaMalloc3D()函数分配内存，用cudaMemcpy2D()和cudaMemcpy3D复制数据，释放时用cudaFree()释放数据。
        以上所有全局内存都是现行内存。CUDA中还有一种内部构造对用户不透明的全局内存，称为CUDA Array。专门用于纹理拾取服务。
        前面介绍的全局内存都是动态的分配内存。在CUDA中允许使用静态全局内存变量，其所占的内存数量是在编译期间就确定的，且必须在设备和主机外部定义，所以还是一种圈觉得静态全局内存变量，这里，第一个全局的含义与C++中的全局变量含义相同，是指从定义开始，一个翻译单元内的所有设备函数均可见。若采用分离编译，可见范围还将进一步扩大。
        静态全局内存变量可由以下方式在任何函数外部定义：
        __device__ T x;//单个变量
        __device__ T y[n];//固定长度数组
其中修饰符__device__说明该变量是设备中的变量，而不是主机中的变量，T是变量的类型；N是一个整形常数。下面代码展示了静态全局内存变量的使用形式
 

#include "error.cuh"
#include<stdio.h>
__device__ int d_x=1;
__device__ int d_y[2];

void __global__ my_kernel(void)
{
    d_y[0]+=d_x;
    d_y[1]+=d_x;
    printf("d_x = %d,d_y[0] = %d,d_y[1]=%d.\n",d_x,d_y[0],d_y[1]);

}
int main(void)
{
    int h_y[2]={10,20};
    CHECK(cudaMemcpyToSymbol(d_y,h_y,sizeof(int)*2));

    my_kernel<<<1,1>>>();
    CHECK(cudaDeviceSynchronize());

    CHECK(cudaMemcpyFromSymbol(h_y,d_y,sizof(int)*2));
    printf("h_y[0]=%d,h_y[1]=%d.\n",h_y[0],h_y[1]);

    return 0;
}

程序将输出:d_x=1,d_y[0]=11,d_y[1]=21.
                   h_y[0]=11,h_y[1]=21
在核函数中可以直接对静态全局内存变量进行访问，并不需要将他们以参数的形式传给核函数。不可在主机函数中访问静态全局内存变量，但可以用cudaMemcpyToSymbol()函数和cudaMemcpyFromSymbol()函数在静态全局内存与主机内存之间传递数据。这两个CUDA运行时API函数的原型如下：
        cudaMemcpyToSymbol( void *symbol,const void *src,size_t count,size_t offset=0,cudaMemcpyKind kink=cudaMemcpyHostToDevice);
            cudaMemcpyFromSymbol( void *src,const void *symbol,size_t count,size_t offset=0,cudaMemcpyKind kink=cudaMemcpyDeviceToHost);
        这两个函数的symbol即可以使静态全局内存变量也可以是常量内存变量的变量名。   在16行中调用cudaMemcpyToSymbol()将全局内存数据h_y复制到静态全局内存数组d_y中，第21行调用cudaMemcpyFromSymbol()函数将静态全局变量复制到主机数组h_y中。

6.2.2 常量内存

        常量内存是有常量缓存的全局内存，数量有限，仅有64KB，可见范围和生命周期与全局内存一样。不同的是常量内存仅可读、不可写。由于有缓存，常量内存的访问速度比全局内存高，但要得到高访问速度的前提是一个线程束中的线程(一个线程块中相邻的32个线程)要读取相同的常量内存数据。
        一个使用常量内存的方法是在核函数外面用__constant__定义变量，并用前面介绍的CUDA运行时函数cudaMemcpyToSymbol()将数据从主机段复制到常量内存后供核函数使用。当计算能力不小于2.0之时，给核函数传递的参数(传值，不是想全局变量那样传递指针)就存放在常量内存中，但给核函数传递参数最多只能在一个核函数中使用4KB常量内存。
        所以其实已经使用过了常量内存。在数组相加的例子中，核函数中的参数const int N就是主机端定义的变量，并通过传值的方式传送给核函数中的线程使用。在核函数的代码段if中，这个参数N就被每一个线程使用了。所以核函数的每一个线程都知道该变量的值，并且对他的访问比对全局内存的访问要快。除给核函数传递单个的变量外，还可以传递结构体，同样也是使用常量内存。结构体中可以定义单个的变量，也可以定义固定长度的数组。

6.2.3 纹理内存和表面内存

        纹理内存和表面内存类似于常量内存，也是一种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读(表面内存也可写)。不同的是，纹理内存和表面内存容量更大，使用方式也和常量内存不一样。
        对于计算能力不小于3.5的GPU而言，将某些只读内存数据用__ldg()函数通过只读数据缓存读取，既可以达到使用纹理内存加速效果，也可以使代码简洁，该函数的原型为
        T __ldg(const T*address);
其中T是需要读取的数据的类型；address是数据的地址。

6.2.4 寄存器

        在核函数中定义的不加任何限定符的变量一般来说存放于寄存器中。核函数中定义的不加任何限定符的数组可能存放于寄存器中，但也有可能放在局部内存中。另外以前提到的各种内建变量如gridDim、blockDim、blockIdx、threadIdx、warpSize都保存在特殊的寄存器中。在核函数中访问这些内建变量是很高效的。
        我们已经使用过寄存器变量，在数组求和的例子中，我们在核函数中有如下const int n=threaIdx.x+blockDim.x*blockIdx.x;
这里的n就是一个寄存器变量。寄存器变量可读可写。上述语句的作用就是定义一个寄存器变量n并将其赋值为threadIdx.x+blockDim.x*blockIdx.x.在稍后的语句中z[n]=x[n]+y[n]中，寄存器变量n被读出。
        寄存器变量仅仅对于一个线程可见。也就说是说每一个线程都有一个变量n的副本。虽然在核函数中用了同一个变量名，但是不同的线程中该寄存器变量的值可以是不同的。每个线程都只能对他的副本进行读写。寄存器的生命周期也与所属线程的生命周期一致，从定义开始，直到线程消失时结束。
        寄存器内存在芯片上(on-chip)，是所有内存访问中最快的，但是数量有限，一个寄存器占有32b(4字节)的内存。所以，一个双精度浮点数将使用两个寄存器，这是估计估计寄存器数量的使用量需要注意的。

6.2.5 局部内存

        当前还没有使用过局部内存，从用法上来看，局部内存和寄存器一样，在核函数中不加任何限定符的变量可能存在寄存器也可能在局部内存。寄存器放不下以及索引不在编译时就确定的数组，都有可能存放在局部内存中。这种判断是由编译器来做的。对于数组相加的const int n这个变量一定存在于寄存器中，因为核函数所用数量远远没有达到上限。
        虽然局部内存用法上类似于寄存器，但从硬件上来讲，局部内存是全局内存的一部分。所以局部内存的延迟也很高，每个线程最多使用512KB的局部内存，过多影响程序性能。

6.2.6 共享内存

        还并没有使用过共享内存。共享内存与寄存器相似，存在于芯片上，具有仅次于寄存器的读写速度，数量也有限。
        不同于寄存器的是，共享内存对整个线程可见，其生命周期也与整个线程块一致，也就是说每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块中不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能访问其他线程块的共享内存变量副本，主要作用是减少全局内存的访问，或者改变访问模式。

6.2.7 L1和L2缓存

        从费米架构开始，有了SM层次的L1缓存(一级缓存)和设备(一个设备有多个SM)层次的L2缓存(二级缓存)。他们主要用来缓存全局内存和局部内存的访问，减少延迟。
        从硬件的角度来看，开普勒架构中的L1缓存和共享内存使用同一块物理片；麦克斯韦和帕斯卡架构中，L1缓存、纹理缓存统一起来，而共享内存是独立的；伏特架构和图灵架构中，L1缓存、纹理缓存与共享内存三者统一。从编程的角度来看，共享内存是可编程的缓存，L1和L2是不可编程的。
        对某些架构而言，可针对单个核函数或者整个程序改变L1缓存和共享内存的比例。

6.3 SM及其占有率

6.3.1 SM的构成

        在前一章讨论并行规模对CUDA程序性能的影响时提到了流多处理器SM，一个GPU由多个SM构成的。一个SM包含如下资源：
        (1)一定数量的寄存器。
        (2)一定数量的共享内存
        (3)常量内存的缓存。
        (4)纹理和表面内存的缓存。
        (5)L1缓存。
        (6)两个(计算能力为6.0)或四个线程束调度器，用于在不同的线程的上下文之间迅速切换，以及为准备就绪的线程束发出执行命令。
        (7)执行核心，包括
                1)若干整数运算核心INT32.
                2)若干单精度浮点数运算核心FP32.
                3)若干双精度浮点数运算的核心FP64。
                4)若干单精度浮点数运算核心 SFUs
                5)若干混合精度的张量核心 tensor cores

6.2.3 SM的占有率

        因为一个SM中的各种计算资源是有限的，那么有些情况下一个SM中驻留的线程数目就有可能达不到理想的最大值，此时说该SM的占有率小于100%，一般来说占有率不小于25%才能获得较高的性能。
        当并行规模较小的时候，有些SM可能没有被利用，占有率为零。一个线程块不论是几个维度，线程数不能超过1024.要分析SM的理论占有率还需要知道两个指标：
        (1)一个SM中最多拥有的线程块数目为Nb=16(开普勒和图灵)或者Nb=32(麦克斯韦、帕斯卡和伏特架构)。
        (2)一个SM中最多拥有的线程个数为Nt=2048(开普勒到伏特)或者Nt=1024(图灵架构)。
        以下通过并行规模足够大的前提下分析SM的理论占有率。
        (1)寄存器和共享内存使用量很少时，此时SM占有率完全由执行配置中的线程块大小决定。之前所使用的线程块大小是128，这是因为SM中线程的执行是以线程束为单位的，所以将线程块大小取为线程束大小的整数倍。若将线程块大小定义为100，那么线程块中将有三个完整的线程束和一个不完整的线程束。在执行核函数的指令时，不完整线程束所花时间与完整线程束一致，造成了计算浪费。所以建议将线程块大小blockSize取值为32的整数倍，在此前提下，任何不小于Nt/Nb且能整除Nt的线程块大小都能达到100%占有率，一般都习惯了128的线程块大小。
        (2)有限的寄存器数量对占有绿的约束情况，对于常见GPU架构(Tesla、GeForce)一个SM中最多能使用的寄存器个数是64K(1024*1024),若想一个SM中驻留最多的线程2048个，核函数中的线程最多只能使用32个寄存器。当每个线程中寄存器个数大于64时，SM占有率将小于50%；当每个线程寄存器数量使用个数大于128时，SM占有率小于25%。
        (3)有限的共享内存对占有率的约束情况，共享内存的数量随着计算能力的上升没有显著变化规律，所以此处针对3.5来进行讨论。若线程块大小为128，那么每个SM要激活16个线程块才能有2048个线程，达到100%占有率，此时一个线程块最多能使用3KB的共享内存。对于50%的占有率最多使用6kb的共享内存，对于要达到25%的线程，一个线程最多使用12KB的共享内存，若线程块使用了48KB的共享内存，核函数直接无法运行。
        以上单独分析了线程块大小、寄存器数量及共享内存对SM占有率的影响，在CUDA工具箱中，有一个名为CUDA_Occupany_Calculator.xls的EXCELL文档，可用来计算各种情况下SM的占有率。
        指的一提的是用编译器选项--ptxas-options=-v可以报道每个核函数的寄存器使用情况。CUDA还提供了核函数__launch_bounds__()修饰符和--maxrregcount=编译选项来让用户对一个核函数和所有核函数中寄存器的使用数量进行控制。

6.4 用CUDA运行时API函数查询设备

        在前面可以使用nvidia-smi程序对设备进行某些方面的查询与设置，本节介绍如何通过CUDA运行时API函数查询所用GPU的规格。

#include "error.cuh"
#include <stdio.h>

int main(int argc, char *argv[])
{
    int device_id = 0;
    if (argc > 1) device_id = atoi(argv[1]);
    CHECK(cudaSetDevice(device_id));

    cudaDeviceProp prop;
    CHECK(cudaGetDeviceProperties(&prop, device_id));

    printf("Device id:                                 %d\n",
        device_id);
    printf("Device name:                               %s\n",
        prop.name);
    printf("Compute capability:                        %d.%d\n",
        prop.major, prop.minor);
    printf("Amount of global memory:                   %g GB\n",
        prop.totalGlobalMem / (1024.0 * 1024 * 1024));
    printf("Amount of constant memory:                 %g KB\n",
        prop.totalConstMem  / 1024.0);
    printf("Maximum grid size:                         %d %d %d\n",
        prop.maxGridSize[0], 
        prop.maxGridSize[1], prop.maxGridSize[2]);
    printf("Maximum block size:                        %d %d %d\n",
        prop.maxThreadsDim[0], prop.maxThreadsDim[1], 
        prop.maxThreadsDim[2]);
    printf("Number of SMs:                             %d\n",
        prop.multiProcessorCount);
    printf("Maximum amount of shared memory per block: %g KB\n",
        prop.sharedMemPerBlock / 1024.0);
    printf("Maximum amount of shared memory per SM:    %g KB\n",
        prop.sharedMemPerMultiprocessor / 1024.0);
    printf("Maximum number of registers per block:     %d K\n",
        prop.regsPerBlock / 1024);
    printf("Maximum number of registers per SM:        %d K\n",
        prop.regsPerMultiprocessor / 1024);
    printf("Maximum number of threads per block:       %d\n",
        prop.maxThreadsPerBlock);
    printf("Maximum number of threads per SM:          %d\n",
        prop.maxThreadsPerMultiProcessor);

    return 0;
}

程序的第十行定义了一个CUDA定义好的cudaDeviceProp得变量prop，在11行利用cudaDeviceProperties()得到编号为device_id的设备的性质，存放在结构体prop中，得到以下输出

选择设备号为0的GPU，若有多块GPU可以选择其他设备，其中cudaSetDevice()是将设备初始化，其中在CUDA工具箱中，有一个名为deviceQuery.cu的程序，可以输出更多的信息。