RX9070XT — Phase A–F を終えて何が分かったか RX9070XT — What Phase A–F Revealed

このページの結論 Page Conclusion

調査環境 — なぜこの構成を anchor にしたか Investigation Environment — Why This Configuration Was Chosen as Anchor

観測ポイント： どの環境・ワークロードを「基準点」として調査するか。

Observation target: What environment and workload to use as the investigation's anchor point.

Q4_K_M を anchor にした理由は、ggml の dispatch 分岐が量子化型と ne11 の組み合わせで決まるため、 1 つの量子化型を固定することで dispatch 閾値の変数を絞れるからだ。 deepseek-r1-distill-qwen-7b:q4_k_m は Q4_K family の代表として選んだ。

Q4_K_M was chosen as anchor because ggml dispatch branching is determined by the combination of quantization type and ne11. Fixing one quantization type reduces the variables in dispatch threshold analysis. deepseek-r1-distill-qwen-7b:q4_k_m was selected as representative of the Q4_K family.

確認できたこと — Phase A–F の成果 What Was Confirmed — Phase A–F Results

観測ポイント： 6 フェーズを通じて、何が「確定した事実」として積み上がったか。

Observation target: Across 6 phases, what accumulated as "confirmed facts"?

• GPU ロード確認（Phase B/C）： offloaded 29/29 layers to GPU を runtime ログで確認。 Flash Attention が auto 設定で有効化された（set to enabled）。 依存チェーン libggml-hip.so → libhipblas.so → librocblas.so → libamdhip64.so → hsa-runtime64 を ldd で確認。
• タイミング基準値（Phase B）： cold generate: load=17.455s, prompt_eval=0.0238s, eval=0.306s。 hot (resident): load=0.083s, prompt_eval=0.0248s, eval=0.149s。 hot の eval が cold の半分になっており、GPU 常駐の効果が数値で現れている。
• gfx1201 custom binary の存在（Phase D/E）： .hip_fatbin（587 MiB）から MMVQ / MMQ / Flash Attention / RoPE の gfx1201 hsaco を isolate・逆アセンブル済み。 binary が「存在しないために動かない」という状況ではないことを確認した。
• BLAS side の存在（Phase E）： rocblas/library/Kernels.so-000-gfx1201.hsaco に Cijk_* family が存在することを逆アセンブリで確認。
• phase proxy probe（Phase F）： prompt_eval_count = 3（MMVQ 圏）・11（MMQ-eligible 圏）・290（境界候補）を取得。 runtime を壊さずに dispatch 圏の近似を読めることを確認した。

• GPU load confirmed (Phase B/C): Runtime log confirmed offloaded 29/29 layers to GPU. Flash Attention enabled automatically (set to enabled). Dependency chain libggml-hip.so → libhipblas.so → librocblas.so → libamdhip64.so → hsa-runtime64 confirmed via ldd.
• Timing baseline (Phase B): Cold generate: load=17.455s, prompt_eval=0.0238s, eval=0.306s. Hot (resident): load=0.083s, prompt_eval=0.0248s, eval=0.149s. Hot eval is roughly half of cold, numerically confirming GPU residency benefit.
• gfx1201 custom binary presence (Phase D/E): Isolated and disassembled gfx1201 hsacos for MMVQ / MMQ / Flash Attention / RoPE from .hip_fatbin (587 MiB). Confirmed that "binary missing" is not the reason any path would fail to run.
• BLAS side presence (Phase E): Cijk_* family confirmed in rocblas/library/Kernels.so-000-gfx1201.hsaco via disassembly.
• Phase proxy probe (Phase F): Obtained prompt_eval_count = 3 (MMVQ zone), 11 (MMQ-eligible), 290 (boundary candidate). Confirmed that dispatch zone approximation is readable without breaking the runtime.

現時点で最も筋のいい読み Strongest Current Reading

観測ポイント： 確認できた事実から、「今この実行で何が起きていたか」をどう読むのが最も自然か。

Observation target: From the confirmed facts, what is the most natural reading of "what was happening during this run"?

• decode 側（small ne11）： prompt_eval_count = 3 は ne11 ≤ 8 の MMVQ 圏に入っていると読むのが最も自然。 custom Q4_K mul_mat_vec_q（bundle_0030）が主経路として機能している可能性が高い。 BLAS 支配を積極的に示す証拠は現時点でない。
• Attention： Flash Attention が runtime で有効化されており、custom flash_attn_ext_f16（bundle_0019）family で処理されているとみるのが自然。 KQ/KQV GEMM が直接見えないことは、Flash Attention ルーティングと整合する。
• prefill 側（medium ne11）： prompt_eval_count = 11 は MMQ-eligible 圏。 bundle_0030（MMVQ）の閾値 ne11 ≤ 8 を超えており、bundle_0096（exact Q4_K mul_mat_q）との整合が強い。
• RoPE： custom gfx1201 family として bundle_0039 に isolate 済み。

• Decode-side (small ne11): prompt_eval_count = 3 most naturally places in the MMVQ zone (ne11 ≤ 8). Custom Q4_K mul_mat_vec_q (bundle_0030) likely serves as the primary path. No positive evidence for BLAS dominance at this stage.
• Attention: Flash Attention enabled at runtime — natural to read this as the custom flash_attn_ext_f16 family (bundle_0019) handling attention. Absence of visible KQ/KQV GEMM names is consistent with Flash Attention routing.
• Prefill-side (medium ne11): prompt_eval_count = 11 is in the MMQ-eligible zone. Crosses the MMVQ threshold (ne11 ≤ 8), aligning more strongly with bundle_0096 (exact Q4_K mul_mat_q).
• RoPE: Isolated as a concrete custom gfx1201 family in bundle_0039.

示せること / 示せないこと What Can and Cannot Be Shown

未確定事項Open Questions

• short live case で BLAS 側 Cijk_* family が実際に入ったかどうか。
• 11 token resident 窓が bundle_0096 を踏んだかどうか（dispatch 帰属なしには確定しない）。
• prompt_eval_count = 290 の境界が MMQ heuristic 超過（BLAS fallback）か internal chunking か。
• decode 側で mul_mat_f dense family など隣接 kernel が混在する割合。

• Whether the BLAS-side Cijk_* family participated in the short live case.
• Whether the 11-token resident window exercised bundle_0096 (not confirmed without dispatch attribution).
• Whether the prompt_eval_count = 290 boundary is MMQ heuristic overflow (BLAS fallback) or internal chunking.
• The fraction of adjacent kernels like mul_mat_f dense family in the decode-side mix.

次の観測点（Phase G 以降） Where to Look Next (Phase G+)

• dispatch-safe observer が得られたら： BLAS 側 Cijk_* の live phase ownership、bundle_0096 の実行確認、prompt_eval_count=290 の切り分けを再試行する。
• 量子化横断 dispatch map： Q4_K_M anchor を維持しつつ、他の tensor type（Q5_K / Q6_K / Q8_0）への dispatch 分岐を source から先にマップする。 file type と tensor type の違いに注意（Q4_K_M ラベルでも内部は混在しうる）。
• このページを読んだ後は： 経路観測ページ（dispatch 分岐と bundle 対応の詳細）、 Observer 制約マップ（なぜ直接確認できないか）。

• When a dispatch-safe observer becomes available: Retry BLAS-side Cijk_* live phase ownership, bundle_0096 execution confirmation, and prompt_eval_count=290 disambiguation.
• Cross-quantization dispatch map: Maintain Q4_K_M anchor while mapping dispatch branches for other tensor types (Q5_K / Q6_K / Q8_0) from source first. Note: file type ≠ tensor type (a Q4_K_M model can internally mix quantization types).
• After this page: Path observation (dispatch branches and bundle correspondence in detail), Observer constraint map (why direct confirmation is blocked).