前言

本课程来自于腾讯课堂《tensorRT从零起步高性能部署》，对一些比较重要的内容做一些记录，便于后续查阅和回顾。

概述

显卡，显卡驱动，nvcc，cuda driver，cudatoolkit，cudnn 到底是什么？

CUDA Driver 是与 GPU 沟通的驱动级别底层 API
对 DriverAPI 的理解，有利于理解后续的 RuntimeAPI
CUDA Driver 随显卡驱动发布，与 cudatoolkit 分开看
CUDA Driver 对应于 cuda.h 和 libcuda.so 文件

各种API的层级

关于 Context

手动管理的context，cuCtxCreate（手动管理，以堆栈方式push/pop）
自动管理的context，cuDevicePrimaryCtxRetain（自动管理，runtime api以此为基础）

关于内存

CPU内存（Host Memory）< Pageable Memory：可分页内存；Page-Locked Memory / Pinned Memory：页锁定内存 >
GPU内存（Device Memory）< Global Memory：全局内存；Shared Memory：共享内存 >

Driver API

cuInit 和返回值检查

cuInit：初始化驱动 API，如果不执行，则所有 API 都将返回错误，全局执行一次即可。

/* 
cuInit(int flags), 这里的flags目前必须给0;
    对于cuda的所有函数，必须先调用cuInit，否则其他API都会返回CUDA_ERROR_NOT_INITIALIZED
    https://docs.nvidia.com/cuda/archive/11.2.0/cuda-driver-api/group__CUDA__INITIALIZE.html
    */
CUresult code=cuInit(0);  //CUresult 类型：用于接收一些可能的错误代码
if(code != CUresult::CUDA_SUCCESS){
    const char* err_message = nullptr;
    cuGetErrorString(code, &err_message);    // 获取错误代码的字符串描述
    // cuGetErrorName (code, &err_message);  // 也可以直接获取错误代码的字符串
    printf("Initialize failed. code = %d, message = %s\n", code, err_message);
    return -1;
}

返回值检查：正确友好的检查cuda函数的返回值，有利于程序的组织结构
使得代码可读性更好，错误更容易发现

#define checkDriver(op)  __check_cuda_driver((op), #op, __FILE__, __LINE__)

bool __check_cuda_driver(CUresult code, const char* op, const char* file, int line){

    if(code != CUresult::CUDA_SUCCESS){    
        const char* err_name = nullptr;    
        const char* err_message = nullptr;  
        cuGetErrorName(code, &err_name);    
        cuGetErrorString(code, &err_message);   
        printf("%s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

CUcontext

context是一种上下文，关联对GPU的所有操作
context与一块显卡关联，一个显卡可以被多个context关联
每个线程都由一个栈结构储存context，栈顶是当前使用的context，对应有push、pop函数操作context的栈，所有api都以当前context为操作目标。

// 创建上下文
CUcontext context = nullptr;
CUdevice device = 0;
checkDriver(cuCtxCreate(&context, CU_CTX_SCHED_AUTO, device));

这样做便于对多个显卡的内存进行管理和使用，但是当频繁对同一个显卡进行操作的时候，显得非常的“丑陋”（每次操作都要携带显卡的 device 指针）。因此 runtimeAPI 自动使用 cuDevicePrimaryCtxRetain 对 context 进行“自动”管理。

Runtime API

Runtime API，与 Driver 最大的区别是懒加载
即，第一个 Runtime API 调用时，才会进行 cuInit 初始化
即，第一个需要 context 的 API 调用时，才会进行 context 关联并创建 context 和设置当前 context，调用 cuDevicePrimaryCtxRetain 实现
对应 cuda_runtime.h 和 libcudart.so
不提供直接管理 context 的 API
随 cuda toolkit 发布

Runtime API 依然需要返回值检查：

#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;
}

Runtime API 的基本特性：

// 获取显卡数量，这里 Runtime API 会自动执行 cuInit(0) 进行 cuda 初始化
// 但是不会自动创建 context，因为cudaGetDeviceCount是一个不需要context的函数
int device_count = 0;
checkRuntime(cudaGetDeviceCount(&device_count));

// Runtime 使用 setdevice 来控制当前上下文，当你要使用不同设备时使用不同的device id
// context是线程内作用的，其他线程不相关的，一个线程一个context stack
// 由于set device函数是“第一个执行的需要 context 的函数”，所以他会自己执行 cuDevicePrimaryCtxRetain 创建 context
int device_id = 0;
checkRuntime(cudaSetDevice(device_id));

Memory

主要理解 pinned memory、global memory、shared memory。其示意图如下：

Device Memory 即 GPU 显存，包括 Global memory 和 Shared memory。

Device Memory

Host Memory 即主机内存，包括 Pageable memory 和 Pinned memory。

Host Memory

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

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

内存方面的原则总结：

GPU可以直接访问pinned memory，称之为（DMA Direct Memory Access）
对于GPU访问而言，距离计算单元越近，效率越高，所以就访问速度而言：PinnedMemory < GlobalMemory < SharedMemory
代码中，由new、malloc分配的，是pageable memory，由 cudaMallocHost 分配的是PinnedMemory，由 cudaMalloc 分配的是GlobalMemory
尽量多用PinnedMemory储存host数据，或者显式处理Host到Device时，用PinnedMemory做缓存，都是提高性能的关键

Stream

流是一种基于context之上的任务管道抽象，一个context可以创建n个流
流是异步控制的主要方式
nullptr表示默认流，每个线程都有自己的默认流

// 创建流
cudaStream_t stream = nullptr;
checkRuntime(cudaStreamCreate(&stream));

// 异步复制操作，主线程不需要等待复制结束才继续
checkRuntime(cudaMemcpyAsync(memory_device, memory_host, sizeof(float) * 100, cudaMemcpyHostToDevice, stream));

// 等待同步流
checkRuntime(cudaStreamSynchronize(stream));

核函数 Kernel

核函数是 cuda 编程的关键
通过 xxx.cu 创建一个 cudac 程序文件，并把 cu 交给 nvcc 编译，才能识别 cuda 语法
__global__ 表示为核函数，由 host 调用。__device__ 表示为设备函数，由 device 调用
__host__ 表示为主机函数，由 host 调用。__shared__ 表示变量为共享变量
host 调用核函数：function<<<gridDim, blockDim, sharedMemorySize, stream>>>(args…);
只有 __global__ 修饰的函数才可以用 <<< >>> 的方式调用
调用核函数是传值的，不能传引用，可以传递类、结构体等，核函数可以是模板
核函数的执行，是异步的，也就是立即返回的
线程layout主要用到 blockDim、gridDim
核函数内访问线程索引主要用到 threadIdx、blockIdx、blockDim、gridDim 这些内置变量

TensorRT 基础

TensorRT 的核心是对模型算子的优化
TensorRT 需要在根据在目标 GPU 上实际运行的方式选择最优的算法和配置
TensorRT 生成的模型只能在特定条件下运行（相同 trt、cuda 版本，相同 GPU 型号）

Hello

包含的 c++ 头文件是 <NvInfer.h> 和 <NvInferRuntime.h>。

使用 createNetworkV2 并指定为 1 （表示显性 batch）
builder 、config 等指针释放，ptr->destroy()
markOutput 表示输出节点，addInput 表示输入节点
workspaceSize 是工作空间大小，layer 直接找 tensorrt 要 workspace 空间，为了内存复用。

// 日志类：通过该方式进行排查和调试
class TRTLogger : public nvinfer1::ILogger{};

//定义 builder
nvinfer1::IBuilder* builder = nvinfer1::createInferBuilder(logger);

//创建一个构建配置，指定Tensorrt应该如何优化模型，tensorrt生成的模型只能在特定配置下运行
nvinfer1::IBuilderConfig* config = builder->createBuilderConfig();

//创建网络定义
nvinfer1::INetworkDefinition* network = builder->createNetworkV2(1);

一定要记住，保存的模型只能适配编译时的 trt 版本、编译时指定的设备。也只能保证在这种配置下是最优的。如果用 trt 跨不同设备执行，有时候可以运行，但不是最优的，也不推荐。

inferrence

本文由困困鱼原创，采用 CC BY-NC-SA 4.0 许可协议，转载请注明出处。

Tensorrt学习笔记

前言

概述