精简CUDA教程——CUDA Runtime API

tensorRT从零起步迈向高性能工业级部署（就业导向）课程笔记，讲师讲的不错，可以去看原视频支持下。

Runtime API 概述

环境

在这里插入图片描述

图中可以看到，Runtime API 是基于 Driver API 之上开发的一套 API。
之前提到过 Driver API 基本都是 cu 开头的，而Runtime API 基本都是以 cuda 开头的。

Runtime API 的特点

Runtime API 与 Driver API 最大的区别是懒加载，即在真正执行功能时才自动完成对应的动作，即：
- 第一个 Runtime API 调用时，会自动进行 cuInit 初始化，避免 Driver API 未初始化的错误；
- 第一个需要 context 的 API 调用时，会创建 context 并进行 context 关联，和设置当前 context，调用 cuDevicePrimaryCtxRetain 实现；
- 绝大部分 api 都需要 context，例如查询当前显卡名称、参数、内存分配释放等
CUDA Runtime 是封装了 CUDA Driver 的更高级别、更友好的 API
Runtime API 使用 cuDevicePrimaryCtxRetain 为每个设备设置 context，不再手动管理 context，并且不提供直接管理 context 的 API（可 Driver API 管理，通常不需要）
可以更友好地执行核函数，.cpp 可以与 .cu 文件无缝对接
Runtime API 对应 cuda_runtime.h 和 libcudart.so
Runtime API 随 cudatoolkit 发布
主要知识点是核函数的使用、线程束布局、内存模型、流的使用
主要是为了实现归约求和、放射变换、矩阵乘法、模型后处理，就可以解决绝大部分问题

错误处理

类似于在介绍 Driver API 时的情况，我们同样提出 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;
}

内存模型 pinned memory

内存模型是 CUDA 中很重要的知识点，主要理解 pinned_memory、global_memory、shared_memory 即可，其他的不太常用。
pinned_memory 属于 host memory，而 global_memory、shared_memory 属于 device memory。

下图是的 Device Memory 的分类

在这里插入图片描述

锁定性和性能

对于主机内存，即整个 host memory 而言，操作系统在逻辑上将其区分为两个大类：

pageable memory，可分页内存
page lock memory (pinned memory)，页锁定内存/锁页内存

可以理解为 page lock memory 是酒店的 vip 房间，锁定给你一个人使用。而 pageable memory 是普通房间，在酒店房间不够的时候，选择性地将你的房间腾出来（交换到硬盘上）给其他人使用，这样就能容纳更多人了。造成房间很多的假象，代价是性能很低。pageable memory 就是常见的虚拟内存的特性。

基于前面的理解，我们总结如下：

锁定性
- pinned memory 具有锁定特性，是稳定不会被交换的，这很重要，相当于每次去这个房间都一定能找到你
- pageable memory 没有锁定特性，对于第三方设备（如 GPU）去访问时，因为无法感知内存是否被交换，可能得到不到正确的数据，相当于每次去房间找你，说不定你的房间正好被交换了
- 因此， GPU 可以直接访问 pinned memory 而不能访问 pageable memory
性能
- pageable memory 的性能比 pinned memory 差，因为我们的 pageable memory 很可能会被交换到硬盘上
- pageable memory 策略能使用内存假象，比如实际只有 8G 内存却能使用 16G（借助 swap 交换），从而提高程序的运行数量
- pinned memory 也不能太多，会导致操作系统整体性能变差（可同时运行的程序变少），而且 8G 内存最多就 8G 锁页内存。

数据传输到GPU

在这里插入图片描述

pinned memory 可以直接传送数据到 GPU
而 pageable memory ，由于并不锁定，需要先传到 pinned memory。

关于内存其他几个点

GPU 可以直接访问 pinned memory，称为 DMA (Direct Memort Access)
对于 GPU 访问而言，距离计算单元越近，效率越高，所以：

SharedMemory > GlobalMemory > PinnedMemory
代码中，
- 由 new/malloc 分配的是 pageable memory
- 由 cudaMallocHost 分配的是 PinnedMemory
- 由 cudaMalloc 分配的是 GlobalMemory
尽量多用 PinnedMemory 储存 host 数据，或者显式处理 Host 到 Device 时，用 PinnedMemory 做缓存，都是提高性能的关键

流 stream

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

生活中的例子

同步（串行）	异步

在这个例子中，男朋友的微信消息，就是任务队列，流的一种抽象
女朋友发出指令之后，她可以做任何事情，无需等待指令执行完毕。即异步操作中，执行的代码加入流的队列之后，立即返回，不耽误时间。
女朋友发的指令被送到流中排队，男朋友根据流的队列，顺序执行。
女朋友选择性，在需要的时候等待所有的执行结果
新建一个流，就是新建一个男朋友，给他发指令就是发微信，可以新建很多个男朋友
通过 cudaEvent 可以选择性等待任务队列中的部分任务是否就绪

注意

要十分注意，指令发出后，流队列中储存的是指令参数，不能在任务加入队列后立即释放参数指针，这会导致流队列执行该指令时指针失效而出错。应当在十分肯定流已经不需要这个指针之后，才进行修改或释放，否则会有非预期行为出现。

就比如，女朋友让男朋友去卖西瓜并转给了他钱，但是却在男朋友买瓜成功前将转账撤了回去，这时就无法知道男朋友在水果店会发生什么，比如会不会跟老板打起来之类的。因此，要保证买瓜行为顺利完成（行为符合预期），在买瓜成功前就不能动买瓜的钱。

核函数

简介

核函数是 cuda 编程的关键
通过 xxx.cu 创建一个 cudac 程序文件，并把 cu 文件交给 nvcc 编译，才能识别 cuda 语法；
__xxx__ 修饰
- __global__ 表示为核函数，由 host 调用；
- __device__ 表示设备函数，由 device 调用；
- __host__ 表示主机函数，由 host 调用；
- __shared__ 表示变量为共享变量。
- 可能存在上述多个关键字修饰同一个函数，如 __device__ 和 __host__ 修饰的函数，既可以设备上调用，也可以在主机上调用
host 调用核函数：
```
function<<<gridDim, blockDim, sharedMemorySize, stream>>>(args, ...)
```
gridDim 和 blockDim 的变量类型为 dim3，是一个三维的值；

function 函数总共启动的线程数目可以这样计算：n_threads = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

详细请参考线程束的相关知识
只有 __global__ 修饰的函数才可以用 <<< >>> 的方式调用s
调用核函数是传值的，不能传引用，可以传递类，结构体等，核函数可以使模板
核函数的返回值必须是 void
核函数的执行是异步的，也就是立即返回的
线程 layout 主要用到 blockDim、gridDim
和函数内访问线程索引主要用到 threadIdx、blockIdx、blockDim、gridDim 这些内置变量

线程索引计算

共涉及四个变量：blockDim、gridDim、threadIdx、blockIdx ，其中前两者可以认为是形状，后两者可以认为是对应的索引。就像我们 PyTorch 中如果一个张量的形状为 $(2, 3)$ ，那么对应的，其两个维度上索引的取值范围就是： $0 - 1, 0 - 2$ 。

在这里插入图片描述

线程索引 id 计算方法：左乘右加，如上图所示。

共享内存

由 __shared__ 关键字修饰
共享内存因为更靠近计算单元，所以访问速度更快
共享内存通常可以作为访问全局内存的缓存使用
可以利用共享内存实现线程间的通信
通常与 __syncthreads 同时出现，这个函数是同步 block 内的所有线程，全部执行到这一行才往下继续执行

如：
```
__shared__ int shared_value1;
__shared__ int shared_value2;if (threadIdx.x == 0) {if (blockIdx.x == 0) {shared_value1 = 123;shared_value2 = 55;}else {shared_value1 = 331;shared_value2 = 8;}__syncthreads();printf("...")
}
```
其他 threadIdx.x 不为 0 的线程不会进到判断语句，但是会卡在 __syncthreads() ，等待 threadIdx.x 为 0 的线程设置好共享内存，再一起继续向下执行。
共享内存使用方式：通常是在线程 id 为 0 的时候从 global memory 取值，然后 __syncthreads ，然后再使用
动态共享内存与静态共享内存
- 动态共享内存的声明需要加 extern 关键字，不需要指定数组大小，如：
```
extern __shared__ char dynamic_shared_memory[];
```
- 静态共享内存的声明需要指定数组大小，如：
```
const size_t static_shared_memory_size = 6 * 1024; // 6KB
__shared__ char static_shared_memory[static_shared_memory_size];
```

warp affine 实战

chapter: 1.6, caption: vector-add, description: 使用cuda核函数实现向量加法
chapter: 1.7, caption: shared-memory, description: 共享内存的操作
chapter: 1.8, caption: reduce-sum, description: 规约求和的实现，利用共享内存，高性能
chapter: 1.9, caption: atomic, description: 原子操作，实现动态数组的操作
chapter: 1.10, caption: warpaffine, description: 仿射变换双线性插值的实现，yolov5的预处理
chapter: 1.11, caption: cublas-gemm, description: 通用矩阵乘法的cuda核函数实现，以及cublasSgemm的调用
chapter: 1.12, caption: yolov5-postprocess, description: 使用cuda核函数实现yolov5的后处理案例

TODO

thrust

相当于 cuda 的 stl，但并不常用

错误处理

若核函数出错，由于它是异步的，立即执行 cudaPeekAtLastError 只会拿到对输入参数校验是否正确的状态，而不会拿到核函数是否正确执行的状态。

需要等待核函数真正执行完毕之后才知道当前核函数是否出错，一般通过设备同步或者流同步进行等待

错误分为可恢复和不可恢复两种

可恢复
- 参数配置错误，例如 block 越界（一般最大值是 1024），shared memory 超出大小范围（一般是 64KB）等
- 通过 cudaGetlastError 可以获取错误代码，同时把当前状态恢复为success
- 该种错误可以在调用核函数之后立即通过 cudaGetLastError / cudaPeekAtLastError 拿到
- 该种错误在下一个函数调用时会覆盖
不可恢复
- 核函数执行错误，例如访问越界等
- 该错误会传递到之后所有的 cuda 操作上
- 错误状态通常需要等到核函数执行完毕才能够拿到，也就是有可能在后续的任何流程中突然异常（因为是异步的）