HGX

一个典型配置就是 8 卡和 4 个 NVSwitch。

HGX 平台上的 GPU 应该只有 NVLink 接口，没有 PCIe 接口，也就是说只能通过 NVLink 连接。

NVIDIA 自己造的超级计算平台。HGX 是一个计算模组，没有 CPU，没有操作系统。

和 JBOG 的关系感觉就是：JBOG 是对于 HGX 上的又一层封装，比如浪潮出过 16 卡的 JBOG（里面包含了两个 HGX）。

HGX 还是包含 BF DPU 的。

GPU Box / JBOG (Just a Bunch of GPUs)

一种新型服务器架构。

JBOG 通过将多个 GPU 模块集中在一个专用机箱中运行，并通过高速互联技术与主机服务器连接，提供了一个灵活且高效的解决方案。这个专用机箱 JBOG，一般叫做机尾；而控制 JBOG 的真正服务器，也就是传统的带有 CPU 和内存的控制节点，叫做机头（Head Node）。

JBOG 的命名方式和设计理念，类似于前几年存储服务节点 JBOD（Just a Bunch of Disks）和 JBOF（Just a Bunch of Flash）。JBOD 是最简单的存储方案，就是将多个硬盘简单地组合在一起，形成一个存储池。随着闪存技术的发展，JBOF 应运而生，它采用同样的理念，但使用闪存替代了传统硬盘，提供更快的存储性能。

和 HGX^ 的区别，

AI算力底座技术：GPU/NPU box是什么？JBOG是什么？什么又是HIB、OAM、OAI、UBB？ - 知乎

GPU 知识体系学习文档

Modal 是一个大模型推理托管平台，他们的 developer 们深入研究了 GPU 相关的各种知识并放到了这个在线文档里面，比较杂，可以一看：GPU Glossary

Nvidia Confidential Computing 学习文档：[[2025-04-17-NVIDIA-GPU-CC#Resources for NVIDIA Confidential Computing]]

nvidia-smi

nvcc

GPU 机间架构

查看 GPU 机内机间拓扑信息 / nvidia-smi topo -m

即使我们有 RDMA 网卡，如果没有启用 GDR 的话，那么输出也不会有 RDMA 网卡，而是仅仅有传统 TCP 网卡？是和 GDR 有关系吗？

nvidia-smi topo -m

一个典型的输出结果：

       GPU0    GPU1    GPU2    GPU3    GPU4    GPU5    GPU6    GPU7    NIC0    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      NV18    NV18    NV18    NV18    NV18    NV18    NV18    SYS     0-87    0               N/A
GPU1    NV18     X      NV18    NV18    NV18    NV18    NV18    NV18    SYS     0-87    0               N/A
GPU2    NV18    NV18     X      NV18    NV18    NV18    NV18    NV18    SYS     0-87    0               N/A
GPU3    NV18    NV18    NV18     X      NV18    NV18    NV18    NV18    SYS     0-87    0               N/A
GPU4    NV18    NV18    NV18    NV18     X      NV18    NV18    NV18    SYS     88-175  1               N/A
GPU5    NV18    NV18    NV18    NV18    NV18     X      NV18    NV18    SYS     88-175  1               N/A
GPU6    NV18    NV18    NV18    NV18    NV18    NV18     X      NV18    SYS     88-175  1               N/A
GPU7    NV18    NV18    NV18    NV18    NV18    NV18    NV18     X      SYS     88-175  1               N/A
NIC0    SYS     SYS     SYS     SYS     SYS     SYS     SYS     SYS      X

• NV#：表示使用 NVLink 传输，可以看到 GPU 之间都是使用 NVLink 直接传输的；
• SYS：跨 NUMA 节点通信，需要通过 CPU 间总线比如 UPI （也就是系统总线 SYS）来通信。所有的 GPU 和网卡通信都是需要通过 SYS 总线；
• NODE: 表示同一 NUMA 节点内的跨 PHB 通信，需要经过 PCIe 总线 + 节点内 PCIe Host Bridge 之间的互连，典型场景：两个 PCIe 设备连接到同一个 CPU 的不同 PHB（如 CPU 的 PCIe x16 和 PCIe x8 插槽）。一个 CPU 上可以有多个 PHB。
• PHB (PCIe Host Bridge): 表示通过单个 PCIe Host Bridge 的通信。设备直接连接到 CPU 内置的 PCIe 控制器，最佳 PCIe 直连情况，延迟约 100ns，带宽可达 PCIe x16 的满速。
• PXB (PCIe Bridge)：表示跨多个 PCIe 交换机的通信，经过 PCIe Switch 层但未到达 PHB。
• PIX：表示单级 PCIe 交换的通信，最多经过一个 PCIe Switch。

NVLink > PIX > PXB > PHB > NODE > SYS。虽然 PHB 和 PXB 都有 Bridge 这个后缀，但是其实表示的是交换机。

更形象的解释：

• GPU8 和 GPU9 通过一跳 PCIe switch 就能相连，这叫 PIX（Connection traversing at most a single PCIe bridge，需要 PLX 等特殊的硬件支持）
• GPU8 和 GPU10 通过两跳 PCIe switch 才能相连，这叫 PXB（Connection traversing multiple PCIe bridges but without traversing the PCIe Host Bridge，需要 PLX 等特殊的硬件支持）
• GPU8 和 GPU12 通过 CPU 根节点 PCIe host bridge 才能相连，这叫 PHB（PHB = Connection traversing PCIe as well as a PCIe Host Bridge，没有特殊硬件支持的话，同一个 CPU 上的设备都要通过 PCIe host bridge 才能相连）
• GPU8 和 GPU0 要跨过 CPU 节点之间的 QPI 才能相连，这叫 SYS（Connection traversing PCIe as well as the SMP interconnect between NUMA nodes, e.g., QPI/UPI）

NIXL

NIXL: NVIDIA Inference Xfer Library，随着 Dynamo 开源出来的。这个重点是在于：

• 优化 AI 推理框架的端到端通信。
• Providing an abstraction over various types of memory (e.g., CPU and GPU) and storage (e.g., file, block and object store)。

NVSHMEM

OpenSHMEM

GPUDirect / GPDR / GDR / GPUDirect RDMA

同一台机器的卡间可以使用 NVLink 通信，那么不同机器的机间可以使用 GDR 来通信。

GDR 技术是基于 PCIe 标准的。GDR 技术其实就是让 RDMA 网卡和 GPU 之间可以通过 PCIe 进行 DMA，不再需要 CPU 侧以及主机内存参与了。

GPUDirect RDMA 的演进与实现 - 知乎

GDR 的合适使用场景是什么呢，比较推荐的场景就是 GPU 与第三方设备在同 PCIe switch 下的场景，这种情况下是存在性能增益的。

技术改变AI发展：RDMA能优化吗？GDR性能提升方案（GPU底层技术系列二）-阿里云开发者社区

如何使用和关闭 GDR？

# 开启 GDR
export NCCL_NET_GDR_LEVEL=SYS
# 关闭 GDR
export NCCL_NET_GDR_LEVEL=LOC

See: Environment Variables — NCCL 2.26.5 documentation

GDR 测试

先 nvcc --version 检查下 CUDA 有没有装。因为 GDR 本身是 CUDA 包的一部分。

GDR 测试依赖 NCCL 吗？

好像是不依赖的，相反，NCCL 本身是依赖 GDR 来进行运行的，因为 NCCL 是一个集合通信库，更偏上层一些。

# 一般来说就是下面这个测试模版
ib_write_bw --use_cuda <gpu index>

# 一打一测试（发送端和接收端唯一的区别在于一个指定了目的地址，一个没有）
# 使用 read 是因为使用了 1-side 测试而不是基于 send recv 的 2-side 测试
# 这个应该是从一边的显存 GDR copy 到另一边的显存当中。
# ib_read_bw 这个命令就是如果指定了 ip 地址就是服务端，否则就是客户端。
# client 端
ib_read_bw -F -d mlx5_0 -x 3 -q 1 -s 1K --run_infinitely --report_gbits -p 10000 --use_cuda 0 192.168.0.24
# server 端
ib_read_bw -F -d mlx5_0 -x 3 -q 1 -s 1K --run_infinitely --report_gbits -p 10000 --use_cuda 0

# 二打二测试
# client 端（可见 client 端也使用了不同的 GPU（但是使用了同一网卡，那么就存在竞争的情况），通过同一 IP 的不同的端口连接服务端）
ib_read_bw -F -d mlx5_0 -x 3 -q 1 -s 64K --run_infinitely --report_gbits -p 10000 --use_cuda 0 192.168.0.24
ib_read_bw -F -d mlx5_0 -x 3 -q 1 -s 64K --run_infinitely --report_gbits -p 10001 --use_cuda 1 192.168.0.24

多机多卡训练问题排查

前备：

• 确保所有机器安装相同版本的 CUDA、NVIDIA 驱动和 NCCL。
• 确保主节点可以免密登录到所有节点（如 ssh hostname 无需密码）。
• 确保所有机器上的 nccl-tests 可执行文件路径一致，或手动同步编译后的二进制文件。

nccl-tests

git clone https://github.com/NVIDIA/nccl-tests.git
cd nccl-tests
make NCCL_HOME=/path/to/nccl  # 指定 NCCL 安装路径（默认 /usr/local/nccl）

# 单机 8 卡测试 AllReduce
./build/all_reduce_perf -b 8 -e 256M -f 2 -g 8

# 多机多卡测试
# host 通常可以包含 master 节点，因为一般来说 master 节点也希望参与到计算中去而不仅仅是调度
mpirun -np 16 -H 192.168.0.24:8,192.168.0.19:8 -x NCCL_DEBUG=INFO -x LD_LIBRARY_PATH --allow-run-as-root ./build/all_reduce_perf -b 8 -e 256M -f 2 -g 8

NCCL

NCCL 同时涉及到 GPU 卡间架构和 GPU 机间架构，但是为了通用性所以放在 GPU 机间架构这一部分。

NCCL 和 CUDA 的区别就在于一个为了使能计算，一个为了使能通信。

NVIDIA Collective Communications^ Library (NCCL pronounced “Nickel”)。NCCL 和其他集合通信库一样，都需要抽象“节点”这个概念，不过 NCCL 里一个节点应该是一个 GPU 而不是一个机器吧？

Collective Communication / Collective Operation 请搜索^。

NVIDIA 的技术 PDF：NCCL

NCCL 下面也依赖不同的物理层通信协议来通信：

• 在绝大多数情况下都可以通过服务器内的 PCIe、NVLink、NVSwitch 等
• 服务器间的 RoCEv2、IB、TCP 网络实现高带宽和低延迟。

因此：NCCL 屏蔽了底层复杂的细节，向上提供 API 供训练框架调用，向下连接机内机间的 GPU 以完成模型参数的高效传输。

相关笔记及文档：

• [[分布式训练#分布式训练通信后端]]

如何查看 NCCL 是否已经安装以及版本？/ NCCL 安装路径 / 如何检查 NCCL 环境已经设置好了？

查看动态链接库是否已经包含了 nccl：

ldconfig -p | grep libnccl

可以问一下 deepseek 看一下有什么输出：

git clone https://github.com/NVIDIA/nccl-tests.git
cd nccl-tests
make

查看 NCCL 版本：

cat /usr/include/nccl.h | grep "NCCL_VERSION_CODE"

NCCL 如何使用 RDMA 进行测试?

NCCL 默认会优先选择高性能协议（如 RDMA/InfiniBand），若不可用则回退到 TCP。也可以通过环境变量强制指定协议。

# 如果要使用 RDMA
export NCCL_DEBUG=INFO # 查看 NCCL 通信详细信息 (debug RDMA)
export NCCL_IB_GID_INDEX=3
export NCCL_IBEXT_DISABLE=0
export NCCL_IB_DISABLE=0
# 使用 GDR
export NCCL_NET_GDR_LEVEL=SYS

# 如果就是不想使用 RDMA
export NCCL_IBEXT_DISABLE=1
export NCCL_IB_DISABLE=1

NCCL Plugin / NCCL-RDMA-SHARP / SHArP (Scalable Hierarchical Aggregation Protocol)

是同一个东西。是提升通信性能的关键工具，它通过优化数据在网络中的传输方式，显著提高了大规模 GPU 集群的通信效率。

将集合操作部分卸载到交换机中进行，进行网内计算。将集合通信（Reduce、Allreduce、Barrier）卸载到网络（交换机）上进行。

GPU 卡间架构

/var/run/nvidia-topologyd

P2P (GPUDirect Peer-to-Peer)

主要用于单机 GPU 间的高速通信，它使得 GPU 可以通过 PCIe 直接访问目标 GPU 的显存。

两个 GPU 通过 PCIe 进行直接 DMA，没有主机 CPU/内存的参与。

NVLink / NVSwitch

本质上都是为了直接能够 P2P，绕过主机内存和 CPU，同时不需要占用 PCIe 带宽（相比于 P2P）。

不要被 NVLink 里的 Link 所迷惑觉得这是一种线材，其实和 PCIe 一样，都是集成在板子上的，因此不存在单独购买 NVLink 然后把两个 GPU 连起来的情况。

使用了 NVSwitch，每一个 GPU 需要连的线数仅仅和 Switch 数量相关，而 Switch 一般不会加很多，所以不会出现八卡 GPU 之间要互连但是每一个卡最多只有四个 NVLink 口的尴尬情况（仅仅比喻）。

NVSwitch 可以形成树形结构吗？

NVSwitch 无法形成级联的结构，无法形成树形的结构，每一个 NVSwitch 连接的不只是 GPU 卡。每个 NVSwitch 芯片还预留了相当数量的端口用于连接到其他 NVSwitch 芯片。这些 NVSwitch 间链路（ISL, Inter-Switch Links）是实现多芯片扩展和构建统一交换网络的关键。

NVLink/NVSwitch 的设计目标是为 GPU 提供超高速、超低延迟、完全对等的互连。多个 NVSwitch 之间的互连通常是高度互联的网状（mesh）或类似拓扑。

NVSwitch 之间需要连接吗？

需要连接，如果不需要连接，为什么还需要多个 NVSwitch 呢？直接一个 NVSwitch 接所有 GPU 就好了。多个 NVSwitch 芯片必须通过高速专用链路互相连接。这样的话，一个 GPU 到另一个 GPU 之间的通信就有多条链路可以走了，因此引申出下面问题：

两个 GPU 之间通信如何选路？

简而言之，有一个硬件实现的自适应路由算法。

GPU 单卡架构

从 CUDA kernels 如何被调度并执行讲起

同一时间，一个 NV GPU 上可以有多少个 kernel 同时执行？ Hyper-Q 之前，kernel 之间是串行的，因此在某一个时刻，只能执行一个 kernel。但是在 Hyper-Q 之后（Kepler），可以同时执行。

同一时间，一个 NV GPU 的 SM 上可以有多少个 kernel 同时执行？ 一个。但是一个 SM 上可以同时执行来自一个 kernel 的多个不同线程块（Block）的线程束（Warp），但这些线程块（Block）都属于同一个 Kernel。 不同 CUDA Kernel 不能同时在一个 SM 上混合执行。

一个 Block 是被分配给一个 SM 的，一个 SM 上可以分配多个 Warp。 这是资源分配的基本单位。一个 Block 内的线程会被组织成一个或多个 Warp。Warp 是 SM 上实际调度和执行的最小单位。一个 Block 必定是在一个 SM 上执行的（因为要共享该内核的资源），故目前来说一个 block 块最多 1024 个线程。并且在这个 SM 内，它的线程会被拆分成多个 Warp（硬件上的拆分）。这些 Warp 会在该 SM 的 Warp 调度器管理下，与其他 Block 的 Warp 交错执行（如果该 SM 上有多个 Block 的话）。

在同一时刻，一个 SM 中的不同 warp 可以同时执行来自不同 block 的任务。

一个 block 被划分出的 warp 不会固定在特定的物理资源上执行，而是由 SM 的硬件调度器根据 warp 的状态动态分配计算资源。

每个时钟周期，如果 Warp Scheduler 发现有 warp 空闲了，那么其就从队列中选择就绪的 warp（例如已完成上一指令、无内存阻塞）进行计算。 调度决策与 warp 所属的 block 无关——可能连续执行同一 block 的多个 warp，也可能交替执行不同 block 的 warp。

一图以蔽之：

https://k48xz7gzkw.feishu.cn/docx/GBa6dc0SFoRhDyxNYA7cNt7knNe?openbrd=1&doc_app_id=501&blockId=doxcn5pdj8kZMqP7epATmMrG5Lh&blockType=whiteboard&blockToken=AT1zwNf0OhUqR2bFbomc8EjjnBe#doxcn5pdj8kZMqP7epATmMrG5Lh

GPU 工作队列（Work Queue）

Nvidia driver 对上提供什么 API

SIMT vs. SIMD

首先明白什么是 SIMD：SIMD 是一种数据级并行的技术。一条指令，多个计算单元同时执行计算。MMX, SSE, AVX (Vector), AMX (Matrix) 都是 SIMD 技术。

比如我们有 4 个数字要加上 4 个数字，那么我们可以用这种 SIMD 的指令来 1 次完成本来要做 4 次的运算。这种机制的问题就是过于死板，不允许每个分支有不同的操作，所有分支必须同时执行相同的指令，必须执行没有例外。

相比之下 SIMT 就更加灵活了，虽然两者都是将相同指令广播给多个执行单元，但是 SIMT 的某些线程可以选择不执行，也就是说同一时刻所有线程被分配给相同的指令，SIMD 规定所有人必须执行，而 SIMT 则规定有些人可以根据需要不执行，这样 SIMT 就保证了线程级别的并行，而 SIMD 更像是指令级别的并行。

SIMT 包括以下 SIMD 不具有的关键特性：

• 每个线程都有自己的指令地址计数器；
• 每个线程都有自己的寄存器状态；
• 每个线程可以有一个独立的执行路径；

而上面这三个特性在编程模型可用的方式就是给每个线程一个唯一的标号 (blckIdx, threadIdx)，并且这三个特性保证了各线程之间的独立。

CUDA执行模型GPU架构 GPU架构是围绕一个流式多处理器（SM）的扩展阵列搭建的。通过复制这种结构来实现GPU的硬件 - 掘金

线程束（Warps）/ Warp scheduler

所有的内容可以在这里找到：

• 1. Introduction — CUDA C++ Programming Guide
• 这里讨论更准确：GPU architecture and warp scheduling - CUDA / CUDA Programming and Performance - NVIDIA Developer Forums

这是 H100 的的架构图，可以看到一个 SM 可以包含多个 Warp（比如 64 个）。注意 wrap scheduler 和 wrap 并不是一一对应的关系（比如 4 个 wrap scheduler 对应到 64 个 warp，每一个 warp 是否绑定到了一个 scheduler 上这个 NVIDIA 没有公开，所以需要讨论）。可以看到一个 warp 有 32 个线程（Warp 的大小固定是 32 个线程）：

CUDA 采用单指令多线程 SIMT 架构管理执行线程，不同设备有不同的线程束大小，但是到目前为止基本所有设备都是维持在 32，也就是说每个 SM 上有多个 block，一个 block 有多个线程（可以是几百个，但不会超过某个最大值），但是从机器的角度，在某时刻，SM 上虽然可以同时执行多个 warp，每一个 warp 内也就是 32 个线程可以同时同步执行，但是线程束中的每个线程执行同一条指令，但是使用的都是私有数据（比如 instruction address counter）和寄存器信息，因此可以有自己的 branch 选择执行与不执行。也就是以不同的数据资源执行相同的指令。线程同时在相同的程序地址启动。 从性能的角度来说，一个 warp 里 diverge 的 thread 越少，那么执行效率越高，所以我们软件编程应当尽量避免 diverge 的情况发生。

一般一个 SM 上最多调度 64 个 thread groups（一个 group 对应一个 warp）。

不同 block 的 thread 不能在同一个 warp 里面，也就是一个 warp 需要整着切 block 而不是混着切。

线程束是 SM 里的硬件概念。

Warp scheduler 调度是硬件行为，一个 SM 上 warp scheduler 会自动感知是否有 warp 是等待的状态，如果是然后切换到另外一个 warp。调度基本上 1ns 发生一次（CPU 上因为要 save/restore context 基本上要 1ms 一次），每次调度，warp 会被 switched out 换入换出。GPU 上的 warp schedule 不需要任何的 save/restore。这是因为每一个线程（指软件线程）都有自己的 register file（从 warp 处申请的，见上图里每一个 warp 里的共享的 register file）。

Warp 和 block，一个是硬件层面的线程集合，一个是逻辑层面的线程集合。 一旦 block 被调度在一个 SM 上，线程块中的线程会被进一步划分为 wrap，划分逻辑非常直白，block 里的所有 thread 按照线程号的连续顺序来划分。因此线程束（wraps）是比线程块（blocks）更小的单位。

了解 warp 可以帮助理解和优化特定 CUDA 设备上 CUDA 应用程序的性能。

SM详解与Warp Scheduler，合理块和线程的数量对GPU利用率非常重要 - 知乎

线程束执行本质 - 知乎

流式多处理器（SM, Streaming Multiprocessor）

Nvidia 特有的硬件概念，每个 GPU 通常有多个 SM，当一个 kernel 的 grid 被启动的时候，多个 block 会被同时分配给可用的 SM 上执行。

当一个 block 被分配给一个 SM 后，他就只能在这个 SM 上执行了，不可能重新分配到其他 SM 上了，多个 block 可以被分配到同一个 SM 上（一个 SM 上可以同时跑多个 blocks 吗，而不仅仅是顺序执行）。因此同一块中的线程能够以不同于不同块之间的线程的方式相互交互（因为 share 同一个 SM）。也就是说，block 和 SM 是多对一的关系。

在 SM 上同一个 block 内的多个线程进行线程级别并行，而同一线程内，指令利用指令级并行将单个线程处理成流水线。

为什么 GPU 要划分 SM？

• 主要目的是提高 GPU 的并行计算能力和资源利用率。GPU 就可以通过将计算任务分解成多个小部分的工作分配给不同的 SM 并行执行，从而加快计算速度。
• 避免不同计算任务之间的资源竞争，提高 GPU 并行性能。

体感认识：Ampere A100 GPU 具有 108 个 SM，每个 SM 有 64 个核心，总共在整个 GPU 中有 6912 个核心。H100 GPU 有 132 个 SM。所以从这种角度来说，一个 SM 更像对应一个 CPU 里一个 core 的概念，比如 AMD CPU 有 256 个 core。

一个 SM 不可以同时被分配不同 CUDA kernel (grid) 的 block，这就是说，如果一块 GPU 需要同时执行多个 kernel，那么应该将它们分配给不同的 SM 来一起执行，而不是在 SM 里面分配不同的 block 来执行。这里面分两部分调度：

• 是由 GPU 硬件调度器来进行而不是软件驱动进行的。

一般不支持乱序执行。

下图区分硬件和软件上的概念：

GPU 如何知道自己应该运行什么程序？

显存 / GPU Memory

显存大小和模型大小的对应关系：如果只是进行推理的话，还是比较容易计算的。目前模型的参数绝大多数都是 float32 类型, 占用 4 个字节。所以一个粗略的计算方法就是，每 10 亿（1B）个参数，占用 4G 显存 (实际应该是 $10^9*4/1024/1024/1024=3.725G$，为了方便可以记为 4G)。

HBM2 / HBM2e / GDDR

因为 DDR^ 是面向内存的，而内存带宽在 GPU 领域里还是不太够。因此如果我们想要以更快的速度访问显存的话，就需要 GDDR 了。

DDR 在下降沿和上升沿都传输，所以叫做 Double Rate，而 GDDR (Graphic DDR) 通过定义了多个沿来传输，实现了四倍甚至更高的速度。

Nvidia GPU 架构演进 / A10

截止到 2024 年的所有 Nvidia GPU 架构：NVIDIA GPU 核心与架构演进史 – 陈少文的网站

Ampere 和 Hopper 之间也是有一个新的架构叫做 Ada Lovelace。是一个全新的架构，应该不属于 Tesla 架构。

区分系列名，架构名和型号名：

• GeForce 等等都是系列名；
• Ampere, Hopper 都是架构名；
• A100, H100 都是型号名。

可以看到，后面出的 CPU 型号名和架构名都是有关联的，比如：

• Ampere 架构：型号名 A100
• Hopper 架构：型号名 H100
• Blackwell 架构：型号名 B100

FYI: 3070ti 显卡用的是 Ampere 架构。

A10: NVIDIA A10 datasheet

• 24GB 显存。

A100: NVIDIA A100 | Tensor Core GPU

• 80GB 显存。

B200 和 GB200 区别

B200 表示的单纯是 GPU，而 GB200 是芯片的“组合”，是通过一个板子将 2 颗 B200 加上一颗 Grace CPU（72 核心的 ARM 架构处理器）组合而成。

H200, H100, H20 区别

H20 是中国特供，H200, H100 GPU core 架构都是一样的，主要区别在于内存容量和带宽上有区别。

H100 PCIe vs. H100 NVL vs. H100 SXM

三种不同的接口，这里能够看到他们的区别：NVIDIA Announces H100 NVL - Max Memory Server Card for Large Language Models

NVIDIA 的 GPU 需要哪些固件

VBIOS

NVIDIA VBIOS 是 NVIDIA 显卡的 固件（Firmware），全称为 Video BIOS 或 显卡 BIOS。它是存储在显卡硬件芯片中的一段程序，负责显卡的底层初始化、硬件控制和与操作系统/驱动程序的通信。

AMD GPU 架构演进

RDNA：架构名字，比如 RDNA, RDNA 2, RDNA 3, RDNA 4。

架构系列：

• Vega 架构：发布于 2017 年，特点是采用高带宽缓存控制器（HBCC），以及下一代计算单元设计，支持高级图形与高性能计算任务。比如 Radeon RX Vega 64/56 显卡。
• RDNA 2/3/4 架构：2019 年推出的第一代 RDNA 架构是 GCN 架构的重大革新，旨在提供更高的能效比，并首次应用于 Radeon RX 5000 系列显卡上，如 RX 5700 XT 和 RX 5700。偏向游戏玩家。
• CDNA 架构：专注于数据中心和高性能计算市场，具备高度优化的计算性能，适用于机器学习、深度学习、科学计算等领域。

主要 GPU 产品线：

• 消费级桌面显卡： • Radeon RX 系列：例如 RX 5000 系列、RX 6000 系列等，为个人电脑玩家和内容创作者提供中高端图形处理能力。
• 移动版显卡： • Radeon RX 移动版：针对笔记本电脑市场，如 Radeon RX 5000M 系列和最新的 RX 6000M 系列，为游戏本和平板电脑提供强劲的图形性能。
• 专业显卡： • Radeon Pro WX 系列：面向专业工作站用户，如 WX 7100、WX 8200 以及基于 Vega 架构的 WX 9100 等型号，满足 CAD、渲染、建模等专业应用需求。

AMD GPU体系知识大全_vega架构和rdna架构-CSDN博客

GPU 图形渲染

着色器（Shader）

是一个软件程序，不是一个硬件单元。

顶点着色器（Vertex Shader）

像素着色器（Pixel Shader）

CUDA

CUDA 是用户态的 Driver，NVIDIA Driver 是内核态的东西。所以一般的玩法是 CUDA 装在业务容器中，Driver 状态宿主机或者虚拟机中： libcuda.so，就是一个二进制库：

ls -l /usr/lib/x86_64-linux-gnu/libcuda.so

Compute Unified Device Architecture.

目前 NVIDIA GPU 在硬件上同一时刻只能运行一个 CUDA Context（对应一个 host 上的进程），多个 CUDA 进程会因为争抢资源而等待。只有一个释放算力后，另一个才能获得算力。因此，为了稳定性，通常一张 GPU 卡只分配给一个容器使用。

CUDA 提供了对其它编程语言的支持，如 C/C++，Python，Fortran 等语言。CUDA 程序中既包含 host 程序，又包含 device 程序，它们分别在 CPU 和 GPU 上运行。

典型的 CUDA 程序的执行流程如下：

• 分配 host 内存，并进行数据初始化；
• 分配 device 内存，并从 host 将代码和数据拷贝到 device 上。因为 CPU 本身代码数据也都位于内存中进行计算，因此 GPU 也可以根据 device 内存（显存）里的代码和数据来进行计算；
• 调用 CUDA 的核函数在 device 上完成指定的运算（应该是类似发送一个信号的方式？）；
• 将 device 上的运算结果拷贝从显存拷贝到 host 内存上；
• 释放 device 和 host 上分配的内存。

可见内存或者说 DMA 是非常重要的，因为 GPU 的输入数据以及数据的输出都需要通过 DMA 的方式从 CPU 拷贝或者直接拷贝到 CPU 侧内存。

一些代码在 CPU 上跑，一些代码在 GPU 上跑，需要区分 host 和 device 上的代码，在 CUDA 中是通过函数类型限定词区别 host 和 device 上的函数主要的三个函数类型限定词如下：

• __global__：在 device 上执行，从 CPU 侧调用（一些特定的 GPU 也可以从 device 上调用），返回类型必须是 void，不支持可变参数参数，不能成为类成员函数。注意用 __global__ 定义的 kernel 是异步的，这意味着 host 不会等待 GPU 上的 kernel 执行完就执行下一步。
• __device__：在 device 上执行，单仅可以从 device 中调用，不可以和 __global__ 限定词同时用。
• __host__：在 host 上执行，仅可以从 host 上调用，一般省略不写，不可以和 __global__ 同时用，但可和 __device__，此时函数会在 device 和 host 都编译。那么是不是不管是 host 上还是 device 上都可以调用了？

好文精读：CUDA编程入门极简教程 - 知乎

CUDA 对 Linux 和 Windows 适配。

查看安装的 CUDA 版本

nvcc -v

CUTLASS

是一个 CUDA C++ 模板抽象的集合，这里可以理解成一个抽象模板库。目的是为实现高性能的矩阵乘法和相关计算。说白了就是简化编写 CUDA 代码的，一些东西都给封装好了。

CUTLASS 的设计初衷是将 GEMM 中一些“可变的部分”分解成若干 C++ 抽象模板实现的基础组件，这种设计可以使开发者轻松的定制到他们自己的 CUDA kernel 中。

异构计算

要了解 CUDA，就要先明白什么是异构计算。为什么 GPU 和硬盘、网卡都是外设，那么 CPU + 硬盘或者说 CPU + 网卡就不叫异构计算，而 CPU + GPU 就叫做异构计算？

我觉的这是因为网卡/硬盘在设计之初就不是为了计算功能，而是为了其他功能（存储/网络）；而 GPU，尤其是 GPGPU 在设计之初就是为了计算，就是为了卸载 CPU 的一部分的计算功能到 GPU 上，因此 CPU + GPU 叫做异构计算。

相比于主机上的计算，显卡可以看作是一个独立的计算生态系统，包含自己的 GPU, 显存，可以自成一体进行计算，而显然网卡和硬盘等等并没有计算的功能，因此，使用显卡外设的可以叫做异构计算但是使用其他的却不行。

CUDA 和 GPU Driver 的区别

In the NVIDIA driver package, there is a libcuda.so.

CUDA driver is libcuda.so which is included in Nvidia driver and used by CUDA runtime api.

CUDA 测试带宽：cuda-samples/Samples/1_Utilities/bandwidthTest/bandwidthTest.cu at master · NVIDIA/cuda-samples

GPU 线程 / CUDA 线程 / CUDA Grid / CUDA Block

Grid, block 和 thread 都是一个软件概念。

Kernel 在 device 上执行时实际上是启动很多线程：

• 一个 kernel 所启动的所有线程称为一个网格（grid），也就是说一个 Kernel 对应一个 Grid，同一个网格上的线程共享相同的全局内存空间，这个模型和 CPU 侧很像，一个网格就像是一个进程，一个进程里的所有线程也是共享全局内存空间的（虚拟内存空间），进程也有自己的页表，那么一个网格里的所有线程难道也是通过页表机制分享全局内存空间的吗？
• Grid 是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程（threads），这是第二个层次；

// 表示一个 grid 有 6 个 block，3 * 2 进行排列
dim3 grid(3, 2);
// 表示一个 block 有 15 个 thread，5 * 3 进行排列
dim3 block(5, 3);
// 因此，一共有多少个线程，可以通过 3 * 2 * 5 * 3 的方式计算出来
kernel_fun<<< grid, block >>>(prams...);

上面代码会生成下面这样的架构：

一个 Grid 可以最多包含三个维度的 block，一个 block 也最多可以包含三个维度的 threads。因此：

• 一个 block 在 Grid 中的位置可以用 (blockIdx.x, blockIdx.y, blockIdx.z) 来表示，
• 一个 thread 在 block 中的位置可以用 (threadIdx.x, threadIdx.y, threadIdx.z) 来表示。

因此有以下九种组合：

1. Grid 为 一维，Block 为一维：int threadId = blockIdx.x *blockDim.x + threadIdx.x;；
2. Grid 为 一维，Block 为二维：int threadId = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
3. Grid 为 一维，Block 为三维：int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z + threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;
4. … 更多请见：初识多线程并行计算 | Notebook

对于一维 Grid 和一维 Block 结构来说，一个线程仅需要两个内置的坐标变量 (blockIdx, threadIdx) 来唯一标识。

一个 Block 一旦被分配给一个特定的 SM，在它的整个生命周期内都将在该 SM 上执行，绝对不会迁移到其他 SM 上。这是 CUDA 执行模型的核心设计原则之一。

Nvidia Shared Memory

下面这个图已经非常明显了。一个 SM 里的所有线程之间会 share memory。

// 推进指针到起始位置
A += cRow * BLOCKSIZE * K;                    // 行=cRow，列=0
B += cCol * BLOCKSIZE;                        // 行=0，列=cCol
C += cRow * BLOCKSIZE * N + cCol * BLOCKSIZE; // 行=cRow，列=cCol

float tmp = 0.0;
// 外部循环推进 A 沿列和 B 沿行，直到我们完全计算出 C 中的结果。
for (int bkIdx = 0; bkIdx < K; bkIdx += BLOCKSIZE) {
  // 每个线程从全局内存加载 A 和 B 中的一个元素到共享内存中。
  // 将 threadCol（=threadIdx.x）设为连续的索引，以允许全局内存访问协同。
  As[threadRow * BLOCKSIZE + threadCol] = A[threadRow * K + threadCol];
  Bs[threadRow * BLOCKSIZE + threadCol] = B[threadRow * N + threadCol];

  // 阻塞本块内的线程，直到缓存完全填充
  __syncthreads();

  // 在当前缓存块上执行点积
  for (int dotIdx = 0; dotIdx < BLOCKSIZE; ++dotIdx) {
    tmp += As[threadRow * BLOCKSIZE + dotIdx] *
            Bs[dotIdx * BLOCKSIZE + threadCol];
  }
  // 在最后需要再次同步，以避免更快的线程在较慢的线程完成之前将下一个块提取到缓存中
  __syncthreads();

  // 推进指针到下一个块
  A += BLOCKSIZE;
  B += BLOCKSIZE * N;
}
C[threadRow * N + threadCol] = tmp;

CUDA kernel / Operator 算子

Kernel 是在 device (GPU) 上线程中并行执行的函数。在调用时需要用 <<<grid, block>>> 来指定 kernel 要执行的线程数量，在 CUDA 中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号 thread ID，这个 ID 值可以通过核函数的内置变量 threadIdx 来获得。

算子其实是从神经网络模型架构的角度来说的，深度学习算法由一个个计算单元组成，我们称这些计算单元为算子（Operator，简称 OP）。在网络模型中，算子对应层中的计算逻辑，例如：卷积层（Convolution Layer）是一个算子；全连接层（Fully-connected Layer， FC layer）中的权值求和过程，是一个算子。再例如：tanh、ReLU 等，为在网络模型中被用做激活函数的算子。

算子其实就是一个函数。Kernel 和 Operator 可能不是一对一的关系，一个 Operator 下可能有多个 Kernel。

算子是实现在哪里？

• GPU 驱动？❌
• CUDA 驱动？❌
• 推理框架层比如 PyTorch/Tensorflow ✅
• 大模型推理服务层比如 vLLM/SGLang（也是通过 torch extension 的方式实现的算子） ✅

因此不同的推理框架里的算子实现是不一样的。

CUDAGraph

在一个多层模型推理的过程中，CPU 和 GPU 之间不同层不同算子的 kernel launch 和 H2D/D2H copy 是一个很大的性能开销，所以我们可以使用 CUDA Graph 让整个模型的推理过程在 GPU 上运行，从而免去了这些开销。

NVIDIA GPU 可观测性

nvml / dgcm

这两个工具之所以放到一起，因为都是 NVIDIA 开发并提供的可观测性工具。两者最重要的区别是一个是单个 GPU 的 metrics，一个是多 GPU 多节点级别的 metrics。

• NVML: NVIDIA Management Library. C-based API.
• DGCM: Data Center GPU Manager. 也是要依赖 NVML 组件的：it builds on NVML but adds higher-level features for cluster-wide management.

nvidia-smi 或许是每一个接触 GPU 编程的人都知道的命令，它能够显示 GPU 的相关信息，经常用于查询 GPU 状态、使用情况等。

有些工具借助 nvidia-smi，对它的输出进行分析，从而获取 GPU 状态；更高级一些的工具，则会使用 nvidia-smi --query-gpu=index --format=csv 等类似的指令，分析结构化的 csv 输出。而本质上，nvidia-smi 的很多信息，其实都是来自于 NVIDIA management library（简称 nvml）。

其次是 dgcm：

What is the difference between NVIDIA DCGM and NVIDIA NVML? - Massed Compute