flydsl-kernel-trace-analysis

カーネルトレース解析

GPU カーネルの ATT トレースをプロファイルして解析し、ストールのホットスポットを特定して最適化計画を作成します。

すべての解析は、同梱の hotspot_analyzer.py（ATT）と pmc_l2_analyzer.py（PMC CSV）を実行して行います。メトリクスを自分で再計算しないでください。 GUI ツールは使用しないでください。

質問が「変更によってレジスタプレッシャー、スピル、または LDS が増えたかどうか」だけの場合は、まず /isa-resource-diff を使用してください。コンパイルのみで、GPU もプロファイラーの実行も不要なため、数秒で回答できます。カーネルが遅い 理由を知りたい場合は、ここを使用してください。カーネルが使用するリソースを 知りたいだけの場合とは異なります。

コンパイラーが出力する内容を変更する 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 件のストールホットスポット、ストールタイプの内訳、占有率 （combined-VGPR-pool モデル。out_kernel_trace.csv から accum/LDS/SGPR を読み込みます）を報告します。
• ${CLAUDE_SKILL_DIR}/scripts/pmc_l2_analyzer.py — rocprofv3 PMC カウンター CSV を読み込み、L2 ヒット率、HBM 32B 部分割合、オーバーフェッチ率を報告します。カーネルがメモリバウンドで、その 理由を知る必要がある場合に使用します（ATT にはキャッシュカウンターがありません）。 Step 5 の「L2 / HBM 効率分析」を参照してください。

---

ワークフロー

モード A：既存のディスパッチディレクトリを解析

ユーザーが --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

Step 5：結果の解釈に進みます。

---

モード B：収集ワークフロー全体

Step 1：カーネルの検出

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

統計 CSV を解析し、次のカーネルテーブルを提示します：

順位	カーネル名	呼び出し回数	合計（us）	平均（us）	GPU 時間に占める割合

明らかでない場合は、トレースするカーネルをユーザーに尋ねます。

利用可能な場合は 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;"

Step 2：input.yaml の設定

以下のジョブを /tmp/trace_input.yaml に書き込み、 kernel_include_regex に Step 1 で見つかったカーネルを設定します：

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 をトレースします
• att_buffer_size：SE あたり 96MB。切り詰められた場合は "0xC000000" に増やします
• att_target_cu: 1：CU を 1 つにすると出力を扱いやすくできます

Step 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

---

Step 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

---

Step 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：すべてのウェーブにわたってこの命令に加算されたサイクル数
• col[3] source_loc："/path/to/file.py:LINE" — snapshots.json で解決される仮想パス
• col[6] exec_count：この命令を実行したウェーブスレッドの数

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	ウェーブ間の同期 — 最も遅いウェーブが全体を支配します
MFMA/FMA	v_mfma_*	MFMA の依存チェーン（RAW ハザード）
LDS	ds_read_*、ds_write_*	LDS アクセスのレイテンシー

よくあるホットスポットパターン

以下のスニペットは模式例です。命令スケジューリングを示すため、ロードは省略形の raw 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 prob を読み込む → 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 アーキテクチャを自動検出し、 すべてのリソース制限要因のうち最小値として occupancy（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

VGPR は、gfx942 と gfx950 の両方で 512 エントリーの共有プールです。 CDNA2（gfx90a） では、arch VGPR ファイル（256）と accum VGPR ファイル（256）が SIMD ごとに 1 つの 512 の予算へ統合されており、gfx942/gfx950 もこれを引き継いでいます。両方の VGPR による occupancy は 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（同じ）
occupancy の計算式（VGPR）	512 / (arch_alloc + accum_alloc)	512 / (arch_alloc + accum_alloc)
割り当て粒度	8 VGPR	8 VGPR
LDS サイズ	64 KB	160 KB
LDS 割り当てブロック	256 bytes	1280 bytes
VMCNT の幅	6 bits（最大 63 件を実行中に保持）	6 bits（最大 63 件を実行中に保持）
LGKMCNT の幅	4 bits（最大 15 件を実行中に保持）	4 bits（最大 15 件を実行中に保持）

CDNA3 と比べた CDNA4 の実際の変更点は、LDS のサイズ（64KB→160KB）と LDS の 割り当て粒度です。VGPR プールのモデルは変更されていません。

実際のカウント値を読み取る。 code.json には（多くの場合、単一 CU の、 多くの場合 vgpr 形式の）逆アセンブル結果しか含まれていないため、accum_vgpr / LDS / SGPR / workgroup size は分かりません。そのため、AGPR 形式を認識できない ISA スキャンでは accum=0 と報告され、occupancy が大きく誤って計算されます。アナライザーは、 正確な値である Accum_VGPR_Count / LDS_Block_Size / SGPR_Count / Workgroup_Size_* を取得するため、out_kernel_trace.csv（dispatch ディレクトリの 隣に配置）を読み取ります。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 decode、gfx942）：arch 144 + accum 136 = 合計 280 → 512//280 = 1 wave/SIMD となり、VGPR が制限要因です（LDS では 5、SGPR では 7 を確保できます）。2 waves を 実現するには、合計を 256 以下にする必要があります。たとえば、VGPR を約 24 個削減します。

注意：maxnreg は削除されました（LLVM に渡されたことはありません）。これを使って accum_vgpr=0 にすると occupancy は 2 倍になりましたが、arch_vgpr を通じて MFMA のスピルが 発生し、GPU が 4.5 倍遅くなることを実測しました。代わりに waves_per_eu を使用してください。

L2 / HBM 効率分析（ATT ではなく PMC）

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

3 つの指標と、3 つの判断：

指標	計算式	わかること
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 バイトは 1 回だけ読み込まれます。GQA（×heads）と MTP（×seq）による再利用はレジスター／LDS で完結し、L2 を再読み込みすることはありません。「L2 ヒット率を改善する」ことは目標ではありません。 ヒット率を上げるのは、実際の KV 再利用である共有プレフィックスのサービングだけです（これはカーネルの変更ではなく、ワークロード／スケジューリングの特性です）。
2. 32B の割合 ≈ 0% → 64B キャッシュラインをすべて使っており、空間的局所性による無駄はありません。ライン単位で修正すべきことはありません。（32B の割合が高い場合は、再構成する価値のある散在／非アラインアクセスが示唆されます。）
3. 過剰フェッチ ≈ 1.0x → カーネルは必要なデータだけを読み込みます。この場合、達成帯域幅（ideal_bytes / kernel_time として計算）がこのアクセスパターンの実際の上限です。ストリーミングが適切でも、理論上の HBM ピークの 50～60％は正常です。部分アクセスが 0％、過剰フェッチが約 1.0x で、ページ分割ギャザー型デコードが約 54％で動作しているなら、これは健全な状態であり、不具合ではありません。

計算例（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 Hang が発生したことがあります。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 or 32	CDNA4 のみ。A または B が FP8 の場合は 32
v_mfma_f32_32x32x64_f8f6f4	32x32x64	32 or 64	CDNA4 のみ。A または B が FP8 の場合は 64
v_mfma_scale_f32_16x16x128_f8f6f4	16x16x128	16 or 32	CDNA4 のみ。ブロック指数スケーリングあり
v_mfma_scale_f32_32x32x64_f8f6f4	32x32x64	32 or 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	HW はこれをチェックしません — ユーザーが待機を追加する必要があります
V_CMPX* が EXEC に書き込み -> V_MFMA*	4	MFMA には EXEC のフォワーディングがありません

「5、8、12、または 20」の待機数は MFMA のバリアントによって異なります。

• 5 回の待機：16x16 の 4 ブロックバリアント（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	デコーダーの .so をインストールします（ステップ 3 を参照）
トレース出力が空	kernel_include_regex が完全に一致しているか確認します
トレースが途中で切れている	att_buffer_size を "0xC000000" に増やします
kernel_iteration_range の不一致	範囲を調整します。"[0, [1-2]]" を試してください
INVALID_SHADER_DATA	aqlprofile とデコーダーのバージョンが不一致です — 両方を更新します
すべての Source loc が ""	FLYDSL_DEBUG_ENABLE_DEBUG_INFO=1 を設定します。コンパイルパイプラインの -g フラグを確認します
最上位のホットスポットがカーネルデコレーターの行	デバッグ情報のアーティファクトです — その行はスキップし、演算の行に注目します
--att フラグのエラー	--att はブール値で、値は不要です。完全な設定には -i input.yaml を使用します
PMC 実行中の GPU ハング／HW 例外	カウンターが多すぎます → 複数パスにします。TCC カウンターが約 4 個以下の単一パスジョブに分割します
カーネルが MFMA を使用しているのに PMC の accum_vgpr=0	VGPR 形式の MFMA では、アキュムレーターはアーキテクチャの VGPR ファイルにあります。VGPR_Count + Accum_VGPR_Count から合計を読み取ります