k. Kumaling / notes
← 所有笔记
GPU 系统 / INTERACTIVE EXPLORATION

深入 DeepGEMM

走进 CTA 与 SM 内部，在交互时间线中理解 DeepGEMM 的计算、数据搬运与流水线。

CUDA 编程模型

四个问题：Host / Device 如何协同，工作如何编号，Threads 如何以 Warp 执行，数据如何共享与同步。

1. Host–Device Execution

先不打开 Kernel 内部。这里只观察 Host 如何提交 GPU 工作、Stream 如何组织顺序， 以及为什么 Host API、Runtime enqueue 与真正的 Device execution 是不同时间段。

Host / Device Stream Async Event
1 先点代码 观察一行 CUDA API 在 trace 上对应哪些事件。
2 再点 Trace 反向定位它来自哪一行 Host code。
3 最后看时间关系 理解 submit、execute、complete 与 synchronize 的区别。
一行 launch 代码，会对应 Host submit 与 Device execution 两段不同时间。
Execution Trace
横轴 = 时间 · 纵轴 = execution context · 点击事件反查源码
Profiler view
Host Source
点击代码行，高亮对应的 Host / Runtime / Device trace
code ↔ trace
basic_async.cu click a line
Host API / CPU work Runtime enqueue Memory operation Kernel execution Wait / dependency
Selected
kernelA<<<...>>>
Mental model
Submit ≠ Complete Host API 返回，通常只说明工作已提交；不表示 GPU 已完成。
Stream 定义 ordering 同一 Stream 中 operation 有确定顺序；不同 Stream 才有潜在 overlap。
Synchronization 显式汇合 Host 或另一个 Stream 可以等待某个 Event / progress，再继续推进。

2. Parallel Work Organization

Grid → CTA / Thread Block → Thread 定义逻辑工作层级；blockIdx / threadIdx 决定每个 Thread 的全局坐标。

Grid CTA / Block Thread
Hierarchy
Grid
→
CTA / Block
→
Thread
Grid · 4 × 3 Blocks
选择一个 CTA
gridDim = (4, 3)
x=0
x=1
x=2
x=3
y=0
y=1
y=2
Selected CTA · blockIdx = (1, 1)
选择一个 Thread
blockDim = (8, 4)
0
1
2
3
4
5
6
7
0
1
2
3

3. SIMT Execution Model

32 个 Thread 构成一个 Warp。 重点看 active lanes、branch divergence，以及 lane 间协作。

Warp = 32 Threads Active Mask Divergence
Warp Lane Viewer
切换模式，观察 32 lanes 的状态变化
Warp 0 all 32 lanes active
active mask
0xffffffff

4. Memory & Cooperation

Memory visibility 与 cooperation scope 对齐理解即可。

Scope
Memory / Data
Cooperation / Sync
Thread
Register / Local thread-private
无跨线程同步
Warp
— 无独立 memory space
shuffle / vote / __syncwarp()
CTA / Block
Shared Memory CTA 内共享
__syncthreads() CTA barrier
Cluster
Distributed Shared Memory Cluster 内 CTA 可互访 SMEM
cluster sync
Device
Global Memory Device-wide
atomic / fence / kernel boundary

CTA ↔ output tile：先看映射，再看任务分配

这一层先把 GEMM 简化成 C = A × B，只讨论二维 output tile 的映射 F(i) → (m,n)。 选择任意 CTA / output tile 时，同时观察它需要的 A[m,*] 与 B[*,n]，从输入侧理解 CTA 之间的复用关系。 Persistent 不是另一种 tile mapping：它只决定固定 CTA / worker p 从线性任务序列里取哪些 i。

1先选映射 F(i)决定第 i 个线性任务对应哪块 C[m,n]。
2再看 CTA 怎么拿 iStatic 与 Persistent 是任务分配，不是新的 mapping。
3再看 A / B 复用同一 wave 中相同 m 复用 A 行，相同 n 复用 B 列。
C = A × B · CTA ↔ output tile 映射
先从 A/B 输入关系理解 CTA 复用，再看 C 的 output-tile mapping
Grouped-M strip
① Mapping · F(i)
② CTA assignment
③ Inspect
第 1 / 8 批
输入 tile 关系
当前 CTA 负责 C[m,n]，因此沿 K 方向读取整条 A[m,*] 与 B[*,n]。
C[0,0] = Σk A[0,k] × B[k,0]
A tiles · M × K相同 m → 复用同一 A row
×
B tiles · K × N相同 n → 复用同一 B column
Unique A bands
4
4 × A[m,*]
Unique B bands
2
2 × B[*,n]
A+B logical working set
6 MiB
4 MiB A + 2 MiB B
示例口径：BM = BN = 128 · K = 4096 · BF16 · 8 concurrent CTAs；统计当前 wave / iteration 的 unique operand bands，不等同于真实 L2 occupancy。
C tiles · M × N选择 output tile 或下方 CTA，A/B 面板会同步高亮
n0
n1
n2
n3
n4
n5
n6
n7
m0
m1
m2
m3
m4
m5
m6
m7
浅灰 = 当前 wave / iteration 的 8 个任务 粗边框 = 当前观察对象 虚线 = persistent CTA 跨 iteration 的轨迹
① MappingF(i) → (m,n)只回答：第 i 个线性任务是哪块 C output tile。
＋
② AssignmentCTA → iStatic：CTA 编号随任务递增；Persistent：固定 CTA 以 stride 反复取 i。
当前 wave 的 CTA → tile Static: CTA i → F(i)

常见 CTA ↔ output tile 映射模式 · 开源实现支持矩阵

表格按“映射 family”归类，而不是按仓库起名字。直接表示源码有明确现成实现；等价表示在注明条件下映射结果相同； 可自定义表示框架能写出该映射，但这里没有把它当成官方现成 preset。Persistent / multicast 不列成 mapping mode。

标记：直接源码有现成实现等价在注明条件下结果同构可自定义框架能表达，但不是这里核对的 preset
Mapping family 核心形状 Triton CUTLASS DeepGEMM Tensile Composable Kernel 关键参数 / 条件
Linear · N-fastC00 → C01 → … → C10 等价 GROUP_SIZE_M=1 等价 AlongN + swizzle=1* 非 Normal 默认 未单列这种 preset 等价 M01=1 最基础 row-major；*CUTLASS 此处指 cluster=1×1
Linear · M-fastC00 → C10 → … → C01 可自定义 等价 AlongM + swizzle=1* Batched 分支可出现 未单列这种 preset 等价 N01=1 最基础 column-major；*CUTLASS 此处指 cluster=1×1
Grouped-M stripM 切成 G 行；组内 M-fast 直接 GROUP_SIZE_M 等价 AlongN + swizzle G* 直接 group on M Blocking 是相关但不同 family 直接 M00_N0_M01 Triton 常见 G=8；DeepGEMM 候选 G=8/16；*cluster=1×1 时 CUTLASS 同构
Grouped-N stripN 切成 G 列；组内 N-fast 可自定义 / 转置 等价 AlongM + swizzle G* 直接 group on N Blocking 是相关但不同 family 直接 N00_M0_N01 DeepGEMM multicast A 时使用 group on N；*cluster=1×1 时 CUTLASS 同构
2D Blocking / box当前活跃 WG 压进局部矩形，再移动 box 可自定义 不是该命名的公开 preset 不是该命名的公开 preset 直接 WorkGroupMappingType=B 有 grouped WGP 变体 Tensile WorkGroupMapping / WGM
Morton / Z-orderbit interleave；2D 空间填充顺序 可自定义 未见现成 preset 未见现成 preset 直接 WorkGroupMappingType=Z 未在本次源码核对中确认 Tensile 源码还注明 Z-order 并不比 blocking 更快
Persistent 是第二个维度：Triton persistent matmul 用 tile_id += NUM_SMS；CUTLASS StaticPersistentTileScheduler 用 linear_idx += total_grid_size；DeepGEMM 用 blockIdx.x + iter × kNumSMs。它们都可以叠在各自的 F(i) → (m,n) 之上，而不是另起一种 output-tile mapping。
核对源码： Triton grouped matmul · Triton persistent matmul · CUTLASS scheduler params · CUTLASS static persistent scheduler · DeepGEMM scheduler · Tensile WorkGroupMapping · Composable Kernel BlockToCTileMap

SM hardware map：先认硬件模块，再讨论 GEMM 数据流

这一层只画相对稳定的硬件结构，不把 TMA/WGMMA 的执行顺序、mbarrier、proxy/fence 画成流程图。 点击硬件模块，右侧再展开它关联的 PTX / CUDA 指令，以及前面讨论过的 register、warp scheduler、scoreboard、TMA、WGMMA、cache、barrier 与 proxy 等知识点。

Hopper / CC 9.0SM90 / GH100 family不是物理 floorplan
1先分清整卡 vs 单个 SM78 SM、96 GB 属于 H20 整卡；64 warps、4 Tensor Cores 属于单个 SM。
2再点硬件模块主图只画硬件；右侧解释它在 GEMM 中的角色和具体上限。
3最后点相关指令把 wgmma、TMA、ldmatrix、mbarrier 等放回对应硬件理解。
整卡 ① GPU SKU profile · 这些数字描述整张卡，不是单个 SM
Enabled SMs
78
整卡启用的 SM 数
FP32 cores
9,984*
78 × 128 · 架构推导
Tensor Cores
312*
78 × 4 · 架构推导
HBM
96 GB
HBM3
HBM bandwidth
4.02 TB/s†
由实机 clock × bus 推导
L2
60 MiB
CUDA Runtime 实机
Compute capability
9.0
Hopper / SM90
H20 实机 CUDA Runtime 校验 · 8 / 8 GPU 属性一致PyTorch 2.10.0+cu128 · CUDA 12.8
Reported memory
97,356 MiB
102,085,361,664 bytes
Memory bus
6144 bit
cudaDeviceProp
Memory clock
2619 MHz
memory_clock_rate
SM clock
1980 MHz
clock_rate
SMEM / SM
228 KiB
233,472 bytes
Opt-in SMEM / block
227 KiB
232,448 bytes
规格来源与校验说明
单个 SM ② CC 9.0 / GH100 SM 的关键上限
H20 / H100 / H200 / H800 的单-SM 心智模型基本一致
32 threads / warpwarp size 固定
64 warps / SM最多 2048 resident threads
32 CTAs / SM单 CTA 最多 1024 threads
4 Tensor Cores / SM第四代 Tensor Core
65,536 registers / SM32-bit registers · 255/thread max
228 KiB SMEM / SM256 KB unified subsystem · 32 banks
上面的 per-SM 数字是 Compute Capability 9.0 / GH100 SM 的架构属性。本页 H20 与 H800 PCIe 都已有 8×GPU CUDA Runtime 实机校验；两组机器的 warp/thread、register、shared-memory 等单-SM 属性一致，差异主要体现在整卡启用 SM 数、L2、显存总量、memory bus/clock 等。Tensor Core 与 FP32 core 总数不是 CUDA Runtime 直接属性，因此仍按 GH100 的每-SM 架构数量 × 启用 SM 数推导，并用 * 标记。
③ H20 · 单个 SM 硬件模块 · 主图只回答“有哪些硬件”
Selected: L1TEX / Shared Memory subsystem
点击规则：选中的模块只加粗，其他模块保持可见。右侧：先看硬件数字与 GEMM 角色，再点指令。动态数据流：留到 Chapter 4。
GPU MEMORY HIERARCHY · 从整卡到单个 SM HBM / Global Memory A · B · C/D · scales · off-chip memory L2 Cache 全 SM 共享 · cache line / sector · CTA 之间复用 单个 SM · CC 9.0 · 逻辑硬件视图 最多 64 warps · 2048 threads · 32 CTAs / SM Warp Scheduler / Dispatch 4 个 SMSP / warp scheduler · ready warp → issue 最多 64 resident warps · scoreboard 跟踪依赖 片上存储 · ON-CHIP MEMORY 寄存器状态 · REGISTER STATE 执行单元 · EXECUTION UNITS L1TEX / Shared Memory Subsystem 256 KB unified L1/Texture/SMEM subsystem L1 Data Cache Shared Memory SMEM up to 228 KB · 32 banks Register File 65,536 × 32-bit registers / SM fragments · accumulators / partial sums addresses · predicates · state max 255 regs/thread · pressure ↔ occupancy/spill Tensor Cores 4 × 4th Gen Tensor Cores / SM WGMMA / MMA execute here CUDA Cores / ALU 128 FP32 · 64 INT32 · 64 FP64 / SM scale · conversion · epilogue 数据搬运 · DATA MOVEMENT 同步 / memory model 放到右侧解释，不在硬件图画流程箭头 Tensor Memory Accelerator (TMA) Hopper TMA · 1D–5D 异步 GMEM ↔ SMEM 搬运 Load / Store Unit 普通 load/store 与 matrix fragment 数据路径 相关机制在右侧 Inspector： mbarrier · async completion · proxy/fence ldmatrix / stmatrix · WGMMA waits
④ H20 / H100 / H200 / H800 · 整卡规格对照（可展开）

四种 Hopper SKU 的整卡差异

这里明确保留具体 form factor，不再把 PCIe 与 SXM 混为一类：H20 / H100 / H200 使用当前页面的 SXM profile，H800 则采用本会话实测的 H800 PCIe。单个 SM 的 CC 9.0 上限基本一致；整卡启用 SM 数、L2、显存接口与时钟会因 SKU / form factor 不同。

GPU profileEnabled SMsFP32 cores4th-gen Tensor CoresHBMHBM bandwidthL2说明 / 证据
H20 SXM5 96GB789,984*312*96 GB marketed
97,356 MiB runtime		4.02 TB/s†60 MiB
62,914,560 B		8×H20 实机 CUDA Runtime 12.8：78 SM、CC 9.0、6144-bit bus、2619 MHz memory clock、60 MiB L2。* core 数按每-SM 架构数量推导；† 带宽由 clock × bus × DDR 推导。
H100 SXM5 80GB13216,89652880 GB HBM33.35 TB/s50 MBNVIDIA 官方 132 SM；每 SM 128 FP32 + 4 Tensor Core。
H200 SXM 141GB13216,896*528*141 GB HBM3e4.8 TB/s50 MB*官方 + 架构推导 NVIDIA CUDA 资料给出 132 SM；带 * 数字由 GH100 每-SM 结构推导。
H800 PCIe11414,592*456*81,079 MiB runtime
85,017,493,504 B		2.04 TB/s†50 MiB
52,428,800 B		8×H800 PCIe 实机 CUDA Runtime 12.8：114 SM、CC 9.0、5120-bit bus、1593 MHz memory clock、50 MiB L2。* core 数按每-SM 架构数量推导；† 带宽由 clock × bus × DDR 推导。
规格来源： NVIDIA Hopper Architecture In-Depth · NVIDIA Hopper Tuning Guide · CUDA Compute Capability 9.0 limits · NVIDIA H100 · NVIDIA H200 · H20: 本会话 8×GPU CUDA Runtime 实测 · H800 PCIe: 本会话 8×GPU CUDA Runtime 实测 · H20 secondary cross-check
Chapter 1 只负责建立“有哪些硬件”的稳定心智模型；不把某个具体 kernel 的 producer / consumer、TMA stage、WGMMA 或 writeback 顺序硬塞进硬件图。动态计算流留给 Chapter 4。

DeepGEMM：CTA 内部执行与 Kernel Configuration

先沿紧凑的 U 型硬件路径观察一个已选定的 BF16 kernel 如何执行：L2 → SMEM → Tensor Core ↓ Registers → smem_c → Global C。下面再结合 Producer / Consumer multi-stage pipeline 理解重叠，最后回到 M×N×K 看 kernel configuration 如何选择。

① CTA Dataflow ② Multi-stage Pipeline ③ Heuristic
CTA Hardware Dataflow · Tile Motion
m0 / m1 在 Tensor Core 中并行消费同一 stage；timeline 继续抽象成一条 WGMMA consumer 链
Output = C
0 / 8 尚未开始 · 等待 3-stage pipeline fill WGMMA · m0 / m1 columns
L2 Cache · A/B
K0
A0
B0
K1
A1
B1
K2
A2
B2
K3
A3
B3
K4
A4
B4
K5
A5
B5
SMEM
S0
empty
empty
S1
empty
empty
S2
empty
empty
smem_c
empty
Tensor Core · WGMMA
Registers · C accumulator
accum = 0
Global Memory · C
C[M,N] · not written
Stage-local Timeline · 6 execution lanes · 时间从左到右
3 stages × (TMA + WGMMA)

TMA：不同 stage 的 async copy 可以同时 in-flight；复用某个 stage 前必须经过 WAIT EMPTY。 WGMMA：先经过 WAIT FULL 等当前 stage ready，再经过必要的 ORDER WAIT；图中 WGMMA K0 → K1 → K2 → … 不重叠。 交互：点击任意 bar，只重放该条 lane 对应的具体 TMA / WGMMA operation。

WAIT FULL: WGMMA 等当前 stage 的 TMA 完成ORDER WAIT: WGMMA 等前序 WGMMAWAIT EMPTY: TMA 等当前 stage 被 WGMMA release
Kernel Configuration Heuristic · 候选从哪里来，为什么选它

核心流程：硬件 / 实现先限定候选范围 → 资源约束过滤 → 对剩余配置估计并行效率与数据搬运成本 → 选择预计 cycles 最低的 configuration。

MNKBF16num_smslayout / GEMM type
1Candidate Space · 参数为什么只能从这些值里选
不是从任意整数里搜索。先由 WGMMA ISA 给出硬件原子，再由 DeepGEMM 把搜索空间缩成一个很小的离散候选集。
Hardware atom · SM90 BF16 WGMMA
64 × N × 16
M=64、K=16 固定；N=8,16,24,…,256。
影响 CTA tileBM 沿 64-row 拼；BN 必须落在支持的 N；BK 沿 16 拼。
DeepGEMM · 手工限定候选
BM: 64 / 128small-M: 16 / 32special: 256*
BN: aligned subset ≤ 256cluster: 1×1 / 2×1 / 1×2G: 8 / 16
这些范围是实现中预先限定的 search space；heuristic 只在这里面比较。
Derived · 不是独立搜索旋钮
BK = 64BF16 基础 K tile
WGMMA = 64 × BN × 16由 BN 选择 instruction shape
num_stages由 BM/BN/BK 的每-stage SMEM 与总 SMEM 决定
Group-M / N由 cluster / multicast 方向派生；G 再从 8/16 选
Resource filter SMEM capacityregister pressurealignment / swizzlecluster legalitypipeline depth
2Performance Model · 对可行候选估分
主要看两件事：SM 是否被充分利用，以及 A/B 数据是否搬得足够少。
Parallelism · 并行度
CTA count整个 GEMM 生成多少 CTA。ceil(M/BM) × ceil(N/BN)
Waves这些 CTA 需要分几轮铺到所有 SM。ceil(CTA / num_sms)
Wave efficiency所有 wave slot 被占用的比例；最后一轮越空，惩罚越大。
Data movement · 数据搬运
L2-side traffic上游 L2 / memory 侧需要为一个 CTA block 提供多少 A/B 数据；multicast 可以降低这一项。
Local / SM-side traffic源码变量名里叫 L1，但这里应理解成 cost model 对 SM 本地数据通路压力的抽象，而不是“TMA 必须经过普通 L1 cache”。
每个 CTA block 的字节估计
L2-side · A/B
bytes_L2_AB = K × (BM / cluster_N + BN / cluster_M) × sizeof(AB)
cluster_N=2 时 A 可以 multicast，所以 A 的上游读取近似减半；cluster_M=2 时对应减少 B。
Local / SM-side · A/B
bytes_local_AB = K × (BM + BN) × sizeof(AB)
每个 CTA 最终仍要在自己的 SMEM 中得到完整 A/B tile，因此 multicast 不会把这一项按 cluster 再除一次。
Tensor Core-side
bytes_TC = K × (max(64,BM) + BN) × sizeof(AB) + BM × BN × sizeof(C)
近似描述 SMEM → WGMMA operand consumption，以及 accumulator / output tile 的本地数据通路压力。
C read / write
bytes_CD = BM × BN × sizeof(C) × (with_accumulation ? 2 : 1)
普通写回算一份；如果带已有 C 的 accumulation，则还要再读一份 C。
L2 / memory side→CTA SMEM / local path→Tensor Core / accumulator
物理含义：TMA 的典型输入路径仍然是 Global/L2 → SMEM。这里源码里的 num_bytes_l1_* 是性能模型的“本地带宽层级”命名；因为它还包含 SMEM→Tensor Core 和结果 staging 流量，所以不能把它解释成真实的 L1 cache transaction 数。
cycles_L2 = (bytes_L2_AB + bytes_CD) × num_blocks / BW_L2
cycles_local = (bytes_local_AB + bytes_TC + bytes_CD) × num_blocks / BW_local
wave_efficiency = num_blocks / (num_waves × num_sms)
score = max(cycles_local, cycles_L2) / wave_efficiency · minimum wins
Trade-off：tile 太大 → CTA 少、复用好，但 wave 可能填不满；tile 太小 → CTA 多、并行度好，但数据搬运和调度开销会上升。
最终 configuration →
Search
BM / BN
Search
cluster M/N
Derived
BK = 64
Derived
num_stages
Derived
WGMMA shape
Scheduler
Group-M/N · G

继续探索

前沿模型训练图谱↗看见 Ring AllReduce↗