CUDA学习入门（二）
前言
本篇接着上一篇，本仍然是对CUDA入门作一个宏观的概述，并不涉及具体编程细节，是关于CUDA基础信息的科普性文章，适用于想简单了解什么是CUDA，以及对CUDA工作流程感兴趣的小伙伴。
上一篇连接：CUDA学习入门（一）（包含CUDA安装和相关基础知识）
上一篇目录如下：
三、CUDA程序的编译过程
CUDA程序编写完成之后，同其他高级语言一样，我们需要将代码进行编译之后生成可执行文件后才能运行。
3.1 区分不同的NVIDIA架构
现在已经到了2024年，英伟达的GPU从2010年的 费米（Fermi） 架构到2022年的 赫柏（Hopper） 架构，历经十多年的发展历程，每一代架构的更迭都会带来新的功能，就会导致不同型号的GPU在硬件配置上会有差异，这就意味着我们在进行CUDA程序编译时需要注意区分不同的GPU架构。
NVIDIA给自己的GPU提供了一套编号来表示GPU的架构，格式为 sm_xy 。x 代表的是该GPU架构的编号， y 代表的是该GPU在本架构中的版本号。新的架构并不一定能兼容老的架构，比如在 Fermi 架构上编译的代码就无法在Kepler GPU 上运行，因为它们的指令集与指令编码都是不一样的，所以在编译时必须确定好你所使用的GPU的架构。
3.2 编译的两个阶段
为了解决CUDA编程的共性和GPU架构的不同所产生的的兼容性问题，CUDA代码编译过程采用了两个阶段：
第一个阶段：
nvcc编译器会将CUDA代码进行预编译，也就是不具体区分不同GPU架构的区别，采用一个虚拟架构对CUDA代码进行编译，将CUDA代码编译成中间过度的PTX文件。
PTX是英文 Parallel Thread Execution 首字母的缩写，是 Nvidia 为 CUDA 设计的一种低级虚拟机和指令架构，它是CUDA 代码的中间表示形式，采用抽象的寄存器和线程模型。
第二个阶段：
nvcc 编译器是包含在 GPU 驱动中，第一阶段生成的PTX文件不依赖于具体的 GPU 架构，可以针对不同的 GPU 生成优化机器代码，可以在不同的 CUDA 运行环境和 GPU设备上执行。PTX 代码在执行前， GPU 驱动则根据真实架构将 PTX 代码编译成实际的二进制代码，也就是编译生成针对特定 GPU 的机器代码，运行在具体的机器上，CUDA代码成功编译完成。
3.3 部分架构代号列表
| 虚拟架构代号 | 真实架构代号 | CUDA支持版本 | 支持架构 | 支持硬件 |
|---|---|---|---|---|
| compute_50 | sm_50 | CUDA 6~11 | Maxwell | Tesla/Quadro M series |
| compute_52 | sm_52 | CUDA 6~11 | Maxwell | GTX-980, GTX Titan X |
| compute_53 | sm_53 | CUDA 6~11 | Maxwell | Tegra TX1, Jetson Nano |
| compute_60 | sm_60 | CUDA 8 | Pascal | Tesla P100 |
| compute_61 | sm_61 | CUDA 8 | Pascal | GTX 1080, GTX1070 |
| compute_62 | sm_62 | CUDA 8 | Pascal | Jetson TX2 |
| compute_70 | sm_70 | CUDA 9 | Volta | Tesla V100 |
| compute_72 | sm_72 | CUDA 9 | Volta | Jetson AGX Xavier |
| compute_75 | sm_75 | CUDA 10 | Turing | RTX 2080, RTX 2070 Tesla T4 |
| compute_80 | sm_80 | CUDA 11.1 | Ampere | A100 |
| compute_86 | sm_86 | CUDA 11.1 | Ampere | RTX 3090 |
| compute_87 | sm_87 | CUDA 11.1 | Ampere | Jetson AGX Orin |
| compute_89 | sm_89 | CUDA 11.8 | Lovelace | RTX 4090 |
| compute_90 | sm_90 | CUDA 12 | Hopper | H100 H200 |
| compute_95 | sm_95 | CUDA 12 | Blackwell | B100 |
我这里介绍的CUDA代码编译过程较为笼统，只是对大的过程进行概括，更加细节的可以参考下面这位大佬的文章：
https://blog.csdn.net/fb_help/article/details/80462853
四、CUDA并行计算基础
4.1 CUDA线程层次
英伟达GPU的架构是围绕一个流式多处理器(Streaming Multiprocessors，SM)的可扩展阵列构建的，通过复制这种架构的构建来实现GPU的硬件并行。
一个GPU中通常有多个SM，每个SM上支持许多个线程并发地执行，CUDA采用单指令多线程(Single-Instruction Multiple-Thread，SIMT)来管理和执行GPU上的众多线程，并提出一个三级的线程层级结构的概念以便组织线程。
由一个内核（Kernel）启动所产生的所有线程统称为一个线程网格（Grid），同一网格中的所有线程共享全局内存空间，一个网格由多个线程块（Block） 组成，一个线程块包含一组线程（Thread），同一线程块内的线程通过同步和共享内存的方式实现协作，不同块内的线程不能协作。
当host通过内核函数启动一个内核网格时，这个内核网格的线程块就被分配到可用的SM上来执行，一个线程块内的多个线程在SM上并发执行，多个线程块可以并发地在一个SM上执行，当线程块终止时，新的线程块又可以在腾出的SM上启动执行。
CUDA线程层次如下图所示：
-
Thread: sequential execution unit
— 所有线程执行相同的核函数
— 并行执行 -
Thread Block: a group of threads
— 执行在一个Streaming Multiprocessor (SM)
— 同 一 个Block 中的线程可以协作 -
Thread Grid: a collection of thread blocks
— 一 个Grid当中的Block 可以在多个SM中执行
4.2 GPU代码示例
这里有一个两个数组相加的GPU代码段示例：
__global__ void VectorAddGPU(const float *a, const float *b,
float *c, const int n) {
int i = blockDim.x * blockIdx.x + threadIdx.x; // 线程ID
if (i < n) {
c[i] = a[i] + b[i]; //每个线程需要做的事情
}
}
可以看到，要想实现两个数组相加，在GPU代码中并不需要像CPU代码（可以自行尝试写出）那样的循环，只是需要一个线程ID来进行索引，并告诉每个线程需要做的事情。线程依靠两个内置变量来进行区分：
-
threadldx.[x y z]
是执行当前kernel函数的线程在block中的索引值 -
blockldx.[x y z]
是指执行当前kernel函数的线程所在block, 在grid中的索引值
在调用内核函数的时候，会在<<< >>>内设置两个参数，分别代表线程网格的维度和线程块的维度。CUDA可以组织三维的线程网格和线程块，它们的维度由下列两个内置变量来决定：
-
blockDim.[x y z]
表示一个block中包含多少个线程 -
gridDim.[x y z]
表示一个grid中包含多少个block
例如我们想要调用上面写的Kernel函数可以写样写：
VectorAddGPU<<<1,4>>>(a,b,c,4);
我们定义一个线程块，每个线程块下一共有四个线程同时运行。
实际上在设备上运行的样子如下图所示，线程0~3，四个线程并行计算：
4.3 CUDA线程层次与GPU硬件的关系
CUDA的线程产生的层级分别其实在GPU硬件结构上也是有所体现的，如下图所示：
我们的线程是在我们的CUDA核心里进行执行的；
一个线程块是在一个SM里执行的；
而一个网格则是在整个Device中执行的。
硬件调度：
- Grid:GPU[GPC]级别的调度单位
- Block(CTA):SM 级别的调度单位
- Threads/Warp:CUDA core级别的调度单位
资源和通信：
- Grid: 共享同样的kernel 和 Context
- Block(CTA): 同一个SM(Streaming Multiprocessor),同一个SM(Shared Memory)
- Threads/Warp: 允许同一个warp中的thread 读取其他 thread 的值
4.4 CUDA的执行流程
-
加载核函数
-
将Grid分配到一个Device
-
根据<<<…>>>内的执行设置的第一个参数，Giga threads engine将 block 分配到 SM 中。 一个 Block 内的线程一定会在同一个 SM 内， 一个 SM 可以有很多个 Block 。
-
根据<<<…>>>内的执行设置的第二个参数，Warp 调度器会将调用线程。
-
Warp调度器为了提高运行效率，会将每32个线程分为一组，称作一 个warp（线程束）。
-
每个Warp会被分配到32个core上运行。
本人是一名学生，也是正在学习Jetson的过程中，如有错误请批评指正！
部分图表信来源：https://mp.weixin.qq.com/s/em309Ho6AaV5f1ogWpajJw
部分代码示例参考链接：https://zhuanlan.zhihu.com/p/440112374
8422
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
