赞
踩
本篇接着上一篇,本仍然是对CUDA入门作一个宏观的概述,并不涉及具体编程细节,是关于CUDA基础信息的科普性文章,适用于想简单了解什么是CUDA,以及对CUDA工作流程感兴趣的小伙伴。
上一篇连接:CUDA学习入门(一)(包含CUDA安装和相关基础知识)
上一篇目录如下:
CUDA程序编写完成之后,同其他高级语言一样,我们需要将代码进行编译之后生成可执行文件后才能运行。
现在已经到了2024年,英伟达的GPU从2010年的 费米(Fermi) 架构到2022年的 赫柏(Hopper) 架构,历经十多年的发展历程,每一代架构的更迭都会带来新的功能,就会导致不同型号的GPU在硬件配置上会有差异,这就意味着我们在进行CUDA程序编译时需要注意区分不同的GPU架构。
NVIDIA给自己的GPU提供了一套编号来表示GPU的架构,格式为 sm_xy
。x
代表的是该GPU架构的编号, y
代表的是该GPU在本架构中的版本号。新的架构并不一定能兼容老的架构,比如在 Fermi 架构上编译的代码就无法在Kepler GPU 上运行,因为它们的指令集与指令编码都是不一样的,所以在编译时必须确定好你所使用的GPU的架构。
为了解决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代码成功编译完成。
虚拟架构代号 | 真实架构代号 | 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
英伟达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中执行
这里有一个两个数组相加的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,四个线程并行计算:
CUDA的线程产生的层级分别其实在GPU硬件结构上也是有所体现的,如下图所示:
我们的线程是在我们的CUDA核心里进行执行的;
一个线程块是在一个SM里执行的;
而一个网格则是在整个Device中执行的。
硬件调度:
资源和通信:
加载核函数
将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
Copyright © 2003-2013 www.wpsshop.cn 版权所有,并保留所有权利。