第3章 GPU 体系结构（下）：片上资源与数据通路

本章导读

上一章我们跟着一次 kernel 提交看清了线程怎么划分、波前怎么落到硬件上执行——编程模型这张地图已经在你手上。这一章补上硬件体系结构的另一半。本章会做四件事：搞清楚 GPU 能同时塞下多少活儿（片上资源 VGPR/SGPR/LDS 与占用率）、把数据从寄存器到显存的内存层级理清楚、弄懂为什么读取地址的排列能让带宽差出好几倍（合并访存与 LDS bank 冲突）、最后认识一下矩阵专用指令 WMMA。

这几样是 Part 2 算子优化的直接地基，每一样都有对应的实战章会展开：合并访存在 第 8 章 Element-Wise、LDS 协作与 bank 在 第 9 章 Reduction、占用率/分块/寄存器累加在 第 11 章 GEMM-Like、WMMA 与融合在 第 12 章 Fusion：融合算子。本章只把概念和直觉立起来，具体怎么优化，留到对应算子章。

本章对应代码在：

code/part0-intro/
├── pyproject.toml
├── uv.lock
├── activate-rocm.sh
└── chapter3/
    ├── global_memory_access.hip   # 选做：合并访存
    ├── lds_bank_conflict.hip      # 选做：LDS bank 冲突
    ├── rdna4_wmma.hip             # 选做：VALU vs WMMA
    └── run_all.sh

每个概念都配了一个只改单一变量的受控实验作为选做（就在上面这些文件里），跟着正文跑一遍能建立更牢的直觉。

3.1 VGPR、SGPR、LDS 与占用率

决定 GPU 能同时塞下多少活儿的，是几种片上资源里谁先被用光。

当一个 wavefront 在等数据（访存）时，硬件不会干等，而是切到另一个已经准备好的 wavefront 去跑——这就是用「同时跑很多 wavefront」把等待的时间藏起来（术语叫延迟隐藏）。所以关键问题不是「线程数占满没有」，而是「硬件还能同时塞下多少工作」。对 gfx1201，ROCm 规格给出的资源盘子是：768 KiB 向量寄存器（VGPR）、32 KiB 标量寄存器（SGPR）和 128 KiB 片上内存（LDS）。ROCm GFX1201 register and LDS specifications 注意，这些只是硬件的总容量；你某个具体 kernel 到底用掉多少寄存器、多少 LDS、有没有溢出到 scratch，得从编译结果或 profiling 里读出来才知道。

把 GPU 想成一座工厂：VGPR/SGPR 是每个工人的私人工具柜，LDS 是车间共用的工作台。一座工厂能同时开工多少条产线（也就是占用率），不取决于工人总数，而取决于哪种资源先被用光——工具柜塞太满，能容纳的工人就少；工作台分太大，能并排的产线就少。

图 3.1　驻留资源图：VGPR、SGPR、LDS 都可能成为上限；图不是从总容量直接计算 占用率 的公式。

这三样东西分工不同：VGPR 存每个 lane 自己的向量临时值；SGPR 存整个 wavefront 可以共用的标量状态；LDS 则是 workgroup 用来互相协作、复用数据的那块片上内存。

所以别指望从「LDS 有 128 KiB」这种总数，直接算出某个 kernel 的占用率（Occupancy，也就是能同时塞下多少 wavefront）；更别以为占用率越高越好。真实的有效并发，是寄存器用量、LDS 的分配粒度、workgroup 的形状、硬件上限、能同时跑几个 workgroup、以及瓶颈到底在哪，这些因素一起决定的。比如硬把寄存器压得很低，反而可能逼出 spill（数据被挤到慢得多的显存里）或额外的指令。正确的顺序永远是：先把算法和访存逻辑写对，再去读编译器报告的资源用量，最后用 profiling 判断它是不是真的在干等访存。

迁移范围： 本节的这些容量数字，只锚定 gfx1201 的规格；「哪种资源先用光，就先限制能塞多少」这个思路可以迁移，但具体的占用率、会不会 spill、最佳的 tile 大小，都必须针对目标 GPU、编译器和 kernel 单独去测。占用率与分块在真实算子里怎么权衡，第 11 章 GEMM-Like 会第一次系统地做。

3.2 内存层级：从寄存器到 GDDR6

这一节最容易混淆，所以我们先把三类完全不同的东西分开。

第一类：编译器有时候会把可以复用的临时值，一直留在 VGPR/SGPR 里——这叫寄存器驻留值（Register-Resident Value），它压根不是一次「去内存取数、然后命中了寄存器」的 load/store。第二类：__shared__/LDS 是由你的程序显式分配、用专门的 DS 指令访问的，它是 workgroup 自己的一块片上存储，不是缓存（cache）。第三类：只有真正的全局访存（global 的 vector/scalar memory operation），才会沿着各自的「缓存→显存」路径去要数据。

再看 gfx1201 的硬件数字：ROCm 规格列出 32 KiB 向量 L0、16 KiB 标量 L0、8 MiB L2 和 64 MiB Infinity Cache；RX 9070 XT 的产品规格则是 16 GB GDDR6 显存、256-bit 位宽、最高约 640 GB/s 的理论板卡带宽。ROCm GFX1201 cache specifications AMD RX 9070 XT 产品规格 AMD Radeon 9000 Quick Reference Guide

图 3.2　gfx1201 的并列概念路径图：寄存器驻留、LDS/DS、vector-global 和 scalar-global 不能串成一条统一访问链。箭头不保证每次访问必经，也不是完整的物理拓扑、固定延迟、带宽或命中率图。

这条数据路径可以用取件来理解：寄存器是你手边的桌面，LDS 是车间里你自己摆好货的货架，L2 和 Infinity Cache 是楼下的周转仓库，GDDR6 则是远处的总仓。越远的地方越大、越慢。但要特别记住：LDS 上的货是你亲手搬上去的（显式分配），仓库里的货才是系统自动帮你缓存的——这正是 LDS 和 cache 的本质区别。

如果把这些层级按「离计算多远」排成一座金字塔（图 3.3），规律就一句话：越往下，容量越大、带宽越低、延迟越高。

图 3.3　gfx1201 / RX 9070 XT 的内存层级金字塔：越往下容量越大、带宽越低、延迟越高。再强调一次：LDS 是你显式分配的片上存储，不是缓存；只有全局访存才会沿 L0→L2→Infinity Cache→GDDR6 这条路径走。容量为规格值。

记住两点：图里这些容量数字不是性能排名；也别把 64 MiB 的 Infinity Cache 误写成 L2。更重要的是把下面三件事分清楚：① 规格给的，只是容量和理论板卡带宽；② 你的程序按算法算出来要读写的字节数，叫逻辑字节；③ 数据真正在 GDDR6 上跑了多少（物理流量），还会受缓存命中、写入方式、硬件事务的影响，必须用专门的计数器或受控实验才能说得清。逻辑字节不等于物理流量——这条区分会在 Part 2 反复用到。

迁移范围： 32 KiB L0、8 MiB L2、64 MiB Infinity Cache、16 GB GDDR6 和 640 GB/s，都只锚定 RX 9070 XT / gfx1201 的规格。某次访问到底命中了哪一级、延迟多少、物理流量多大，必须在目标机器上实测。

3.3 全局内存访问与合并访存

写入的地方一模一样，只是读取的顺序不同，速度就能差出好几倍。

想象一排 32 个 lane 同时去读数据：如果它们读的是连续的下标（lane 0 读 0、lane 1 读 1……），一次访存就能喂饱整排——这叫合并访存（Coalesced Access）。如果它们读的下标是离散的（lane 0 读 0、lane 1 读 257……），就得牵出多得多的 cache line，带宽利用率直线下降。

图 3.4　同样 32 个 lane、同样写 output[tid]，差别只在读地址的排布。连续下标一次访存就能喂饱整排 lane，离散下标要牵出多得多的 cache line。这里画的是逻辑地址的排列方式，不是实测的物理事务数。

我们用一个只改读取步长（stride）的受控实验量了这个差距：每个线程照样写连续的 output[tid]，只把读取下标换成 input[(tid * Stride) & (n - 1)]。结果，把下标从连续（stride-1）打散到 stride-257，逻辑有效带宽掉到了约 1/8。再强调一次：这是按算法字节（2 × N × sizeof(float)）反推的逻辑带宽，不是硬件计数器数出来的物理 GDDR6 流量。选做实验：读取步长 → 逻辑有效带宽

核心 kernel 只有一个变量 Stride：每个线程照样写连续的 output[tid]，但读取下标变成 (tid * Stride) & (n - 1)。Stride=1 时相邻 lane 读相邻地址（合并访存），Stride 越大地址越离散。

代码：global_memory_access.hip（核心 kernel）

template <unsigned Stride>
__global__ void gather_copy(
    const float* input, float* output, std::size_t n
) {
    std::size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        std::size_t source = (tid * Stride) & (n - 1);
        output[tid] = input[source];
    }
}

完整可运行版本位于 code/part0-intro/chapter3/global_memory_access.hip。

怎么运行（原生 Ubuntu 实验机，已 uv sync 并 source ./activate-rocm.sh）：

cd code/part0-intro/chapter3
hipcc --offload-arch=gfx1201 -O3 -std=c++17 global_memory_access.hip -o global_memory_access
# --implementation all 会依次跑 stride-1 / stride-17 / stride-257
./global_memory_access --implementation all --size 16777216 --warmup 10 --repeat 50

参数说明：

• --implementation：选择要运行的读取步长。all 会依次运行 stride-1、stride-17 和 stride-257，也可以单独指定其中一种。
• --size：输入和输出数组包含的 FP32 元素数量，结果中记为 N。
• --warmup：正式计时前的预热次数，不计入结果。
• --repeat：正式计时的重复次数，程序报告这些测量结果的中位数。

结果（RX 9070 XT + ROCm 7.13，N=16,777,216 FP32，逻辑有效带宽按 2 × N × sizeof(float) / time）：

读取 stride	中位数时间（ms）	逻辑有效带宽（GB/s）	相对 stride-1
1	0.235040 [0.231941, 0.237641]	571.042	基线
17	0.771721 [0.770461, 0.772321]	173.920	-69.54%
257	1.887641 [1.884983, 1.923662]	71.103	-87.55%

为什么连续下标能喂饱带宽、在真实算子里怎么保证合并访存（含 Grid-Stride Loop、向量化、尾部处理），第 8 章 Element-Wise 会系统地讲。

迁移范围： 「先假设用连续相邻的下标」通常是个值得试的起点；但具体的步长曲线、缓存的影响、最佳的数据布局，都必须用目标的数据类型、规模、编译器和 GPU 重新测。

3.4 LDS bank 冲突

LDS 快不快，还要看同一个 wavefront 里的线程，访问的地址是怎么排布的。

LDS 内部可以并行访问的小单元叫存储体（Bank）。用收银台来理解最直观：LDS 好比一排并行的收银台，一个 wavefront 的 32 个 lane 同时来结账。stride-1 时一人一台，瞬间结完；如果 32 人全挤向同一个台（stride-32），只能排长队，于是慢了好几倍；把每人错开一个台（stride-33），队伍又散了。

图 3.5　LDS 地址映射示意：它显示实验的索引模式，而不是声称 gfx1201 固定的 bank 数、bank 公式或每周期服务规则。

我们的受控实验测得：在这套 wave32 / 共享内存索引方式下，stride-32 比 stride-1 慢约 5 倍；stride-33 错开一个台后基本回到基线。需要强调：我们测出了这个惩罚，但没有把「bank 有几个、怎么取模」当成 gfx1201 的官方结论——那要靠文档或计数器确认。选做实验：共享数组索引步长 → bank 冲突

核心 kernel 先把一块 __shared__ 填好，再让每个 lane 按 wavefront * kRegion + lane * Stride + (iteration & 31) 反复读取。Stride 改变的是同一 wavefront 内 32 个 lane 落到哪些 bank 上。

代码：lds_bank_conflict.hip（核心 kernel）

template <unsigned int Stride>
__global__ void lds_kernel(const float* input, float* output, std::size_t n) {
    constexpr unsigned int kSharedElements = kWavesPerBlock * kRegion;
    volatile __shared__ float shared[kSharedElements];
    volatile float* shared_pointer = shared;
    for (unsigned int offset = threadIdx.x; offset < kSharedElements;
         offset += kBlockSize) {
        shared[offset] = shared_value(offset);
    }
    __syncthreads();

    const std::size_t tid =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    if (tid >= n) {
        return;
    }
    unsigned wavefront = threadIdx.x / kWaveSize;
    unsigned lane = threadIdx.x % kWaveSize;
    float accumulator = input[tid];
    for (unsigned int iteration = 0; iteration < kReadIterations;
         ++iteration) {
        unsigned index = wavefront * kRegion + lane * Stride +
                                 (iteration & 31);
        accumulator += shared_pointer[index];
    }
    output[tid] = accumulator;
}

完整可运行版本位于 code/part0-intro/chapter3/lds_bank_conflict.hip。

怎么运行：

cd code/part0-intro/chapter3
hipcc --offload-arch=gfx1201 -O3 -std=c++17 lds_bank_conflict.hip -o lds_bank_conflict
# --implementation all 会依次跑 stride-1 / stride-32 / stride-33
./lds_bank_conflict --implementation all --size 16777216 --warmup 10 --repeat 50

参数说明：

• --implementation：选择 LDS 访问步长。all 会依次运行 stride-1、stride-32 和 stride-33，也可以单独指定其中一种。
• --size：输入和输出数组包含的 FP32 元素数量。
• --warmup 和 --repeat：含义同上一实验。

结果（RX 9070 XT + ROCm 7.13，256-thread block、每线程 256 次 LDS 读取）：

LDS stride	中位数时间（ms）	对照读法
1	8.075865 [8.065751, 8.096146]	基线
32	42.218658 [41.267262, 52.018600]	5.23× 基线时间
33	8.093651 [6.570887, 8.099028]	≈基线，但进程范围更宽

LDS 的协作加载、同步与 bank 布局在真实算子里怎么用（以及 Wave Shuffle 这种更细的协作），第 9 章 Reduction 会结合归约系统地讲。

迁移范围： 「先看同一个 wavefront 的共享地址怎么排，再动手测」这个方法可以迁移；但本节的具体结果，只属于当前这个 wave32、这种数组布局、这个循环、这个编译器和 gfx1201 的受控场景。

3.5 矩阵指令 WMMA

算矩阵乘法时，让整排线程协作的专用指令，比普通算法快多少？

普通 VALU（向量算术逻辑单元）路径，是每个 lane 各自独立算出一个输出元素。RDNA 4 还提供了一类专门的波前矩阵乘加指令（WMMA，Wave Matrix Multiply-Accumulate）：由整个 wave32 一起调用，32 个 lane 分工协作，合力拼装一块 16×16 的矩阵。打个比方：普通路径像 32 个人各算各的答案；WMMA 则让这 32 人组成一条流水线，每人只负责搬运其中的一小块。分工协作让矩阵乘这类规整计算快得多，代价是积木的尺寸（16×16×16）和摆放方式（fragment layout）都是 gfx12 定死的，不能随意改。

图 3.6　两个路径计算相同的小矩阵任务，但 WMMA 的 fragment layout 与 intrinsic 是 gfx12 专属接口；图不表示生产 GEMM 的 tile、流水或 library 实现。

按 AMD GPUOpen 的说明，这条指令 __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12 由整个 wave32 调用；每个 lane 负责给一块矩阵碎片（fragment）装入/存出 8 个元素，其中 A 按转置（column-major）解释，B、C、D 按行优先（row-major）。GPUOpen：RDNA 4 WMMA intrinsic 与 fragment layout

我们的受控实验测得：对一批独立的 16×16×16（FP16→FP32）小矩阵乘，WMMA 的教学吞吐约为普通 VALU 路径的 2.3 倍。但务必注意：这是教学级微基准，没有大矩阵 tiling、双缓冲、生产 epilogue 或库级调度，不是 rocBLAS 基准，也远不是 RX 9070 XT 的理论矩阵峰值。选做实验：普通 VALU vs gfx12 WMMA

两条路径算同一批独立的 16×16×16（FP16→FP32）矩阵乘：valu_kernel 让每个线程各算一个输出元素；wmma_kernel 由整个 wave32 调用一条 WMMA 指令，32 个 lane 分工装载 fragment、协作完成矩阵乘。

代码：rdna4_wmma.hip（两条路径的核心 kernel）

// 普通 VALU 路径：一线程算一个 C[row, col]
__global__ void valu_kernel(const _Float16* matrix_a_batches,
                            const _Float16* matrix_b_batches,
                            float* matrix_c_batches) {
    const std::size_t batch = blockIdx.x;
    const _Float16* matrix_a = matrix_a_batches + batch * kMatrixElements;
    const _Float16* matrix_b = matrix_b_batches + batch * kMatrixElements;
    float* matrix_c = matrix_c_batches + batch * kMatrixElements;
    const int row = threadIdx.x / 16;
    const int column = threadIdx.x % 16;
    float accumulator = 0.0f;
    for (int k = 0; k < 16; ++k) {
        accumulator += static_cast<float>(matrix_a[row * 16 + k]) *
                       static_cast<float>(matrix_b[k * 16 + column]);
    }
    matrix_c[row * 16 + column] = accumulator;
}

// WMMA 路径：整个 wave32 协作，一条指令算一块 16×16×16
__global__ void wmma_kernel(const _Float16* matrix_a_batches,
                            const _Float16* matrix_b_batches,
                            float* matrix_c_batches) {
    const std::size_t batch = blockIdx.x;
    const _Float16* matrix_a = matrix_a_batches + batch * kMatrixElements;
    const _Float16* matrix_b = matrix_b_batches + batch * kMatrixElements;
    float* matrix_c = matrix_c_batches + batch * kMatrixElements;
    Half8 a_frag;
    Half8 b_frag;
    Float8 c_frag{};
    const int laneWrapped = threadIdx.x % 16;
    const int laneGroup = threadIdx.x / 16;
    for (int ele = 0; ele < kFragmentElements; ++ele) {
        a_frag[ele] = matrix_a[16 * laneWrapped + (ele + laneGroup * 8)];
        b_frag[ele] = matrix_b[16 * (ele + laneGroup * 8) + laneWrapped];
    }
    c_frag = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(
        a_frag, b_frag, c_frag);
    for (int ele = 0; ele < kFragmentElements; ++ele) {
        matrix_c[16 * (ele + laneGroup * 8) + laneWrapped] = c_frag[ele];
    }
}

完整可运行版本（含 CPU FP32 参考校验）位于 code/part0-intro/chapter3/rdna4_wmma.hip。

怎么运行：

cd code/part0-intro/chapter3
hipcc --offload-arch=gfx1201 -O3 -std=c++17 rdna4_wmma.hip -o rdna4_wmma
# --implementation all 会依次跑 valu 和 wmma 两条路径；--size 是矩阵批数
./rdna4_wmma --implementation all --size 4096 --warmup 10 --repeat 50

参数说明：

• --implementation：选择矩阵乘路径。all 会依次运行普通向量计算的 valu 和使用波前矩阵指令的 wmma，也可以单独指定其中一种。
• --size：独立矩阵乘法的批次数；每一批都是一次固定的 16×16×16（FP16 输入、FP32 累加与输出）矩阵乘，不是矩阵边长。
• --warmup 和 --repeat：含义同前述实验。

结果（RX 9070 XT + ROCm 7.13，batch 4,096 个独立 16×16×16 乘，正确性对 CPU FP32 参考，smoke 最大绝对误差 2.98e-08）：

路径	中位数时间（ms）	吞吐（TFLOPS）	相对
valu	0.037200 [0.036080, 0.038320]	0.902001	基线
wmma	0.016160 [0.016121, 0.016160]	2.076388	2.30×

WMMA 在真实算子里怎么用（含分块、fragment 排布、与融合的配合），第 11 章 GEMM-Like 和 第 12 章 Fusion：融合算子 会展开。

迁移范围： 这条指令、wave32 的碎片宽度和布局，只适用于 gfx12/RDNA 4；「先验证布局和对数，再比较两条语义相同的路径」这个方法可以迁移，但任何真正的大 GEMM，都得单独做自动调参（autotune）并和库实现对照。

3.6 写 Kernel 前的硬件决策清单

我们把上下两章的要点，压缩成一份真正能照着做的检查清单。第一次写 kernel，不需要每条都答对；但要做到：每一条结论，都能在源码、编译报告、实验记录或 profiler 里找到证据。

1. 启动（launch）： 输出的逻辑分块是怎么切的？grid、block、tid 加上越界保护，是不是正好覆盖全部、且只覆盖一次？
2. 工作组/波前： 目标机器实际的波前大小是多少？组内的协作和分支，有没有跨越那些不该混在一起的 lane？
3. 放在哪执行（placement）： 哪些正确性假设，依赖了「同一个 workgroup 共享 LDS/同步」？你是不是错误地假定了固定的 CU/WGP 结构？
4. 控制流： 判断条件会不会让同一个 wavefront 的 lane 走进不同的长路径？能不能靠调整数据布局来改善一致性，又不损害语义？
5. 资源占用（residency）： 编译产物用了多少 VGPR、SGPR、LDS、scratch？几种候选 tile/block 的资源变化，记录了吗？
6. 数据路径： 每个数组在哪段范围里被复用？逻辑字节和真实硬件流量，有没有分清？
7. 全局地址： 同一个 wavefront 的 lane，一轮里读的逻辑下标是不是相邻的？每种步长/布局，越界处理都对吗？
8. LDS： 把「lane → 共享地址」列出来之后，改变填充/步长，同步和数值结果还保持一致吗？用对照实验测过吗？
9. 矩阵路径： 如果用了 WMMA，目标是不是 gfx12？碎片 A/B/C/D 的布局、波前大小、CPU 参考，是不是全都对得上？
10. 证据： 硬件、软件、源码版本、规模、预热/重复次数，固定了吗？报告的是中位数和多个独立进程的范围，而不是单次最好值吗？

结论类型	例子	下一步
可迁移	先做边界正确性，再分离逻辑指标与物理归因	在新 kernel 中继续采用这套验证顺序
需重测	stride、LDS layout、VGPR/LDS 资源、占用率、分支代价	更换 GPU、dtype、shape、编译器或算法后重新 benchmark/profile
仅 gfx12	__builtin_amdgcn_wmma_*_w32_gfx12、WMMA fragment layout	仅在 gfx12 编译/运行守卫下使用；其他 target 查其对应文档

本章小结

• 一个 wavefront 等访存时，硬件会切到别的 wavefront 去跑（延迟隐藏）；能同时塞下多少 wavefront，取决于 VGPR/SGPR/LDS 哪种资源先用光，而不是线程数。占用率不是越高越好。
• 内存层级越往下越大、越慢；LDS 是你显式分配的片上存储，不是缓存；逻辑字节不等于物理流量。
• 让相邻 lane 读相邻下标（合并访存）能显著抬高带宽；LDS 要注意同一 wavefront 的地址排布避免 bank 冲突；WMMA 让整排 lane 协作算矩阵、比通用 VALU 快，但布局是 gfx12 定死的。
• 这些概念的具体优化，分别在 Part 2 的第 8/9/11/12 章结合真实算子展开。至此入门篇的硬件心智模型就完整了；下一章我们跑通第一个真正的 GPU 程序 vector add，并建立 Roofline 心智模型。

自我检验

读完本章，你应该能：

1. 能说清 VGPR、SGPR、LDS 各自存什么，以及为什么占用率取决于「哪种资源先用光」、为什么不是越高越好。
2. 能区分寄存器驻留值、显式 LDS 访问和全局访存这三类操作，并说明为什么 LDS 不是 cache。
3. 能区分规格容量、逻辑字节和物理 GDDR6 流量这三个概念。
4. 能解释为什么合并访存（连续下标）能抬高带宽，以及把下标打散为什么会掉到约 1/8。
5. 能说明 LDS bank 冲突为什么发生（同一 wavefront 挤同一个 bank），以及 stride 错开为什么能缓解。
6. 能说出 WMMA 相对普通 VALU 快在哪里、代价是什么，以及它为什么不能等同于 rocBLAS 或理论峰值。

延伸阅读

• ROCm GPU specifications：核对 gfx1201 的寄存器、LDS 与各级缓存容量。
• AMD Radeon RX 9070 XT 产品规格：显存容量、位宽、理论板卡带宽与理论计算规格。
• GPUOpen: Using the Matrix Cores of AMD RDNA 4 architecture GPUs：仅 gfx12 的 WMMA fragment layout 与 intrinsic 示例。
• 选做实验：code/part0-intro/chapter3/（global_memory_access / lds_bank_conflict / rdna4_wmma 受控对照）。
• 实战展开：第 8 章 Element-Wise、第 9 章 Reduction、第 11 章 GEMM-Like、第 12 章 Fusion：融合算子。