第 2 章 AI InfraGPU分布式训练

第 2 章 GPU 微结构

第2章 GPU 微结构

本章从 SIMT 执行模型出发，逐层拆解 NVIDIA GPU 的微架构：架构代际演进、SM 内部结构与 Warp 调度、内存层次与数据通路、Tensor Core 与 MMA 指令、Block 调度与 Occupancy，以及 GPU 间互联与拓扑/MIG 运维实践。CUDA 语言细节与分布式训练算法不在本章展开。

2.1 SIMT vs SIMD 计算范式

GPU之所以适合AI计算，根源在于其SIMT（Single Instruction, Multiple Thread）执行模型。理解SIMT与CPU的SIMD （Single Instruction, Multiple Data）之间的本质差异，是掌握GPU编程和性能优化的基础。

2.1.1 CPU SIMD 范式

CPU的SIMD指令集（x86的AVX-512、ARM的SVE/NEON）将多个数据打包到一个宽寄存器中，用一条指令同时处理。例如，AVX-512使用512位寄存器，可同时处理16个FP32或8个FP64运算。 // CPU SIMD example (AVX-512 intrinsic) #include <immintrin.h> void vector_add_simd(const float* a, const float* b, float* c, int n) { for (int i = 0; i < n; i += 16) { __m512 va = _mm512_loadu_ps(&a[i]); // Load 16 floats __m512 vb = _mm512_loadu_ps(&b[i]); __m512 vc = _mm512_add_ps(va, vb); // 1 instruction: 16 additions _mm512_storeu_ps(&c[i], vc); } } CPU SIMD的核心限制在于：所有数据通路共享同一个程序计数器（PC），分支会导致严重的性能退化。这使得CPU擅长处理控制流密集的应用，但在大规模数据并行场景下效率较低。

2.1.2 GPU SIMT 范式

NVIDIA GPU的SIMT模型将大量线程组织为 Warp（含32个线程）。一个Warp内的32个线程在相同的SM（Streaming Multiprocessor）上同时执行同一条指令，但各自拥有独立的寄存器和程序计数器。SIMT与SIMD的关键区别如表2-1所示。表2-1 SIMT与SIMD对比特性 CPU SIMD GPU SIMT 执行单元宽度 4-16 (AVX-512) 32 (Warp) 程序计数器共享（一个）每线程一个（32个）分支处理需手动mask 硬件自动Warp Divergence处理寄存器独立性无（共享向量寄存器）每线程独立寄存器文件切换代价高（OS上下文切换）零（硬件Warp调度器）

2.1.3 Warp 与线程分歧

Warp是GPU调度的最小执行单元。当Warp内的32个线程执行到分支语句时： global void divergent_kernel(float* data, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { if (idx % 2 == 0) { data[idx] = data[idx] * 2.0f; // Even threads execute path A } else { data[idx] = data[idx] + 1.0f; // Odd threads execute path B

         }
     }
 }

当Warp中一半线程走路径A、一半走路径B时，SM会串行化这两条路径：先让走A的线程执行（B的线程闲置），再让走B 的线程执行（A的线程闲置）。这称为 Warp Divergence，会导致实际吞吐减半，其时序如图2-1所示。 SM Warp (32 Threads) Even Threads Odd Threads Issue branch instruction Execute path A (Even threads active) Idle (Masked) Execute path B (Odd threads active) Idle (Masked) Issue reconvergence instruction SM Warp (32 Threads) Even Threads Odd Threads 图2-1 Warp Divergence的串行执行时序

2.1.4 GPU 为何适合 AI 计算

GPU在AI计算中表现卓越的核心原因有四。

极高的算术密度矩阵乘法（GEMM）是AI计算的核心操作，其计算访存比（Operational Intensity）为 O(N³) / O(N²) = O(N)。对于大矩阵（M,N,K ≥ 4096），每从HBM读取一个元素，可执行数百次乘加运算。这使得GPU的高吞吐计算单元能得到充分利用。
隐藏延迟的能力 GPU通过 Fine-Grained Multithreading 隐藏内存访问延迟。每个SM上常驻数千个线程（多个Warps），当Warp A等待全局内存数据时，Warp调度器在零额外开销下切换到Warp B继续执行。只要SM上有足够多的就绪Warp，计算单元就能保持忙碌。
内存层次的精心设计寄存器文件（每SM 256KB）→ 共享内存/L1（每SM 228KB，Hopper架构）→ L2 Cache（50MB，H100）→ HBM3 （80GB，H100），每一层都针对不同的访问模式和复用策略优化。
专门硬件单元 Tensor Core（Volta起）专为矩阵乘累加设计，单指令可执行4×4矩阵乘加（FP16），吞吐远高于CUDA Core的标量计算。

2.1.5 编程模型解析

// CUDA programming model core concepts dim3 grid(256, 1, 1); // Grid: 256 Blocks dim3 block(256, 1, 1); // Each Block: 256 Threads (= 8 Warps) // Total: 65536 Threads // Kernel launch my_kernel<<<grid, block, shared_mem_bytes, stream>>>(args);

 // Built-in variables
 // threadIdx.x   : Thread index within Block (0 ~ blockDim.x-1)
 // blockIdx.x    : Block index within Grid (0 ~ gridDim.x-1)
 // blockDim.x    : Number of Threads per Block
 // gridDim.x     : Number of Blocks in Grid

理解SIMT的执行模型是后续学习Tensor Core编程、Shared Memory优化和Occupancy调优的基础。

2.2 NVIDIA GPU 架构演进

NVIDIA GPU架构从前深度学习时代的Fermi（2010）到最新的Rubin（2026），经历了十余年的持续迭代。每一次架构升级都带来了关键创新，且这些创新越来越聚焦于AI/ML工作负载。

2.2.1 演进总览

NVIDIA GPU 的主要架构世代如图2-2所示。 NVIDIA GPU Architecture Generations 2010 2012 2014 2016 2017 2020 2022 2024 2025 2026 图2-2 NVIDIA GPU架构世代演进时间线

2.2.2 各代关键创新

Fermi 2010 首个完整的GPU计算架构。首次引入L1/L2缓存层次、支持ECC纠错、统一地址空间，理论双精度性能515 GFLOPS。
Kepler 2012 SMX架构，每SM 192个CUDA Core。引入Dynamic Parallelism（GPU Kernel可启动子Kernel），解耦CPU-GPU同步。 GK110芯片拥有7.1B晶体管（当时最大），为后续大GPU奠定工艺基础。
Pascal 2016 首次采用HBM2（高带宽存储器），GP100拥有732 GB/s带宽。NVLink 1.0（160 GB/s双向）取代PCIe，实现多GPU高速互联。统一内存（Unified Memory）做到CPU和GPU虚拟地址空间统一，Page Fault自动迁移数据。深度学习时代的真正起点，DGX-1深度学习工作站于2016年发布，内置8块P100。
Volta 2017 这是NVIDIA GPU历史上的转折点，Volta引入了 Tensor Core，首个专为矩阵乘累加（D = A × B + C）设计的硬件单元。第一代Tensor Core每个SM含8个单元，每单元每时钟执行4×4×4 FP16 FMA，即64次FMA/时钟 = 128 FLOPS/时钟； V100（80 SM）的FP16峰值约125 Tensor TFLOPS。NVSwitch 1.0于2018年随DGX-2发布，支持16块V100的全互联拓扑（DGX-2: 16×V100）。
Ampere 2020 A100（GA100芯片）引入三大关键创新，成为大模型时代的标配： •第三代Tensor Core：支持TF32（TensorFloat-32），在FP32范围内使用19位精度，训练中无需修改代码即可获得 FP16级别性能且保持FP32训练稳定性；同时增加BF16和FP64 Tensor Core支持 •MIG（Multi-Instance GPU）：将A100物理切分为最多7个独立GPU实例，每个实例拥有独立的内存、缓存和计算资源，使推理场景可以安全地多租户共享GPU •结构化稀疏（2:4）：每4个连续值中恰好有2个为零，可将Tensor Core吞吐翻倍 A100规格：6912 CUDA Core, 432 Tensor Core (3rd gen), 80GB HBM2e (2.0 TB/s), NVLink 3.0 (600 GB/s双向)。
Hopper 2022 H100（GH100芯片）是首个显式针对Transformer架构优化的GPU： •FP8 Transformer Engine：在训练中动态调整FP8的数据格式范围，自动混合E4M3和E5M2格式，使H100 FP8算力达 1979 TFLOPS（稠密）/3958 TFLOPS（稀疏） •TMA（Tensor Memory Accelerator）：硬件张量内存访问单元，支持异步全局到共享内存的2D/3D拷贝，释放寄存器与指令供计算使用 •Thread Block Clusters：跨SM的Thread Block协作，共享内存可跨Block访问 •DPX指令：动态编程加速指令，用于生物信息学等特定负载 Hopper引入的CUDA 12.0+中的Cluster功能，允许Block在L1层面共享数据，这对FlashAttention等需要频繁跨Block数据共享的算子有显著加速效果。
Blackwell 2024-2025 B200和GB200代表了NVIDIA在AI领域的最大一次架构跃升，关键规格对比如表2-2所示。表2-2 B200与GB200规格对比特性 B200 GB200 (NVL2 Superchip) 晶体管数 208B (2×104B die) 208B FP16 稠密算力 1125 TFLOPS 2250 TFLOPS FP16 稀疏算力 2250 TFLOPS 4500 TFLOPS FP8 稠密算力 2250 TFLOPS 4500 TFLOPS FP8 稀疏算力 4500 TFLOPS 9000 TFLOPS FP4 稠密算力 4500 TFLOPS 9000 TFLOPS FP4 稀疏算力 9000 TFLOPS 18000 TFLOPS HBM 192GB HBM3e (8 TB/s) 384GB HBM3e (16 TB/s) NVLink 5.0 (1.8 TB/s/卡) 5.0 (3.6 TB/s/Superchip) TDP 1000W — 工艺 TSMC 4NP TSMC 4NP GB200 NVL72 将 72 颗 B200 与 36 颗 Grace CPU 整合为单一 NVLink 域（fabric 带宽 130 TB/s），FP8 稠密算力 720 PFLOPS，可支撑万亿参数模型的单机架训练。
Blackwell Ultra 2025 B300/GB300 在 Blackwell 基础上演进，FP4 算力提升 1.5×，注意力层加速 2×，单机架 HBM3e 容量由 GB200 的约 13.5 TB 增至 20 TB（单卡 288 GB），FP4 稠密算力由 720 提升至 1080 PFLOPS，FP4 稀疏算力保持 1440 PFLOPS。8 卡 HGX B300 平台的 FP4 稠密 108 PFLOPS、总 NVLink 带宽 14.4 TB/s 等规格，完整列于附录 B。数据来源：NVIDIA 官网（2026-01）
Rubin 2026 R100 GPU采用HBM4内存，单卡338亿晶体管，TDP约2300W；Vera Rubin NVL72机柜配套第六代NVLink和ConnectX-9 SuperNIC（1600 Gb/s端口带宽）。NVIDIA官网HGX Rubin NVL8系统规格如下： •8颗Rubin GPU，第六代NVLink，NVLink Switch总带宽28.8 TB/s •NVFP4推理400 PFLOPS（稀疏，系统级）/NVFP4训练280 PFLOPS（稠密） •GPU内存2.3 TB HBM4，带宽176 TB/s •相对HGX B200：token工厂吞吐10×，训练所需GPU数量减少4×，内存带宽3× 数据来源：NVIDIA 官网（2026-01）

2.2.3 Tensor Core 代的算力跃迁

各代GPU的Tensor Core算力与带宽对比如表2-3所示。表2-3 Tensor Core算力与带宽对比 GPU 年 FP16 Tensor 稠密 FP16 Tensor 稀疏 FP8 Tensor 稠密 FP8 Tensor 稀疏 HBM带 NVLink带份 (TFLOPS) (TFLOPS) (TFLOPS) (TFLOPS) 宽宽 V100 201 125 N/A N/A N/A 900 300 GB/s 7 GB/s A100 2020 312 624 N/A N/A 2039 600 GB/s GB/s H100 2022 989 1979 1979 3958 3352 900 GB/s GB/s H200 2024 989 1979 1979 3958 4800 900 GB/s GB/s B200 2024 1125 2250 2250 4500 8000 1800 GB/s GB/s B300 2025 约2250 约4500 约4500 约9000 8000 1800 GB/s GB/s R100 202 TBD TBD TBD TBD HBM4 NVLink 6 (Rubin) 6 TBD TBD 说明：B300行数据来源于NVIDIA HGX B300系统规格（8卡聚合FP16 36 PFLOPS、FP8 72 PFLOPS反推单卡），对应单卡 FP4稠密约13.5 PFLOPS，HBM带宽8000 GB/s对应单卡288 GB HBM3E。从V100到B200，Tensor Core FP16 稀疏算力增长了约18倍，而内存带宽增长了约9倍，带宽增速远低于算力增速，这一 “计算-带宽缺口”深刻影响了算法设计，催生了FlashAttention等IO-aware算法。到B300/Rubin时代，HBM4的引入有望部分缓解这一瓶颈，但根本性的算力-带宽剪刀差预计仍将持续。

2.3 SM 微架构与 Warp 调度

Streaming Multiprocessor（SM）是NVIDIA GPU的核心计算单元，相当于CPU的“核”。理解SM的微观结构，是编写高性能CUDA Kernel的基础。

2.3.1 SM 总体架构

H100（GH100）GPU包含132个SM，每个SM的组成如下： H100 SM (per SM): ├── 4 Warp Schedulers (4 per SM) │ └── Each scheduler manages 16 Warps (64 Warps/SM total) ├── 128 FP32 CUDA Cores (32 per scheduler) ├── 64 FP64 Cores (16 per scheduler) ├── 4 Tensor Cores (4th gen, 1 per scheduler) ├── 256KB Register File (max 255 32-bit registers per thread) ├── 228KB L1 Data Cache/Shared Memory (256KB total; up to 228KB configurable as shared memory) ├── 256KB L0 Instruction Cache └── 1 TMA Unit 总计：132 SM × 128 CUDA Core = 16896 CUDA Core。加上Tensor Core和专用硬件，GH100总晶体管数高达80B。

2.3.2 Warp 调度器的工作原理

Warp Scheduler是SM内部的指令发射单元。每时钟周期，Warp调度器从就绪（Operands Ready）的Warp中选择一个，发射其下一条指令。其内部结构如图2-3所示。 SM (H100) Warp Scheduler 0 INT Cores Warp 0 Warp 1 Warp 2 Warp 3 FP32 Cores Warp Scheduler 1 Warp 4-7 Dispatch Units (one per sc heduler) Warp Scheduler 2 Warp 8-11 Warp Scheduler 3 FP64 Cores Warp 12-15 Tensor Cores 图2-3 H100 SM内部的Warp调度器架构 H100每个Warp Scheduler配备一个Dispatch Unit，每时钟周期从选中Warp发射一条指令到对应执行单元（FP32/FP64/Tensor Core/SFU）。不同Warp Scheduler可同时各自发射，SM每周期最多可并发发射4条来自4个不同 Warp的指令。（注：双发射是Kepler时代的特性，从Volta起已移除。）

2.3.3 指令流水线

GPU指令执行的关键延迟指标（H100为例）如表2-4所示。表2-4 GPU指令延迟操作类型延迟 (cycles) 说明寄存器访问 0 零延迟共享内存访问约20-30 不含Bank Conflict L1缓存命中约30-40 含Tag查询和TEX L2缓存命中约200-300 跨SM共享 HBM访问约400-800 取决于负载和频率 Tensor Core MMA 约100-200 取决于MMA形状 FP32 FMA 约4 全流水线化 GPU隐藏这些延迟的方法是多Warp并发。当SM上有32个活跃Warp时，即使每个Warp有400个cycle的HBM延迟，只要每12.5个cycle有Warp完成等待并发射新指令，计算单元就能满负荷运行。隐藏延迟所需的活跃Warp数可用公式估算：活跃 Warp 数 = 延迟（cycles）/ 每时钟发射指令数 / Warp 宽度。以 HBM 延迟 400 cycles、每时钟 4 条指令、Warp 宽度 32 计算，理论下界约为 3 个 Warp；实际通常需要 8-16 个 Warp 以应对突发延迟。

2.3.4 寄存器文件与 Bank 组织

寄存器是GPU上最快的内存。每个H100 SM拥有256KB寄存器文件（65536个32位寄存器），按Bank组织以支持多端口并发访问。关键约束： •每个Thread最多使用255个32位寄存器（H100） •每个SM的寄存器总数 = 65536 •每个SM的Warp槽位数 = 64（Warp Limit） •若每个Thread使用N个寄存器，则SM可容纳的最大Thread数为 65536 / N 寄存器分配示例： •32 寄存器/线程：SM 容量 65536/32 = 2048 线程 = 64 Warps，满 Occupancy •128 寄存器/线程：SM 容量 65536/128 = 512 线程 = 16 Warps，Occupancy 25% •255 寄存器/线程（上限）：SM 容量约 256 线程 = 8 Warps，Occupancy 12.5% Register Spilling：当某个Kernel需要超过255个寄存器时，编译器会将溢出数据放入L1缓存或HBM（Local Memory），导致严重的性能退化（延迟从0 cycle变为约400 cycles）。

2.3.5 共享内存与 L1 缓存

共享内存（Shared Memory）以Bank为单位组织，每个SM的共享内存分为32个Bank（每个Bank 4字节宽）。同一Warp 内的32个Thread如果访问不同Bank（理想情况），可在1个cycle内完成；如果多个Thread访问同一Bank，则产生 Bank Conflict，需串行化处理。 // No Bank Conflict: sequential access global void no_conflict(float* data) { shared float sdata[256]; int tid = threadIdx.x; float val = sdata[tid]; // Thread i -> Bank i%32

     // 32 Threads access 32 different Banks, completes in 1 cycle
 }
 // 2-way Bank Conflict: strided access

global void bank_conflict(float* data) { shared float sdata[256]; int tid = threadIdx.x; float val = sdata[tid * 2]; // Thread 0 -> Bank0, Thread 1 -> Bank2, Thread 16 -> Bank0 // Thread 0 and 16 access the same Bank, requires 2 cycles } Padding技巧可消除Bank Conflict： // Add padding to avoid Bank Conflict shared float sdata[256 + 1]; // +1 shifts address, breaks Bank alignment // Or use padding dimension: shared float sdata[256][32+1]; 从Volta架构起，L1数据缓存和共享内存共用同一块片上SRAM（H100上为256KB总容量，其中共享内存最多可配置为 228KB），比例可通过API配置： // Set L1/shared memory allocation ratio cudaFuncSetAttribute(myKernel, cudaFuncAttributePreferredSharedMemoryCarveout, cudaSharedmemCarveoutMaxShared); // Max shared memory mode // Or cudaSharedmemCarveoutMaxL1 // Max L1 cache mode 在AI训练的矩阵乘法Kernel中，通常需要最大化共享内存以存放Tile（如128×128的FP16矩阵Tile需要32KB共享内存），因此在调用矩阵乘法Kernel时通常配置为最大共享内存模式。

2.4 GPU 内存层次与通路

GPU的内存层次设计是决定AI训练和推理性能的关键因素。每一层内存都在容量（Capacity）、带宽（Bandwidth）和延迟（Latency）之间做出权衡，形成经典的内存金字塔。

2.4.1 内存金字塔

H100的内存层次结构如图2-4所示，从片上寄存器到片外HBM逐层容量增大、带宽降低。 GPU Memory Pyramid H100 CPU DRAM (via PCIe) ~1TB BW: 64 GB/s (PCIe Gen5 x1 Register File Shared Memory/L1 HBM3 6) 256KB/SM × 132 = 33.0MB 228KB/SM × 132 = ~30MB L2 Cache 80GB on H100 BW: ~100 TB/s (SM aggreg BW: ~64 TB/s (SM aggregat 50MB (H100) BW: 3.35 TB/s (H100), 4.8 T ated) ed) BW: ~12 TB/s B/s (H200) Local NVMe SSD Multi-TB BW: ~7 GB/s 图2-4 H100内存层次与带宽示意

2.4.2 各层详细分析

寄存器文件 •容量：256KB/SM（65536 × 32-bit），H100共132 SM •带宽：0 cycle延迟（访问即完成） •作用：保存Kernel内的中间计算结果、循环变量、指针等 •限制：每个Thread最多255个32位寄存器；寄存器Spilling至Local Memory会导致约100×的性能退化
共享内存 •容量：228KB/SM（H100，可配置L1/共享内存分配比例；最大228KB共享内存） •带宽：每SM每时钟读写约128字节（H100） •延迟：约20-30 cycles（不含Bank Conflict） •作用：Block内Thread间通信、矩阵乘法的Tile缓存、归约操作的中间结果 •访问规则：32个Bank，同一Warp内无Bank Conflict时1 cycle完成
L2缓存 •容量：50MB（H100，A100为40MB） •带宽：约12 TB/s（与HBM带宽之比约3.6:1） •作用：跨SM的数据共享、减少HBM访问、缓存Texture/Constant数据 •策略：LRU或自适应替换策略；对内存合并访问（Coalesced Access）友好 L2 Cache的驻留策略是近年来的优化热点。通过 cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size) 配合 cudaStreamSetAttribute 和 cudaAccessPolicyWindow （CUDA 11.0引入），可为特定数据在L2中预留持久化区域（Persisting L2 Cache），减少频繁数据的HBM往返。在大模型推理的KV Cache管理中，将热KV Block固定在L2中是提升吞吐的有效手段。
HBM HBM（High Bandwidth Memory）是GPU的全局内存，通过硅中介层（Silicon Interposer）以2.5D封装与GPU Die连接。各代GPU的HBM规格如表2-5所示。表2-5 各代GPU的HBM规格 GPU型号 HBM规格容量带宽 HBM Stack数量 A100 HBM2e 40/80GB 1.6/2.0 TB/s 5 H100 HBM3 80GB 3.35 TB/s 5 H200 HBM3e 141GB 4.8 TB/s 6 B200 HBM3e 192GB 8.0 TB/s 8 HBM访问的关键优化技术是内存合并访问（Memory Coalescing）：当同一Warp内的32个Thread访问连续的32个地址（128字节对齐）时，这些访问会合并为一次或少量内存事务（Transaction）。若访问模式不规则（如随机索引），则内存事务数会膨胀32倍，浪费大量带宽。 // Coalesced access: 1 x 128-Byte Transaction float val = global_array[threadIdx.x + blockIdx.x * blockDim.x]; // Non-coalesced (Strided): 32 x 32-Byte Transactions (severe HBM BW waste) float val = global_array[threadIdx.x * 1024]; // Random access (Gather): 32 x 32-Byte Transactions (worst case) float val = global_array[indices[threadIdx.x]];

2.4.3 数据通路与带宽计算

在实际AI训练中，理解数据在内存层次间的流动路径至关重要。以Transformer的Self-Attention计算为例（假设 batch=2、seq_len=4096、num_heads=32、head_dim=128、FP16）： •Q、K、V 各为 [2,32,4096,128] 张量，即 33.6M 元素 × 2 字节 = 67.1 MB •分数矩阵 S = QK^T 为 [2,32,4096,4096]，即 1.07G × 2 字节 = 2.15 GB •QK^T 的 FLOPs 为 2×32×4096×4096×128×2 = 549.8 GFLOPs，SV 与之相同朴素 attention 的 HBM 流量（读+写）为：读 Q+K+V 共 201.3 MB，写 S 为 2.15 GB，softmax 读 S 写 P 为 4.30 GB，读 P+V 写 O 为 2.28 GB，合计约 6.8 GB。计算访存比为 549.8 GFLOPs / 6.8 GB = 80.9 FLOPS/Byte。H100 FP16 的 ridge point 为 989 TFLOPS / 3.35 TB/s = 295 FLOPS/Byte。由于 80.9 < 295，朴素 attention 属于 memory-bound，可达性能约 80.9 × 3.35 TB/s = 271 TFLOPS（约为峰值的 27%）。这就是FlashAttention诞生的背景——通过算法重构（Tiling + Recomputation），将QKV从全局内存到共享内存的搬运量大幅减少，从而将计算访存比提升至硬件能高效执行的范围。

2.4.4 内存带宽实战验证

Test HBM bandwidth using CUDA

git clone https://github.com/NVIDIA/cuda-samples.git cd cuda-samples/Samples/1_Utilities/bandwidthTest make ./bandwidthTest --mode=shmoo

2.4.5 统一内存与超额分配

 # Typical output (H100):
 # Host to Device Bandwidth: ~51 GB/s (PCIe Gen5 limited)
 # Device to Host Bandwidth: ~51 GB/s
 # Device to Device Bandwidth: ~3300 GB/s (close to HBM3 3.35 TB/s peak)

CUDA的统一内存（Unified Memory, UM）允许CPU和GPU共享同一个虚拟地址空间，Page Fault时自动迁移数据。在推理场景中，当模型超过单GPU显存时，可通过UM实现“超额分配（Oversubscription）”——不过这会付出约50×的性能代价（HBM → PCIe）。 // Unified Memory simple example float *data; cudaMallocManaged(&data, 10ULL * 1024 * 1024 * 1024); // 10GB UM // Automatic Page Fault migration: CPU access migrates to DRAM, GPU access migrates to HBM 实际推理部署中，模型并行（Tensor/Pipeline Parallelism）或KV Cache Offloading比统一内存超额分配更常用。

2.5 Tensor Core 与 MMA 指令

Tensor Core是NVIDIA GPU中最强大的矩阵计算单元，自Volta架构（2017年）引入以来，它已成为AI训练和推理算力的核心来源。本节剖析Tensor Core的硬件架构、编程模型和性能特性。

2.5.1 Tensor Core 硬件规格演进

Tensor Core的硬件规格演进如表2-6所示。表2-6 Tensor Core硬件规格演进架构 Tensor Core 精度支持性能 (TFLOPS，各GPU自身峰值) 关键特性代数 Volta (V100) 第1代 FP16 125 基4×4×4 MMA Turing (T4) 第2代 INT8/INT4/FP16 130 (INT8) 增加整数MMA Ampere 第3代 TF32/BF16/FP64/INT8/INT4 312 (TF32) TF32自动加速、2:4 (A100) 稀疏 Hopper 第4代 FP8/FP16/BF16/TF32/FP64/IN 989 (FP16/BF16, 稠密), 1979 (FP8, FP8、WGMMA、 (H100) T8 稠密) TMA Blackwell 第5代 FP4/FP6/FP8/FP16/BF16/TF32 1125 (FP16/BF16, 稠密), 2250 FP4、微张量缩放 (B200) /FP64 (FP8, 稠密)

2.5.2 MMA 指令的矩阵形状

MMA（Matrix Multiply-Accumulate）指令是Tensor Core编程的基础。不同架构支持的MMA形状如下：

2.5.3 WMMA API 编程示例

 // MMA shapes defined in PTX ISA (SM80=Ampere, SM90=Hopper)
 // SM80 (A100): m16n8k16 (FP16 input, FP16/FP32 accumulate)
 // Meaning: A[16,16] x B[16,8] -> C[16,8] (16x8 result matrix)
 // Each Thread provides: A's 2 fragments (4 FP16x2), B's 1 fragment (2 FP16x2)
 // SM90 (H100): supports larger shapes
 // m16n8k16 (FP16/BF16)
 // m16n8k8 (TF32)
 // m16n8k32 (FP8 E4M3)
 // Key MMA features (H100 WGMMA):
 // 1. Async execution: can overlap with CUDA Core instructions
 // 2. Shared Memory direct input: no Thread-level data distribution needed
 // 3. Larger K dimension: k32 for FP8, k256 for FP4

WMMA（Warp Matrix Multiply-Accumulate）是CUDA C++层面访问Tensor Core的高级API： #include <cuda_fp16.h> #include <mma.h> using namespace nvcuda; // Matrix multiply 16x16x16 using WMMA API global void wmma_gemm(

     const half* A, // M×K = 16×16
     const half* B, // K×N = 16×16
     float* C)       // M×N = 16×16
 {

wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag; wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag; wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag; // Initialize accumulator to 0 wmma::fill_fragment(c_frag, 0.0f); // Load matrix tiles into WMMA Fragments (auto-distributed to Thread registers) wmma::load_matrix_sync(a_frag, A, 16); // LDM = 16 wmma::load_matrix_sync(b_frag, B, 16); // LDM = 16 // Execute MMA on Tensor Core: C += A x B wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); // Store back to global memory wmma::store_matrix_sync(C, c_frag, 16, wmma::mem_row_major);

2.5.4 MMA 数据布局与 Fragment

} 理解WMMA Fragment的数据布局是编写正确MMA代码的关键。以 matrix_a, 16, 16, 16, half, row_major 为例： •Fragment大小： 16×16 ÷ 32 threads = 8个half值/Thread •数据分配：32个Thread按照Warp内规则分配A矩阵的元素（非等距分配，受Tensor Core硬件固定调度规则约束） WMMA API的完整数据流动路径如图2-5所示：矩阵先经 wmma::load_matrix_sync 从全局内存载入共享内存，再分发到各线程的 Fragment（寄存器），最后由 wmma::mma_sync 在 Tensor Core 上完成矩阵乘累加。 Global Memory (HBM) Matrix A (MxK) Matrix B (KxN) wmma::load_matrix_sync wmma::load_matrix_sync Shared Memory A Tile (16x16) B Tile (16x16) Distribute to 32 Threads Distribute to 32 Threads Register File A Fragment B Fragment (8 half per thread) (8 half per thread) wmma::mma_sync wmma::mma_sync C Fragment (8 float per thread) wmma::store_matrix_sync Matrix C (MxN) 图2-5 WMMA API的数据流动路径

2.5.5 Tensor Core 编程最佳实践

选择合适的输入数据类型

 // Tensor Core throughput comparison across precisions (H100)
 // FP16 (m16n8k16): 989 TFLOPS (with FP32 accumulate)
 // BF16 (m16n8k16): 989 TFLOPS
 // TF32 (m16n8k8):   494 TFLOPS
 // FP8 (m16n8k32): 1979 TFLOPS (with FP32 accumulate)
 // INT8 (m16n8k32): 1979 TFLOPS
 // Best practice: inference with FP8/INT8 (accuracy loss <0.1% on LLM), training with BF16 (numerically stable)

充分利用流水线并行

 // Async copy (cp.async) + Tensor Core execution pipeline
 for (int k_tile = 0; k_tile < K / K_TILE; ++k_tile) {
     // Stage 1: async copy from global memory to shared memory (non-blocking compute)

cp_async(&As[stage][0], &A[k_tile * ...]); cp_async(&Bs[stage][0], &B[k_tile * ...]); // Stage 2: execute previous Tile compute on Tensor Core (overlap with copy) if (k_tile > 0) { wmma::mma_sync(c_frag, a_frag_prev, b_frag_prev, c_frag); } cp_async_commit(); cp_async_wait<0>();

     // Stage 3: switch current Stage
     // Prepare data for next iteration load
 }

利用2:4结构化稀疏 Ampere架构引入的2:4稀疏模式要求每4个连续值中恰好有2个为零。启用此功能后：

 # Use 2:4 sparsity in PyTorch (requires PyTorch 2.0+)
 from torch.sparse import to_sparse_semi_structured
 # Convert dense weights to 2:4
 weight = torch.randn(4096, 4096, device='cuda', dtype=torch.float16)
 weight_sparse = to_sparse_semi_structured(weight)
 # Sparse matrix multiply inference throughput
 output = torch.nn.functional.linear(input, weight_sparse)

最佳实践要点： •Sparse GEMM的实际加速比通常在1.5-1.6×（相对于理论2×），因为有稀疏表示的额外开销 •最适合推理场景（权重固定，可提前转换；训练中稀疏模式的动态变化会带来额外开销） •迁移学习微调时可尝试对预训练权重做稀疏化 4) 用CUTLASS Profiler验证性能

Use CUTLASS Profiler to compare

git clone https://github.com/NVIDIA/cutlass.git cd cutlass && mkdir build && cd build cmake .. -DCUTLASS_NVCC_ARCHS=90

Run GEMM performance test

./tools/profiler/cutlass_profiler \

2.5.6 WGMMA 原生接口

   --kernels=cutlass_tensorop_f16_s16816gemm_f16 \
   --m=4096 --n=4096 --k=4096
 # Typical output (H100):
 # Problem: 4096x4096x4096
 # TFLOPs: 645.2 (~65% of peak 989 TFLOPS for FP16 dense)
 # Compare to FP32 SGEMM: ~120 TFLOPS (CUDA Core only)

Hopper（H100）引入了 WGMMA（Warp Group Matrix Multiply-Accumulate）作为第4代 Tensor Core 的主要编程接口，其与旧版 WMMA 存在根本性差异，理解这一区别对于充分利用 H100 算力至关重要。核心差异如表2-7所示。表2-7 WMMA与WGMMA对比特性 WMMA（Volta~Hopper 均支持） WGMMA（仅 Hopper SM90 原生）操作粒度单 Warp（32 线程） Warp Group（4 个 Warp，128 线程）输入来源寄存器文件共享内存（Shared Memory）直接输入执行方式阻塞式（等待完成）完全异步（可与 CUDA Core 指令重叠）典型 K 维 k=16（FP16） k=16（FP16/BF16），k=32（FP8）每次 MMA 结果 16×16（M×N） 64×N（更大 Tile） WGMMA 直接从共享内存读取矩阵 A/B，绕过了 WMMA 中每个 Thread 先从共享内存加载 Fragment 到寄存器、再传入 Tensor Core 的中间步骤，减少了寄存器压力，同时异步执行使得数据加载与矩阵计算可以完全流水线化。在 FlashAttention-3、cuBLAS 和 CUTLASS 3.x 内部，H100 路径全部使用 WGMMA 而非 WMMA，这也是 H100 实际能逼近 989 TFLOPS（FP16 稠密）峰值的关键之一——使用 WMMA 的旧式写法在 H100 上通常只能达到 400-600 TFLOPS。

 // WGMMA conceptual sketch (PTX level, simplified)
 // wgmma.mma_async.sync.aligned.m64n128k16.f32.f16.f16
 // meaning: Warp Group MMA, async, output 64x128, FP32 accumulate, FP16 inputs
 // In CUTLASS 3.x this is abstracted via cute::wgmma()
 // Production code should prefer CUTLASS 3.x or cuBLAS over hand-written PTX

对于需要最大化 H100 Tensor Core 利用率的自定义算子，应使用 CUTLASS 3.x（提供 WGMMA 封装）或 Triton（自动生成 WGMMA 指令）；避免在 H100 上使用仅支持 WMMA 的旧版 Kernel。

2.5.7 FP8 与 Transformer Engine

FP8 并非单一格式，而是包含两种不同的指数/尾数配置，适用于不同的数值分布场景，如表2-8所示。表2-8 FP8两种格式对比格式指数位尾数位动态范围精度推荐用途 E4M3 4位 3位 ±448 较高前向激活值、权重（数值分布集中） E5M2 5位 2位 ±57344 较低反向梯度（分布范围更宽，需更大动态范围）两种格式与 BF16/FP16 相比动态范围均大幅压缩，因此不能直接替换，必须配合动态量化缩放（Dynamic Quantization Scaling）使用。 H100 Transformer Engine 的自动缩放机制在每次矩阵乘法前自动执行以下步骤：

计算缩放因子：统计当前 Tensor 的最大绝对值 amax ，推导使该值恰好落在 FP8 表示范围内的缩放因子 scale = FP8_MAX / amax
量化：将 FP16/BF16 输入乘以 scale，转换为 FP8 存储
执行 MMA：Tensor Core 以 FP8 精度计算，结果以 FP32 累加
反量化：将结果除以 scale_A × scale_B ，恢复到 FP16/BF16 在 PyTorch 中启用 FP8 训练只需：

 import transformer_engine.pytorch as te
 # Replace standard Linear with TE
 layer = te.Linear(1024, 4096, bias=True)
 # Wrap forward pass with FP8

with te.fp8_autocast(enabled=True, fp8_recipe=te.recipe.DelayedScaling()): output = layer(input) # E4M3 for forward, E5M2 for backward DelayedScaling 是 Transformer Engine 的默认策略：每隔若干步更新一次 amax 历史记录并重新计算缩放因子，避免了每步都调用 torch.max() 的额外开销。精度影响：在主流 LLM（LLaMA、GPT、BERT 等）的训练中，启用 FP8 混合精度通常导致最终精度损失 < 0.2%（以 perplexity 衡量），同时在 H100 上可将训练吞吐提升约 1.5-2×（相比 BF16）。

2.6 Block 调度与 Occupancy

Occupancy是衡量GPU计算资源利用效率的核心指标。它定义为SM上活跃Warp数与该SM可支持的最大Warp数的比值。本节系统讲解Block调度机制、Occupancy计算方法及其性能影响。

2.6.1 Block 到 SM 的调度策略

NVIDIA GPU使用 Round-Robin策略将Thread Block分配给SM。当用户启动一个Kernel时：

检查每个SM是否有足够的剩余资源（寄存器、共享内存、Warp槽位）容纳新的Block
在资源充足的SM中，按Round-Robin顺序分配Block
Block一旦被调度器分配给某个SM，就在该SM上运行直至完成
同一Block内的线程不能跨SM执行（Block是调度的原子单元）关键推论：如果Block配置过大（过多Threads/寄存器/共享内存），每个SM只能容纳很少的Block，剩余资源被浪费。

2.6.2 Occupancy 计算公式

Occupancy 的定义与计算： •Occupancy = SM 活跃 Warp 数 / SM 最大 Warp 数（H100 为 64） •活跃 Warp 数 = min(常驻 Block 数 × 每 Block Warp 数, SM 最大 Warp 数) •常驻 Block 数 = floor(min(SM 最大 Block 数（H100 为 32）, SM 寄存器总数 / 每 Block 寄存器数, SM 共享内存总量 / 每 Block 共享内存用量)) Occupancy的三个限制因素： •寄存器限制（Register Limiter）：最常遇到。当每个Thread使用大量寄存器时，SM可容纳的Thread/Block大幅减少。 •共享内存限制（Shared Memory Limiter）：当Block使用大量共享内存（如矩阵乘法的Tile、FlashAttention的工作缓冲区）。 •Block/Warp槽位限制（Slot Limiter）：Block数量太少（每个Block含极多Threads）或Block数量太多但每个Block Threads太少。

2.6.3 FP16 矩阵乘 Occupancy

以 FP16 GEMM、Tile 128×128 的 Kernel 为例（共享内存用量 A_tile 128×128×2=32KB + B_tile 32KB = 64KB，C_tile 存于寄存器）： •Config A（256 线程/Block，8 Warps，48 寄存器/线程）：每 Block 寄存器 256×48=12288，SM Block 容量 min(32, 65536/12288≈5.3, 228KB/64KB≈3.5)=3，活跃 Warp 3×8=24，Occupancy 24/64=37.5% •Config B（128 线程/Block，4 Warps，64 寄存器/线程）：每 Block 寄存器 128×64=8192，SM Block 容量 8，活跃 Warp 32，Occupancy 50% •Config C（512 线程/Block，16 Warps，32 寄存器/线程）：每 Block 寄存器 512×32=16384，SM Block 容量 4，活跃 Warp 64，Occupancy 100%

2.6.4 Occupancy 与性能的权衡

高Occupancy不一定意味着高性能，这是许多初学者常见的误区。不同 Kernel 类型的最优 Occupancy 如下： •访存密集 Kernel（如逐元素操作、归约）：需要高 Occupancy 以通过 Warp 切换隐藏内存延迟，最优 Occupancy 为 50%-100% •计算密集 Kernel（如 GEMM、卷积）：更多寄存器意味着每个线程每周期做更多工作，但每线程寄存器多会导致 SM 容纳的 Block 减少、Occupancy 下降，最优 Occupancy 为 25%-50%（寄存器饱和优先于 Occupancy） •延迟敏感 Kernel（如顺序依赖操作）：需要最高 Occupancy 隐藏指令延迟，最优 Occupancy 尽量接近 100% 经验法则：对于Tensor Core密集的计算（如矩阵乘法、卷积），让寄存器使用饱和（每Thread接近255个寄存器）通常比追求高Occupancy更重要，因为Tensor Core的计算能力远高于HBM带宽，数据复用远比延迟隐藏关键。

2.6.5 TMA 与 Thread Block Cluster

Hopper（H100/SM90）引入了两项重要特性来缓解Occupancy与寄存器压力的矛盾。

TMA TMA（Tensor Memory Accelerator）是一个专用的硬件单元，支持异步2D/3D张量从全局内存到共享内存的拷贝，完全在CUDA Core之外执行：

 // TMA usage example (simplified)
 // Traditional method (occupies CUDA Cores and registers):
 // Each Thread computes source address, executes load, writes shared memory -- consumes ~20-30 registers
 // TMA method:
 // 1 Thread calls cp.async.bulk to copy the entire Tile -- 0 extra register cost
 // TMA executes on dedicated hardware, freeing CUDA Cores for concurrent compute

TMA将数据加载操作从CUDA Core卸载到专用硬件，释放了宝贵的寄存器资源，使得相同Occupancy下Kernel能执行更多计算，这对于FlashAttention类需要大量数据搬运的算子尤为有利。 2) Thread Block Cluster Cluster允许一个GPC内的Block共享共享内存的分布式访问： // Thread Block Cluster programming cluster_dims(2, 1, 1) // Declare Block Cluster as 2x1x1 global void cluster_kernel() {

     // Can access across Blocks via distributed_shared_memory
 }
 // Launch Cluster

cudaLaunchKernelEx(&config, cluster_kernel, nullptr); // config.clusterDim = dim3(2, 1, 1); Cluster对于FlashAttention等需要跨Block数据共享的算子提供了低延迟的数据交换通道。

2.6.6 Occupancy 分析工具

Nsight Compute

 # Analyze Occupancy with Nsight Compute
 ncu --set full --kernel-name my_gemm_kernel \
     --section Occupancy \
     --section SpeedOfLight \

./my_program

 # Key metrics:
 # - achieved_occupancy: actual Occupancy observed during kernel execution
 # - theoretical_occupancy: theoretical upper bound given resource constraints
 # - occupancy_limit_reason: limiting factor (register/shmem/warp/block)
 # - sm__warps_active.avg: average number of active Warps per SM

CUDA Occupancy API

#include <cuda_runtime.h> #include <cuda_occupancy.h> int min_grid_size, block_size; cudaOccupancyMaxPotentialBlockSize( &min_grid_size, &block_size, my_kernel, 0, // dynamic_shared_memory_size 0 // block_size_limit (0=use max) ); printf("Suggested block size: %d\n", block_size);

2.7 GPU 间互联技术

GPU间互联是万卡集群的“神经网络”，其带宽和拓扑直接决定了分布式训练的可扩展性上限。NVIDIA构建了从芯片级到机架级的完整GPU专有互联体系，涵盖NVLink、NVSwitch、NVLink-C2C等核心互联技术。

2.7.1 互联层次全景

NVIDIA 的互联体系覆盖节点内、机架与集群三个层次，如图2-6所示。 Node-Level Interconnect NVLink: GPU-GPU direct H100: 18 Links x 50 GB/s = Rack-Level Interconnect Cluster-Level Interconnect 900 GB/s bidirectional B200: NVLink 5.0 = 1.8 TB/ NVLink domain: 72 GPU (G s bidirectional B200 NVL72) PCIe: CPU-GPU/Host-GPU NVLink-Crossbar: 130 TB/s Gen5 x16: 64 GB/s unidirec full bidirectional tional

       NVLink-C2C: CPU-GPU 900                Gen6 x16: 128 GB/s (2025
                                                         +)                     NVSwitch: Full mesh insid
                  GB/s                                                                     e node

H100 NVSwitch 3: 64 port s, 3.2 TB/s full-duplex/GPU Cross-domain: NVSwitch c onnected via NVLink 图2-6 GPU互联技术层次体系

2.7.2 NVLink 技术演进

NVLink是NVIDIA专有的高带宽GPU-GPU互联技术，每一代性能翻倍，演进历程如表2-9所示。表2-9 NVLink技术演进版本年代单Link带宽 GPU Link数总双向带宽代表GPU NVLink 1.0 2016 20 GB/s 4 160 GB/s P100 NVLink 2.0 2017 25 GB/s 6 300 GB/s V100 NVLink 3.0 2020 50 GB/s 12 600 GB/s A100 NVLink 4.0 2022 50 GB/s 18 900 GB/s H100 NVLink 5.0 2024 100 GB/s 18 1.8 TB/s B200 NVLink的实际应用场景： •张量并行（Tensor Parallelism）：每层前向/反向传播的AllReduce通信在NVLink域内完成，要求低延迟高带宽 •节点内 AllReduce：8卡服务器的AllReduce操作通过NVSwitch实现全带宽

2.7.3 NVSwitch 架构

NVSwitch是GPU服务器的核心交换芯片，提供节点内所有GPU间的全互联。

架构与规格第3代NVSwitch（H100/H200配套）： •64端口，每端口50 GB/s全双工 •每个GPU连接全部4个NVSwitch（全互联设计） •总全双工带宽：8 GPU × 900 GB/s = 7.2 TB/s •物理形态：4个NVSwitch芯片，集成在HGX基板上
AllReduce硬件加速 NVSwitch通过硬件组播（Multicast）加速AllReduce。传统 Ring-AllReduce 中，N 个 GPU 完成一次 AllReduce 需要传输 2(N-1)/N 倍数据量，以 8 卡、800MB 梯度、400 GB/s NVLink 计算约需 3.5ms；NVSwitch 的硬件组播将传输量降至约 1 倍（800MB），按 7.2 TB/s 交换机容量理论耗时约 0.11ms，实际受限于 NVSwitch 内部缓冲与调度，约 0.2-0.3ms。 NVLS（NVLink Sharp）是Hopper引入的集合通信硬件加速特性，允许NVSwitch在芯片内部完成Reduce/AllReduce操作（而非简单的数据转发），进一步降低了延迟和带宽消耗。

2.7.4 NVLink-C2C 封装

NVLink-C2C是NVIDIA针对CPU-GPU互联的Die-to-Die互连技术，首次出现在Grace-Hopper超级芯片（GH200）中，并在GB200中进一步扩展： •GH200（Grace CPU + H100 GPU）：NVLink-C2C提供900 GB/s的双向带宽（约为PCIe Gen5 ×16的7倍） •GB200（Grace CPU + 2×B200 GPU）：通过NVLink-C2C和NVLink-Crossbar连接，实现130 TB/s的全双向带宽这种封装创新使CPU和GPU可以在同一个统一内存地址空间中高效协作，消除了传统的cudaMemcpy开销，这对内存受限的大模型推理（如405B参数模型的KV Cache管理）具有重大意义。

2.7.5 GPU 拓扑与通信特性

常用 GPU 拓扑的通信特性对比如表2-10所示。表2-10 GPU拓扑与通信特性拓扑典型设备直径对分带宽适用场景 All-to-All (NVSwitch) DGX H100 (8卡) 1 全带宽张量并行，密集AllReduce Ring (8-GPU) 某些4/8卡服务器 N-1 2×单Link 数据并行（Ring-AllReduce） Hierarchical 8×8拓扑（64卡） 3 有限合理混合TP/PP/DP Rail-Optimized H100万卡集群 2 非阻塞大规模数据并行AllReduce Rail-Optimized拓扑是万卡集群的关键设计：将同一位置的GPU（如所有Node0的GPU0）连接到同一交换机，使得 AllReduce的Ring/Recursive Halving/Doubling操作中的每个Step都是同构的（同一牌号交换机、等长光纤），从而将带宽利用率提升至理论峰值的90%-95%。

2.8 GPU 拓扑与 MIG 配置

本节提供可直接在GPU服务器上执行的运维技能，涵盖拓扑检测、NUMA绑定和MIG多租户配置三项核心操作。

2.8.1 GPU 拓扑检测与可视化

 # 1. Basic GPU info
 nvidia-smi
 # Output: GPU model
 # 2. GPU topology matrix
 nvidia-smi topo -m

典型8卡H100服务器输出解读： GPU0 GPU1 GPU2 GPU3 GPU4 GPU5 GPU6 GPU7 GPU0 X NV18 NV18 NV18 NV18 NV18 NV18 NV18 GPU1 NV18 X NV18 NV18 NV18 NV18 NV18 NV18 ... GPU7 NV18 NV18 NV18 NV18 NV18 NV18 NV18 X Legend: NV : NVLink (NV18 = 18 NVLink lanes x 50 GB/s per link bidirectional = 900 GB/s bidirectional) PIX : PCIe (under the same PCIe Switch) PHB : PCIe Host Bridge SYS : Cross-NUMA node (via CPU/memory, slowest) 关键分析： •所有GPU之间均为NV18，这意味着8卡全互联，张量并行无跨节点瓶颈 •如果看到SYS标识，表示GPU对跨NUMA节点，CUDA Kernel的数据访问需经过CPU内存

2.8.2 NUMA 亲和性绑定

 # 3. Detect NUMA topology and CPU affinity
 nvidia-smi topo -c cpu
 # Output: GPU0 -> CPU0 (NUMA node 0)
 # 4. Check NVLink interconnect status
 nvidia-smi nvlink -s
 # Output: Link status
 # 5. List GPU-supported PCIe
 nvidia-smi -q -d CLOCK | grep -A2 "PCIe"
 # PX16: PCIe Gen4 x16
 # 6. Get detailed GPU bus
 nvidia-smi --query-gpu=index,name,pci.bus_id,gpu_bus_id \
     --format=csv,noheader

在多NUMA节点系统中，CPU进程与GPU之间的NUMA跨节点访问会导致30%-50%的性能退化：

Check NUMA topology

lscpu | grep NUMA numactl --hardware

Typical output: 2 NUMA nodes

使用numactl绑定训练进程：

 # Check GPU0 NUMA node
 cat /sys/bus/pci/devices/0000:17:00.0/numa_node   # GPU0 PCI BDF
 # Output: 0 (NUMA0)
 # Method 1: numactl + GPU visibility
 NUMA0_GPUS="0,1,2,3" # GPUs on NUMA0
 NUMA1_GPUS="4,5,6,7" # GPUs on NUMA1
 # Process 1: bind to NUMA0
 CUDA_VISIBLE_DEVICES=$NUMA0_GPUS numactl --cpunodebind=0 --membind=0 \

python train.py --gpu 0,1,2,3 &

Process 2: bind to NUMA1

CUDA_VISIBLE_DEVICES=$NUMA1_GPUS numactl --cpunodebind=1 --membind=1
python train.py --gpu 0,1,2,3 &

Method 2: use taskset (CPU affinity binding)

CUDA_VISIBLE_DEVICES=0,1,2,3 taskset -c 0-31 python train.py & PyTorch中检测NUMA亲和性：

 import torch
 import os
 # Get GPUs available to current process
 ngpus = torch.cuda.device_count()
 for i in range(ngpus):
     props = torch.cuda.get_device_properties(i)
     print(f"GPU {i}: {props.name}, "

f"total memory: {props.total_memory / 1024**3:.1f} GB")

2.8.3 MIG 配置实战

 # Set environment variables at startup
 os.environ['CUDA_VISIBLE_DEVICES'] = '0,1,2,3'
 # Coordinate with numactl binding at process launch

MIG（Multi-Instance GPU）将一张物理GPU切分为多个独立实例，每实例拥有独立的内存、缓存和计算单元。MIG最适用于多租户推理场景（多个小模型共享一张GPU，彼此互不干扰）。各GPU的MIG支持情况如表2-11所示。表2-11 MIG支持矩阵 GPU MIG 最大实例数 MIG Profile粒度 A100-40GB 支持 7 GPU Instances 1g.5gb 约 7g.40gb A100-80GB 支持 7 1g.10gb 约 7g.80gb H100 支持 7 1g.10gb 约 7g.80gb A30 支持 4 1g.6gb 约 4g.24gb V100 不支持 N/A 不支持 RTX 4090 不支持 N/A 消费级不支持

MIG 管理命令

 # 1. Enable MIG mode (requires GPU reset)
 nvidia-smi -i 0 -mig 1
 # Or enable on all GPUs: nvidia-smi -mig 1
 # 2. Check MIG capability
 nvidia-smi mig -lgip
 # Output:
 # GPU 0 Profile ID 19 Placements: {0}:1      (1g.10gb)
 # GPU 0 Profile ID 14 Placements: {0}:2      (2g.20gb)
 # GPU 0 Profile ID 9 Placements: {0}:4       (3g.40gb)
 # GPU 0 Profile ID 5 Placements: {0}:7       (7g.79gb)
 # 3. Create GPU instances
 # Plan A: 2x 2g.20gb + 1x 3g.40gb = 7g
 nvidia-smi mig -cgi 14,14,9 -C
 # Plan B: 7x 1g.10gb (full partitioning into 7 equal instances)
 nvidia-smi mig -cgi 19,19,19,19,19,19,19 -C
 # 4. Create compute instances
 nvidia-smi mig -cci -i 0
 # 5. View MIG instances
 nvidia-smi
 # Output will show MIG 1g.10gb instances as separate GPU entries
 # 6. Delete MIG config
 nvidia-smi mig -dci # Delete all compute instances
 nvidia-smi mig -dgi # Delete all GPU instances
 # 7. Disable MIG mode
 nvidia-smi -i 0 -mig 0
 nvidia-smi -r # Or reset GPU

Profile 选择策略 H100 80GB 的常见切分方案： •7g.79gb：整卡单实例 •4g.40gb：半卡，适合大型推理 •3g.40gb：单实例，可与 2 个 2g.20gb 组合 •2g.20gb：中等推理，最多 3 个实例（2+2+3） •1g.10gb：小型推理，最多 7 个实例例如部署 7 个 INT4 量化的小型模型（如 Llama-7B 权重约 4GB），每个实例 1g.10gb 约 10GB HBM 足够： nvidia-smi mig -cgi 19,19,19,19,19,19,19
Kubernetes 部署

Pod YAML requesting MIG device

apiVersion: v1 kind: Pod spec: containers:

name: inference resources: limits: nvidia.com/mig-1g.10gb: 1 # Request 1x 1g.10gb MIG instance

env: - name: CUDA_VISIBLE_DEVICES value: "0" 4) 性能验证

 # Verify compute after creating MIG instances
 # Run mini inference test on each MIG instance
 for mig_device in $(nvidia-smi -L | grep MIG | awk -F':' '{print $1}'); do
      CUDA_VISIBLE_DEVICES=$mig_device python -c "
 import torch
 x = torch.randn(1024, 1024, device='cuda')
 t = torch.mm(x, x.T)
 print(f'Device {torch.cuda.current_device()}: {t.sum():.2f}, OK')

" done 5) 局限性

CUDA 11.4之前MIG与MPS（Multi-Process Service）互斥；CUDA 11.4+起，每个MIG Compute Instance内可独立启动MPS Server，但跨MIG实例的MPS共享仍不支持
MIG实例间不能P2P通信（无NVLink，即使同一物理GPU的不同MIG实例之间）
MIG实例不支持动态调整（需删除重建）
仅Ampere及后续架构的数据中心专业GPU支持（A100、H100、A30等；V100等Volta架构及消费级GPU不支持） MIG改变了每个实例的SM数量，从而影响Occupancy上限。Kubernetes 通过 Device Plugin 与 MIG 策略实现此类资源的调度。