RX9070XT rDNA4

RX9070XT（rDNA4 / gfx1201）の経路観測・最適化調査の入口です。MI25 由来の多層観測方法論を移植する C14 pilot として、Phase A–F（ソース追跡〜ランタイム相関）を完了しています。

Entry point for RX9070XT (rDNA4 / gfx1201) path observation and optimization study. Phase A–F (source trace through runtime correlation) of the C14 pilot transplanting MI25 methodology are complete.

gfx1201 / ROCm 7.2.0 Ollama 0.17.4 Phase A–F 完了 Phase A–F Complete

観測済み事実Observed Facts

plain ollama serve で inference compute library=ROCm compute=gfx1201 を確認。offloaded 29/29 layers to GPU・Flash Attention was auto, set to enabled を取得しました（2026-03-27）。
runtime bundle は /usr/local/lib/ollama/rocm/ が先行し、libggml-hip.so → libhipblas.so → librocblas.so → libamdhip64.so の依存チェーンを確認しました。
逆アセンブリで gfx1201 向け custom hsaco（Q4_K MMVQ・MMQ・Flash Attn・RoPE）と BLAS 側 Cijk_* family の両方を isolate 済みです。
resident phase-proxy probe: prompt_eval_count = 3（MMVQ 圏）・11（MMQ-eligible 圏）・290（境界候補）。
rocprofv3 wrap は bootstrap を壊す（CPU fallback に落ちる）ことを確認。rocprofv3 --attach は ptrace_scope=1 でブロック。dispatch-safe observer はまだ確保できていません。

Plain ollama serve confirmed inference compute library=ROCm compute=gfx1201, offloaded 29/29 layers to GPU, and Flash Attention was auto, set to enabled (2026-03-27).
Runtime bundle uses /usr/local/lib/ollama/rocm/ first; confirmed dependency chain: libggml-hip.so → libhipblas.so → librocblas.so → libamdhip64.so.
Disassembly isolated both gfx1201 custom hsacos (Q4_K MMVQ/MMQ, Flash Attn, RoPE) and the BLAS-side Cijk_* family.
Resident phase-proxy probe: prompt_eval_count = 3 (MMVQ-compatible), 11 (MMQ-eligible), 290 (boundary candidate).
Confirmed that rocprofv3 wrap breaks bootstrap (CPU fallback). rocprofv3 --attach blocked by ptrace_scope=1. A dispatch-safe observer has not yet been secured.

このセクションのページPages in This Section

未確定事項Open Questions

short live case での BLAS 側 Cijk_* family の参加有無は未確定です。
prompt_eval_count = 290 境界が MMQ heuristic 超過（fallback）か internal chunking かは未確定です。
dispatch-safe observer なしでは、per-kernel attribution は推論にとどまります。

Whether the BLAS-side Cijk_* family participated in the short live case is unresolved.
Whether the prompt_eval_count = 290 boundary is MMQ heuristic overflow (fallback) or internal chunking is unresolved.
Without a dispatch-safe observer, per-kernel attribution remains inference only.

掲載情報は観測記録に基づきます。未確定事項は各ページの「未確定事項」欄に明示し、確定次第更新します。 Content is grounded in observation logs. Uncertain items are listed under "Open Questions" on each page and updated as findings solidify.