flydsl-kernel-trace-analysis

Kernel 轨迹分析

分析 GPU kernel 的 ATT 轨迹，找出停顿热点并制定优化计划。

所有分析都通过运行随附的 hotspot_analyzer.py（ATT）和 pmc_l2_analyzer.py（PMC CSV）完成；切勿自行重新计算这些指标。 不要使用 GUI 工具。

如果问题只是“我的改动是否增加了寄存器压力、溢出或 LDS？”，请先使用 /isa-resource-diff——它仅需编译，无需 GPU 或分析器运行，几秒内即可给出答案。当你需要了解 kernel 为什么运行缓慢，而不是使用了哪些资源时，再使用此工具。

如果要选择能改变编译器输出的 LLVM 选项——编译提示、函数属性或 -mllvm 标志——并验证其是否生效，请使用 /llvm。它同样仅需编译。

参数

参数	说明
<CMD>	要分析的命令。示例：python bench_pa.py --batch 32
--dir <path>	跳过采集；分析现有的 ui_output_agent_*_dispatch_* 目录
--topk N	显示前 N 个热点（默认：15）

---

分析器脚本

• ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py — 读取 ui_output_agent_*_dispatch_* ATT 目录；报告前 K 个停顿热点、停顿类型细分和占用率 （组合 VGPR 池模型，从 out_kernel_trace.csv 读取 accum/LDS/SGPR）。
• ${CLAUDE_SKILL_DIR}/scripts/pmc_l2_analyzer.py — 读取 rocprofv3 PMC 计数器 CSV；报告 L2 命中率、HBM 32B 部分数据占比和过量读取比率。用于 kernel 受内存带宽限制且需要了解原因时 （ATT 没有缓存计数器）。 参见第 5 步中的“L2 / HBM 效率分析”。

---

工作流

模式 A：分析现有 dispatch 目录

如果用户提供了 --dir <path> 或已有 ui_output_agent_*_dispatch_* 目录：

# Both analyzers ship with this skill -- do not reimplement them.
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --kernel <kernel> --topk 15 --mode both --detail --context 4

跳至第 5 步：解读结果。

---

模式 B：完整采集工作流

第 1 步：发现 kernel

touch /tmp/trace_ts
rocprofv3 --stats --kernel-trace -f csv -- <CMD> 2>&1
find . -maxdepth 3 -name "*stats*" -newer /tmp/trace_ts -type f 2>/dev/null

解析 stats CSV 并展示 kernel 表格：

排名	Kernel 名称	调用次数	总计（us）	平均（us）	GPU 时间占比

如果目标不明确，请询问用户要跟踪哪个 kernel。

如果有 results.db，优先使用它——使用 sqlite3 执行结构化查询：

sqlite3 results.db "
SELECT ks.KernelName, COUNT(*) calls,
       ROUND(AVG(kd.end-kd.start)/1000.0,1) avg_us
FROM rocpd_kernel_dispatch kd
JOIN rocpd_info_kernel_symbol ks ON kd.kernel_symbol_id=ks.id
GROUP BY ks.KernelName ORDER BY avg_us DESC LIMIT 20;"

第 2 步：配置 input.yaml

使用下面的任务写入 /tmp/trace_input.yaml，将 kernel_include_regex 设置为第 1 步中找到的 kernel：

jobs:
   -
       kernel_include_regex: <KERNEL_NAME_PATTERN>
       kernel_iteration_range: "[1, [3-4]]"
       output_file: out
       output_directory: kernel_trace_output
       output_format: [csv]
       truncate_kernels: true
       sys_trace: true
       advanced_thread_trace: true
       att_target_cu: 1
       att_shader_engine_mask: "0xf"
       att_simd_select: "0xf"
       att_buffer_size: "0x6000000"

要点：

• kernel_iteration_range："[1, [3-4]]" 跳过预热，跟踪第 3–4 次 dispatch
• att_buffer_size：每个 SE 为 96MB；如果数据被截断，将其增大为 "0xC000000"
• att_target_cu: 1：使用单个 CU，以控制输出规模

第 3 步：采集 ATT 轨迹

FLYDSL_DEBUG_ENABLE_DEBUG_INFO=1 rocprofv3 -i /tmp/trace_input.yaml -- <CMD> 2>&1
find . -type d -name "ui_output_agent_*" -newer /tmp/trace_ts 2>/dev/null

如果缺少 rocprof-trace-decoder 库，请安装与 /opt/rocm/.info/version 中 ROCm 版本匹配的发行版。在运行下面的代码块之前，将 RTD_VERSION 设置为该发行版。build-rocm-image 会为它构建的镜像写死自己的 RTD_VERSION，也不会据此推导；如果两者不一致，以正在分析的机器上的 ROCm 版本为准：

RTD_VERSION=0.1.6   # <-- set to the release matching your ROCm version
RTD_INSTALLER="rocprof-trace-decoder-manylinux-2.28-${RTD_VERSION}-Linux.sh"
wget -q "https://github.com/ROCm/rocprof-trace-decoder/releases/download/${RTD_VERSION}/${RTD_INSTALLER}"
chmod +x "${RTD_INSTALLER}"
"./${RTD_INSTALLER}" --skip-license --prefix=/tmp/rtd-install
find /tmp/rtd-install -name '*.so*' -exec cp -a {} /opt/rocm/lib/ \;
ldconfig

输出结构：

ui_output_agent_<PID>_dispatch_<N>/
├── code.json          ← PRIMARY: per-instruction stall/cycle data
├── snapshots.json     ← source file path mapping (virtual → local filename)
├── source_0_*.py      ← embedded source files
├── filenames.json     ← wave file index
├── occupancy.json     ← occupancy timeline
└── se*_sm*_sl*_wv*.json  ← per-wave raw traces

---

第 4 步：运行 hotspot_analyzer.py

此技能随附该分析器——不要重新实现。运行：

# Full report
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --topk 15 --mode both

# Source-level with code context (best for optimization)
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --topk 5 --mode src --detail --context 4

# ASM-only for instruction-level detail
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --mode asm --topk 20

---

第 5 步：解读结果

code.json 字段参考

code["code"] 中的每一行如下：

[asm, _, pc_index, source_loc, _, pc_addr, exec_count, total_cycles, stall_cycles, issue_cycles]
  0   1     2          3       4     5          6            7              8             9

• col[8] stall_cycles：指令因阻塞而无法发射的周期数——主要热点指标
• col[7] total_cycles：所有 wave 中计入此指令的总周期数
• col[3] source_loc："/path/to/file.py:LINE"——通过 snapshots.json 解析的虚拟路径
• col[6] exec_count：执行此指令的 wave 线程数

snapshots.json：解析源文件路径

snapshots.json 编码了一个嵌套字典树，将虚拟路径映射到本地文件名：

{"/": {"FlyDSL": {"kernels": {"pa_decode_sw_fp8_ps.py": "source_0_pa_decode_sw_fp8_ps.py"}}}}

递归展平：/FlyDSL/kernels/pa_decode_sw_fp8_ps.py → source_0_pa_decode_sw_fp8_ps.py

停顿类型分类

类型	指令	根本原因
VMEM-load	buffer_load_*、global_load_*	加载本身停顿（VMEM 队列已满，或缺少可用于隐藏延迟的计算，导致反压）
VMEM-wait	s_waitcnt vmcnt(N)	等待尚未完成的 VMEM 加载
LDS/SMEM-wait	s_waitcnt lgkmcnt(N)	等待 LDS 或 SMEM 操作
barrier	s_barrier	跨 wave 同步——最慢的 wave 决定整体速度
MFMA/FMA	v_mfma_*	MFMA 依赖链（RAW 冒险）
LDS	ds_read_*、ds_write_*	LDS 访问延迟

常见热点模式

以下代码片段仅为示意——展示的是指令调度，因此加载使用了简写的原始 intrinsic 形式。在当前接口中，这些操作通过 make_buffer_tensor 视图上的 fx.copy 实现；无论采用哪种形式，调度建议都不变。

模式 1：MFMA 循环内的 V/K 加载 → 停顿率极高（80–95%）

# BAD: load and MFMA alternate — only 1 MFMA of hiding time
for k_step in range_constexpr(QKHELOOP * 2):
    if k_step % 2 == 0:
        v_data = buffer_ops.buffer_load(...)   # stall_rate ~92%
    acc = rocdl.mfma_f32_16x16x32_fp8_fp8(...)

# GOOD: batch all loads before the MFMA loop
for td in range_constexpr(TLOOP):
    v_prefetch[td] = [buffer_ops.buffer_load(...) for _ in range_constexpr(QKHELOOP)]

for td in range_constexpr(TLOOP):
    for k_step in range_constexpr(QKHELOOP * 2):
        acc = rocdl.mfma_f32_16x16x32_fp8_fp8(...)   # entire QK MFMA hides VMEM latency
    v_results[td] = v_prefetch[td]   # already in registers

模式 2：没有计算的顺序加载 → VMEM 队列饱和

# BAD: all loads back-to-back, no compute interleaved
for td in range_constexpr(TLOOP):
    for qkhe in range_constexpr(QKHELOOP):
        k4 = buffer_ops.buffer_load(k_rsrc, ka_dw, ...)   # queue fills up

# GOOD: prefetch next tile's K loads during current tile's MFMA computation

模式 3：PV MFMA 前立即读取 LDS 概率值 → lgkmcnt 停顿

# BAD: LDS reads and MFMA in same loop
for vhe in ...:
    for vt in ...:
        p_i64 = lds_read(...)    # issued here
        tmp = mfma(v_i64, p_i64, ...)   # immediately consumed → lgkmcnt stall

# GOOD: batch all LDS reads first, then all MFMAs
for vhe in ...:
    for vt in ...:
        p_i64s.append(lds_read(...))    # all LDS reads issued first

for vhe in ...:
    for vt in ...:
        tmp = mfma(v_i64s[...], p_i64s[...], ...)   # LDS data already ready

模式 4：缩放值加载位置离使用点过近

# BAD: scale load and usage separated by only TLOOP MFMAs
for td in range_constexpr(TLOOP):
    k_scale = buffer_ops.buffer_load(ks_rsrc, ...)   # issued here
# ... small compute gap ...
    result = acc * k_scale   # used too soon → stall

# GOOD: issue scale loads at the very beginning of the block,
# before K loads, to maximise latency hiding distance

模式 5：热点归因到内核入口行

当 @flyc.kernel / 内核装饰器行显示为首要热点，且混合出现 VMEM-wait 与 barrier 停顿类型时，这是调试信息聚合伪象。MLIR/编译器生成的指令（地址运算、cndmask、序言设置）会映射到最外层作用域所在行。忽略这一行；关注包含明确用户操作的行。

寄存器压力检查（架构感知）

hotspot_analyzer.py 会根据 ISA 指令模式自动检测 GPU 架构，并将占用率（waves/SIMD）计算为所有资源限制因素中的最小值：

occupancy = min(vgpr_limit, lds_limit, sgpr_limit, hw_max=8)
  vgpr_limit = 512 // (arch_vgpr_alloc + accum_vgpr_alloc)              # per SIMD
  lds_limit  = (LDS_total // lds_per_wg) * waves_per_wg // 4_SIMDs      # per SIMD
  sgpr_limit = 800 // sgpr_alloc                                        # per SIMD

在 gfx942 和 gfx950 上，VGPR 共用一个包含 512 个条目的池。 CDNA2（gfx90a）将 arch（256）和 accum（256）VGPR 文件合并为每个 SIMD 一个 512 的预算，gfx942/gfx950 也沿用了这一设计。在两者上，VGPR 占用率均为 512 / (arch + accum)，而不是 256 / max(arch, accum)。（独立池的 256/max 模型仅适用于 gfx908 / CDNA1；当时 accum VGPR 是独立的文件，仅 MFMA 可访问。）

属性	CDNA3（gfx942）	CDNA4（gfx950）
VGPR 池	共用 512 个（256 个 arch + 256 个 accum，统一预算）	共用 512 个（相同）
占用率公式（VGPR）	512 / (arch_alloc + accum_alloc)	512 / (arch_alloc + accum_alloc)
分配粒度	8 个 VGPR	8 个 VGPR
LDS 大小	64 KB	160 KB
LDS 分配块	256 字节	1280 字节
VMCNT 位宽	6 位（最多 63 个在途项）	6 位（最多 63 个在途项）
LGKMCNT 位宽	4 位（最多 15 个在途项）	4 位（最多 15 个在途项）

CDNA4 相较于 CDNA3 的实际变化是 LDS 大小（64KB→160KB）和 LDS 分配粒度，而不是 VGPR 池化模型。

读取真实计数。 code.json 只包含反汇编内容（通常是单个 CU，且通常为 vgpr 形式），因此无法反映 accum_vgpr / LDS / SGPR / 工作组大小；无法识别 AGPR 形式的 ISA 扫描会报告 accum=0，导致占用率计算严重错误。分析器会读取 out_kernel_trace.csv（暂存于 dispatch 目录旁），获取权威的 Accum_VGPR_Count / LDS_Block_Size / SGPR_Count / Workgroup_Size_*。arch_vgpr 取 max(ISA_scan, CSV)，因此即使 CSV 的 VGPR_Count 字段异常偏低，也不会导致报告值偏低。如果找不到 CSV，则会回退到仅使用 ISA，并打印警告。

自动检测：gfx950 特有指令（v_mfma_scale_f32_*、v_mfma_f32_16x16x128_*、v_mfma_f32_32x32x64_*）表示 CDNA4；若不存在，则表示 CDNA3。

sqlite3 results.db "
SELECT ks.KernelName, ki.arch_vgpr_count, ki.accum_vgpr_count, ki.lds_size
FROM rocpd_kernel_dispatch kd
JOIN rocpd_info_kernel_symbol ks ON kd.kernel_symbol_id=ks.id
JOIN rocpd_info_kernel ki ON kd.kernel_id=ki.id LIMIT 5;"

示例（PA 解码，gfx942）：arch 144 + accum 136 = 合计 280 → 512//280 = 1 wave/SIMD，受 VGPR 限制（LDS 允许 5，SGPR 允许 7）。要达到 2 个 wave，合计值需要 ≤ 256，例如释放约 24 个 VGPR。

注意：maxnreg 已被移除（它从未传递到 LLVM）。使用它强制设置 accum_vgpr=0 虽然使占用率翻倍，却会通过 arch_vgpr 导致 MFMA 溢出；实测 GPU 速度降低了 4.5 倍。请改用 waves_per_eu。

L2 / HBM 效率分析（PMC，而非 ATT）

当 ATT 热点主要是高停顿率的 VMEM-load（例如停顿占比为 40-50%，每次加载约 94%）时，内核受内存限制，接下来要问的是为什么——而 ATT 无法回答（它没有缓存计数器）。 捕获 PMC 计数器（参见 capture-kernel-trace 的“PMC Mode”），并使用 ${CLAUDE_SKILL_DIR}/scripts/pmc_l2_analyzer.py 进行分析：

python ${CLAUDE_SKILL_DIR}/scripts/pmc_l2_analyzer.py \
    /tmp/pmc_out/pass_1/pmc_l2_counter_collection.csv \
    /tmp/pmc_ea_out/pass_1/pmc_ea_counter_collection.csv \
    --kernel <kernel> --ideal-gb <bytes_per_dispatch_GB> --ea-channels 2

三个指标，三个判断：

指标	公式	含义
L2 命中率	TCC_HIT/(TCC_HIT+TCC_MISS)	是否存在可利用的时间复用？
32B 占比	TCC_EA0_RDREQ_32B/TCC_EA0_RDREQ	空间局部性／缓存行浪费
过量读取	est_HBM_bytes / (ideal_GB × dispatches)	冗余读取

内存受限解码内核的决策树：

1. L2 命中率 < 5% → 纯流式访问，没有复用。对于每个序列独立分页的 KV 解码，这符合预期且正确——每个 KV 字节只读取一次；GQA（×heads）和 MTP（×seq）的复用发生在寄存器／LDS 中，不会再次读取 L2。“提高 L2 命中率”并非目标。 唯一能提高它的是真正的 KV 复用，即共享前缀服务（这是工作负载／调度属性，而非内核改动）。
2. 32B 占比 ≈ 0% → 使用完整的 64B 缓存行，没有空间局部性浪费。无需在线级别进行优化。（较高的 32B 占比则表明访问分散／未对齐，值得重构。）
3. 过量读取 ≈ 1.0x → 内核读取的数据恰好满足需求。此时达到的带宽（按 ideal_bytes / kernel_time 计算）就是这种访问模式的实际上限。即使是干净的流式访问，达到理论 HBM 峰值的 50-60% 也很正常；分页聚合解码达到约 54%，且部分读取为 0%、过量读取约为 1.0x，属于正常表现，并非缺陷。

示例（PA 解码，gfx942，bs=16，ctx=131072，batch=256）： L2 命中率为 1.7%，32B 占比为 0%，过量读取为 1.04x，带宽为 2.85 TB/s（峰值的 54%）。结论：内存子系统运行正常；KV 加载已无可优化之处——这一点也通过测试 block_size 从 16→64（性能下降 7.8%）得到验证，并确认 CDNA3 上不存在 dwordx8（dwordx4 / 16B 是单次向量加载的最大值）。

计数器采集注意事项：每个 PMC 任务的 TCC 计数器应不超过约 4 个（单次硬件遍历）。多遍采集曾导致 gfx942 上的 GPU 挂起——参见 capture-kernel-trace。

MFMA 延迟参考（周期数 = 流水线深度）

指令	变体	周期数	备注
v_mfma_f32_*_f16 / _bf16	16x16x16	16
v_mfma_f32_*_f16 / _bf16	32x32x8	32
v_mfma_f32_*_fp8_fp8	16x16x32	16	适用于 CDNA3+CDNA4
v_mfma_f32_*_fp8_fp8	32x32x16	32	适用于 CDNA3+CDNA4
v_mfma_f32_16x16x128_f8f6f4	16x16x128	16 或 32	仅适用于 CDNA4；A 或 B 为 FP8 时为 32
v_mfma_f32_32x32x64_f8f6f4	32x32x64	32 或 64	仅适用于 CDNA4；A 或 B 为 FP8 时为 64
v_mfma_scale_f32_16x16x128_f8f6f4	16x16x128	16 或 32	仅适用于 CDNA4；使用块指数缩放
v_mfma_scale_f32_32x32x64_f8f6f4	32x32x64	32 或 64	仅适用于 CDNA4；使用块指数缩放
v_mfma_f32_*_f32	16x16x4	32
v_mfma_f32_*_f32	32x32x2	64
v_mfma_f64_16x16x4_f64	16x16x4	64

MFMA 依赖 NOP（CDNA4，摘自 ISA 参考手册表 38）

MFMA 结果生成后到被使用前所需的最少独立指令数（或 s_nop 数量）。具体数值因 MFMA 变体而异：

依赖模式	所需等待数	说明
XDL 写入 -> 同一 XDL 读取 SrcC（累加，完全相同的 vDst）	0-2	转发路径；可背靠背累加
XDL 写入 -> VALU/VM/LDS/FLAT 读取结果（RAW）	5、8、12 或 20	不支持转发；必须等待 MFMA 提交到 VGPR
XDL 写入 -> MFMA 读取 SrcA 或 SrcB	5、8、12 或 20	不存在转发路径
非 DLops VALU 写入 -> MFMA 读取	2	不支持 4/8 周期转发路径
VALU 写入 SGPR -> VMEM 读取该 SGPR	5	硬件不会检查此项——用户必须添加等待
V_CMPX* 写入 EXEC -> V_MFMA*	4	MFMA 不支持 EXEC 转发

“5、8、12 或 20”的等待数取决于 MFMA 变体：

• 5 次等待：16x16 的 4-block 变体（8 周期 MFMA）
• 8 次等待：16x16x16 F16/BF16 等（16 周期 MFMA）
• 12 次等待：32x32x8、16x16x4 F32 等（32 周期 MFMA）
• 20 次等待：32x32x4 F32、32x32x2 F32（64 周期 MFMA）

---

步骤 6：优化计划

运行 hotspot_analyzer.py --detail 后，生成一份按优先级排序的计划：

## Stall Summary
- Total stalls: X cycles (Y% of kernel)
- Top stall type: VMEM-load (Z%)

## Hotspot Analysis

### #1 :LINE  stall=XK (N%)  VMEM-load  stall_rate=92%
Root cause: buffer_load inside QK MFMA loop — only 1 MFMA of hiding time.
Fix: Move all V loads before the QK MFMA loop.
Estimated gain: ~20% kernel cycle reduction (hypothesis; verify correctness and re-time).

### #2 :LINE  stall=XK (N%)  VMEM-load  stall_rate=80%
Root cause: K loads sequential with no compute interleaved.
Fix: Prefetch next tile's K during current tile's MFMA (double-buffer pattern).
See /prefetch-data-load skill.

### #3 ...

## Priority Order
1. [HIGH]  Fix V-load position (24% of all stalls, easy refactor)
2. [HIGH]  K-load cross-tile prefetch (8% of stalls, needs _process_block restructure)
3. [MED]   Move scale loads earlier (8% of stalls, trivial move)
4. [LOW]   Batch LDS reads before PV MFMA (4% of stalls, loop split)

---

错误处理

错误	修复方法
rocprof-trace-decoder library path not found	安装 decoder .so（参见步骤 3）
跟踪输出为空	检查 kernel_include_regex 是否完全匹配
跟踪被截断	将 att_buffer_size 增大到 "0xC000000"
kernel_iteration_range 不匹配	调整范围；尝试 "[0, [1-2]]"
INVALID_SHADER_DATA	aqlprofile/decoder 版本不匹配——同时更新两者
源位置全为 ""	设置 FLYDSL_DEBUG_ENABLE_DEBUG_INFO=1；检查编译流水线中的 -g 标志
顶部热点是内核装饰器行	这是调试信息问题——跳过该行，聚焦于 op 行
--att 标志错误	--att 是布尔值，不需要参数；使用 -i input.yaml 加载完整配置
PMC 期间 GPU 挂起 / 硬件异常	计数器过多 → 多次采集。拆分为单次采集作业，每个作业使用 ≤ 约 4 个 TCC 计数器
PMC 中 accum_vgpr=0，但内核使用了 MFMA	vgpr-form MFMA：累加器位于架构 VGPR 文件中；从 VGPR_Count + Accum_VGPR_Count 读取总数