HGX

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

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

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

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

HGX 还是包含 BF DPU 的。

PTX (Parallel Thread Execution) / SASS (Streaming Assembly) / NVRTC

https://k48xz7gzkw.feishu.cn/docx/GBa6dc0SFoRhDyxNYA7cNt7knNe#share-OZ1mdKvNEo0Gk1xXHnbcJnoMnSh

cuda - Is NVIDIA's JIT compilation cache used when you don't use NVCC? - Stack Overflow

PTX 是什么：CUDA 平台的汇编语言，需要在虚拟机（这个虚拟机 JIT 是包含在驱动中的）上执行，虚拟机再使用真正的操作硬件的 CPU 指令来运行。

  • PTX 是一种中间语言，可以在不同的 GPU 上运行，
  • 而 SASS 是一种特定的汇编语言，只能在特定的 GPU 上运行。

PTX 代码示例：

.visible .entry test(int&)(
        .param .u64 test(int&)_param_0
)
{

        ld.param.u64    %rd1, [test(int&)_param_0];
        cvta.to.global.u64      %rd2, %rd1;
        mov.u32         %r1, %ctaid.x;
        st.global.u32   [%rd2], %r1;
        ret;
}

SASS 代码示例：

test(int&):
    MOV R1, c[0x0][0x20]
    MOV R2, c[0x0][0x140]
    S2R R0, SR_CTAID.X

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？ - 知乎

NVIDIA GPU 加速库

下图很直观地展示了各个加速库（cuDNN, CUTLASS, TENSORRT, cuBLAS 等等）的位置，位于 Pytorch 之下，CUDA 之上。

  • cudnn、cuBLAS 这样的基础算子原语库在常见的卷积层上性能表现很好，通常都能够满足用户的需求，但是在面对用户高度定制化的算法时，基础算子库往往并不能充分发挥硬件的性能。这是由于算子优化的长尾问题引起的，基础算子库引入了许多卷积优化的通用策略，但是这些优化的策略并不能覆盖所有的情况，实际算法中的卷积层有可能并不能从通用的优化策略中获得收益，从而无法充分发挥硬件的性能。基础算子库的另一个问题是用户无法对这些基础算子进行定制化开发，当算法开发人员想为卷积算子添加一种新的激活函数，或者想添加一种特殊的卷积算子 (比如：LocalConv) 时，就会变得束手无策。
  • cutlass 是 NVIDIA 推出的一款线性代数模板库，它定义了一系列高度优化的算子组件，开发人员可以通过组合这些组件，开发出性能和 cudnn、cublas 相当的线性代数算子。但是 cutlass 仅支持矩阵乘法运算，不支持卷积算子，从而难以直接应用到计算机视觉领域的推理部署中。
  • TensorRT 是一款非常强大的深度学习推理部署框架，在 CUDA 平台上性能表现非常优秀，而且目前已经比较成熟，用户使用起来比较方便。然而 TensorRT 也存在着一些问题，对于开发人员来说，TensorRT 是一个黑盒，用户没有办法细粒度控制 TensorRT 内部的实现细节。例如：在部署量化网络时，开发人员无法控制 TensorRT 底层的量化细节，有可能会出现部署和训练的精度对不齐的问题。再比如：TensorRT 在推理部署时，用户无法精细的控制算子的显存使用情况，有时 TensorRT 在运行网络时耗费了大量的显存，而用户却没有特别好的办法对此进行优化。

(48 封私信 / 82 条消息) cudnn、cublas、cutlass、TensorRT - 知乎

cuDNN

GPU 加速库。提供了深度学习算法中常见算子的高效实现。

它能将模型训练的计算优化之后，再通过 CUDA 调用 GPU 进行运算，也就是位于 CUDA 的上层，算子也是位于 CUDA 上层（所以叫做 CUDA 算子开发）。

主要是为了加速训练过程。

Pytorch(GPU)配环境原理：cuda+cudnn+pytorch配环境的每一步到底干了些什么？_pytorch cudnn-CSDN博客

TensorRT / TensorRT-LLM

也是加速库，只负责模型的推理过程，一般不用 TensorRT 来训练模型的，而是用于部署时加速模型运行速度。

CUTLASS / cuBLAS

CUTLASS 是开源的，cuBLAS 是闭源的。

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

这个是工作在 CUDA 上层，的一般来说调用顺序是这样的：

torch.nn.functional.linear
    // pytorch
    libtorch_cpu.so
        libtorch_cuda.so
            // cuBLAS
            libcublas.so
                // CUDA
                libcuda.so

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

GPU 知识体系学习文档

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

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

nvidia-smi

查看一个 GPU 特定的信息：

nvidia-smi -q -i 0

给 GPU 锁频：

nvidia-smi -lgc 2490,2490

nvcc / nvrtc

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

nvidia_peermem.ko / nv_peer_mem.ko

这是一个内核模块。

The nvidia_peermem module is a drop-in replacement for nv_peer_mem

GDR 毕竟要涉及到 DMA 等等底层操作，而 NVIDIA 并不打算开源他们的驱动，所以搞了个可以安装的内核模块。

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

nccl-tests/nccl 本身没有办法在多机之间进行通信。

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-tests 和 MPI 库之间的关系

一图以蔽之：

NCCL 和 MPI 一个控制 GPU 之间的通信，一个控制 CPU 之间的通信。 两者没有半毛钱关系。

NCCL can be easily used in conjunction with MPI. NCCL collectives are similar to MPI collectives, therefore, creating a NCCL communicator out of an MPI communicator is straightforward. It is therefore easy to use MPI for CPU-to-CPU communication and NCCL for GPU-to-GPU communication.

nccl-tests 直接 include 了 mpi.h，同时也使用了 NCCL。所以当没有 GDR/NVLink/GPU P2P 时，NCCL 基本上就退化成 MPI 通信的方式了。注意 NCCL 本身并没有 import mpi.h，只是方式退化成了和 mpi.h 类似的通信方式：

  • 先从 GPU 显存拷贝到 CPU（这一步显然 MPI 库是做不到的）；
  • 然后 CPU 再通过网络发送到对端网卡，对端网卡再到 CPU（这一步 MPI 库可以做到）；
  • 然后对端 CPU 再发送到 GPU 显存上（这一步显然 MPI 库也是做不到的）。

有没有可能中间第二步的时候，也是使用了 MPI 库的呢？并没有，可以 clone 下来 NCCL 仓库然后 grep 一下 mpi.h，并不能找到对应代码引入了这个头文件。所以 CPU 通过网络发送并接收的这一步应该也是 NCCL 自己实现的。

实际上，NCCL 的网络层是自己实现的，NCCL 有一个插件机制，可以引入一个网络插件，NCCL 的自研网络层：NCCL 实现了自己的网络抽象层 (ncclNet 插件)。对于标准的以太网（TCP/IP）和 InfiniBand/ RoCE (使用 Verbs API)，NCCL 提供了内置的、不依赖 MPI 的实现：

  • TCP/IP：NCCL 直接使用 OS 提供的 socket API (socket, connect, bind, listen, accept, send, recv 等) 来实现基于 TCP 的网络通信。
  • InfiniBand / RoCE：NCCL 直接使用 libibverbs 库提供的底层 Verbs API (ibv_post_send, ibv_post_recv, ibv_poll_cq 等) 来实现 RDMA 通信。

因为 nccl-tests include 了 mpi.h，所以 nccl-tests 其实是能够感知到那些 MPI 的环境变量设置的，因此我们可以通过 mpirun 来运行 nccl-tests 的程序。

如何查看 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）卸载到网络（交换机）上进行。

NCCL 环境变量解读

NCCL_SOCKET_IFNAME

主要用于指定 NCCL 控制面 通信（TCP/IP 套接字）所使用的网络接口（网卡）。它通常不直接控制数据面（高性能数据传输，如 RDMA/IB）所使用的网卡。

NCCL 的双层通信机制：

  • 控制面 (Control Plane)：负责节点间的握手、交换地址信息、协调集合操作、错误处理等管理任务。这部分通信通常使用标准的 TCP/IP 套接字。
  • 数据面 (Data Plane)：负责实际的、高性能的集合通信数据传输（如 AllReduce, AllGather, Broadcast 等）。这部分通信会优先使用最高性能的传输方式，如：
    • InfiniBand Verbs (IB/RoCE)：通过 RDMA 实现零拷贝、低延迟、高带宽传输。
    • NVIDIA GPUDirect RDMA：允许 GPU 内存直接与网卡进行 RDMA 操作，绕过 CPU 和系统内存。
    • TCP/IP：当高性能传输不可用时作为后备方案（性能较差）。

NCCL_IB_HCA

控制 GPU 之间通信所使用的网卡。

这个参数和 btl_openib_if_include 之间的区别可以看 RDMA 集合通信测试^。简单来说，一个控制 GPU <-> GPU 之间的通信，一个控制 CPU <-> CPU 之间的通信。

GPU 卡间架构

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

CUDA Core / Tensor Core

CUDA cores are general-purpose processors in NVIDIA GPUs. While, Tensor cores are built specifically to accelerate deep learning tasks like training and inference.

CUDA cores can do “more” than the tensor cores.

GPU 工作队列（Work Queue）

NVIDIA driver

有广义（full driver）有狭义：

  • 广义 NVIDIA driver：包含了 CUDA library 等等用户态的 driver；
  • 狭义 NVIDIA driver：仅仅包含内核模块。

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

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

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

CUDA 采用单指令多线程 SIMT 架构管理执行线程，不同设备有不同的 warp 大小，但是到目前为止基本所有设备都是维持在 warp size == 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 显存 (实际应该是 109∗4/1024/1024/1024=3.725G10^9*4/1024/1024/1024=3.725G，为了方便可以记为 4G)。

int8 量化模型更简单，1B 对应 1G 显存。

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）

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

GPU XID

The Xid message is an error report from the NVIDIA driver that is printed to the kernel log or event log.

Nsight

Nsight 的安装

Installation Guide — nsight-systems 2025.3 documentation

上面文档里的下面几行：

apt update
apt install -y --no-install-recommends gnupg
echo "deb http://developer.download.nvidia.com/devtools/repos/ubuntu$(source /etc/lsb-release; echo "$DISTRIB_RELEASE" | tr -d .)/$(dpkg --print-architecture) /" | tee /etc/apt/sources.list.d/nvidia-devtools.list
apt-key adv --fetch-keys http://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/7fa2af80.pub
apt update
apt install nsight-systems-cli

Nsight 的使用

# 结束进程时，会显示：
# Generating '/tmp/nsys-report-01cd.qdstrm'
# [1/1] [========================100%] report1.nsys-rep
# Generated:
#        /root/p/virt_pgo/honor/report1.nsys-rep
# 这些表示 profile 完成了。
nsys profile <进程启动命令>

拿到 .nsys-rep 文件后，我们在我们的客户端（工作机）打开 Nsight Sytem，选择 file -> load 这个文件来查看即可，非常简单。

Nsight 的解读

‌‍⁡⁡⁢⁣⁢‬‬⁣‬​‬‌​‌‍⁣‍⁡‌‌​‬⁤​‬⁡⁣​⁣⁢​‬‍​⁣‬‌​‍⁤‬Nsight System 输出解读 - 飞书云文档

NVTX (NVIDIA Tools Extension SDK)

一个调试工具，具体功能这里说的比较清楚：探索NVIDIA NVTX：强大的性能分析工具库-CSDN博客

NVIDIA GPU 性能

GPU P-State

在 NVIDIA GPU 上，“P-State”（Performance State）用来表示显卡当前的性能/功耗等级，P0 代表最高性能。

如何查看当前 P-State？

nvidia-smi -q -d PERFORMANCE

里面有 Performance State。

GPU d2h h2d 性能问题

nvidia-smi -q -i 0

可以看 PCI -> GPU Link Info 这部分，能看到有 PCIe Gen 和 Link Width 两部分，都可以作为排查依据。

How to install NVIDIA Bunch of things?

How to install and check NVIDIA driver?

可以看一下这个页面：CUDA Toolkit 12.9 Update 1 Downloads | NVIDIA Developer 可以直接把 driver 和 CUDA 都安装了。

NVIDIA 最近开源了一版驱动：NVIDIA/open-gpu-kernel-modules: NVIDIA Linux open GPU kernel module source version 575.64。从 515 开始，NVIDIA 会 release 专用的和开源的两版驱动。目前的话，开源的已经是默认安装方式了。

The NVIDIA Linux GPU Driver contains several kernel modules:

  • nvidia.ko
  • nvidia-modeset.ko
  • nvidia-uvm.ko
  • nvidia-drm.ko
  • nvidia-peermem.ko

How to install and check CUDA?

See [[#How to install and check NVIDIA driver?]]

How to uninstall CUDA?

How to install and check cuDNN?

How to install and check GDR?

Installing GPUDirect RDMA - NVIDIA Docs

How to install and check NCCL?

Installation Guide NVIDIA Deep Learning NCCL Documentation

NCCL 下载还需要登录？

# CUDA 和 Driver 安装好后，好像直接执行下面这个就可以了。
sudo apt install libnccl2 libnccl-dev

How to install and check NIXL?

NVIDIA Container Toolkit / How to install NVIDIA Container Toolkit

Installing the NVIDIA Container Toolkit — NVIDIA Container Toolkit

需要在 host 上而不是在 container 中安装，同时注意安装完了还要配置（里面有配置的章节）。

作用是能够让：docker run --gpus all 用上 GPU Driver，否则会报错：

docker: Error response from daemon: could not select device driver "" with capabilities: [[gpu]]

同时 docker run --runtime=nvidia 能够跑通。

CUDA Runtime API