大家好,我是Tim。
在Nvidia机器中GPU与CUDA一直是双宿双飞的,CUDA是其中的平行处理平台,CUDA Core则是GPU内部的处理单元。
CUDA 是 Compute Unified Device Architecture 的缩写,是GPU并行编程处理和直接访问Nvidia GPU指令集API的总称,它适用于流行的编程语言C、C++,方便用户调用。
CUDA Core则为GPU中的处理单元,如果我们将GPU处理器比作玩具工厂,那么CUDA Core就是其中的流水线,如果你想生产更多的玩具,就需要使用更多的流水线。
我们来看一张图灵102架构中SM的图示:

由上图我们可以看到,通用的GPU中有三种核心,分别是Cude Core, Tensor Core,RT Core。
如果我们把每种Core都比作流水线的话,那么通用的GPU中就存在三类流水线,他们各自的特性是什么呢?
快谈GPU架构的进化
Nvidia显卡从 Tesla 架构开始,即“出道”开始,所有 Nvidia GPU 都带有有 CUDA Core,但Tensor Core和RT Core确并非都具有。
在Fermi架构之前,GPU的处理核心一直被叫做Processor Core(SPs),随着GPU中处理核心的增加,直到2010年Nvidia的Fermi架构它被换了一个新名字叫做CUDA Core。

如上图所示,在Fermi架构中其计算核心由16个SM(Stream Multiprocesser)组成,每个SM包含2个线程束(Warp),16组加载存储单元(LD/ST)和4个特殊函数单元(SFU)组成。
最核心的是,每个线程束Warp包含16个Cuda Core组成,每一个Cuda Core由1个浮点数单元FPU和1个逻辑运算单元ALU组成。
我们先来看看CUDA Core里最初有什么吧。
可以看出,一个cuda core 包含了一个整数运算单元(INT Unit)和一个浮点运算单元(FPU)。
然后,这个core能进行一种原子性的fused multiply-add (FMA)的操作,通俗一点就是一个加乘操作的融合。
你一定记得这个运算:
y = W * X + b
在常规的运算中,先把寄存器里面的数据送入乘法器,然后把结果送回寄存器,然后再把寄存器的数据送入加法器。而这个cuda core,可以实现单指令完成。
而这些CUDA core 在显卡里面是并行运算,就是说大家分工计算。从逻辑上说,那么CUDA core越多,算力也就相应的会越强。
所以说Fermi这个架构确立了英伟达GPU整体的发展方向,2012年的Kepler架构和2014年的Maxwell架构,都在这个基础上疯狂加倍增加Cuda Core。
到了2016年的Pascal架构,老黄注意到了深度学习这个“大金矿”,Nvidia GPU也开始往深度学习方向进行演进。

由上图可以看出,在Pascal SM的内部,除了以往支持单精度的FP32 Cuda Core,还增加了支持双精度的DP Unit,而DP Unit实际上是FP64的Cuda Core。
一个SM由64个FP32 Cuda Cores和32个 FP64 Cuda Cores(DP Unit)组成,此外,FP32 Cuda Core也具备处理半精度FP16的能力,以满足当时行业开始对低精度计算的需求。
此外,NVLink也是这个时候开始引入的。
到了2017年的Volta架构,Nvidia GPU 已经深入深度学习进行优化。

由上图可以看出,在Volta架构的SM中,在FP64 Cuda Cores和FP32 Cuda Core基础上增加了INT32 Cuda Core,意味着可以执行INT32的操作。
更重要的是,引入了张量核Tensor Core模块,用于执行融合乘法加法。
在 Tensor Core 发布之前,CUDA Core 一直是加速深度学习的硬件。
由于Cuda Core每个时钟周期只能进行一次计算,而CUDA Core数量和每个Core时钟速度还存在一定的限制,一定程度上制约了深度学习计算性能的提升。
在这个背景下,NVIDIA 开发了 Tensor Core。

Tensor Core 是支持混合精度训练的专用核心。第一代Tensor Core内核是通过融合乘加计算来实现这一点,它允许两个 4 x 4 FP16 矩阵相乘并添加到 4 x 4 FP16 或 FP32 矩阵中。
混合精度计算之所以如此命名,是因为虽然输入矩阵可以是低精度 FP16,但最终输出将是 FP32,且输出中的精度损失极小,大大加速了模型的训练。
Tensor Core的针对深度学习的训练速度是非常牛掰的,它使得 Volta 提供了比 Pascal 高 3 倍的训练和推理性能。
与 NVIDIA Pascal 相比,用于训练的峰值 teraFLOPS (TFLOPS) 性能提升了高达 12 倍,用于推理的峰值 TFLOPS 性能提升了高达 6 倍。这项关键功能使 Volta 提供了比 Pascal 高 3 倍的训练和推理性能。
在Volta基础上,2018年英伟达发布Turing架构,对Tensor Core进行了升级,增加了对INT8、INT4、Binary(INT1)的计算能力,这使得混合精度训练操作能够将 GPU 的性能吞吐量提高 Pascal GPU 的 32 倍!
此外,在图灵架构中还配备了RT Core(专用光线追踪处理器),能够高速对光线和声音进行渲染。
在2020年发布Ampere 架构,对Tensor Core又进行了升级,增加了TF32和BF16两种数据格式的支持,也增加了对稀疏矩阵计算的支持,同时也带来第二代RT Core光线追踪核心。
这些 CUDA 核心的有趣之处在于它可以处理整数和浮点运算。这意味着 Ampere 架构中的每个 CUDA 核心每个时钟周期可以处理两个 FP32 或一个 FP32 和一个 INT 操作。
随着大模型的爆火,深度学习进入加速时代。在2022年NVIDIA GTC上,老黄介绍了基于全新Hopper架构,完全深度学习芯片。

Hopper架构的新一代流式多处理器中,引入了 FP8张量核心(Tensor Core)来加速 AI 训练和推理。同时去掉了RT Core来给深度学习腾挪空间, 此外引入了Transformer引擎,所以说Hopper架构H100是深度学习专用芯片。
在Hopper架构中,每个SM包括128个FP32 CUDA 核心,和4个第4代张量核心(Tensor Core)。进入SM单元的指令首先存入L1指令缓存(L1 Instruction Cache),然后再分发到L0指令缓存(L1 Instruction Cache)。与L0缓存配套的线程束排序器(Wrap Scheduler)和调度单元(Dispatch Unit)来为CUDA核心和张量核心分配计算任务。
Hopper的张量核心支持 FP8、FP16、BF16、TF32、FP64 和 INT8 MMA 数据类型。这一代张量核心的关键点是引入了Transformer引擎。
什么影响CUDA Core性能
在继续讲解Tensor Core之前,我们需要先来了解下GPU的性能受什么影响。
GPU 核心不能获取或解码指令,而只能执行计算。
每一个CUDA Core是一个“工厂流水线”的话,由于每一个CUDA Core都可以并行执行,那么是否可以通过CUDA Core的数量来判断任何显卡的性能。
显然是不可能,虽然说“工厂流水线”越多“工厂”的性能可能会越好,但这也受限于每个流水线的生产效率,生产设备的架构,生产存储资源的能力等等。
反应到GPU上,也就是我们需要考虑,显卡架构、时钟速度、CUDA 核心数量,内存带宽、内存速度、VRAM等很多因素。
如果比较同一代的 GPU,CUDA 核心的数量可以作为性能的良好指标,只要不存在其他性能瓶颈因素,CUDA 数量越高意味着同代 GPU 的性能越好。
但如果比较不同代的GPU, CUDA 核心的数量则比一定能反应性能的好坏,我们以Maxwell 架构GTX 980 Ti和Pascal 架构GTX 1080进行对比举例。

上表中我们可以看到,两款显卡的 晶体管/CUDA核心数量 仅仅存在微小差异,那么这两个显卡的性能相似吗?
答案是否定的。
在Maxwell架构中,为了使制造商能够在小型芯片上容纳更多的晶体管,使用了更小的晶体管,而更小的晶体管尺寸降低了总体功耗,使得其时钟的速度较低。
而Pascal 架构虽然其Cuda Core的数量上略少,但由于其时钟速度较快,使得其整体的性能更好。
下面是一个假设方程,显示 GPU 的性能如何取决于各种因素。
GPU性能 = 核心数 * 时钟频率 * 架构系数;
现在让我们总结一下CUDA 核心性能相关。
更多 CUDA 核心数量意味着可以并行处理更多数据。
更高的时钟速度意味着单个核心可以执行得更快。
GPU 在新一代和架构中变得更好,因此具有更多 CUDA 核心的显卡不一定比具有较少 CUDA 核心的显卡更强大。
正是由于时钟速度和CUDA 核心数量的限制,Nvidia GPU才针对深度学习推出了Tensor Core。
Tensor 就是矩阵的意思,也就是说他在tensor矩阵场景下算得更快。
Tensor Core原理
首先,Tensor Core的计算速度比 CUDA Core快得多。CUDA Core 每个时钟周期执行一次操作,而Tensor Core 内核每个时钟周期可以执行多个操作。
相比于 FP32 Core 一次只能对两个数字进行计算,Tensor Core 能一次对两个 4×4 的 FP16 Tensor 进行矩阵乘计算并累加到另一个 4×4 的 Tensor 上,即 D = A * B + C,这也是其取名为 Tensor Core 的原因。
通过硬件上的特殊设计,Tensor Core 理论上可以实现 8 倍于 FP32 Core 的计算吞吐量,并且没有明显的占用面积和功耗增加。
其次,混合精度也是利用 Tensor Core 的这一特性,才能够实现训练加速。
混合精度并不是网络层面既有 FP16 又有 FP32,它指的是在底层硬件算子层面,使用半精度(FP16)作为输入和输出,使用全精度(FP32)进行中间结果计算从而不损失过多精度的技术,这个底层硬件其实指的就是 Tensor Core。
如下图所示,Tensor Core 对 FP16 输入数据进行操作,并进行 FP32 累积。FP16 乘法会产生一个全精度结果,该结果在 FP32 运算中与 4x4x4 矩阵乘法的给定点积中的其他乘积一起累加。

在程序执行期间,通过完整的执行Warp同时使用多个 Tensor Core。warp 内的线程提供更大的 16x16x16 矩阵运算以供 Tensor Core 处理。
CUDA 将这些操作公开为 CUDA C++ WMMA API 中的扭曲级矩阵操作。这些 C++ 接口提供专门的矩阵加载、矩阵乘法和累加以及矩阵存储操作,以在 CUDA C++ 程序中有效利用 Tensor Core。
此外,目前许多深度学习框架,包括 Tensorflow、PyTorch都支持Tensor Core,可以参考混合精度训练指南。而对于深度学习推理,TensorRT 3 也支持了 Tensor Core。
既然Tensor Core这么好,为什么Nv不生产一个全部是Tensor Core的深度学习专用GPU架构?
虽然说Tensor Core是专门为加速深度学习和 AI 工作负载(例如矩阵运算)而设计,但目前深度学习也不能离开Cuda Core。
CUDA Core 针对各种并行计算任务进行了优化,更适合于通用并行计算任务。
首先,深度学习任务其不仅仅是矩阵运算,还有很多的并行计算,这就看瓶颈在哪里了。如果瓶颈时在并行计算,那么这种类型的深度学习任务可能更适合CUDA Core。
其次,Tensor Core的使用是有些限制的,对于GEMM计算其效果很好,其次其输入和输出数据类型需要是半精度或单精度,矩阵的维度最好是 8 的倍数。
当然Tensor Core也可用于通用并行计算任务,但它们可能不如 CUDA Core 高效。
总结
通用的 GPU 中包含三种核心,分别是 CUDA Core、Tensor Core 和 RT Core,这三种核心各自具有不同的特性和功能。
CUDA Core:CUDA Core 是用于通用并行计算任务的计算核心,可以执行单精度和双精度浮点运算,以及整数运算。它在处理广泛的并行计算任务方面非常高效。
Tensor Core:Tensor Core 是针对深度学习和 AI 工作负载而设计的专用核心,可以实现混合精度计算并加速矩阵运算,尤其擅长处理半精度(FP16)和全精度(FP32)的矩阵乘法和累加操作。Tensor Core 在加速深度学习训练和推理中发挥着重要作用。
RT Core:RT Core 是专门用于光线追踪处理的核心,能够高速进行光线和声音的渲染,对于图形渲染和光线追踪等任务具有重要意义。
Tensor Core 作为混合精度训练赖以加速的底层硬件支持,在Volta 架构之后才开始支持。
NVIDIA Tensor Core 是完全可编程的。Warp级别的 Tensor Core 编程 API 已在命名空间mma.h下的标头中声明nvcuda::wmma。
然而Tensor Core的加速也并不是万能的。
首先,深度学习任务其不仅仅是矩阵运算,还有很多的并行计算,这就看瓶颈在哪里了。如果瓶颈时在并行计算,那么这种类型的深度学习任务可能更适合CUDA Core。
其次,Tensor Core的使用是有些限制的,对于GEMM计算其效果很好,其他则性能不一定很好,其次其输入和输出数据类型需要是半精度或单精度,矩阵的维度最好是 8 的倍数。