大模型算力

摩尔线程

2022年，摩尔线程就推出了GPU统一系统架构MUSA，发布并量产“苏堤”和“春晓”两颗全功能GPU芯片，这也是国内采用现代GPU架构

主流 AI 芯片配置

厂商	型号	图形处理器	架构	显存	FP16 算力	BF16 算力	INT8 算力	FP32算力	TF32 算力	FP8算力	CUDA Core	Tensor Core
英伟达	RTX 3090	GA102-300-A1	Ampere	24GB（GDDR6X）	35.58 TFLOPS	-	-	35.58 TFLOPS	-	不支持	10496	328
英伟达	RTX 3090 Ti	GA102-350-A1	Ampere	24GB（GDDR6X）	40.00 TFLOPS	-	-	40.00 TFLOPS	-	不支持	10752	336
英伟达	RTX 4090	AD102-300-A1	Ada Lovelace	24GB（GDDR6X）	369.7 TFLOPS（Tensor Core）82.58 TFLOPS	369.7 TFLOPS（Tensor Core）	739.4 TFLOPS（Tensor Core）	82.58 TFLOPS	-	-	16384	512
英伟达	RTX 4090 Ti	AD102-400-A1	Ada Lovelace	24GB（GDDR6X）	93.24 TFLOPS	-	-	93.24 TFLOPS	-	-	18176	568
英伟达	RTX 4090D-特供-消费级	AD102-250-A1	Ada Lovelace	24GB（GDDR6X）	329.3 TFLOPS（Tensor Core）73.54 TFLOPS	329.3 TFLOPS（Tensor Core）	658.6 TFLOPS（Tensor Core）	73.54 TFLOPS	-	-	14592	456
英伟达	L20（PCIe）-特供-推理（PCIe）	AD102	Ada Lovelace	48GB（GDDR6）	119.5 TFLOPS（Tensor Core）	119.5 TFLOPS（Tensor Core）	239 TOPS（Tensor Core）	59.8 TFLOPS	59.8 TFLOPS（Tensor Core）	239 TFOPS（Tensor Core）	11776	368
英伟达	H20-特供-训练（PCIe、Nvlink）	-	Hopper	96GB（HBM3）	148 TFLOPS（Tensor Core）	148 TFLOPS（Tensor Core）	296 TOPS（Tensor Core）	44 TFLOPS	74 TFLOPS（Tensor Core）	296 TFOPS（Tensor Core）	-	-
英伟达	A800（PCIe）	GA100	Ampere	80GB（HBM2e）	312 TFLOPS（Tensor Core）77.97 TFLOPS	312 TFLOPS（Tensor Core）	624 TOPS（Tensor Core）	19.5 TFLOPS	156 TFLOPS（Tensor Core）	不支持	6912	432
英伟达	H800（ SXM）	GH100	Hopper	80GB（HBM3）	1,979 TFLOPS（Tensor Core）	1,979 teraFLOPS（Tensor Core）	3,958 TOPS（Tensor Core）	67 teraFLOPS	989 teraFLOPS （Tensor Core）	3,958 teraFLOPS（Tensor Core）	18,432	640
昇腾	Atlas 800T A2训练（910B3-HCCS）	-	达芬奇	64GB（HBM2e）	313 TFLOPS	313 TFLOPS	640 TOPS	75 TFLOPS	141 TFLOPS（HF）	不支持	-	-
昇腾	Atlas 800I 推理（910B4）	-	达芬奇	32GB（HBM2e）	280 TFLOPS	280 TFLOPS	550 TOPS	75 TFLOPS	141 TFLOPS（HF）	不支持	-	-

CPU vs GPU

首先明确什么是 GPU ，以及 GPU 和 CPU 的主要区别是什么？

现在探讨一下 CPU 和 GPU 在架构方面的主要区别， CPU 即中央处理单元（Central Processing Unit），负责处理操作系统和应用程序运行所需的各类计算任务，需要很强的通用性来处理各种不同的数据类型，同时逻辑判断又会引入大量的分支跳转和中断的处理，使得 CPU 的内部结构异常复杂。

GPU 即图形处理单元（Graphics Processing Unit），可以更高效地处理并行运行时复杂的数学运算，最初用于处理游戏和动画中的图形渲染任务，现在的用途已远超于此。两者具有相似的内部组件，包括核心、内存和控制单元。

GPU 和 CPU 在架构方面的主要区别包括以下几点：

并行处理能力： CPU 拥有少量的强大计算单元（ALU），更适合处理顺序执行的任务，可以在很少的时钟周期内完成算术运算，时钟周期的频率很高，复杂的控制逻辑单元（Control）可以在程序有多个分支的情况下提供分支预测能力，因此 CPU 擅长逻辑控制和串行计算，流水线技术通过多个部件并行工作来缩短程序执行时间。GPU 控制单元可以把多个访问合并成，采用了数量众多的计算单元（ALU）和线程（Thread），大量的 ALU 可以实现非常大的计算吞吐量，超配的线程可以很好地平衡内存延时问题，因此可以同时处理多个任务，专注于大规模高度并行的计算任务。
内存架构： CPU 被缓存 Cache 占据了大量空间，大量缓存可以保存之后可能需要访问的数据，可以降低延时； GPU 缓存很少且为线程（Thread）服务，如果很多线程需要访问一个相同的数据，缓存会合并这些访问之后再去访问 DRMA，获取数据之后由 Cache 分发到数据对应的线程。 GPU 更多的寄存器可以支持大量 Thread。
指令集： CPU 的指令集更加通用，适合执行各种类型的任务； GPU 的指令集主要用于图形处理和通用计算，如 CUDA

和 OpenCL。
功耗和散热： CPU 的功耗相对较低，散热要求也相对较低；由于 GPU 的高度并行特性，其功耗通常较高，需要更好的散热系统来保持稳定运行。

因此，CPU 更适合处理顺序执行的任务，如操作系统、数据分析等；而 GPU 适合处理需要大规模并行计算的任务，如图形处理、深度学习等。在异构系统中， GPU 和 CPU 经常会结合使用，以发挥各自的优势。

GPU 起初用于处理图形图像和视频编解码相关的工作。 GPU 跟 CPU 最大的不同点在于， GPU 的设计目标是最大化吞吐量（Throughput），相比执行单个任务的快慢，更关心多个任务的并行度（Parallelism），即同时可以执行多少任务；CPU 则更关心延迟（Latency）和并发（Concurrency）。

CPU 优化的目标是尽可能快地在尽可能低的延迟下执行完成任务，同时保持在任务之间具体快速切换的能力。它的本质是以序列化的方式处理任务。 GPU 的优化则全部都是用于增大吞吐量的，它允许一次将尽可能多的任务推送到 GPU 内部。然后 GPU 通过大数量的 Core 并行处理任务。

处理器带宽（Bandwidth）、延时（Lantency）和吞吐（Throughput）

带宽：处理器能够处理的最大的数据量或指令数量，单位是 Kb、Mb、Gb；
延时：处理器执行指令或处理数据所需的时间，传送一个数据单元所需要的时间，单位是 ms、s、min、h 等；
吞吐：处理器在一定时间内从一个位置移动到另一个位置的数据量，单位是 bps（每秒比特数）、Mbps（每秒兆比特数）、Gbps（每秒千比特数），比如在第 10s 传输了 20 bit 数据，因此在 t=10 时刻的吞吐量为 20 bps。

解决带宽相比较解决延时更容易，线程的数量与吞吐量成正比，吞吐量几乎等于带宽时说明信道使用率很高，处理器系统设计所追求的目标是提高带宽的前提下，尽可能掩盖传送延时，组成一个可实现的处理器系统。
并发与并行

并行和并发是两个在计算机科学领域经常被讨论的概念，它们都涉及到同时处理多个任务的能力，但在具体含义和应用上有一些区别。

并行（Parallelism）

并行指的是同时执行多个任务或操作，通常是在多个处理单元上同时进行。在计算机系统中，这些处理单元可以是多核处理器、多线程、分布式系统等。并行计算可以显著提高系统的性能和效率，特别是在需要处理大量数据或复杂计算的情况下。例如，一个计算机程序可以同时在多个处理器核心上运行，加快整体计算速度。

2. 并发（Concurrency）

并发指的是系统能够同时处理多个任务或操作，但不一定是同时执行。在并发系统中，任务之间可能会交替执行，通过时间片轮转或事件驱动等方式来实现。并发通常用于提高系统的响应能力和资源利用率，特别是在需要处理大量短时间任务的情况下。例如，一个 Web 服务器可以同时处理多个客户端请求，通过并发处理来提高系统的吞吐量。

因此并行和并发的主要区别如下：

并行是指同时执行多个任务，强调同时性和并行处理能力，常用于提高计算性能和效率。
并发是指系统能够同时处理多个任务，强调任务之间的交替执行和资源共享，常用于提高系统的响应能力和资源利用率。

在实际应用中，并行和并发通常结合使用，根据具体需求和系统特点来选择合适的技术和策略。同时，理解并行和并发的概念有助于设计和优化复杂的计算机系统和应用程序。在实际硬件工作的过程当中，更倾向于利用多线程对循环展开来提高整体硬件的利用率，这就是 GPU 的最主要的原理。

以三款芯片为例，对比在硬件限制的情况下，一般能够执行多少个线程，对比结果增加了线程的请求（Threads required）、线程的可用数（Threads available）和线程的比例（Thread Ration），主要对比到底需要多少线程才能够解决内存时延的问题。从表中可以看到几个关键的数据：

GPU（NVIDIA A100）的时延比 CPU （AMD Rome 7742，Intel Xeon 8280）高出好几个倍数；
GPU 的线程数是 CPU 的二三十倍；
GPU 的可用线程数量是 CPU 的一百多倍。计算得出线程的比例，GPU 是 5.6， CPU 是 1.2~1.3，这也是 GPU 最重要的一个设计点，它拥有非常多的线程为大规模任务并行去设计。

AMD Rome 7742	Intel Xeon 8280	NVIDIA A100

Memory B/W(GB/sec) 204 143 1555
DRAM Latency(ns) 122 89 404
Peak bytes per latency 24,888 12,727 628,220
Memory Efficiency 0.064% 0.13% 0.0025%
Threads required 1,556 729 39,264
Threads available 2048 896 221,184
Thread Ration 1.3X 1.2X 5.6X

CPU 和 GPU 的典型架构对比可知 GPU 可以比作一个大型的吞吐器，一部分线程用于等待数据，一部分线程等待被激活去计算，有一部分线程正在计算的过程中。GPU 的硬件设计工程师将所有的硬件资源都投入到增加更多的线程，而不是想办法减少数据搬运的延迟，指令执行的延迟。

相对应的可以把 CPU 比喻成一台延迟机，主要工作是为了在一个线程里完成所有的工作，因为希望能够使用足够的线程去解决延迟的问题，所以 CPU 的硬件设计者或者硬件设计架构师就会把所有的资源和重心都投入到减少延迟上面，因此 CPU 的线程比只有一点多倍，这也是 SIMD
（Single Instruction, Multiple Data）和 SIMT

（Single Instruction, Multiple Threads）架构之间最大的区别。 CPU 不是通过增加线程来去解决问题，而是使用相反的方式去优化线程的执行速率和效率，这就是 CPU 跟 GPU 之间最大的区别，也是它们的本质区别。
SIMD (Single Instruction, Multiple Data) 和 SIMT (Single Instruction, Multiple Threads)
SIMD 架构是指在同一时间内对多个数据执行相同的操作，适用于向量化运算。例如，对于一个包含多个元素的数组，SIMD 架构可以同时对所有元素执行相同的操作，从而提高计算效率。常见的 SIMD 架构包括 SSE (Streaming SIMD Extensions) 和 AVX (Advanced Vector Extensions)。
SIMT 架构是指在同一时间内执行多个线程，每个线程可以执行不同的指令，但是这些线程通常会执行相同的程序。这种架构通常用于 GPU (Graphics Processing Unit) 中的并行计算。CUDA (Compute Unified Device Architecture) 和 OpenCL 都是支持 SIMT 架构的编程模型。
SIMD 适用于数据并行计算，而 SIMT 适用于任务并行计算。在实际应用中，根据具体的计算需求和硬件环境选择合适的架构可以提高计算性能。

GPU 工作原理
基本工作原理

首先通过$AX+Y$这个加法运算的示例了解 GPU 的工作原理，$AX+Y$ 的示例代码如下：

void demo(double alpha, double *x, double *y)
{
int n = 2000;
for (int i = 0; i < n; ++i)
{
y[i] = alpha * x[i] + y[i];
}
}

示例代码中包含 2 FLOPS 操作，分别是乘法（Multiply）和加法（Add），对于每一次计算操作都需要在内存中读取两个数据，$x[i]$ 和 $y[i]$，最后执行一个线性操作，存储到 $y[i]$ 中，其中把加法和乘法融合在一起的操作也可以称作 FMA（Fused Multiply and Add）。

在 O(n) 的时间复杂度下，根据 n 的大小迭代计算 n 次，在 CPU 中串行地按指令顺序去执行 $AX+Y$ 程序。以 Intel Exon 8280 这款芯片为例，其内存带宽是 131 GB/s，内存的延时是 89 ns，这意味着 8280 芯片的峰值算力是在 89 ns 的时间内传输 11659 个比特（byte）数据。$AX+Y$ 将在 89 ns 的时间内传输 16 比特（C/C++中 double 数据类型所占的内存空间是 8 bytes）数据，此时内存的利用率只有 0.14%（16/11659），存储总线有 99.86% 的时间处于空闲状态。

GPU 缓存机制

在 GPU 工作过程中希望尽可能的去减少内存的时延、内存的搬运、还有内存的带宽等一系列内存相关的问题，其中缓存对于内存尤为重要。NVIDIA Ampere A100 内存结构中 HBM Memory

的大小是 80G，也就是 A100 的显存大小是 80G。

其中寄存器（Register）文件也可以视为缓存，寄存器靠近 SM（Streaming Multiprocessors）执行单元，从而可以快速地获取执行单元中的数据，同时也方便读取 L1 Cache
缓存中的数据。此外 L2 Cache

更靠近 HBM Memory，这样方便 GPU 把大量的数据直接搬运到 cache 中，因此为了同时实现上面两个目标， GPU 设计了多级缓存。80G 的显存是一个高带宽的内存，L2 Cache 大小为 40M，所有 SM 共享同一个 L2 Cache，L1 Cache 大小为 192kB，每个 SM 拥有自己独立的 Cache，同样每个 SM 拥有自己独立的 Register，每个寄存器大小为 256 kB，因为总共有 108 个 SM 流处理器，因此寄存器总共的大小是 27MB，L1 Cache 总共的大小是 20 MB。

NVIDIA Ampere A100 内存结构

GPU 和 CPU 内存带宽和时延进行比较，在 GPU 中如果把主内存（HBM Memory）作为内存带宽（B/W, bandwidth）的基本单位，L2 缓存的带宽是主内存的 3 倍，L1 缓存的带宽是主存的 13 倍。在真正计算的时候，希望缓存的数据能够尽快的去用完，然后读取下一批数据，此时时候就会遇到时延（Lentency）的问题。如果将 L1 缓存的延迟作为基本单位，L2 缓存的延迟是 L1 的 5 倍，HBM 的延迟将是 L1 的 15 倍，因此 GPU 需要有单独的显存。

假设使用 CPU 将 DRAM（Dynamic Random Access Memory）中的数据传入到 GPU 中进行计算，较高的时延（25 倍）会导致数据传输的速度远小于计算的速度，因此需要 GPU 有自己的高带宽内存 HBM（High Bandwidth Memory），GPU 和 CPU 之间的通信和数据传输主要通过 PCIe 来进行。
NVIDIA Ampere A100 存储延迟对比

DRAM 动态随机存取存储器（Dynamic Random Access Memory）
一种计算机内存类型，用于临时存储计算机程序和数据，以供中央处理器（CPU）快速访问。与静态随机存取存储器（SRAM）相比，具有较高的存储密度和较低的成本，但速度较慢。它是计算机系统中最常用的内存类型之一，用于存储操作系统、应用程序和用户数据等内容。
DRAM 的每个存储单元由一个电容和一个晶体管组成，电容负责存储数据位（0 或 1），晶体管用于读取和刷新数据。由于电容会逐渐失去电荷，因此需要定期刷新（称为刷新操作）以保持数据的正确性，这也是称为“动态”的原因，用于临时存储数据和程序，提供快速访问速度和相对较低的成本。

存储类型 结构 工作原理 性能 应用
DRAM（Dynamic Random Access Memory） 一种基本的内存技术，通常以单层平面的方式组织，存储芯片分布在一个平面上 当读取数据时，电荷被传递到输出线路，然后被刷新。当写入数据时，电荷被存储在电容中。由于电容会逐渐失去电荷，因此需要周期性刷新来保持数据 具有较高的密度和相对较低的成本，但带宽和延迟相对较高 常用于个人电脑、笔记本电脑和普通服务器等一般计算设备中
GDDR（Graphics Double Data Rate） 专门为图形处理器设计的内存技术，具有较高的带宽和性能 在数据传输速度和带宽方面优于传统的 DRAM，适用于图形渲染和视频处理等需要大量数据传输的应用 GDDR 与标准 DDR SDRAM 类似，但在设计上进行了优化以提供更高的数据传输速度。它采用双倍数据速率传输，即在每个时钟周期传输两次数据，提高了数据传输效率 主要用于高性能图形处理器（GPU）和游戏主机等需要高带宽内存的设备中
HBM（High Bandwidth Memory） 使用堆叠设计，将多个 DRAM 存储芯片堆叠在一起，形成三维结构 堆叠设计允许更短的数据传输路径和更高的带宽，同时减少了功耗和延迟。每个存储芯片通过硅间连接（Through Silicon Via，TSV）与其他存储芯片通信，实现高效的数据传输 具有非常高的带宽和较低的延迟，适用于高性能计算和人工智能等需要大量数据传输的领域 主要用于高端图形处理器（GPU）、高性能计算系统和服务器等需要高带宽内存的设备中

不同存储和传输的带宽和计算强度进行比较，假设 HBM 计算强度为 100，L2 缓存的计算强度只为 39，意味着每个数据只需要执行 39 个操作，L1 的缓存更少，计算强度只需要 8 个操作，这个时候对于硬件来说非常容易实现。这就是为什么 L1 缓存、L2 缓存和寄存器对 GPU 来说如此重要。可以把数据放在 L1 缓存里面然后对数据进行 8 个操作，使得计算达到饱和的状态，使 GPU 里面 SM 的算力利用率更高。但是 PCIe 的带宽很低，整体的时延很高，这将导致整体的算力强度很高，算力利用率很低。
DataLocation Bandwidth(GB/sec) ComputeIntensity Latency(ns) Threads Required
L1 Cache 19,400 8 27 32,738
L2 Cache 4,000 39 150 37,500
HBM 1,555 100 404 39,264
NVLink 300 520 700 13,125
PCIe 25 6240 1470 2297

在带宽增加的同时线程的数量或者线程的请求数也需要相对应的增加，这个时候才能够处理并行的操作，每个线程执行一个对应的数据才能够把算力利用率提升上去，只有线程数足够多才能够让整个系统的内存处于忙碌的状态，让计算也处于忙碌的状态，因此看到 GPU 里面的线程数非常多。
GPU 线程原理

GPU 整体架构和单个 SM（Streaming Multiprocessor）的架构，SM 可以看作是一个基本的运算单元，GPU 在一个时钟周期内可以执行多个 Warp，在一个 SM 里面有 64 个 Warp，其中每四个 Warp 可以单独进行并发的执行，GPU 的设计者主要是增加线程和增加 Warp 来解决或者掩盖延迟的问题，而不是去减少延迟的时间。
GPU 整体架构与 SM 架构

为了有更多的线程处理计算任务，GPU SMs 线程会选择超配，每个 SM 一共有 2048 个线程，整个 A100 有 20 多万个线程可以提供给程序，在实际场景中程序用不完所有线程，因此有一些线程处于计算的过程中，有一些线程负责搬运数据，还有一些线程在同步地等待下一次被计算。很多时候会看到 GPU 的算力利用率并不是非常的高，但是完全不觉得它慢是因为线程是超配的，远远超出大部分应用程序的使用范围，线程可以在不同的 Warp 上面进行调度。
Pre SM A100
Total Threads 2048 221,184
Total Warps 64 6,912
Active Warps 4 432
Waiting Warps 60 6,480
Active Threads 128 13,824
Waiting Threads 1,920 207,360
小结

本节首先从架构层面分析了 CPU 和 GPU 的主要区别，因为 CPU 的设计目标是尽可能在低的延迟下执行任务，GPU 的设计目标是最大化吞吐量，因此决定了 CPU 适合处理顺序执行的任务，GPU 适合处理大规模并行计算的任务。

以 $AX+Y$ 为例讲解了并发和并行的区别以及和串行的区别，在串行计算时内存利用率很低，当把程序用并发的方式展开，并发操作的多流水线会受到 CPU 架构的限制，当程序用并行的方式循环展开时，程序的执行则只受到线程数量和内存请求的约束，因此 CPU 跟 GPU 的本质区别是并行的问题而不是并发的问题，GPU 通过大量的线程提供并行的能力。

为了提供并行的能力，GPU 通过多级缓存、多级流水、多级 Cache 提供并行的机制，同时可以尽可能减少内存的时延。为了将数据充分利用起来，引入了 GPU 线程的原理，GPU 里面提供了大量超配的线程去完成对不同层级数据的搬运和计算。
GPU 硬件基础概念

A100 GPU 架构中 GPC（Graphic Processing Cluster）表示图像处理簇，一共有 8 个。共有两个 L2 Cache 并且可以互相实现数据同步，通过 Memory Controller 实现与高带宽存储器 HBM2（High Bandwidth Memory）进行数据交换。
每个 GPC 中包含 TPC（Texture processing cluster）表示纹理处理簇，每个处理簇被分为多个 SM（Streaming Multiprocessors）流处理器，SM 中包含多个 CUDA core 和 Tensor Core，用于处理图形图形和 AI 张量计算。
SM（Streaming Multiprocessors）称作流式多处理器，核心组件包括 CUDA 核心、共享内存、寄存器等。SM 包含很多为线程执行数学运算的 core，是英伟达 GPU 的核心，在 CUDA 中可以执行数百个线程、一个 block 上线程放在同一个 SM 上执行，一个 SM 有限的 Cache 制约了每个 block 的线程数量。

SM 主要组成如表所示，以英伟达 GP 100 为例，一共有 64 个 CUDA Core，Register File 存储大小为 256 KB，Shared Memory 内存大小为 64 KB，Active Thread 总线程数量是 2048，Active Block 数量是 32，Active Grid 数量是 8。
CUDA Core 向量运算单元 FP32-FPU、FP64-DPU、INT32-ALU
Tensor Core 张量运算单元 FP16、BF16、INT8、INT4
Special Function Units 特殊函数单元 超越函数和数学函数，例如反平方根、正余弦等
Warp Scheduler
线程束调度器 XX Thread/clock
Dispatch Unit 指令分发单元 XX Thread/clock
Multi Level Cache 多级缓存 L0/L1 Instruction Cache、L1 Data Cache & Shared Memory
Register File 寄存器堆
Load/Store 访问存储单元 LD/ST，负责数据处理

SP（Streaming Processor）流处理器是最基本的处理单元，最后线程具体的指令和任务都是在 SP 上进行处理的，GPU 在进行并行计算时就是很多个 SP 同时处理。在 Fermi 架构之后，SP 被改称为 CUDA Core，通过 CUDA 来控制具体的指令执行。

SP（Streaming Processor）流处理器是最基本的处理单元，最后线程具体的指令和任务都是在 SP 上进行处理的，GPU 在进行并行计算时就是很多个 SP 同时处理。在 Fermi 架构之后，SP 被改称为 CUDA Core，通过 CUDA 来控制具体的指令执行。

在 Fermi 架构中，通过 CUDA 来控制具体的指令执行，是最小的运算执行单元。所以对于现在的 NVIDIA GPU 架构来讲，流处理器的数量就是 CUDA Core 的数量。一个 SM 中包含了 2 组各 16 个 CUDA Core，每个 CUDA Core 包含了一个整数运算单元 ALU（Arthmetic Logit Unit）和一个浮点运算单元 FPU（Floating Point Unit）。

Volta 架构取消 CUDA core，变为单独的 FP32 FPU 和 INT32 ALU，因为 FP32:INT32 是 1:1 的关系，因此还是可以将它们合并起来一起称为原来的 CUDA Core，这样做的好处是每个 SM 现在支持 FP32 和 INT32 的并发执行，同时新增了光线追踪 RT Core。

Warp 是线程束，逻辑上所有 Thread 并行执行，但是从硬件的角度讲并不是所有的 Thread 能够在同一时刻执行，因此引入 Warp。Warp 是 SM 基本执行单元，一个 Warp 包含 32 个并行 Thread（warp_size=32），这 32 个 Thread 执行 SIMT（Single Instruction Multiple Thread）指令模式。

也就是说，所有的 Thread 以锁步的方式执行同一条指令，但是每个 Thread 会使用各自的 Data 执行指令分支。如果在 Warp 中没有 32 个 Thread 需要工作，那么 Warp 虽然还是作为一个整体运行，但这部分 Thread 是处于非激活状态。此外，Thread 是最小的逻辑单位，Warp 是硬件执行单位。
CUDA 基本概念

2006 年 11 月，NVIDIA 推出 CUDA（Compute Unified Device Architecture），通用并行计算架构（Parallel Computing Architecture）和编程模型（Programming Model），利用 GPU 的并行处理能力，将 GPU 用作通用并行计算设备，以加速各种计算任务，而不仅限于图形处理。

CUDA 编程模型
允许开发人员在 GPU 上运行并行计算任务，基于 LLVM 构建了 CUDA 编译器，开发人员可以使用 CUDA C/C++语言编写并行程序，通过调用 CUDA API 将计算任务发送到 GPU 执行。CUDA 编程模型包括主机（CPU）和设备（GPU）之间的协作，此外还提供了对其它编程语言的支持，比如 C/C++，Python，Fortran 等语言，支持 OpenCL 和 DirectCompute 等应用程序接口。
CUDA 在软件方面由一个 CUDA 库、一个应用程序编程接口（API）及其运行库（Runtime）、两个较高级别的通用数学库，即 CUFFT 和 CUBLAS 组成。CUDA TOOLKIT 包括编译和 C++核，CUDA DRIVER 驱动 GPU 负责内存和图像管理。CUDA-X LIBRARIES 主要提供了机器学习（Meachine Learning）、深度学习（Deep Learning）和高性能（High Performance Computing）计算方面的加速库，APPS & FRAMEWORKS 主要对接 Tensorflow 和 Pytorch 等框架。
CUDA 线程层次结构

CUDA 最基本的执行单位是线程（Thread），图中每条曲线可视为单个线程，大的网格（Grid）被切分成小的网格，其中包含了很多相同线程数量的块（Block），每个块中的线程独立执行，可以通过本地数据共享实现数据交换同步。因此对于 CUDA 来讲，就可以将问题划分为独立线程块，并行解决的子问题，子问题划分为可以由块内线程并行协作解决。
CUDA 引入主机端（host）和设备（device）概念，CUDA 程序中既包含主机（host）程序也包含设备（device）程序，host 和 device 之间可以进行通信，以此来实现数据拷贝，主机负责管理数据和控制程序流程，设备负责执行并行计算任务。在 CUDA 编程中，Kernel 是在 GPU 上并行执行的函数，开发人员编写 Kernel 来描述并行计算任务，然后在主机上调用 Kernel 来在 GPU 上执行计算。

代码 cuda_host.cpp 是只使用 CPU 在 host 端实现两个矩阵的加法运算，其中在 CPU 上计算的 kernel 可看作是加法运算函数，代码中包含内存空间的分配和释放。

#include
#include <math.h>
#include <sys/time.h>

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y[i] = x[i] + y[i];
}

int main(void)
{
int N = 1<<25; // 30M elements

float *x = new float[N];
float *y = new float[N];

// initialize x and y arrays on the host
for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
}

struct timeval t1,t2;
double timeuse;
gettimeofday(&t1,NULL);

// Run kernel on 30M elements on the CPU
add(N, x, y);

// Free memory
delete [] x;
delete [] y;

return 0;

}

在 CUDA 程序架构中，host 代码部分在 CPU 上执行，是普通的 C 代码。当遇到数据并行处理的部分，CUDA 会将程序编译成 GPU 能执行的程序，并传送到 GPU，这个程序在 CUDA 里称做核(kernel)。device 代码部分在 GPU 上执行，此代码部分在 kernel 上编写(.cu 文件)。

kernel 用global符号声明，在调用时需要用<<>>来指定 kernel 要执行及结构。代码 cuda_device.cu 是使用 CUDA 编程实现 GPU 计算，代码涉及到 host（CPU）和 device（GPU）相关计算，使用global 声明将 add 函数转变为 GPU 可执行的 kernel。

#include
#include <math.h>

// Kernel function to add the elements of two arrays
// global 变量声明符，作用是将 add 函数变成可以在 GPU 上运行的函数
// global 函数被称为 kernel
global
void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y[i] = x[i] + y[i];
}

int main(void)
{
int N = 1<<25;
float *x, *y;

// Allocate Unified Memory – accessible from CPU or GPU
// 内存分配，在 GPU 或者 CPU 上统一分配内存
cudaMallocManaged(&x, Nsizeof(float));
cudaMallocManaged(&y, Nsizeof(float));

// initialize x and y arrays on the host
for (int i = 0; i < N; i++) {
x[i] = 1.0f;
y[i] = 2.0f;
}

// Run kernel on 1M elements on the GPU
// execution configuration, 执行配置
add<<<1, 1>>>(N, x, y);

// Wait for GPU to finish before accessing on host
// CPU 需要等待 cuda 上的代码运行完毕，才能对数据进行读取
cudaDeviceSynchronize();

// Free memory
cudaFree(x);
cudaFree(y);

return 0;
}

因此 CUDA 编程流程总结为：

编写 Kernel 函数描述并行计算任务。
在主机上配置线程块和网格，将 Kernel 发送到 GPU 执行。
在主机上处理数据传输和结果处理，以及控制程序流程。

为了实现以上并行计算，对应于 GPU 硬件在进行实际计算过程时，CUDA 可以分为 Grid，Block 和 Thread 三个层次结构：

线程层次结构Ⅰ-Grid：kernel 在 device 上执行时，实际上是启动很多线程，一个 kernel 所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid 是线程结构的第一层次。
线程层次结构Ⅱ-Block：Grid 分为多个线程块（block），一个 block 里面包含很多线程，Block 之间并行执行，并且无法通信，也没有执行顺序，每个 block 包含共享内存（shared memory），可以共享里面的 Thread。
线程层次结Ⅲ-Thread：CUDA 并行程序实际上会被多个 threads 执行，多个 threads 会被群组成一个线程 block，同一个 block 中 threads 可以同步，也可以通过 shared memory 通信。

算力峰值计算

GPU 的算力峰值是指 GPU 在理想情况下能够达到的最大计算能力，通常以浮点运算每秒（FLOPS）为单位进行衡量，GFLOPS（每秒十亿次浮点运算），TFLOPS（每秒万亿次浮点运算）。计算 GPU 的算力峰值可以帮助开发人员了解其在理论上的最大性能，并进行性能预测和优化，更好地评估大模型训练过程中的算力利用率。

GPU 的算力峰值通常由以下几个因素决定：

CUDA 核心数量：每个 CUDA 核心可以执行一个线程，GPU 的算力峰值与 CUDA 核心数量成正比。
核心频率：GPU 的核心频率越高，每个核心每秒钟能够执行的指令数就越多。
每个核心的计算能力：不同型号的 GPU 具有不同的计算能力，通常以每个核心每个时钟周期能够执行的浮点指令数（FLOPS）为单位进行衡量。
并行度：GPU 的并行度决定了其能够同时执行的线程数量，从而影响了算力峰值。

计算 GPU 的算力峰值可以使用以下公式：

其中，

$F_{\text{clk}}$：GPU 时钟周期内指令执行数 (FLOPS/Cycle)
$N_{\text{SM}}$：SM（Streaming Multiprocessor）数量
$F_{\text{req}}$：Tensor Core 核心运行频率(GHz)

以 NVIDIA A100 为例，其中 FP32 Tensor Core 指令吞吐 64 FLOPS/Cycle ，核心运行频率为 1.41GHz ，SM 数量为 108 ，因此 GPU 的算力峰值是，19,491 GFLOPS，大约为 1.95 TFLOPS：

NVIDIA A100 GPU 算力峰值
Peak FP641 9.7 TFOPS
Peak FP64 Tensor Core1 19.5 TFOPS
Peak FP321 19.5 TFOPS
Peak FP161 78 TFOPS
Peak BF161 39 TFOPS
Peak FP32 Tensor Core1 156 TFOPS | 312 TFOPS2
Peak FP16 Tensor Core1 312 TFOPS | 624 TFOPS2
Peak BF16 Tensor Core1 312 TFOPS | 624 TFOPS2
Peak INT8 Tensor Core1 624 TFOPS | 1,248 TFOPS2
Peak INT4 Tensor Core1 1,248 TFOPS | 2,496 TFOPS2
1 - Peak rates are based on GPU Boost Clock.
2 - Effective TFLOPS/TOPS using the new Sparsity feature
小结

本节主要对 NVIDIA GPU 硬件相关的基础概念进行了讲解，以 A100 GPU 为例，GPU 架构包括 GPC（图形处理簇，Graphics Processing Clusters）、GPC 包含 TPC（纹理处理簇，Texture Processing Clusters）、TPC 包含 SM（流多处理器，Stream Multiprocessors），SM 又包含 SP（流处理器，Streaming Processor，在 Fermi 架构之后，SP 被改称为 CUDA Core），Wrap 是 GPU 执行程序时的调度单位，SM 的基本执行单元。

CUDA 提出的通用并行计算架构和编程模型将 GPU 用作通用并行计算设备，加速各种计算任务，CUDA 编程模型包括主机（CPU）和设备（GPU）之间的协作，支持多种编程语言，在软件方面由 CUDA 库、应用程序编程接口（API）及其运行库（Runtime）和通用数学库构成。

CUDA 利用 Grid，Block 和 Thread 三个线程层次结构，将并行计算任务分为多个线程，这些线程被组织为块，块可以进一步组织为网格，每个线程都可以独立执行计算任务。在 CUDA 编程中，Kernel 是在 GPU 上并行执行的函数，开发人员编写 Kernel 来描述并行计算任务，然后在主机上调用 Kernel 来在 GPU 上执行计算。软件方面的线程 Thread 对应硬件 CUDA Core，Thread Block 对应硬件 SM，多个 Block 组成的 Grid 组成了 TPC 和 GPC 硬件，从而实现了并行计算。

GPU 的算力峰值与 CUDA 核心数量、核心频率、每个核心的计算能力和并行度相关，计算 GPU 的算力峰值可以帮助开发人员了解其在理论上的最大性能，并进行性能预测和优化。