tensorRT从零起步迈向高性能工业级部署（就业导向） 课程笔记，讲师讲的不错，可以去看原视频支持下。
环境
cu 开头的，而Runtime API 基本都是以 cuda 开头的。Runtime API 的特点
cuInit 初始化，避免 Driver API 未初始化的错误；cuDevicePrimaryCtxRetain 实现；cuDevicePrimaryCtxRetain 为每个设备设置 context，不再手动管理 context，并且不提供直接管理 context 的 API（可 Driver API 管理，通常不需要）.cpp 可以与 .cu 文件无缝对接cuda_runtime.h 和 libcudart.so
类似于在介绍 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;
}
下图是的 Device Memory 的分类
对于主机内存，即整个 host memory 而言，操作系统在逻辑上将其区分为两个大类：
可以理解为 page lock memory 是酒店的 vip 房间，锁定给你一个人使用。而 pageable memory 是普通房间，在酒店房间不够的时候，选择性地将你的房间腾出来（交换到硬盘上）给其他人使用，这样就能容纳更多人了。造成房间很多的假象，代价是性能很低。pageable memory 就是常见的虚拟内存的特性。
基于前面的理解，我们总结如下：
pinned memory 可以直接传送数据到 GPU
而 pageable memory ，由于并不锁定，需要先传到 pinned memory。
GPU 可以直接访问 pinned memory，称为 DMA (Direct Memort Access)
对于 GPU 访问而言，距离计算单元越近，效率越高，所以：
SharedMemory > GlobalMemory > PinnedMemory
代码中，
new/malloc 分配的是 pageable memorycudaMallocHost 分配的是 PinnedMemorycudaMalloc 分配的是 GlobalMemory尽量多用 PinnedMemory 储存 host 数据，或者显式处理 Host 到 Device 时，用 PinnedMemory 做缓存，都是提高性能的关键
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
)
(2,3)
(2,3) ，那么对应的，其两个维度上索引的取值范围就是：
0
−
1
,
0
−
2
0-1,0-2
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];
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
相当于 cuda 的 stl，但并不常用
若核函数出错，由于它是异步的，立即执行 cudaPeekAtLastError 只会拿到对输入参数校验是否正确的状态，而不会拿到核函数是否正确执行的状态。
需要等待核函数真正执行完毕之后才知道当前核函数是否出错，一般通过设备同步或者流同步进行等待
错误分为可恢复和不可恢复两种
cudaGetlastError 可以获取错误代码，同时把当前状态恢复为successcudaGetLastError / cudaPeekAtLastError 拿到