写在前面（废话，请跳过）
本来自己的脑子里是不存在 CUDA 编程这个东西的，没错，就是压根儿没听说过。
之所以了解到这个东西，是因为最近开始做 AVM（Around View Monitor），或者说 SVS（Surrounding View System），利用 4 个广角相机，实现车辆周围的 360° 环视功能。
这里面涉及到很多的坐标转换和矩阵乘法，众所周知这些运算是非常耗费资源的，平时 1-2 个相机通过 TensorRT 做模型推理还好，大部分工作被 GPU 做了，留给 SLAM、控制等模块的 CPU 算力勉强还够，但是同时处理 4 路相机，以我目前用的 Jetson Orin AGX 的 8 核 ARM CPU 来说，留给其他功能的 CPU 资源已经不多了……
偶然发现一个同学做的环视，使用 CUDA 加速竟然能达到 10000 FPS，虽然我用她的代码并不能跑出这么快的速度，但是花里胡哨的 CUDA Kernel 函数仍然给了我一个小小的震撼，以及大大的精神冲击（为啥我同学啥都会？我咋啥都不会？）。
既然自己也是做视觉感知的，那就开学捏！
并行性
对于一段代码，如果是有执行次序的，那么就必须串行执行；若没有执行次序，那么则可以并发执行。如果一个任务处理的是另一个任务的输出，即这两个任务间存在相关性。相关性是限制并行性的主要因素之一。
程序中有两种基本并行类型：
- 任务并行：多个任务或函数能独立地并行执行，重点在于多核系统对任务的分配。
- 数据并行：指能够同时处理许多数据，重点在于多核系统对数据的分配。
CUDA 更适合解决 数据并行 计算的问题。
数据并行的第一步是把数据依据 线程 进行划分。一般使用两种方法：block partitioning 和 cyclic partitioning。
- 块划分：一组连续的数据组成一个数据块，一个线程在同一时间通常只处理一个数据块。数据块以任意次序安排给一个线程。
- 周期划分：一个数据块可能只包含一组连续数据中的一部分，数据块按周期排列，相邻线程处理相邻数据块，每个线程可以处理多个数据块。线程选择一个新块意味着要跳过和现有线程一样多的数据块。
 |
 |
NOTE：实际上数据是在一维空间中存储的。我们熟知的 H×W×C 的图像，本质上也是一维数据，因此会有 HWC 和 CHW 的区别。
程序性能对块的大小比较敏感，数据的划分方式与计算机架构有关。
异构计算
在一个异构计算节点中，GPU 不是一个独立运行的平台，而是 CPU 的协处理器。GPU 必须通过 PCIe 总线和基于 CPU 的主机进行操作。因此，GPU 被称为 Device，CPU 被称为 Host。
一个异构应用程序包含两部分：
- Host 代码：在 CPU 上运行
- Device 代码：在 GPU 上运行
描述 GPU 容量通常用如下指标：
- CUDA Core 数量：CUDA Core 内部有分别处理 int 和单精度 float 的单元。TOPS 通常对应 INT8 整型运算的峰值性能，而 TFLOPS 通常对应浮点运算（如 FP32/FP16）的峰值性能，二者描述的是不同数据类型的计算速度。
- Tensor Core 数量：专门为深度学习矩阵运算设计，可以单次完成小尺寸矩阵乘加运算，此外 Tensor Core 支持的数据类型也更多。
- 内存（显存）大小：显存在训练模型阶段，主要影响其
Batch_Size，此参数会直接决定你的训练时长。
描述 GPU 性能通常用如下指标：
- 峰值计算性能：注意区分 TOPS、FLOPS 和 FLOPs。其中 OPS 指的是每秒处理次数，即形容 速度，但通常是默认对 INT8 整型数的处理次数，加上 FL 后 FLOPS 指的是对 FP32 的处理次数。TOPS 即 Tera，每秒几万亿次的计算。FLOPs 指的是深度学习模型自身的计算量，是形容 规模 的。
- 内存带宽：一般单位 GB/s，指从内存中读取/写入数据的速度。
此外，还有一个“计算能力”，即 compute capability，这个看起来也是描述性能的，但是实际上描述的是 GPU 加速器的硬件版本。
Hello World CUDA 版
CUDA 源文件的后缀是 .cu，使用 nvcc 进行编译。当然这并不意味着 cmake 不能编译它。
要注意 .cu 文件的核函数内部不支持 C++ 的 std::cout，而只能使用 C 语言函数 printf。另外，核函数内部的日志输出会大幅增加运行时间，因此调试完毕后需要注释掉。
#include <stdio.h>
__global__ void helloworldFromGPU(void){
printf("Hello world from GPU! \n");
}
int main(void){
printf("Hello world from CPU! \n");
helloworldFromGPU<<<1,10>>>();
cudaDeviceSynchronize();
return 0;
}
在《Professional CUDA C Programming》一书中，使用 cudaDeviceReset() 函数而非 cudaDeviceSynchronize() 函数。虽然 cudaDeviceReset() 也会让所有输出流在程序结束前输出，但所有 CUDA 上下文 context 也会被摧毁，虽然其具有同步数据的功能，但是目前 CUDA 建议使用 cudaDeviceSynchronize()。
此外，cudaDeviceSynchronize() 可以在程序的任何需要同步的地方调用，而 cudaDeviceReset() 通常只在程序结束时调用一次。
编译命令 nvcc hello.cu -o hello，会生成名为 hello 的可执行文件，运行即可得到如下结果：
Hello world from CPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
Hello world from GPU!
在 NVIDIA 官方网站 NVCC Command Options 中，有 nvcc 编译选项的全部内容。
评论