第六天 CUDA内存管理

操作内存分配和数据复制过程概述

1. 在gpu上开辟一块空间，并把地址记录在mem_device上
2. 在cpu上开辟一块空间，并把地址记录在mem_host上，并修改了该地址所指区域的第二个值
3. 把mem_host所指区域的数据都复制到mem_device的所指区域
4. 在cpu上开辟一块空间,并把地址记录在mem_page_locked上
5. 最后把mem_device所指区域的数据又复制回cpu上的mem_page_locked区域

内存模型

1. 内存分为

    - 主机内存：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 pinned
       • 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来释放

// 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 devicefloat* memory_host = new float[100];memory_host[2] = 520.25;checkRuntime(cudaMemcpy(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice)); // 返回的地址是开辟的device地址，存放在memory_devicefloat* memory_page_locked = nullptr;checkRuntime(cudaMallocHost(&memory_page_locked, 100 * sizeof(float))); // 返回的地址是被开辟的pin memory的地址，存放在memory_page_lockedcheckRuntime(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;
}