RX9070XT Path Worklog

調査タイムラインInvestigation Timeline

• 2026-03-27
  Phase A — ソース footing 確立Phase A — Source Footing

  一次コードソースを ollama-upstream/（commit 4fda6980）に固定。 runtime truth を /usr/local/lib/ollama/rocm/、参照ツリーを ROCm-repos/（rocBLAS / hipBLAS / hipBLASLt / rocm-systems）に設定しました。 rocminfo で gfx1201 / device id 0x7550 / 16.7 GiB pool を確認。

  Fixed primary code source to ollama-upstream/ (commit 4fda6980). Set runtime truth as /usr/local/lib/ollama/rocm/ and reference trees as ROCm-repos/ (rocBLAS / hipBLAS / hipBLASLt / rocm-systems). Confirmed gfx1201 / device id 0x7550 / 16.7 GiB pool via rocminfo.
• 2026-03-27
  Phase B — ソース経路追跡Phase B — Source Path Trace

  Ollama → ggml HIP backend の呼び出しチェーンを追跡し、GGML_OP_MUL_MAT dispatch 分岐を確定しました。 Q4_K_M ワークロードに対する候補経路：
  (1) mul_mat_vec_q（ne11 ≤ 8 / MMVQ）
  (2) mul_mat_q（ne11 ≤ 256 for Q4_K / MMQ）
  (3) hipBLAS/rocBLAS fallback
  Flash Attention が有効の場合は ggml_flash_attn_ext に分岐することも確認。

  Traced Ollama → ggml HIP backend call chain and fixed the GGML_OP_MUL_MAT dispatch branches. Candidate paths for Q4_K_M workload:
  (1) mul_mat_vec_q (ne11 ≤ 8 / MMVQ)
  (2) mul_mat_q (ne11 ≤ 256 for Q4_K / MMQ)
  (3) hipBLAS/rocBLAS fallback
  Confirmed that Flash Attention enabled routes to ggml_flash_attn_ext.
• 2026-03-27
  Phase B 実機確認 — plain ROCm path 到達Phase B Live — Plain ROCm Path Reached

  plain ollama serve（port 18568）で deepseek-r1-distill-qwen-7b:q4_k_m を実行し、ROCm path を確認： inference compute library=ROCm compute=gfx1201・offloaded 29/29 layers to GPU・Flash Attention was auto, set to enabled・graph splits = 2。 cold generate: load 17.455s / prompt_eval 0.0238s / eval 0.306s。 hot（resident）generate: load 0.083s / prompt_eval 0.0248s / eval 0.149s。 rocprofv3 wrap は bootstrap を壊し CPU fallback に落ちることを確認。rocprofv3 --attach は ptrace_scope=1 でブロック。

  Ran deepseek-r1-distill-qwen-7b:q4_k_m via plain ollama serve (port 18568), confirming the ROCm path: inference compute library=ROCm compute=gfx1201, offloaded 29/29 layers to GPU, Flash Attention was auto, set to enabled, graph splits = 2. Cold generate: load 17.455s / prompt_eval 0.0238s / eval 0.306s. Hot (resident) generate: load 0.083s / prompt_eval 0.0248s / eval 0.149s. Confirmed that rocprofv3 wrap breaks bootstrap (CPU fallback). rocprofv3 --attach blocked by ptrace_scope=1.
• 2026-03-27
  Phase C — 共有ライブラリ依存追跡Phase C — Shared Library Map

  ldd で runtime 依存チェーンを確定： libggml-hip.so → libhipblas.so.2.3.60303 → librocblas.so.4.3.60303 → libamdhip64.so.6 → hsa-runtime64。 rocblas/library/ 下に gfx1201 対象ファイル 56 本を確認（Kernels.so-000-gfx1201.hsaco を含む）。

  Fixed runtime dependency chain via ldd: libggml-hip.so → libhipblas.so.2.3.60303 → librocblas.so.4.3.60303 → libamdhip64.so.6 → hsa-runtime64. Confirmed 56 gfx1201-targeted files under rocblas/library/ (including Kernels.so-000-gfx1201.hsaco).
• 2026-03-27
  Phase D — コードオブジェクト目録Phase D — Code Object Inventory

  libggml-hip.so の .hip_fatbin セクション（587 MiB）を調査。 custom 側のシンボル群：ggml_cuda_op_mul_mat_q・ggml_cuda_op_mul_mat_vec_q・mul_mat_q_case<Q4_K/Q5_K/Q6_K>・Flash Attention variants・RoPE など。 external rocBLAS 側に Cijk_* contraction family が共存することも確認。

  Surveyed the .hip_fatbin section (587 MiB) in libggml-hip.so. Custom-side symbols include: ggml_cuda_op_mul_mat_q, ggml_cuda_op_mul_mat_vec_q, mul_mat_q_case<Q4_K/Q5_K/Q6_K>, Flash Attention variants, RoPE, and others. Also confirmed that external rocBLAS ships a coexisting Cijk_* contraction family.
• 2026-03-27
  Phase E — 逆アセンブリによる絞り込みPhase E — Disassembly-Assisted Narrowing

  llvm-objdump（ROCm 7.2.0 付属）と clang-offload-bundler を使用し、.hip_fatbin から 28+ 個の gfx1201 hsaco を抽出。 主な binary anchor を確定：bundle_0030（Q4_K MMVQ, 1019 KB）・bundle_0096（exact Q4_K MMQ, 918 KB）・bundle_0037（MMQ repack）・bundle_0039（RoPE）・bundle_0019（Flash Attn）・bundle_0012（Q4_K dequant）。 現行 inspect ウィンドウでは v_mfma 命令は検出されていません。

  Used llvm-objdump (bundled with ROCm 7.2.0) and clang-offload-bundler to extract 28+ gfx1201 hsacos from .hip_fatbin. Fixed key binary anchors: bundle_0030 (Q4_K MMVQ, 1019 KB), bundle_0096 (exact Q4_K MMQ, 918 KB), bundle_0037 (MMQ repack), bundle_0039 (RoPE), bundle_0019 (Flash Attn), bundle_0012 (Q4_K dequant). No v_mfma instructions detected in the current inspection window.
• 2026-03-28
  Phase F — ランタイム・逆アセンブリ相関（初回パス）Phase F — Runtime–Disassembly Correlation (First Pass)

  Phase B–E の証拠を統合し、初回相関パスを実施： (1) decode 側は custom Q4_K MMVQ と最も強く整合 / (2) Attention は custom Flash family で自然に説明可 / (3) RoPE は custom gfx1201 family として確定 / (4) prompt 側（11〜12 token）は MMQ-eligible。 また resident phase-proxy probe を追加し、prompt_eval_count = 3（MMVQ）・11（MMQ-eligible）・290（境界候補）を記録。 short resident case が pure-MMVQ ではなく mixed-by-phase だと判明。

  Integrated evidence from Phases B–E and ran the first correlation pass: (1) Decode-side aligns most strongly with custom Q4_K MMVQ. (2) Attention is naturally explained by the custom Flash family. (3) RoPE is confirmed as a custom gfx1201 family. (4) Prompt-side (~11–12 tokens) is MMQ-eligible. Also added the resident phase-proxy probe: recorded prompt_eval_count = 3 (MMVQ), 11 (MMQ-eligible), 290 (boundary candidate). Found that the short resident case is mixed-by-phase, not pure-MMVQ.
• 2026-04-01
  Follow-up — bounded cluster 再点検Follow-up — Bounded Cluster Recheck

  既知の bundle_0027 / 0030 / 0037 / 0039 / 0082 / 0096 だけを再点検。 bundle_0030 の Q4_K mul_mat_vec_q は 1..8、bundle_0096 の exact Q4_K mul_mat_q は 8..128、bundle_0082 は IQ1_S と整理できた。 これにより、short resident case の prompt 側 11-12 窓は、bounded cluster 内では bundle_0096 とより強く整合すると判断した。

  Rechecked only the known bundle_0027 / 0030 / 0037 / 0039 / 0082 / 0096 cluster. Fixed that bundle_0030 carries Q4_K mul_mat_vec_q for 1..8, bundle_0096 carries exact Q4_K mul_mat_q for 8..128, and bundle_0082 is an IQ1_S neighbor. This makes the prompt-side 11-12 window in the short resident case align more strongly with bundle_0096 within the bounded cluster.

未確定事項（現時点）Open Questions (Current)

• short resident case が bundle_0096 自体を踏んだか、それとも別の同型 Q4_K MMQ bundle だったか。
• BLAS-side Cijk_* family の live phase ownership（short case での参加有無）。
• prompt_eval_count = 290 境界：MMQ heuristic 超過（fallback）か internal chunking か。

• Whether the short resident case exercised bundle_0096 itself or another structurally similar Q4_K MMQ bundle.
• Live phase ownership of the BLAS-side Cijk_* family (whether it participated in the short case).
• The prompt_eval_count = 290 boundary: MMQ heuristic overflow (fallback) or internal chunking?

追記予定Planned

• Phase G 以降：non-distorting observer 確保後に BLAS side / bundle_0096 live ownership を再確認。
• long prompt boundary（fallback vs chunking）の切り分け。

• Phase G+: Re-check BLAS-side and bundle_0096 live ownership once a non-distorting observer is available.
• Separate long-prompt boundary: fallback vs. chunking.