cuda 内存模型

1.关于内存模型

请参照 https://www.bilibili.com/video/BV1jX4y1w7Um
内存大局上分为
主机内存：Host Memory，也就是CPU内存，内存
设备内存：Device Memory，也就是GPU内存，显存

对于整个Host Memory而言，操作系统将其区分为两个大类（这里指的逻辑分区即不是真实存在的，物理上是同一个东西）：
1.Pageable memory，可分页内存
2.Page lock memory/Pinned memory，页锁定内存
我们可以将Page lock memory理解为vip房间，锁定给你一个人用。而Pageable memory是普通房间，在酒店房间不够时，选择性得把你的房间腾出来给其他人交换用，你就到硬盘去待着，这就可以容纳更多人了。造成房间很多的假象，代价是性能降低。因为就是内存被挪到硬盘上去，性能肯定会降低呀。所以Pinned memory是具有锁定特性，是稳定不会被交换的。性能要比Pageable memory好，因为Pageable memory会降低你程序的优先级，把内存交换给别人用。 GPU可以直接访问Pinned memory而不能访问Pageable memory，因为Pageable memory没有锁定特性，对于第三方设备如GPU访问它会造成错误，每次去房间找时，说不准你的房间被人交换了

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

2.对于GPU访问而言，距离计算单元越近，计算效率越高，所以PinnedMemory（host memory）

SharedMemory离GPU最近，其实pageable memory是最远的，但因为GPU不能直接访问pageable memory，所以就比较起来也没有意义。

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

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

4.代码中，有new或者malloc分配的是pageable memory（host memory）；由cudaMallocHost分配的是PinnedMemory（host memory）；由cudaMalloc分配的是GlobalMemory（device memory）。

尽量多用PinnedMemory储存host数据，或者显式处理Host到Device时，可以用PinnedMemory做缓存即先传给PinnedMemory再传给GPU，都是提高性能的关键。因为如果直接用pageable memory传给device的话，在这中间它自己会去分配一次，这样不断分配释放就会很浪费时间。

5.cudaMemcpy

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

6.释放内存记住一个原则：先创建的后释放

使用cuda API来分配内存的一般都有自己对应的释放内存方法；而使用new来分配的使用delete来释放，具体见 7.代码示例

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

7.代码示例

下面代码做了的事情：
在gpu上开辟一块空间，并把地址记录在mem_device上
在cpu上开辟一块空间，并把地址记录在mem_host上，并修改了该地址所指区域的第二个值
把mem_host所指区域的数据都复制到mem_device的所指区域
在cpu上开辟一块空间,并把地址记录在mem_page_locked上
最后把mem_device所指区域的数据又复制回cpu上的mem_page_locked区域

// CUDA运行时头文件
#include 

#include 
#include 

#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)
//cuda返回值检查，即检查cuda函数的返回值，打印报错信息
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));

    //分配global memory（device memory）
    float* memory_device = nullptr;
    checkRuntime(cudaMalloc(&memory_device, 100 * sizeof(float))); // pointer to device

    float* memory_host = new float[100]; //pageable memory（host memory）
    memory_host[2] = 520.25;
    //把memory_host的数据复制到memory_device上
    checkRuntime(cudaMemcpy(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice)); // 返回的地址是开辟的device地址，存放在memory_device

    //pinned memory也叫page locked memory
    float* memory_page_locked = nullptr;
    //关于为什么cudaMallocHost传入的是指针的地址，可以参见这个链接：https://blog.csdn.net/gdjason/article/details/51123978
    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]);
    //释放掉page locked memory
    checkRuntime(cudaFreeHost(memory_page_locked));
    //释放掉pageable memory
    delete [] memory_host;
    //释放掉global memory
    checkRuntime(cudaFree(memory_device)); 

    return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48