3.3.cuda运行时API-内存的学习，pinnedmemory，内存效率问题

前言

杜老师推出的 tensorRT从零起步高性能部署 课程，之前有看过一遍，但是没有做笔记，很多东西也忘了。这次重新撸一遍，顺便记记笔记。

本次课程学习精简 CUDA 教程-内存模型，pinned memory，内存效率问题

课程大纲可看下面的思维导图

1. Memory

内存模型是 CUDA 中很重要的知识点，通常和高性能有关系。

主要理解 pinned memory(Host Memory)、global memory(Deivce Memory）、shared memory(Device memory)即可，其它不常用

图1-1 CUDA内存模型

下图展示了显卡中各种内存所处位置，这是一个不太严谨但是辅助理解的图

图1-2 显卡各种内存位置

其中 shared memory 为片上内存，global memory 为片外内存

显卡的内存即 GPU 内存和主机的内存即 CPU 内存在电脑主板中的位置如下

图1-3 电脑主板中的GPU和CPU内存

Host Memory 即主机内存如下图所示：

图1-4 Host Memory

对于整个 Host Memory 内存条而言，操作系统区分为两个大类（逻辑区分，物理上是同一个东西）：

1. Pageable Memory，可分页内存
2. Page lock Memory，页锁定内存

你可以理解为 Page lock Memory 是 VIP 房间，锁定给你一个人用。而 Pageable Memory 是普通房间，在酒店房间不够的时候，选择性的把你的房间腾出来给其他人交换用(交换到硬盘上)，这就可以容纳更多人了。造成房间很多的假象，代价是性能降低

基于前面的理解，我们总结如下：

1. pinned memory 具有锁定特性，是稳定不会被交换的（这点很重要，相当于每次去这个房间都一定能找到你）
2. pageable memory 没有锁定特性，对于第三方设备（比如GPU），去访问时，因为无法感知内存是否被交换，可能得不到正确的数据（每次去房间找，说不准你的房间被人交换了）
3. pageable memory 的性能比 pinned memory 差，很可能降低你程序的优先级然后把内存交换给别人用
4. pageable memory 策略能使用内存假象，实际是 8GB 但是可以使用 15GB，提高程序运行数量（不是速度）
5. pinned memory 太多，会导致操作系统整体性能降低（程序运行数量减少），8GB 就只能用 8GB
6. GPU 可以直接访问 pinned memory 而不能访问 pageable memory

不同的 Host Memory 数据传输到 GPU 上的方式不同，具体如下图所示：

图1-5 Pinned Memory和Pageable Memory数据传输到GPU

GPU 要访问 Pageable Memory 的数据必须通过 Pinned Memory，GPU 可以直接访问 Pinned Memory 数据，其轨迹主要是通过 PICE 接口到主板再到内存条。

内存方面知识点总结

原则：

1. GPU 可以直接访问 pinned memory，称之为 DMA(Direct Memory Access) 技术
2. 对于 GPU 访问而言，距离计算单元越近，效率越高，所以 Pinned Memory<Global Memory<Shared Memory
3. 代码中，由 new、malloc 分配的是 pageable Memory，由 cudaMallocHost 分配的是 Pinned Memory（C语言函数分配也行），由 cudaMalloc 分配的是 Global Memory
4. 尽量多用 Pinned Memory 储存 host 数据或者显式处理 Host 到 Device 时用 Pinned Memory 做缓存都是提高性能的关键

内存模型案例代码如下：

// CUDA运行时头文件
#include <cuda_runtime.h>

#include <stdio.h>
#include <string.h>

#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)

bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){
    if(code != cudaSuccess){    
        const char* err_name = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

int main(){

    int device_id = 0;
    checkRuntime(cudaSetDevice(device_id));

    float* memory_device = nullptr;
    checkRuntime(cudaMalloc(&memory_device, 100 * sizeof(float))); // pointer to device

    float* memory_host = new float[100];
    memory_host[2] = 520.25;
    checkRuntime(cudaMemcpy(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice)); // 返回的地址是开辟的device地址，存放在memory_device

    float* memory_page_locked = nullptr;
    checkRuntime(cudaMallocHost(&memory_page_locked, 100 * sizeof(float))); // 返回的地址是被开辟的pin memory的地址，存放在memory_page_locked
    checkRuntime(cudaMemcpy(memory_page_locked, memory_device, sizeof(float) * 100, cudaMemcpyDeviceToHost)); // 

    printf("%f\n", memory_page_locked[2]);
    checkRuntime(cudaFreeHost(memory_page_locked));
    delete [] memory_host;
    checkRuntime(cudaFree(memory_device)); 

    return 0;
}

运行效果如下：

图1-6 Memory案例

代码演示了操作内存分配和数据复制，整个流程可用下面的示意图来说明。先创建三个指针(mem_device、mem_host、mem_page_locked)，然后开辟三块空间，将 mem_host 的数据复制到 GPU 上，再将 GPU 上 mem_deivce 的数据复制回 mem_page_locked 上(绿色箭头代表拷贝方向)

图1-7 数据传输示意图

关于内存模型及其知识点(from 杜老师)

1.关于内存模型，请参照https://www.bilibili.com/video/BV1jX4y1w7Um

• 内存大局上分为
• 主机内存：Host Memory，也就是 CPU 内存，内存
• 设备内存：Device Memory，也就是 GPU 内存，显存
  • 设备内存又分为：
    • 全局内存(3)：Global Memory(√)
    • 寄存器内存(1)：Register Memory
    • 纹理内存(2)：Texture Memory
    • 共享内存(2)：Shared Memory(√)
    • 常量内存(2)：Constant Memory
    • 本地内存(3)：Local Memory
  • 只需要知道，谁距离计算芯片近，谁速度越快，空间越小，价格越贵
    • 清单的括号数字表示到计算芯片的距离

2.通过 cudaMalloc 分配 GPU 内存，分配到 setDevice 指定的当前设备上

3.通过 cudaMallocHost 分配 page locked memory，即 pinned memory，页锁定内存

• 页锁定内存是主机内存，CPU 可以直接访问
• 页锁定内存也可以被 GPU 直接访问，使用 DMA（Direct Memory Access）技术
• 注意这么做的性能会比较差，因为主机内存距离 GPU 太远，隔着 PCIE 等，不适合大量数据传输
• 页锁定内存是物理内存，过度使用会导致系统性能地低下（导致虚拟内存等一系列技术变慢）

4.cudaMemcpy

• 如果 host 不是页锁定内存，则：
• Device To Host 的过程，等价于：
  • pinned = cudaMallocHost
  • copy Device to Host
  • copy pinned to Host
  • free pinned
• Host To Device 的过程，等价于：
  • pinned = cudaMallocHost
  • copy Host to pinned
  • copy pinned to Device
  • free pinned
• 如果 host 是页锁定内存，则：
• Device To Host 的过程，等价于
  • copy Device to Host
• Host To Device 的过程，等价于
  • copy Host to Device

对于分配的内存一般来说建议先分配的先释放

checkRuntime(cudaFreeHost(memory_page_locked));
delete [] memory_host;
checkRuntime(cudaFree(memory_device));

使用 CUDA API 来分配内存的一般都有自己对应的释放内存方法；而使用 new 来分配的则使用 delete 来释放

总结

本次课程简单的学习了下各种内存，其中主机内存(即 CPU 内存)需要了解 pinned memory 以及 page locked memory 两大类，设备内存(即 GPU 内存)需要知道 Global memory 和 shared memory。同时页锁定内存可以直接被 GPU 访问(DMA技术)，cudaMallocHost 分配的是页锁定内存(pinned memory)，new、malloc 分配的是可分页内存(pageable memory)，cudaMallo 分配的是 Global Memory。
两大类，设备内存(即 GPU 内存)需要知道 Global memory 和 shared memory。同时页锁定内存可以直接被 GPU 访问(DMA技术)，cudaMallocHost 分配的是页锁定内存(pinned memory)，new、malloc 分配的是可分页内存(pageable memory)，cudaMallo 分配的是 Global Memory。