TensorRT笔记二（RuntimeAPI）

一、Runtime API与Driver API

• Driver API更为底层一些，需要做显式的device初始化，以及context、module的管理，但是使用更为灵活，提供一些Runtime API所没有的功能
• Runtime API是封装了Driver API的更高级更友好的API，且进行懒加载，第一个Runtime API调用时灰进行culnit初始化，第一个需要context的API调用时，会进行context关联并创建context并设置当前context，通过调用cuDevicePrimaryCtxRetain实现
• Runtime API更友好执行核函数，.cpp和.cu文件无缝对接
• Runtime API对应cuda_runtime.h和libcudart.so，随cuda toolkit发布；Driver API对应cuda.h和libcuda.so，随显卡驱动发布

二、CUDA内存模型

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
    - 距离芯片越近，速度越快，空间越小，价格越高(括号中数字表示到计算芯片的距离)
  

对于整个Host Memory，操作系统分为两个大类（逻辑分区）：

• Pageable Memory，可分页内存
• Page lock Memory（Pinned Memory），页锁定内存

2. Pinned Memory(Page lock Memory，页锁定内存)

• Pinned Memory具有锁定特性，稳定不会被交换
• Pageable Memory没有锁定特性，对于第三方设备（GPU）去访问时，因为无法感知内存是否被交换，可能得不到正确的数据
• Pageable Memory性能低于Pinned Memory，可能降低程序优先级然后把内存交换给别人用
• Pageable Memory策略能使用内存假象（虚拟内存），如实际8GB但可以使用16GB，提高程序运行数量
• GPU可以直接访问Pinned Memory，而不是Pageable Memory,通过使用DMA（Direct Memory Access）技术
• Pegaable Memory要传给GPU，先传给Pinned Memory，再到GPU

3.CUDA内存分配

• 对于GPU访问来说，距离计算单元越近，效率越高，故访问速度PinnedMemory < GlobalMemory < ShareMemory
• 通过new、malloc分配的是Pageable Memory
• 通过cudaMalloc分配的是GPU内存（GlobalMemory），分配到setdevice指定的设备上
• 通过cudaMallocHost分配的是Page Lock Memory（即PinnedMemory，页锁定内存）
• 尽量多使用PinnedMemory来储存Host数据，或者显式处理Host到Device时，用PinnedMemory做缓存，提高性能

4. CUDA内存复制

（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

三、stream-流

1.流介绍

• 流是一种基于context之上的任务管道抽象，可以当作一个队列，一个context可以创建n个流
• 流是异步控制的主要方式
• nullptr表示默认流，每个线程都有自己的默认流
• 指令发出后，流队列储存的是指令参数，不能加入队列后立即释放参数指针，否则会导致执行该指令时指针失效而出错

2.cudaStream

//使用流复制内存
int device_id = 0;
checkRuntime(cudaSetDevice(device_id));

cudaStream_t stream = nullptr;
checkRuntime(cudaStreamCreate(&stream));

//在GPU上开辟Pageable Memory空间
float* memory_device = nullptr;
checkRuntime(cudaMalloc(&memory_device, 100 * sizeof(float)));

//在CPU上开辟空间并且存放数据，并把数据复制到GPU
float* memory_host = new float[100];
memory_host[2] = 520.25;
//Async使用流
checkRuntime(cudaMemcpyAsync(memory_device, memory_host, 100 * sizeof(float)), cudaMemcpyHostToDevice, stream); //异步复制操作，主线程不需要等待复制结束才继续

//在CPU上开辟Pinned Memory空间，并将GPU上的数据复制回来
float* memory_page_locked = nullptr;
checkRuntime(cudaMallocHost(&memory_page_locked, 100 * sizeof(float)));
checkRuntime(cudaMemcpyAsync(memory_page_locked, memory_device, 100 * sizeof(float), cudaMemcpyDeviceToHost, stream));//异步复制操作，主线程不需要等待复制结束才继续

//通过cudaStreamSynchronize函数，等待stream中所有指令执行完毕，也就是队列为空
checkRuntime(cudaStreamSynchronize(stream));

//释放内存
checkRuntime(cudaFreeHost(memory_page_locked));
checkRuntime(cudaFree(memory_device));
checkRuntime(cudaStreamDestroy(stream));
delete [] memory_host;

3.Event

向stream中加入Event，监控是否到达某个检查点

• cudaEventCreate，创建事件
• cudaEventRecord，记录事件，即stream中加入某个事件，队列执行到该事件修改其状态
• cudaEventQuery，查询事件当前状态
• cudaEventElapsedTime，计算两个事件间执行时间间隔，如统计某些核函数执行时间
• cudaEventSynchronize，同步某个事件，等待事件到达
• cudaSreamWaitEvent，等待流中的某个事件