TensorRT课程笔记（2） | 暴躁小何的笔记

发布日期: 2024-05-31

作者: LSJune

文章字数: 3.4k

阅读时长: 13 分

阅读次数:

写在前面

本篇文章程序链接：https://github.com/Shao-junliang/TensorRT-Course-Notes/tree/master/cuda-runtime-api

1. 上下文

CUDA Runtime是封装了CUDA Driver的高级别更友好的API。
cudaruntime需要引入cudart这个so文件。
上下文管理：
- 使用cuDevicePrimaryCtxRetain为每个设备设置context，不再手工管理context，并且不提供直接管理context的API。
- 任何依赖CUcontext的API被调用时，会触发CUcontext的创建和对设备的绑定。
  - 此后任何API调用时，会以设备id为基准，调取绑定好的CUcontext。
  - 这种模式被称为懒加载模式，避免了手动维护CUcontext的麻烦。
cuda的状态返回值，都是cudaError_t类型，通过check宏捕获状态并处理是一种通用方式。
- 官方案例采用宏，并非本例程序里的函数加宏。
- 函数加宏具有更加好的便利性。

2. 内存

关于内存模型，请参考：https://www.bilibili.com/video/BV1jX4y1w7Um
- 内存大体可分为
  - 主机内存：Host Memory，也就是CPU内存，内存；
  - 设备内存：Device Memory，也就是GPU内存，显存；
    - 设备内存又分为（括号里的数字表示到计算芯片的距离）：
      - 寄存器内存(1)：Register Memory
      - 纹理内存(2)：Texture Memory
      - 共享内存(2)：Shared Memory
      - 常量内存(2)：Constant Memory
      - 全局内存(3)：Global Memory
      - 本地内存(3)：Local Memory
    - 距离计算芯片越近，计算速度就越快，空间越小，价格越贵。
通过cudaMalloc分配GPU内存，分配到setDevice指定的当前设备上。
通过cudaMallocHost分配page locked memory，即pinned memory，页锁定内存。
- 页锁定内存是主机内存，CPU可以直接访问。
- 页锁定内存也可以被GPU直接访问，使用DMA（Direct Memory Access）技术。
  - 注意这么做的性能会比较差，因为主机内存距离GPU太远，隔着PCIE等，不适合大量数据传输。
- 页锁定内存是物理内存，过度使用会导致系统性能低下（导致虚拟内存等一系列技术变慢）。

内存拷贝：cudaMemcpy

如果host不是页锁定内存，则：

Device To Host的过程，等价于：

pinned = cudaMallocHost // 创建锁页内存
copy Device to pinned	// 拷贝Device数据到pinned
copy pinned to Host		// 拷贝pinned数据到Host
free pinned				// 释放内存

Host To Device的过程，等价于：

pinned = cudaMallocHost	// 创建锁页内存
copy Host to pinned		// 拷贝Host数据到pinned
copy pinned to Device	// 拷贝pinned数据到Device
free pinned				// 释放内存

如果host是页锁定内存，则：
- Device To Host的过程，等价于：copy Device to Host
- Host To Device的过程，等价于：copy Host to Device

3. 流

stream是一个流句柄，可以当做是一个队列。
- cuda执行器从 stream 中一条条的读取指令并执行；如cudaMemcpyAsync函数等同于向stream这个队列中加入一个cudaMemcpy指令并排队；
- 使用到stream的函数，在向stream中加入指令后立即返回，并不会等待指令执行结束；
- 通过cudaStreamSynchronize函数，等待stream中所有指令执行完毕，也就是队列为空，即同步操作。
当使用stream时，注意：由于异步函数会立即返回，因此传递进入的参数要考虑其生命周期，应确认函数调用结束后再做释放。
还可以向stream中加入Event，用以监控是否到达了某个检查点，具体流程：
- cudaEventCreate，创建事件;
- cudaEventRecord，记录事件；在stream中加入某个事件，当队列执行到该事件后，修改其状态;
- cudaEventQuery，查询事件当前状态；
- cudaEventElapsedTime，计算两个事件之间经历的时间间隔，若要统计某些核函数执行时间，请使用这个函数，能够得到最准确的统计；
- cudaEventSynchronize，同步某个事件，等待事件到达；
- cudaStreamWaitEvent，等待流中的某个事件；
默认流，对于cudaMemcpy等同步函数，其等价于执行了：
- cudaMemcpyAsync(… 默认流) 加入队列；
- cudaStreamSynchronize(默认流) 等待执行完成；
- 默认流与当前设备上下文类似，是与当前设备进行的关联。因此，如果大量使用默认流，会导致性能低下。

4. 核函数

cu文件可以当做正常cpp写即可，它是cpp的超集，兼容支持cpp的所有特性。

cu文件中引入的一些新的符号和语法：
- __global__标记，核函数标记，调用方必须是host，返回值必须是void；
```
__global__ void kernel(const float* pdata, int ndata)

// 调用方式，"<<< >>>"参数类型是：<<<dim3 gridDim, dim3 blockDim, size_t bytesSharedMemorySize, cudaStream_t stream>>>
kernel<<<gridDim, blockDim, bytesSharedMemorySize, stream>>>(pdata, ndata);
```
  - 其中gridDim, blockDim, bytesSharedMemory, stream是线程layout参数；
  - gridDim 与 blockDim 都是 dim3 类型，dim3有构造函数dim3(int x, int y=1, int z=1)，因此当直接赋值为int时，实则定义了dim.x = value, dim.y = 1, dim.z = 1
  - 如果指定了stream，则把核函数加入到stream中异步执行；pdata和ndata则是核函数的函数调用参数；
  - 核函数调用参数必须传值，不能传引用，参数可以是类类型等；
  - 核函数执行时无论stream是否为nullptr，都将是异步执行；
  - 因此在核函数中进行printf操作时，必须进行等待，例如cudaDeviceSynchronize、或者cudaStreamSynchronize，否则将无法看到打印的信息。
- __device__标记，设备调用的函数，调用方必须是device。
- __host__标记，主机调用函数，调用方必须是主机；也可以__device__ __host__两个标记同时有，表明该函数可以设备也可以主机。
- __constant__标记，定义常量内存。
- __shared__标记，定义共享内存。
- 通过cudaPeekAtLastError/cudaGetLastError函数，捕获核函数是否出现错误或者异常。
- 内存索引的计算公式
```
int position = 0
for i in range(6):
    position *= dims[i]
    position += indexs[i]
```
buildin变量，内置变量。
- 所有核函数都可以访问，其取值由执行器维护和改变；
- gridDim[x, y, z]：网格维度，线程布局的大小，在核函数启动时指定；
- blockDim[x, y, z]：块维度，线程布局的大小，在核函数启动时指定；
- blockIdx[x, y, z]：块索引，对应最大值为gridDim，由执行器根据当前执行的线程进行赋值，核函数内访问时已经被配置好；
- threadIdx[x, y, z]：线程索引，对应最大值是blockDim，由执行器根据当前执行的线程进行赋值，核函数内访问时已经被配置好；
- Dim是固定的，启动后不会改变，并且是Idx的最大值；
- 每个Dim都具有x、y、z三个维度，分别以z、y、x为高低顺序。
thread, grid, block 和 threadIdx
- 如果假设一个Grid相当于一个立方体，如图所示:
  
  图为一个gridDim[2, 2, 2]的girds，每一个小立方体为一个block，图中的block为blockDim[6, 6, 3]，其中绿色位置的小方块则代表一个block中的thread；
- threadIdx表示的含义为thread的1D idx，所以先确定知道在第几个block里，再确定在这个block里的第几个thread。
- 一个核函数只能有一个grid，一个grid可以有很多个block，每个block可以有很多的线程。

5. layout

layout是设置核函数执行的线程数，包括最大线程数、每个块的最大线程数、以及warp的大小。
- maxGridSize对应gridDim的取值最大值；
- maxThreadsDim对应blockDim的取值最大值；
- warpSize对应线程束中的线程数量；
- maxThreadsPerBlock对应blockDim元素乘积最大值；
layout的4个主要变量的关系：
- gridDim是layout维度，其对应的索引是blockIdx，blockIdx最大值为 [0, gridDim-1]。
- blockDim是layout维度，其对应的索引是threadIdx，threadIdx最大值为 [0, blockDim-1]，blockDim维度乘积必须小于等于maxThreadsPerBlock。
- gridDim、blockDim为维度，启动核函数后是固定的；
- blockIdx、threadIdx为索引，启动核函数后，枚举每一个维度值，不同线程取值不同。
核函数启动时，<<<>>>的参数分别为：<<<gridDim, blockDim, shraed_memory_size, cudaStream_t>>>，shared_memory_size为共享内存（shared memory）大小。
对于一维数组时，通常只定义layout的x维度，若处理的是二维，则可以考虑定义x、y维度，如单通道图像。

6. 共享内存

sharedMemPerBlock 为block中最大可用的共享内存，可以利用共享内存使得 block 内的threads相互通信。
共享内存是片上内存，更靠近计算单元，因此比globalMem速度更快，通常可以充当缓存使用。
- 数据先读入到sharedMem，做各类计算时，使用sharedMem而非globalMem
demo_kernel<<<1, 1, 12, nullptr>>>()；其中第三个参数12，是指定动态共享内存dynamic_shared_memory的大小，单位是bytes，也就是说可以安全存放3个float。
- dynamic_shared_memory变量必须使用 extern __shared__开头；
- 并且定义为不确定大小的数组 []；
- 变量在函数外部与内部定义效果都一样；
- 其指针由cuda调度器执行时赋值；
static_shared_memory作为静态分配的共享内存。
- 不加extern，以__shared__开头；
- 定义时需要明确数组的大小；
- 静态分配的地址比动态分配的地址低；
动态共享变量，无论定义多少个，地址都一样。
静态共享变量，定义几个地址随之叠加。
如果配置的各类共享内存总和大于sharedMemPerBlock，则核函数执行出错，Invalid argument
- 不同类型的静态共享变量定义，其内存划分并不一定是连续的；
- 中间会有内存对齐策略，使得第一个和第二个变量之间可能存在空隙；
- 因此如果变量之间存在空隙，可能小于全部大小的共享内存时就会报错。

7. atomic

atomicAdd是原子加法，同类型原子操作有很多，如原子减法等；
输出的数组需要预先分配空间，如[cuda-runtime-api-9-atomic.cpp][]的：output_capacity；
最后统计结果的时候记得取min，int output_size = min(output_host[0], output_capacity)；
- 因为核函数中，add的次数可能会超过capacity大小，导致后续访问越界的发生；
输出的数组可能会乱序，即不同执行时刻结果不同。这是因为cuda是并行的，调度顺序不确定；
- 因此capacity一定要比预期的最大值大一些，才可能保证结果不会丢失；
- 可以通过储存index，然后在cpu上执行一次排序，实现有序输出；
此类型的代码实现是模型后处理的关键，是处理动态数组的关键。

float *parray = nullptr;
checkRuntime(cudaMalloc(&output_device, sizeof(float) + 1000 * 7 * sizeof(float)));
// parray 是这样的数据格式：[count, box1, box2, ……]，count 为 int，box1 为 yolo 检测结果的结构体；
// atomicAdd(parray, 1) 的操作相当于是 count += 1, 但是返回的是没有 +1 之前的 old_count;
int idx = atomicAdd(parray, 1);

8. 仿射变换

仿射变换+双线性插值，在CV场景下，解决图像预处理是非常合适的。如Yolo的letterbox，实则是边缘填充；CenterNet的居中对齐；
在仿射核函数里，循环的次数为dst.width * dst.height，以dst为参照集；因此，无论src多大，dst固定的话，计算量也是固定的。
仿射变换的公式推导。
- 旋转变换：
  $$
  \begin{bmatrix} x’ \\ y’ \end{bmatrix} = \begin{bmatrix} cos(\theta) & sin(\theta) \\ -sin(\theta) & cos(\theta) \end{bmatrix} \begin{bmatrix} x \\ y \end{bmatrix}
  $$
- 缩放变换：
  $$
  \begin{bmatrix} x’ \\ y’ \end{bmatrix} = \begin{bmatrix} scale_x & 0 \\ 0 & scale_y \end{bmatrix} \begin{bmatrix} x \\ y \end{bmatrix}
  $$
- 平移变换：
  $$
  \begin{bmatrix} x’ \\ y’ \end{bmatrix} = \begin{bmatrix} 1 & 0 \\ 0 & 1 \end{bmatrix} \begin{bmatrix} x \\ y \end{bmatrix} + \begin{bmatrix} \Delta x \\ \Delta y \end{bmatrix}
  $$
- 旋转 + 缩放：
  $$
  \begin{bmatrix} x’ \\ y’ \end{bmatrix} = \begin{bmatrix} cos(\theta)*scale_x & sin(\theta)*scale_y \\ -sin(\theta)*scale_x & cos(\theta)*scale_y \end{bmatrix} \begin{bmatrix} x \\ y \end{bmatrix}
  $$
- 旋转 + 缩放 + 平移：
  $$
  \begin{bmatrix} x’ \\ y’ \\ z’ \end{bmatrix} = \begin{bmatrix} cos(\theta)*scale_x & sin(\theta)*scale_y & \Delta x \\ -sin(\theta)*scale_x & cos(\theta)*scale_y & \Delta y \\ 0 & 0 & 1 \end{bmatrix} \begin{bmatrix} x \\ y \\ 1 \end{bmatrix}
  $$

9. thrust特性

thrust是cuda开发的，基于cuda的stl库。
对于thrust中的lambda表达式，需要增加__device__标记表明函数可以被核函数调用；由于使用到了device vector，因此编译环境需要修改为nvcc编译，且cpp文件要改为cu文件；此时需要在xmake中添加nvcc工具链，具体可参考自定义工具链
内存的复制和分配，被cuda封装。

10. error

由于cuda核函数是异步执行，如果cuda核函数出错，立即执行cudaPeekAtLastError时只会拿到对输入参数校验是否正确的状态，而不会拿到核函数是否执行正确的状态，所以此时拿到的状态码为：no error。
因此一般等待核函数执行完毕后，才可以知道当前核函数是否出错，一般通过设备同步或者流同步进行等待。
错误分为可恢复和不可恢复两种：
- 可恢复：
  - 参数配置错误等，例如block越界（一般最大值是1024），shared memory大小超出范围（一般是48KB）；
  - 通过cudaGetLastError可以获取错误代码，同时把当前状态恢复为success；
  - 该错误在调用核函数后可以立即通过cudaGetLastError/cudaPeekAtLastError拿到；
  - 该错误在下一个函数调用的时候会覆盖；
- 不可恢复：
  - 核函数执行错误，例如访问越界等等异常；
  - 该错误则会传递到之后的所有cuda操作上；
  - 错误状态通常需要等到核函数执行完毕才能够拿到，也就是有可能在后续的任何流程中突然异常（因为是异步的）。