NVIDIA A100：从芯片到 SIMT¶

A100 是理解现代数据中心 GPU 的合适起点：它把大规模 SIMT 计算、第三代 Tensor Core、统一 L1/共享内存、40 MB L2、HBM 和 NVLink 放在同一套层次结构中。本页依次画出 整体架构、SM 微架构、访存架构与 SIMT 编程模型，最后用一个 batch=1 的 W8A8 Llama 线性层把 Kernel、编译产物与 Tensor Core 执行串起来。

1. 整体架构¶

下图把包含、调度、片上互连与完整访存层次放在同一张系统图中：外框表示包含，数量标签与末端摘要卡表示同类单元重复，紫色虚线表示 Host 发起 Kernel 后的控制/调度路径，青色双向箭头表示数据通路。每个 SM 内嵌自己的 Unified L1/Texture Cache + Shared Memory；L2、Crossbar 和内存控制器位于 GPC 之外，由所有启用的 SM 共享。展开的 GPC 0 / TPC 0 是结构代表，不表示只有它包含下一级。

1.1 从完整 Die 到 A100 产品¶

这里最重要的不是背诵 108，而是看清可扩展层次：GPC 包含多个 TPC，TPC 包含两个 SM，SM 是执行 CUDA Thread Block 的核心计算单元。

2. SM 微架构¶

一个 A100 SM 可以理解为四个对称处理分区：每个分区都有自己的 L0 指令缓存、Warp Scheduler/Dispatch、寄存器文件切片和执行单元；四个分区共同连接 SM 级 L1 指令缓存与统一 L1/共享内存。四张分区卡片明确表示物理重复关系，不把面积比例当成真实版图。

四个分区合计得到每 SM 的 64 个 FP32 Core、64 个 INT32 Core 与 4 个 Tensor Core。Tensor Core 执行矩阵乘加；普通 CUDA Core、整数单元、Load/Store 与特殊函数单元共同覆盖地址计算、标量运算和其他指令。

2.1 从 Warp 指令到五类执行管线¶

每个处理分区的 Warp Scheduler 从已就绪 Warp 中选择一条指令；Dispatch 传递的是同一条 Warp 指令及其 Active Mask，各 Lane 的操作数和结果仍分布在寄存器文件中。下图用紫色虚线表示发射/控制，用实线表示操作数或数据流，避免把“谁控制谁”和“数据往哪里走”画成同一种箭头。

flowchart TB
    WS["Warp Scheduler<br/>选择 Ready Warp"] -. "发射 + Active Mask" .-> DP["Dispatch"]
    DP -.-> PIPE["Execution Pipes / Partition<br/>FP32 ×16 · INT32 ×16 · Tensor ×1<br/>LD/ST ×8 · SFU ×4"]
    RF["Register File Slice<br/>16K × 32-bit"] <--> PIPE
    PIPE <-->|"仅 LD/ST 数据路径"| L1["Unified L1 / Shared"]
    L1 <--> L2["Shared L2"]

    linkStyle 0,1 stroke:#7c4dff,stroke-width:2px,stroke-dasharray:5 4
    linkStyle 2,3,4 stroke:#008ea3,stroke-width:1.5px

四个分区合起来就是图中的 64 FP32 + 64 INT32 + 4 Tensor Core + 32 LD/ST + 16 SFU。单元数量描述的是并行资源，不等于每条指令的固定延迟；吞吐还会受数据相关、发射冲突、Active Mask、Cache Miss 和 Bank Conflict 影响。

2.2 CUDA Core：Warp 共享控制的标量向量流水线¶

这里的 CUDA Core 特指白皮书计数口径中的 FP32 算术 Lane，不是一个能独立取指的“小 CPU”。其功能路径可以分成四步：

1. Warp Scheduler 选择一条就绪的 FP32 Warp 指令，例如 SASS FADD、FMUL 或 FFMA。
2. Dispatch 将同一操作码和 Active Mask 发给所在分区；各活跃 Lane 从寄存器文件切片读取自己的源操作数。
3. 每分区 16 个 FP32 Lane 执行逐 Lane 运算，四个分区在峰值吞吐口径下合计产生 64 个 FP32 结果/clk/SM。若是 FFMA，一个结果包含一次乘和一次加，因此性能表按 FLOP 计数时是 128 FLOP/clk/SM。
4. 结果写回各 Thread 的架构寄存器；Scoreboard 在依赖满足前阻止后续相关指令发射。吞吐数字不能直接当成单条指令延迟。

A100 还在每分区配置 16 个独立 INT32 Lane。地址更新、循环计数等整数运算可与 FP32 运算并行，这正是白皮书所说的“FP32 与 INT32 同时执行”；但是否真的重叠取决于编译出的指令序列、依赖关系和 Scheduler 当时可发射的 Warp。Thread 只是状态与 Lane 语义，执行期间会先后使用 FP32、INT32、LD/ST、SFU 等不同管线，不能永久绑定到某一个 CUDA Core。

2.3 Tensor Core：Warp 级矩阵指令和分布式 Fragment¶

Tensor Core 与 CUDA Core 共用 Scheduler、寄存器和存储层次，但消费的是 Warp 级 MMA 指令。一个典型高吞吐数据通路如下：

flowchart TB
    GM["HBM / L2"] -->|"coalesced load 或 cp.async"| SMEM["Shared Memory Tile"]
    SMEM -->|"ldmatrix / warp load"| FRAG["32 Lane 寄存器中的<br/>A/B Fragment"]
    FRAG --> MMA["mma.sync / HMMA · IMMA · DMMA<br/>+ C Accumulator<br/>+ 可选 2:4 稀疏元数据"]
    MMA --> OUT["D / Accumulator Fragment"]
    OUT --> NEXT["后续 MMA 或 store"]

• 协作粒度是一个 Warp。 A、B、C/D 的 Fragment 分散在 32 个 Lane 的寄存器中；某个 Thread 不拥有完整矩阵。mma.sync 要求参与 Lane 共同执行，Scheduler 发射一次 Warp 级矩阵指令。
• API Tile、PTX 指令和物理 Tensor Core 不是同一层。 WMMA 常见 Tile 如 m16n16k16；Ampere PTX/SASS 可使用更细的 m16n8k16 等 MMA 形状；白皮书给出的物理吞吐示例则是每个 Tensor Core 每周期完成一个 8×4×8 的 FP16/FP32 混合精度矩阵乘加。不要把其中任一形状直接解释成电路阵列的永久二维尺寸。
• 每分区一个、每 SM 四个。 每个第三代 Tensor Core 可做 256 个 FP16 FMA/clk，四个合计 1024 FP16 FMA/clk/SM；同时支持 BF16、TF32、FP64、INT8、INT4 与 Binary 等模式，不同模式的输入、累加类型和吞吐不同。
• 结构化稀疏是额外模式。 2:4 稀疏 MMA 携带非零值与元数据，硬件跳过规定位置的零，在满足格式与软件支持时把有效 Tensor 运算吞吐提高到最多 2 倍；它不是对任意稀疏矩阵自动加速。

第三代 Tensor Core 的重要改进不只是“单 Core 更宽”：A100 让数据在 Warp 的 32 个 Thread 间共享（Volta 为 8 Thread 范围），并增大单条指令的 K 维工作量，从而减少寄存器读取和指令数。官方没有公开其 Operand Collector、乘法阵列与归约树的完整流水级，因此这里不画臆测的内部门级结构。

2.4 Load/Store 与 SFU：地址、合并访问和特殊函数¶

Shared Memory 与 L1 共用 192 KB 物理块，但编程语义不同：Shared Memory 是软件显式寻址的 32-Bank Scratchpad，每个 Bank 每周期 32 bit；L1 是硬件管理的 Cache/Coalescing Buffer。32 Banks × 32-bit/clk = 128 B/clk 只描述无 Bank Conflict 时的 Shared Memory Bank 聚合口径，不能改写成“A100 L1 的公开物理总线位宽”。

3. 访存架构¶

3.1 层次、作用域与物理位置¶

flowchart TB
    T["Thread<br/>标量状态"] --> R["Register<br/>线程私有 · SM 内"]
    R --> S["Shared Memory / L1 Data Cache<br/>Block 协作 / 硬件缓存 · SM 内"]
    S --> L2["L2 Cache · 40 MB<br/>所有 SM 共享"]
    L2 --> HBM["Global Memory · HBM<br/>设备级持久数据"]
    HBM <-->|PCIe 4.0| HM["Host Memory"]
    L2 <-->|NVLink 3.0| PM["Peer GPU Memory"]

    LM["Local Memory<br/>线程私有的地址空间"] -. 寄存器溢出等访问最终落到显存并经过缓存 .-> L2

3.2 A100 产品带宽不要混用¶

NVLink 的 600 GB/s 是 12 条链路的双向聚合 GPU 互连能力，不应直接当作 CPU 主存带宽；PCIe Gen4 x16 的 64 GB/s 也是双向聚合标称值，单方向理论量级约为 32 GB/s。

3.3 一次 Tile 计算的数据路径¶

sequenceDiagram
    participant H as HBM
    participant L as L2 / L1
    participant S as Shared Memory
    participant R as Register
    participant C as CUDA / Tensor Core
    H->>L: 读取输入 Tile
    L->>S: 合并并异步搬运
    S->>R: Warp 读取操作数片段
    R->>C: 发射标量或矩阵指令
    C->>R: 累加局部结果
    R->>S: 需要时交换/重排
    S->>H: 经 Cache 写回输出 Tile

A100 引入了从 Global Memory 到 Shared Memory 的硬件加速异步拷贝，可绕开中间寄存器占用，并允许软件把数据搬运与计算重叠。这正是“软件调度 × 存储层次 × 计算单元”协同的典型例子。

4. SIMT 编程模型¶

4.1 软件层次如何映射到硬件¶

先看软件对象的包含关系，再看 Block 如何被调度到 SM。两者是不同关系，不用同一组箭头混画。

上图回答“谁包含谁”；下面回答“运行时谁被交给谁”。这两组层次不是四级一一对齐：Grid → Block → Warp → Thread 是软件组织，GPC → TPC → SM → 执行流水线 是硬件组织，真正稳定的调度契约落在 Block → SM。

把这张图读成四句话即可：

1. Grid 面向整个 CUDA 设备。 A100 的工作分发逻辑把 Grid 中尚未执行的 Block 分派到可用 SM；基础 CUDA 启动语法不能指定 GPC、TPC 或 SM，具体分发次序也不是程序语义的一部分。
2. GPC/TPC 是硬件分组，不是软件层次的“配对项”。 A100 中 GPC → TPC → SM 最终提供 108 个 SM；普通 Grid 可以使用跨多个 GPC 的 SM，但并不存在 1 Grid = 1 GPC 或 1 Block = 1 TPC 的关系。
3. Block 是不可拆分的驻留单位。 一个 Block 的全部线程都由同一个 SM 执行；因此它的所有 Warp 也留在该 SM 内。反过来，只要线程数、寄存器、Shared Memory 等资源允许，一个 SM 可同时常驻多个 Block。
4. Warp 才是 SM 内的指令发射单位。 Scheduler 发射的是“某个 Warp 的下一条指令 + Active Mask”，不是逐个发射 32 个独立 Thread；Thread 保留各自的数据与控制状态，当前指令再决定使用哪类执行单元。

4.2 Block 如何分批进入 SM¶

Grid 可以包含远多于物理 SM 容量的 Block。尚未获得 SM 资源的 Block 只是待执行工作；某个 SM 有资源空位时，完整接纳一个 Block，为它的全部线程分配寄存器、为整个 Block 分配 Shared Memory，并把它拆成常驻 Warp。Block 完成后资源释放，后续 Block 再补入，因此大 Grid 会以许多动态“波次”推进。

A100（Compute Capability 8.0）每个 SM 的硬上限是 2048 Thread、64 Warp、32 Block；但真实常驻 Block 数取各项约束中最紧的一项。可把它粗略理解为：

其中 \(N_{\text{register}}\) 和 \(N_{\text{shared}}\) 分别表示寄存器容量、Shared Memory 容量能支持的并发 Block 数。这两项还受分配粒度影响，实际应以编译器资源报告、CUDA Occupancy API 或 Nsight Compute 为准。尤其要区分：常驻（Resident）只表示执行上下文和资源已经在 SM 上，并不表示这些 Warp 正在同一个周期同时做运算。

4.3 Warp 如何在 SM 内真正执行¶

一个 A100 SM 有四个处理分区，每个分区包含 Warp Scheduler/Dispatch、寄存器文件切片和一组执行流水线。常驻 Warp 会处于“可发射”或“等待数据/依赖/同步”等不同状态；每个调度器从自己负责的常驻 Warp 中选择一个下一条指令已就绪的 Warp 发射。Warp 的寄存器和控制状态一直保存在片上，因此从一个 Ready Warp 换到另一个不需要像 CPU 线程那样保存、恢复上下文。

以 y[i] = a * x[i] + y[i] 为例，进入 if 的一个 Warp 大致会经历下列硬件路径；这是理解 SIMT 比“一个线程对应一个 Core”更准确的方式：

当 Warp A1 因 HBM 访问等待数百周期时，Scheduler 可以发射已经 Ready 的 Warp F3。延迟隐藏就来自大量常驻 Warp 中持续找到可做的工作，而不是让一个 Warp 的 32 个 Thread 随意跑到其他 SM。高 Occupancy 会增加可选择的 Warp，但如果它们都在等待同一个瓶颈，也不保证更高性能。

4.4 分支发散如何发生¶

flowchart LR
    PC["同一 Warp · 同一条指令"] --> P{"每个线程的条件"}
    P -->|偶数 Lane| A["执行 Path A<br/>Mask: 1010…"]
    P -->|奇数 Lane| B["执行 Path B<br/>Mask: 0101…"]
    A --> R[Reconverge]
    B --> R
    R --> NEXT[继续执行共同路径]

当同一 Warp 内线程走不同分支时，硬件用 Active Mask 分别执行所需路径，未参与该路径的 Lane 被屏蔽，因此有效利用率下降。不同 Warp 则可以独立推进，Warp Scheduler 会在可运行 Warp 之间切换以隐藏访存或流水线延迟。

__global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];
    }
}

这个 kernel 中，相邻线程访问相邻元素，有利于合并 Global Memory 访问；边界处最后一个 Warp 可能只有部分 Lane 有效。实践中通常让 Block 线程数为 32 的倍数，但最终仍需结合寄存器、Shared Memory、指令延迟和实际 occupancy 判断。

5. 贯通示例：Llama batch=1 W8A8 Tensor Core GEMM 从 CUDA 到 A100¶

这一章不再分别介绍名词，而是跟踪一次 W8A8 线性层从量化公式、CUDA C++、PTX/SASS，一直走到 A100 的 GPC、SM、Warp Scheduler、Tensor Core 和存储层次。示例刻意保持完整、可推演：它明确发出 INT8 WMMA 指令，但仍是教学 Kernel，不是替代 cuBLAS、CUTLASS、TensorRT-LLM 或专用量化库的生产实现。

5.1 先把 W8A8 的“整数计算、最后反量化”说清楚¶

在单 batch、单 token 的 Llama 风格线性层中，逻辑运算仍是：

从数学形状看，它是 M=1 的 GEMM，也就是 GEMV；本章称其为 Tensor Core GEMM，是因为实现仍把它组织成 Warp 级矩阵 Tile，而不是改回逐元素标量点积。

本章采用一种易于追踪的对称 W8A8 口径：Activation 使用当前 token 的 per-tensor scale \(s_x\)，Weight 使用每个输出通道的 per-channel scale \(s^W_o\)，Zero Point 均为 0。权重在离线阶段量化；Activation 在运行时由上一算子或单独的量化步骤产生：

真正送进 Tensor Core 的两个乘数都是 INT8；乘积累加到 INT32，完整点积结束后才在 Epilogue 中反量化：

即使按最保守的 \(|q|\le128\) 估算，\(|z_o|\le4096\times128^2=67{,}108{,}864<2^{31}\)，本例不会溢出 INT32。若线性层带 Bias，可在反量化后的浮点 Epilogue 中相加；若下一层仍接收 INT8，也可把 Epilogue 改为 Requantize，而不必写出 FP16。

显存中的 qweight[N,K] 按输出通道行主序存放；数学上使用的是它的逻辑转置 [K,N]，运行时不需要真的转置整张权重矩阵。

5.2 教学 Kernel：4 个 Warp 用 Tensor Core 计算 16 个输出¶

WMMA 的 INT8 m16n16k16 运算以一个 Warp 为协作单位，计算 \(16\times16\) 输出 Tile。逻辑 Activation 只有一行，因此下面把它放到 A Tile 的第 0 行，其余 15 行补 0；每个 Block 的 4 个 Warp 分别处理 1024 个 K 元素，最后把四份 INT32 部分和相加。

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89

#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <mma.h>

namespace wmma = nvcuda::wmma;
using int8 = signed char;

constexpr int K = 4096;
constexpr int N = 4096;
constexpr int WMMA_M = 16;
constexpr int WMMA_N = 16;
constexpr int WMMA_K = 16;
constexpr int WARPS_PER_BLOCK = 4;
constexpr int THREADS = WARPS_PER_BLOCK * 32;
constexpr int K_PER_WARP = K / WARPS_PER_BLOCK;

__global__ void w8a8_tc_gemv_kernel(
    const int8* __restrict__ qact,       // [K], INT8
    const int8* __restrict__ qweight,    // [N, K], row-major INT8
    float act_scale,                     // per-token s_x
    const float* __restrict__ w_scale,   // [N], per-output s_W
    half* __restrict__ y) {              // [N], FP16

    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int lane = tid & 31;
    const int out_base = blockIdx.x * WMMA_N;
    const int k_begin = warp * K_PER_WARP;
    const int k_end = k_begin + K_PER_WARP;

    // 每个 Warp 各有对齐的 A/B Tile 与 INT32 部分结果。
    __shared__ __align__(32) int8 a_tile[WARPS_PER_BLOCK][WMMA_M * WMMA_K];
    __shared__ __align__(32) int8 b_tile[WARPS_PER_BLOCK][WMMA_K * WMMA_N];
    __shared__ __align__(32) int c_tile[WARPS_PER_BLOCK][WMMA_M * WMMA_N];

    wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K,
                   int8, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K,
                   int8, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K,
                   int> acc_frag;
    wmma::fill_fragment(acc_frag, 0);

    // 每个 Warp 沿自己的 K 区间循环 1024 / 16 = 64 次。
    for (int k0 = k_begin; k0 < k_end; k0 += WMMA_K) {
        for (int i = lane; i < WMMA_M * WMMA_K; i += 32) {
            // 逻辑 M=1：第 0 行有效，其余 15 行补 0。
            a_tile[warp][i] = (i < WMMA_K) ? qact[k0 + i] : int8{0};

            // B 逻辑形状为 [K,N]；qweight[N,K] 行主序恰好是其列主序存储。
            const int n = i / WMMA_K;
            const int kk = i % WMMA_K;
            b_tile[warp][i] = qweight[(out_base + n) * K + k0 + kk];
        }
        __syncwarp();

        wmma::load_matrix_sync(a_frag, a_tile[warp], WMMA_K);
        wmma::load_matrix_sync(b_frag, b_tile[warp], WMMA_K);
        wmma::mma_sync(acc_frag, a_frag, b_frag, acc_frag); // s8 × s8 → s32
        __syncwarp(); // 确保 Shared Tile 不会在仍被读取时覆写
    }

    wmma::store_matrix_sync(
        c_tile[warp], acc_frag, WMMA_N, wmma::mem_row_major);
    __syncthreads();

    // A Tile 只有第 0 行有效；前 16 个 Thread 合并四个 K Slice 并反量化。
    if (tid < WMMA_N) {
        int z = 0;
        #pragma unroll
        for (int w = 0; w < WARPS_PER_BLOCK; ++w) {
            z += c_tile[w][tid];
        }
        const int o = out_base + tid;
        y[o] = __float2half(static_cast<float>(z) * act_scale * w_scale[o]);
    }
}

void launch_w8a8_tc_gemv(
    const int8* qact,
    const int8* qweight,
    float act_scale,
    const float* w_scale,
    half* y,
    cudaStream_t stream) {
    // qact 与 act_scale 已由上游 Quantize / Requantize 步骤产生。
    w8a8_tc_gemv_kernel<<<N / WMMA_N, THREADS, 0, stream>>>(
        qact, qweight, act_scale, w_scale, y);
}

这段代码与编程模型的对应关系是：

这里不能再把一个 Thread 解释成“计算一个输出元素”。WMMA Fragment 在 32 个 Lane 的寄存器间按硬件规定分布，其元素到 Lane 的具体映射不是 CUDA C++ API 的稳定接口；只有整个 Warp 的 32 个 Lane 都执行 load_matrix_sync、mma_sync 和 store_matrix_sync，一个矩阵 Tile 才完整。

5.3 从 CUDA C++ 编译成 A100 的 INT8 Tensor Core 指令¶

编译 A100（Compute Capability 8.0）版本时，可保留不同层级的产物，并检查最终机器码是否真的出现整数矩阵乘加：

# 常用构建：生成可链接的 Host Object，并内嵌 A100 Device Code
nvcc -O3 -std=c++17 -arch=sm_80 -lineinfo -c w8a8_tc_gemv.cu -o w8a8_tc_gemv.o

# 分别观察虚拟 ISA 与 A100 原生机器码
nvcc -O3 -std=c++17 -arch=compute_80 -ptx w8a8_tc_gemv.cu -o w8a8_tc_gemv.ptx
nvcc -O3 -std=c++17 -arch=sm_80 -cubin w8a8_tc_gemv.cu -o w8a8_tc_gemv.cubin
cuobjdump -sass w8a8_tc_gemv.cubin

要抓住三个边界：

1. WMMA 是 CUDA C++ 的 Warp 级矩阵 API。 Fragment 是分布在 32 个 Lane 寄存器中的协作对象，不是某一个 Thread 的普通数组。
2. PTX 是虚拟 ISA。 本例应能看到 wmma.load、wmma.mma.sync...s32.s8.s8.s32、wmma.store 等语义；A100 不直接执行 PTX 文本。
3. SASS/Cubin 才是目标机器码。 ptxas 为 sm_80 完成寄存器分配、指令选择和调度；最终用 cuobjdump -sass 检查 IMMA-family 指令。

源代码与指令族大致有如下对应。确切寄存器、指令数量和次序随 CUDA 版本与优化而变化，应以本机产物为准：

5.4 Launch 之后：一个 Block 在 A100 上经历什么¶

Host 侧的 <<<256, 128, 0, stream>>> 只描述 Grid、Block、动态 Shared Memory 和 Stream，不指定 GPC/TPC/SM。Kernel Launch 对 Host 通常异步；首次模块装载或 PTX JIT 仍可能带来一次性准备开销。

下面冻结一个 Warp 的一次 K-Tile 迭代，再接上 Block Epilogue：

代码中的关键构造可落到如下硬件层次：

5.5 把规模数字代回 A100¶

• Grid 有 4096 / 16 = 256 个 Block；每 Block 128 Thread，也就是 4 Warp。
• 每个 Warp 处理 1024 个 K 元素，共发出 1024 / 16 = 64 次 Warp 级 WMMA；每 Block 256 次，整个 Grid 共 65,536 次 wmma::mma_sync 调用。
• 逻辑问题仍只有 4096 × 4096 = 16,777,216 次有用 INT8 乘加。一次 m16n16k16 WMMA 覆盖 4096 次物理乘加；由于 batch=1 只使用 16 行中的第 0 行，本教学映射会做约 16 倍矩阵算术，其中 15 行是 Padding。这正是 M=1 的 Tensor Core Tile 利用率问题。
• A100 每 SM 上限为 2048 Thread / 64 Warp，因此仅从 Thread/Warp 上限看，最多可同时容纳 2048/128 = 16 个本例 Block；约 6 KiB Shared Memory、Fragment 寄存器分配以及每 SM 的 Block 上限还会进一步约束实际常驻量，应查看 nvcc -Xptxas=-v 与 Nsight Compute。
• 256 个 Block 足以在理想条件下覆盖 108 个 SM，并留下尾部 Block；实际准入仍是空位即补入，不存在全芯片整齐同步的“波次栅栏”。

从数据量看，INT8 qweight[4096,4096] 约 16 MiB，qA[4096] 仅 4 KiB，FP32 per-output Weight Scale 为 16 KiB，FP16 输出为 8 KiB。每生成一个 token，绝大多数权重仍只使用一次；W8A8 把两个乘数都变成 INT8 并启用高吞吐 Tensor Core，却没有改变 batch=1 时权重流量占主导、数据复用很低的事实。因此它通常仍更接近带宽受限，而不是接近 A100 的峰值 INT8 TOPS。

5.6 用一条因果链复述整个执行流¶

上游把当前 Activation 量化成 qA + sₓ → Driver 准备 sm_80 Cubin → Host 发起 <<<256,128>>> → Block 按资源空位落到 108 个 SM → 每 Block 的 4 个 Warp 分担四段 K → Lane 协作把 INT8 A/B Tile 从 HBM/L2/L1 搬入 Shared → wmma.load 形成分布式寄存器 Fragment → Warp Scheduler 发射 IMMA → Tensor Core 完成 s8 × s8 → s32 → 四个 Warp 的 INT32 K Slice 经 Shared Memory 合并 → Epilogue 乘 sₓ × sᵂ[o] 反量化 → Thread 0–15 写 16 个 FP16 输出 → 全部 256 个 Block 完成后得到 y[1,4096]。

如果这条链能够不看图复述出来，就已经把 W8A8 量化、CUDA SIMT、Tensor Core 编程、编译栈与 A100 硬件架构真正连在了一起。

6. 校验问题¶

读完后应能独立回答：

1. 完整 GA100 的 128 SM 与 A100 的 108 SM 为什么不能混写？
2. 为什么不能把 Grid → Block → Warp → Thread 与 GPC → TPC → SM → CUDA Core 逐级一一对应？
3. 一个大 Grid 为什么会分多个波次执行？Block、Warp 与 SM 的驻留关系分别是什么？
4. 为什么 32 Threads / Warp 不代表每个 Thread 永久对应一个 CUDA Core？
5. 192 KB Unified L1/Shared 为什么不代表 Block 可申请 192 KB Shared Memory？
6. Local Memory 为什么可能比 Shared Memory 慢很多？
7. 分支发散、访存不合并和低 Occupancy 分别损失了什么资源？
8. 80 GB SXM 的 HBM 带宽、NVLink 带宽与 PCIe 带宽各描述哪条路径？
9. 对称 W8A8 中，为什么 INT8 × INT8 可以先完整累加到 INT32，再统一乘 sₓ × sᵂ[o] 反量化？
10. 教学 Kernel 中，Block、Warp 与 Lane 分别负责 N Tile、K Slice 和哪一类协作？
11. 为什么不能把 WMMA Fragment 的某个输出元素永久对应到某个 Thread？
12. wmma::mma_sync、PTX wmma.mma...s32.s8.s8.s32 与 A100 IMMA-family SASS 分别处在哪一层？
13. 为什么 batch=1 即使使用 INT8 Tensor Core，仍会同时面临 M Tile 利用率低和 HBM 权重带宽瓶颈？

参考资料¶

• 参考页面：HiHuo《GPU 架构基础》
• NVIDIA A100 Tensor Core GPU Architecture Whitepaper
• NVIDIA A100 Tensor Core GPU Datasheet
• CUDA Programming Guide：Thread Blocks and Grids
• CUDA Programming Guide：Warps and SIMT
• CUDA Programming Guide：Hardware Multithreading
• CUDA Programming Guide：Writing SIMT Kernels
• CUDA Programming Guide：Coalesced Global Memory Access
• CUDA Programming Guide：Warp Shuffle Functions
• CUDA Programming Guide：Warp Matrix Functions / WMMA
• PTX ISA：Warp-level Matrix Multiply-Accumulate
• CUDA Programming Guide：NVCC
• NVIDIA CUDA Compiler Driver：Compilation Trajectory
• NVIDIA CUDA Binary Utilities：cuobjdump
• NVIDIA Developer Blog：Faster Parallel Reductions on Kepler
• NVIDIA CUDA 11 Features：A100 Multi-precision Tensor Cores
• NVIDIA Matrix Multiplication Background：Tensor Core Tile 与性能
• NVIDIA TensorRT：INT8 Quantization Schemes
• NVIDIA Ampere GPU Architecture Tuning Guide：Streaming Multiprocessor
• NVIDIA CUDA Best Practices Guide：Coalesced Access 与 Shared Memory Banks
• NVIDIA CUDA Binary Utilities：Ampere SASS Opcode（FFMA / HMMA / IMMA / MUFU 等）
• Demystifying the Nvidia Ampere Architecture through Microbenchmarking and Instruction-level Analysis
• Dissecting Tensor Cores via Microbenchmarks: Latency, Throughput and Numeric Behaviors