CUDA ---- Memory Model

灰太狼 2022-04-14 06:57 246阅读 0赞

## Memory ##

kernel性能高低是不能单纯的从warp的执行上来解释的。比如之前博文涉及到的，将block的维度设置为warp大小的一半会导致load efficiency降低，这个问题无法用warp的调度或者并行性来解释。根本原因是获取global memory的方式很差劲。

众所周知，memory的操作在讲求效率的语言中占有极重的地位。low-latency和high-bandwidth是高性能的理想情况。但是购买拥有大容量，高性能的memory是不现实的，或者不经济的。因此，我们就要尽量依靠软件层面来获取最优latency和bandwidth。CUDA将memory model unit分为device和host两个系统，充分暴露了其内存结构以供我们操作，给予用户充足的使用灵活性。

## Benefits of a Memory Hierarchy ##

一般来说，程序获取资源是有规律的，也就是计算机体系结构经常提到的局部原则。其又分为时间局部性和空间局部性。 相信大家对计算机内存方面的知识都很熟悉了，这里就不多说了，只简单提下。

![在这里插入图片描述][092148005358606.png]  
GPU和CPU的主存都是用DRAM实现，cache则是用lower-latency的SRAM来实现。GPU和CPU的存储结构基本一样。而且CUDA将memory结构更好的呈现给用户，从而能更灵活的控制程序行为。

## CUDA Memory Model ##

对于程序员来说，memory可以分为下面两类：

*  Programmable：我们可以灵活操作的部分。
 *  Non-programmable：不能操作，由一套自动机制来达到很好的性能。

在CPU的存储结构中，L1和L2 cache都是non-programmable的。对于CUDA来说，programmable的类型很丰富：

Registers  
Shared memory  
Local memory  
Constant memory  
Texture memory  
Global memory

下图展示了memory的结构，他们各自都有不用的空间、生命期和cache。  
![在这里插入图片描述][092148546606175.png]  
其中constant和texture是只读的。最下面这三个global、constant和texture拥有相同的生命周期。

## Registers ##

寄存器是GPU最快的memory，kernel中没有什么特殊声明的自动变量都是放在寄存器中的。当数组的索引是constant类型且在编译期能被确定的话，就是内置类型，数组也是放在寄存器中。

寄存器变量是每个线程私有的，一旦thread执行结束，寄存器变量就会失效。寄存器是稀有资源。在Fermi上，每个thread限制最多拥有63个register，Kepler则是255个。让自己的kernel使用较少的register就能够允许更多的block驻留在SM中，也就增加了Occupancy，提升了性能。

使用nvcc的\-Xptxas -v,-abi=no（这里Xptxas表示这个是要传给ptx的参数，不是nvcc的，v是verbose，abi忘了，好像是application by interface）选项可以查看每个thread使用的寄存器数量，shared memory和constant memory的大小。如果kernel使用的register超过硬件限制，这部分会使用local memory来代替register，即所谓的register spilling，我们应该尽量避免这种情况。编译器有相应策略来最小化register的使用并且避免register spilling。我们也可以在代码中显式的加上额外的信息来帮助编译器做优化：

__global__ void
    __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
    kernel(...) {
        // your kernel body
    }

**maxThreadsPerBlock**指明每个block可以包含的最大thread数目。**minBlocksPerMultiprocessor**是可选的参数，指明必要的最少的block数目。

我们也可以使用-maxrregcount=32来指定kernel使用的register最大数目。如果使用了\_\_ launch\_bounds \_\_，则这里指定的32将失效。

## Local Memory ##

有时候，如果register不够用了，那么就会使用local memory来代替这部分寄存器空间。除此外，下面几种情况，编译器可能会把变量放置在local memory：  
  
编译期无法决定确切值的本地数组。  
较大的结构体或者数组，也就是那些可能会消耗大量register的变量。  
任何超过寄存器限制的变量。

local memory这个名字是有歧义的：在local memory中的变量本质上跟global memory在同一块存储区。所以，local memory有很高的latency和较低的bandwidth。在CC2.0以上，GPU针对local memory会有L1（per-SM）和L2（per-device）两级cache。

## Shared memory ##

用\_\_ shared\_\_修饰符修饰的变量存放在shared memory。因为shared memory是on-chip的，他相比localMemory和global memory来说，拥有高的多bandwidth和低很多的latency。他的使用和CPU的L1cache非常类似，但是他是programmable的。

按惯例，像这类性能这么好的memory都是有限制的，shared memory是以block为单位分配的。我们必须非常小心的使用shared memory，否则会无意识的限制了active warp的数目。

不同于register，shared memory尽管在kernel里声明的，但是他的生命周期是伴随整个block，而不是单个thread。当该block执行完毕，他所拥有的资源就会被释放，重新分配给别的block。

shared memory是thread交流的基本方式。同一个block中的thread通过shared memory中的数据来相互合作。获取shared memory的数据前必须先用\_\_syncthreads()同步。L1 cache和shared memory使用相同的64KB on-chip memory，我们也可以使用下面的API来动态配置二者：

cudaError\_t cudaFuncSetCacheConfig(const void\* func, enum cudaFuncCachecacheConfig);

func是分配策略，可以使用下面几种：

cudaFuncCachePreferNone: no preference (default)

cudaFuncCachePreferShared: prefer 48KB shared memory and 16KB L1 cache

cudaFuncCachePreferL1: prefer 48KB L1 cache and 16KB shared memory

cudaFuncCachePreferEqual: Prefer equal size of L1 cache and shared memory, both 32KB

Fermi仅支持前三种配置，Kepler支持全部，注意，在Maxwell之后，L1被舍弃了，所以这64KB就完全属于shared Memory了，也就没有了上面这个分配一说。

## Constant Memory ##

Constant Memory驻留在device Memory，并且使用专用的constant cache（per-SM）。该Memory的声明应该以**connstant**修饰。constant的范围是全局的，针对所有kernel，对于所有CC其大小都是64KB。在同一个编译单元，constant对所有kernel可见。

kernel只能从constant Memory读取数据，因此其初始化必须在host端使用下面的function调用：

cudaError\_t cudaMemcpyToSymbol(const void\* symbol, const void\* src,size\_t count);

这个function拷贝src指向的count个byte到symbol的地址，symbol指向的是在device中的global或者constant Memory。

当一个warp中所有thread都从同一个Memory地址读取数据时，constant Memory表现最好。例如，计算公式中的系数。如果所有的thread从不同的地址读取数据，并且只读一次，那么constant Memory就不是很好的选择，因为一次读constant Memory操作会广播给所有thread知道。

## Texture Memory ##

texture Memory驻留在device Memory中，并且使用一个只读cache（per-SM）。texture Memory实际上也是global Memory在一块，但是他有自己专有的只读cache。这个cache在浮点运算很有用（具体还没弄懂）。texture Memory是针对2D空间局部性的优化策略，所以thread要获取2D数据就可以使用texture Memory来达到很高的性能，D3D编程中有两种重要的基本存储空间，其中一个就是texture。

## Global Memory ##

global Memory是空间最大，latency最高，GPU最基础的memory。“global”指明了其生命周期。任意SM都可以在整个程序的生命期中获取其状态。global中的变量既可以是静态也可以是动态声明。可以使用\_\_device\_\_修饰符来限定其属性。global memory的分配就是之前频繁使用的cudaMalloc，释放使用cudaFree。global memory驻留在devicememory，可以通过32-byte、64-byte或者128-byte三种格式传输。这些memory transaction必须是对齐的，也就是说首地址必须是32、64或者128的倍数。优化memory transaction对于性能提升至关重要。当warp执行memory load/store时，需要的transaction数量依赖于下面两个因素：

Distribution of memory address across the thread of that warp 就是前文的连续  
Alignment of memory address per transaction 对齐  
一般来说，所需求的transaction越多，潜在的不必要数据传输就越多，从而导致throughput efficiency降低。

对于一个既定的warp memory请求，transaction的数量和throughput efficiency是由CC版本决定的。对于CC1.0和1.1来说，对于global memory的获取是非常严格的。而1.1以上，由于cache的存在，获取要轻松的多。

## GPU Cache ##

跟CPU的cache一样，GPU cache也是non-programmable的。在GPU上包含以下几种cache，在前文都已经提到：  
  
L1  
L2  
Read-only constant  
Read-only texture

每个SM都有一个L1 cache，所有SM共享一个L2 cache。二者都是用来缓存local和global memory的，当然也包括register spilling的那部分。在Fermi GPus 和 Kepler K40或者之后的GPU，CUDA允许我们配置读操作的数据是否使用L1和L2或者只使用L2。

在CPU方面，memory的load/store都可以被cache。但是在GPU上，只有load操作会被cache，store则不会。

每个SM都有一个只读constant cache和texture cache来提升性能。

## CUDA Variable Declaration Summary ##

下表是之前介绍的几种memory的声明总结：

![在这里插入图片描述][092211035983208.png]  
![在这里插入图片描述][092211350827158.png]

## Static Global Memory ##

下面的代码介绍了怎样静态的声明global variable（之前的博文其实都是global variable）。大致过程就是，先声明了一个float全局变量，在checkGlobal-Variable中，该值被打印出来，随后，其值便被改变。在main中，这个值使用cudaMemcpyToSymbol来初始化。最终当全局变量被改变后，将值拷贝回host。

#include <cuda_runtime.h>
    #include <stdio.h>
    __device__ float devData;
    __global__ void checkGlobalVariable() {
        // display the original value
        printf("Device: the value of the global variable is %f\n",devData);
        // alter the value
        devData +=2.0f;
    }
    
    int main(void) {
        // initialize the global variable
        float value = 3.14f;
        cudaMemcpyToSymbol(devData, &value, sizeof(float));
        printf("Host: copied %f to the global variable\n", value);
        // invoke the kernel
        checkGlobalVariable <<<1, 1>>>();
        // copy the global variable back to the host
        cudaMemcpyFromSymbol(&value, devData, sizeof(float));
        printf("Host: the value changed by the kernel to %f\n", value);
        cudaDeviceReset();
        return EXIT_SUCCESS;
    }

编译运行：

$ nvcc -arch=sm_20 globalVariable.cu -o globalVariable
    $ ./globalVariable

熟悉了CUDA的基本思想后，不难明白，尽管host和device的代码是写在同一个源文件，但是他们的执行却在完全不同的两个世界，host不能直接访问device变量，反之亦然。

我们可能会反驳说，用下面的代码就能获得device的全局变量：

cudaMemcpyToSymbol(devD6ata, &value, sizeof(float));

但是，我们应该还注意到下面的几点：  
  
该函数是CUDA的runtime API，使用的GPU实现。  
devData在这儿只是个符号，不是device的变量地址。  
在kernel中，devData被用作变量。

而且，cudaMemcpy不能用&devData这种方式来传递变量,正如上面所说，devData只是个符号，取址这种操作本身就是错误的：

cudaMemcpy(&devData, &value, sizeof(float),cudaMemcpyHostToDevice); // It’s wrong!!!

不管怎样，CUDA还是为我们提供了，利用devData这种符号来获取变量地址的方式：

cudaError\_t cudaGetSymbolAddress(void\*\* devPtr, const void\* symbol);

获取地址之后，就可以使用cudaMemcpy了

float *dptr = NULL;
    cudaGetSymbolAddress((void**)&dptr, devData);
    cudaMemcpy(dptr, &value, sizeof(float), cudaMemcpyHostToDevice);

我们只有一种方式能够直接获取GPU memory，即使用pinned memory，下文将详细介绍。

## Memory Management ##

CUDA非常接近C的编程风格，以便能够快速上手掌握，在内存管理这点上，CUDA区别于C最明显的操作就是在device和host之间不停的传递数据。很麻烦的一个过程，不过Unified Memory出现后，程序编写就没那么复杂了，但是目前，Unified Memory的使用并未普及，我们还是要关注Memory的显式的操作过程：

*  Allocate and deallocate device Memory
 *  Transfer data between the host and device

为了达到最好的性能，CUDA提供了五花八门的接口供程序员显式的在device和host之间传递数据。

## Memory Allocation and Deallocation ##

前面的博文已经提到一部分内存分配函数了，在分配global Memory时，最常用的就是下面这个了：

cudaError\_t cudaMalloc(void \*\*devPtr, size\_t count);

如果分配出错则返回cudaErrorMemoryAllocation。分配成功后，就得对该地址初始化值，要么从host调用cudaMemcpy赋值，要么调用下面的API初始化：

cudaError\_t cudaMemset(void \*devPtr, int value, size\_t count);

释放资源就是：

cudaError\_t cudaFree(void \*devPtr);

device资源分配是个非常昂贵的操作，所以，device Memory应该尽可能的重用，而不是重新分配。

## Memory Transfer ##

一旦global Memory分配好后，如果不用cudaMemset就得用下面这个：

cudaError\_t cudaMemcpy(void \*dst, const void \*src, size\_t count,enum cudaMemcpyKind kind);

这个大家应该也很熟悉了，kind就是下面这几种：  
  
cudaMemcpyHostToHost  
cudaMemcpyHostToDevice  
cudaMemcpyDeviceToHost  
cudaMemcpyDeviceToDevice

下图是CPU和GPU之间传输关系图，可以看出来，CPU和GPU之间传输速度相对很差（NVLink技术能提高5~10倍），GPU和on-board Memory传输速度要快得多，所以对于编程来说，要时刻考虑减少CPU和GPU之间的数据传输。

![在这里插入图片描述][102323351601184.png]

## Pinned Memory ##

Host Memory的分配默认情况下是pageable的，也就是说，我们要承受因pagefault导致的操作，，这个操作要将host virtual Memory的数据转移到由OS决定的不物理位置。GPU无法安全的获取host的pageable Memory，因为GPU没有办法控制host OS物理上转移数据的时机。因此，当将pageable host Memory数据送到device时，CUDA驱动会首先分配一个临时的page-locked或者pinned host Memory，并将host的数据放到这个临时空间里。然后GPU从这个所谓的pinned Memory中获取数据，如下左图所示：

![在这里插入图片描述][102324218941406.png]  
左图是默认的过程，我们也可以显式的直接使用pinned Memory，如下：

cudaError\_t cudaMallocHost(void \*\*devPtr, size\_t count);

由于pinned Memory能够被device直接访问（不是指不通过PCIE了，而是相对左图我们少了pageable Memory到pinned Memory这一步），所以他比pageable Memory具有相当高的读写带宽，当然像这种东西依然不能过度使用，因为这会降低pageable Memory的数量，影响整个虚拟存储性能，我们不能因小失大。

cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes);
    if (status != cudaSuccess) {
        fprintf(stderr, "Error returned from pinned host memory allocation\n");
        exit(1);
    }

Pinned Memory的释放也比较特殊：

cudaError\_t cudaFreeHost(void \*ptr);

Pinned Memory比pageable Memory的分配操作更加昂贵，但是他对大数据的传输有很好的表现。还有就是，pinned Memory效果的高低也是跟CC有关的。

将许多小的传输合并到一次大的数据传输，并使用pinned Memory将降低很大的传输消耗。这里提及下，数据传输的消耗有时候是可以被kernel的执行覆盖的。

## Zero-Copy Memory ##

一般来说，host和device是不能直接访问对方的数据的，前文也有提到，但是Zero-Copy Memory是个特例。

该Memory是位于host的，但是GPU thread可以直接访问，其优点有

*  当device Memory不够用时，能够利用host Memory。
 *  避免device和host之间显式的数据传输。
 *  提高PCIe传输效率。  
    当使用zero-copy来共享host和device数据时，我们必须同步Memory的获取，否则，device和host同时访问该Memory会导致未定义行为。

Zero-copy本身实质就是pinned memory并且被映射到了device的地址空间。下面是他的分配API：  
cudaError\_t cudaHostAlloc(void \*\*pHost, size\_t count, unsigned int flags);

其资源释放当然也是cudaFreeHost，至于flag则是下面几个选项：

*  cudaHostAllocDefault
 *  cudaHostAllocPortable
 *  cudaHostAllocWriteCombined
 *  cudaHostAllocMapped  
    当使用cudaHostAllocDefault时，cudaHostAlloc和cudaMallocHost等价。cudaHostAllocPortable则说明，分配的pinned memory对所有CUDA context都有效，而不是单单执行分配此操作的那个context或者说线程。cudaHostAllocWriteCombined是在特殊系统配置情况下使用的，这块pinned memory在PCIE上的传输更快，但是对于host自己来说，却没什么效率。所以该选项一般用来让host去写，然后device读。最常用的是cudaHostAllocMapped，就是返回一个标准的zero-copy。可以用下面的API来获取device端的地址：

cudaError\_t cudaHostGetDevicePointer(void \*\*pDevice, void \*pHost, unsigned int flags);

flags是保留参数，留待将来使用，目前必须设置为零。

使用zero-copy memory来作为device memory的读写很频繁的那部分的补充是很不明智的，pinned这一类适合大数据传输，不适合频繁的操作，究其根本原因还是GPU和CPU之间低的可怜的传输速度，甚至，频繁读写情况下，zero-copy表现比global memory也要差不少。

int main(int argc, char **argv) {
    // part 0: set up device and array
    // set up device
    int dev = 0;
    cudaSetDevice(dev);
    // get device properties
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    // check if support mapped memory
    if (!deviceProp.canMapHostMemory) {
    printf("Device %d does not support mapping CPU host memory!\n", dev);
    cudaDeviceReset();
    exit(EXIT_SUCCESS);
    }
    printf("Using Device %d: %s ", dev, deviceProp.name);
    // set up date size of vectors
    int ipower = 10;
    if (argc>1) ipower = atoi(argv[1]);
    int nElem = 1<<ipower;
    size_t nBytes = nElem * sizeof(float);
    if (ipower < 18) {
    printf("Vector size %d power %d nbytes %3.0f KB\n", nElem,\
    ipower,(float)nBytes/(1024.0f));
    } else {
    printf("Vector size %d power %d nbytes %3.0f MB\n", nElem,\
    ipower,(float)nBytes/(1024.0f*1024.0f));
    }
    // part 1: using device memory
    // malloc host memory
    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A = (float *)malloc(nBytes);
    h_B = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef = (float *)malloc(nBytes);
    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);
    // add vector at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    // malloc device global memory
    float *d_A, *d_B, *d_C;
    cudaMalloc((float**)&d_A, nBytes);
    cudaMalloc((float**)&d_B, nBytes);
    cudaMalloc((float**)&d_C, nBytes);
    // transfer data from host to device
    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);
    // set up execution configuration
    int iLen = 512;
    dim3 block (iLen);
    dim3 grid ((nElem+block.x-1)/block.x);
    // invoke kernel at host side
    sumArrays <<<grid, block>>>(d_A, d_B, d_C, nElem);
    // copy kernel result back to host side
    cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
    // check device results
    checkResult(hostRef, gpuRef, nElem);
    // free device global memory
    cudaFree(d_A);
    cudaFree(d_B);
    free(h_A);
    free(h_B);
    // part 2: using zerocopy memory for array A and B
    // allocate zerocpy memory
    unsigned int flags = cudaHostAllocMapped;
    cudaHostAlloc((void **)&h_A, nBytes, flags);
    cudaHostAlloc((void **)&h_B, nBytes, flags);
    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);
    // pass the pointer to device
    cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0);
    cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0);
    // add at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    // execute kernel with zero copy memory
    sumArraysZeroCopy <<<grid, block>>>(d_A, d_B, d_C, nElem);
    // copy kernel result back to host side
    cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
    // check device results
    checkResult(hostRef, gpuRef, nElem);
    // free memory
    cudaFree(d_C);
    cudaFreeHost(h_A);
    cudaFreeHost(h_B);
    free(hostRef);
    free(gpuRef);
    // reset device
    cudaDeviceReset();
    return EXIT_SUCCESS;
    }

编译执行

$ nvcc -O3 -arch=sm_20 sumArrayZerocpy.cu -o sumZerocpy
    $ nvprof ./sumZerocpy
    Using Device 0: Tesla M2090 Vector size 1024 power 10 nbytes 4 KB
    Time(%) Time Calls Avg Min Max Name
    27.18% 3.7760us 1 3.7760us 3.7760us 3.7760us sumArraysZeroCopy
    11.80% 1.6390us 1 1.6390us 1.6390us 1.6390us sumArrays
    25.56% 3.5520us 3 1.1840us 1.0240us 1.5040us [CUDA memcpy HtoD]
    35.47% 4.9280us 2 2.4640us 2.4640us 2.4640us [CUDA memcpy DtoH]

下表是尝试不同数组长度后的结果：

./sumZerocopy

![在这里插入图片描述][102330017549909.png]

因此，对于共享host和device之间的一小块内存空间，zero-copy是很好的选择，因为他简化的编程而且提供了合理的性能。

## Unified Virtual Addressing ##

在CC2.0以上的设备支持一种新特性：Unified Virtual Addressing（UVA）。这个特性在CUDA4.0中首次介绍，并被64位Linux系统支持。如下图所示，在使用UVA的情况下，CPU和GPU使用同一块连续的地址空间：

![在这里插入图片描述][102330542697099.png]  
在UVA之前，我们需要分别管理指向host memory和device memory的指针。使用UVA之后，实际指向内存空间的指针对我们来说是透明的，我们看到的是同一块连续地址空间。

这样，使用cudaHostAlloc分配的pinned memory获得的地址对于device和host来说是通用的。我们可以直接在kernel里使用这个地址。回看前文，我们对于zero-copy的处理过程是：

*  分配已经映射到device的pinned memory。
 *  根据获得的host地址，获取device的映射地址。
 *  在kernel中使用该映射地址。  
    使用UVA之后，就没必要来获取device的映射地址了，直接使用一个地址就可以，如下代码所示：

// allocate zero-copy memory at the host side
    cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);
    cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped);
    // initialize data at the host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    // invoke the kernel with zero-copy memory
    sumArraysZeroCopy<<<grid, block>>>(h_A, h_B, d_C, nElem);

可以看到，cudaHostAlloc返回的指针直接就使用在了kernel里面，编译指令;

$ nvcc -O3 -arch=sm_20 sumArrayZerocpyUVA.cu -o sumArrayZerocpyUVA

修改后的代码执行效率和之前的效率是相差无几的，大家可以自己动手试试。

转自：[http://www.cnblogs.com/1024incn/p/4564726.html][http_www.cnblogs.com_1024incn_p_4564726.html]

[092148005358606.png]: /images/20220414/e6a01e90c1fd426ebbda59b614bf38d2.png
[092148546606175.png]: /images/20220414/e73f053089094203a94936f6664ce3dd.png
[092211035983208.png]: /images/20220414/29ed1ae023704a31b34097f57e216e90.png
[092211350827158.png]: /images/20220414/d28d3bfe37fa48deadf19850e257abd7.png
[102323351601184.png]: /images/20220414/de507454a67e4b33b032e862fcfd328c.png
[102324218941406.png]: /images/20220414/48c0bb68c6fb4145a7aaade41b9ac308.png
[102330017549909.png]: /images/20220414/2bb067453a95422cbda41c32058fb6af.png
[102330542697099.png]: /images/20220414/5d1f0c18c9034f8eb824dbfebb2550f4.png
[http_www.cnblogs.com_1024incn_p_4564726.html]: http://www.cnblogs.com/1024incn/p/4564726.html