なぜこの調査は難しく、なぜ面白いのか Why This Investigation Is Hard — and Why It's Interesting

このページで得られる理解：RX9070XT 調査が「何を問い、何が見えて、何が見えないのか」——データページを読む前に、この文脈を掴む。

What you'll gain here: what the RX9070XT investigation is asking, what can be seen, and what cannot — context to grasp before reading the data pages.

この調査（C14 portability pilot）は、MI25（AMD の古い世代の GPU）で確立した観測方法論を RDNA4（新しい世代）に移植する試みだ。 ところが、MI25 で当然のように使えた観測ツールが RDNA4 では軒並み機能しない。 その「機能しない」という事実が、そのまま研究の成果になる。
RDNA4 / gfx1201 が特に面白いのは、GPU 世代としてちょうど「本命経路（最適化された専用の計算ルート）が通るかどうか」の境界にいるからだ。 fallback（本命が使えないとき、別ルートに切り替えること）だけの世代でも、本命だけの世代でもない——観測しなければ分からない。

This investigation (C14 portability pilot) is an attempt to transplant the observation methodology established for MI25 (an older AMD GPU generation) onto RDNA4 (a newer generation). But the observation tools that worked naturally for MI25 fail across the board on RDNA4. That fact — "it doesn't work" — is itself the research finding.
RDNA4 / gfx1201 is particularly interesting because, as a GPU generation, it sits right at the boundary of "does the primary path (the optimized, dedicated compute route) actually run?" Neither pure-fallback (switching to an alternate route when the primary is unavailable) nor pure-primary — you have to observe to know.

C14 portability pilot gfx1201 / RDNA4 Phase A–F 完了 IQ1_S direct import 確認 IQ1_S direct import confirmed Q5_K_M direct import 確認 Q5_K_M direct import confirmed

まず3行で——この調査は何をしているのか In Three Lines — What This Investigation Is Doing

① GPU（グラフィックスカード）の中で「どの仕事が動いているか」を知りたい。
② でも直接のぞこうとすると、のぞいた瞬間に機械が止まる。
③ だから、外に出てくる数字やヒントを集めて、中で何が起きているかを推理する。 ① We want to know which "jobs" are running inside the GPU (graphics card).
② But if we try to look directly, the machine stops the moment we look.
③ So instead, we collect clues and numbers that emerge from outside, and reason about what's happening inside.

もう少し具体的に言うと——GPU の中では、行列同士のかけ算（GEMM: LLM の推論の大半を占める演算）が どの「仕事の担当カーネル（kernel: GPU の中で動く個々の仕事の単位）」に回されているかを、 実行時にしか確認できない。 ソースコードに経路が書いてあっても、実際にその経路を通るかどうかは動かしてみないと分からない。

To be more specific — inside the GPU, which "kernel (the individual job unit running inside the GPU)" handles the matrix multiplication (GEMM: the operation dominating most of LLM inference) can only be confirmed at runtime. Even if the source code describes a path, whether it's actually taken cannot be known until it runs.

たとえ話： 工場の機械が動いている。仕事の量によって「担当ライン A」か「担当ライン B」かを自動で選ぶ。 どちらのラインが動いているか知りたい。でも、確認しようとして機械に手を触れると止まってしまう。 だから「機械から出てくるメーターの数字」を読んで、どちらが動いているかを推理する——これがこの調査の本質だ。

An analogy: A factory machine is running. Depending on workload, it automatically selects "Line A" or "Line B." We want to know which line is active. But touching the machine to check makes it stop. So we read the meter numbers that come out from the machine, and reason about which line is running — that is the essence of this investigation.

このページの結論（最初にお読みください） Page Conclusion (read this first)

① gfx1201 は「本命経路が通るかどうか」を観測で確認する必要がある唯一の世代——それがこの調査の研究的意義。
② MI25 で使えた観測ツールは RDNA4 では全滅。「失敗」ではなく「観測可能性の地図」として記録する。
③ 消去法の末に残るのは phase proxy（response JSON）だけ。多層証拠の積み上げがこの調査の方法論。
④ 「見えない」には2種類ある：observer が壊れているのか、実際に起きていないのか——この切り分けが調査の核心的な知的作業。
⑤ 比較枝としての `IQ1_S` と `Q5_K_M` も、small single-file GGUF の direct import なら current Ollama で `100% GPU` resident case まで到達する。 ① gfx1201 is the one generation where observation is required to determine whether the primary path actually runs — that is this investigation's research significance.
② Observation tools that worked for MI25 fail entirely on RDNA4. Documented not as failure, but as a map of observability.
③ By elimination, only the phase proxy (response JSON) remains. Multi-layer evidence accumulation is the methodology.
④ "Invisible" has two causes: a broken observer, or something genuinely not happening — distinguishing these is the core intellectual work.
⑤ The comparative `IQ1_S` and `Q5_K_M` branches both reach a `100% GPU` resident case on current Ollama when imported as small single-file GGUFs.
追加で見えてきたこと：`Q4_K_M` main case の prompt-side MMQ owner 候補を囲む local bundle neighborhood は Q4_1 -> Q4_K -> Q5_0 -> Q5_1 -> Q5_K -> Q6_K -> Q8_0 -> Flash と並ぶ。
だからいま残る疑問は「近傍に別の `Q4_K` owner があるか」ではなく、 live `9+` window が既知の `Q4_K` artifact を本当に dispatch したか になっている。 One more thing is now visible: the local bundle neighborhood around the prompt-side MMQ owner candidate in the `Q4_K_M` main case reads as Q4_1 -> Q4_K -> Q5_0 -> Q5_1 -> Q5_K -> Q6_K -> Q8_0 -> Flash.
So the remaining question is no longer "is there another nearby `Q4_K` owner?" but whether the live `9+` window actually dispatched the known `Q4_K` artifact.
source 側ももう一段詰まっていて、RDNA4 では `ggml_cuda_should_use_mmq(...)` が `Q4_K` で source-side true になり、 non-split branch order は MMVQ -> MMQ -> BLAS fallback になる。
つまり `Q4_K_M` main case の `9+` window は、いまや「MMQ-eligible」より少し強く、 source 上も `mul_mat_q` 側へ寄る と読むのが自然になっている。 The source side is now tighter too: on RDNA4, `ggml_cuda_should_use_mmq(...)` is source-side true for `Q4_K`, and the non-split branch order becomes MMVQ -> MMQ -> BLAS fallback.
In other words, the `9+` window in the `Q4_K_M` main case is now a bit stronger than merely "MMQ-eligible": the source itself leans it toward `mul_mat_q`.
さらに、current Ollama path は llama.cpp の default LLAMA_SPLIT_MODE_LAYER をそのまま使っていて、 ROW-mode split buffer を source 上 request していない。
だから runtime の graph splits = 2 は、そのまま source-side split-buffer の証拠にはならない。 いま残る問いは「split だったか」ではなく、 live 9+ window が既知の Q4_K owner を本当に dispatch したか になっている。 In addition, the current Ollama path keeps llama.cpp at its default LLAMA_SPLIT_MODE_LAYER and does not request ROW-mode split buffers in source.
So runtime graph splits = 2 should not be read as direct evidence of source-side split-buffer mode. The remaining question is no longer "was it split?" but whether the live 9+ window actually dispatched the known Q4_K owner.

なぜ gfx1201 が研究として面白いのか — 世代ごとの「観測の世界観」 Why gfx1201 Is Interesting as Research — The Worldview of Each GPU Generation

この研究と観測を理解するために： RX9070XT (gfx1201) は GPU 世代の中でどんな立ち位置にあるのか？

To understand this investigation and observations: Where does RX9070XT (gfx1201) sit among GPU generations?

GPU 世代を横断してこの調査を理解するには、単なる「サポート状況の違い」としてではなく、 各世代が持つ「観測の世界観」の違いとして捉える必要がある。

To understand this investigation across GPU generations, the difference is not simply "support status" — it is the difference in what each generation's world of observation looks like.

世代Generation GFX 観測の世界観Observation Worldview 研究上の役割Research Role
Vega / MI25Vega / MI25 gfx900 古い世代のため、fallback が主戦場。本命経路は通らない前提で観測する。fallback の構造を正確に記録することが研究の価値。 Fallback is the main stage. Observe under the assumption the primary path won't run. Accurately mapping the fallback structure is the research value. fallback 観測Fallback observation
RDNA4 / RX9070XTRDNA4 / RX9070XT gfx1201 新しいが、民生用GPUであるため、通るものと通らないものが混在。観測しないと分からない。fallback でも本命でもない——"分岐そのもの"を観測する世代。研究として最も面白い。 Mixed: some paths run, some don't. Must observe to know. Neither pure-fallback nor pure-primary — this generation observes "the branching itself." Most interesting as research. 分岐観測 ← 今ここBranch observation ← here
MI300XMI300X gfx942 現行フラッグシップであるためほぼ本命が通ると見られる（要確認）。正規ルートの確認が主。fallback 観測は副次的。ただし MI400X 登場も噂されており、陳腐化は不可避。 Primary path expected to run (needs confirmation). Verifying the canonical path is the main work. Fallback is secondary. However, MI400X rumors make obsolescence inevitable. 本命確認（予定）Primary confirmation (planned)
gfx1201 は「fallback でも本命でもない」——"分岐そのもの"を観測する世代。 MI25 では fallback が答えだった。MI300X では本命が答えのはずだ。gfx1201 は「どちらが答えか」を実行時に確認しなければならない唯一の世代。 だからこそ観測が難しく、だからこそ研究として面白い。 gfx1201 is "neither fallback nor primary" — the generation that observes the branching itself. For MI25, fallback was the answer. For MI300X, primary is expected. gfx1201 is the only generation where "which one?" must be confirmed at runtime. That is why observing it is hard — and why it's interesting as research.
GPU generation worldview — gfx900 fallback / gfx1201 mixed / gfx942 primary

クリックで拡大 · gfx900 / gfx1201 / gfx942 の観測世界観の違い

Click to enlarge · Observation worldview differences across gfx900 / gfx1201 / gfx942

C14 portability pilot とは何か — 「移植」ではなく「方法論の検証」 What Is the C14 Portability Pilot — Not "Transplant" but "Method Validation"

観測ポイント： MI25 で確立した方法論は RDNA4 でそのまま使えるのか？

Observation target: Does the methodology established for MI25 work directly on RDNA4?

MI25 調査では「observe first, then optimize（まず観測し、それから最適化する）」の手順で観測クラスを確立した。 3本の柱——rocprofv3（AMD GPU の標準観測ツール）による kernel（GPU の仕事の単位）追跡・ROCBLAS_LAYER ログ・プロセスマップ観測——で「何が実行されたか」を直接確認できた。 C14 pilot はこの方法論を RX9070XT へ持ち込む試みだ。

The MI25 investigation established an observation class using an "observe first, then optimize" procedure. Three pillars — rocprofv3 (AMD's standard GPU observation tool) kernel tracing, ROCBLAS_LAYER logging, and process map observation — allowed direct confirmation of what was executed. C14 pilot is the attempt to bring this methodology to RX9070XT.

問いが変わる： 「MI25 でできたことを RDNA4 でもできるか」ではない。 「MI25 で当然だった観測窓が使えない状況で、何が観測可能で何が観測不可能かを切り分けられるか」——これが C14 pilot が問うていること。 答えが「できない」でも、「何がなぜ観測できないのか」を特定できれば、それは調査の失敗ではなく調査の成果。 The question shifts: Not "can we do in RDNA4 what we did in MI25?" But "in a situation where MI25's standard observation windows are unavailable, can we distinguish what is observable from what is not?" — that is what C14 asks. Even if the answer is "no," identifying what cannot be observed and why is a research finding, not a failure.

Phase A–F はその切り分けのための設計。 ソース追跡で理論的な分岐を固め（A/B）、ライブラリ（プログラムの部品）とバイナリ（実際に動く実行コード）の存在を確認し（C/D）、逆アセンブリ（実行コードを人間が読める形に変換すること）で実体を掴み（E）、runtime proxy（直接観測の代わりに外から読める値を手がかりにする手法）で近似相関を取る（F）。 → 次の問い： しかし、この手順を RDNA4 で実行しようとすると、何が壊れるのか？

Phase A–F is the design for this disambiguation. Fix theoretical branches in source (A/B), confirm library and binary existence (C/D), grasp substance via disassembly (E), get approximate correlation via runtime proxy (F). → Next question: But when we try to execute this procedure on RDNA4, what breaks?

Phase A–F investigation flow — source footing through runtime correlation

クリックで拡大 · Phase A–F の調査フロー（現在の frontier: Phase F）

Click to enlarge · Phase A–F investigation flow (current frontier: Phase F)

調査の流れ — 何が知りたくて、何をして、どうなったか The Investigation Flow — What We Wanted to Know, What We Did, and What Happened

調査の出発点は単純な疑問だ：「RX9070XT 上で Ollama が Q4_K_M モデルを動かすとき、GPU の中で何が起きているのか？」 これを追うために Phase A〜F を設計した。各フェーズで「何が知りたかったか」「何をしたか」「期待と結果のギャップ」を平易にまとめる。

The investigation starts from a simple question: "When Ollama runs a Q4_K_M model on RX9070XT, what is happening inside the GPU?" Phases A–F were designed to pursue this. Below is a plain-language summary of what each phase wanted to know, what was done, and where expectation and reality diverged.

Phase 何が知りたかったかWhat we wanted to know 何をしたかWhat was done 期待 vs 結果Expected vs. Actual
A ソースの dispatch 分岐はどこにあるか？どんな条件で kernel が切り替わるか？ Where are the dispatch branches in source? Under what conditions do kernels switch? ggml / Ollama のソースを追跡し、ggml_cuda_mul_mat() の ne11 閾値（≤ 8 / ≤ 256）を特定した。 Traced ggml / Ollama source and identified ne11 thresholds (≤ 8 / ≤ 256) in ggml_cuda_mul_mat(). 期待通り確定Confirmed as expected
分岐はソースで明確に読めた。 Branches were clearly readable in source.
B Ollama から GPU kernel までの呼び出しチェーンはどうなっているか？ What does the call chain from Ollama to GPU kernel look like? Ollama → ggml → hipBLAS / custom kernel の経路をソースで追跡した。 Traced Ollama → ggml → hipBLAS / custom kernel path through source. 期待通り確定Confirmed as expected
呼び出しチェーンをソースで完全に追えた。 Full call chain confirmed through source.
C 実際に動く Ollama プロセスはどのライブラリをロードしているか？ Which libraries does the running Ollama process actually load? /proc/<runner>/maps を確認し、12 ライブラリのロードを確認した。 Checked /proc/<runner>/maps and confirmed 12 libraries loaded. 確定（ただし load ≠ dispatch）Confirmed (but load ≠ dispatch)
ロードは確認できた。しかし「ロードされている = 使われている」ではないことも同時に発見。 Loading confirmed. Also discovered simultaneously that "loaded ≠ dispatched."
D gfx1201 向けのバイナリは実際に fatbin に入っているか？どの kernel が存在するか？ Are gfx1201 binaries actually present in the fatbin? Which kernels exist? libggml-hip.so の .hip_fatbin（587 MiB）を解析し、gfx1201 向け hsaco を列挙した。 Analyzed the .hip_fatbin (587 MiB) in libggml-hip.so and enumerated gfx1201 hsacos. 期待通り確定Confirmed as expected
MMVQ / MMQ / BLAS の gfx1201 binary はすべて存在した。 gfx1201 binaries for MMVQ / MMQ / BLAS all exist.
E Q4_K を担う具体的な kernel の register footprint・命令パターンは？ What are the register footprints and instruction patterns of the specific Q4_K kernels? Q4_K 候補 bundle を isolate し逆アセンブルした（bundle_0030 / bundle_0096 等）。 Isolated and disassembled Q4_K candidate bundles (bundle_0030 / bundle_0096, etc.). 期待通り確定Confirmed as expected
gfx1201 固有の wavefront_size=32 を確認。v_mfma は検出されず。 gfx1201-specific wavefront_size=32 confirmed. No v_mfma detected.
F live run で実際にどの dispatch 圏に入っているか？runtime を壊さずに確認できるか？ Which dispatch zone does a live run actually enter? Can this be confirmed without breaking the runtime? rocprofv3・--attach・ROCBLAS_LAYER を試みたが全滅。消去法で phase proxy（response JSON）を採用し probe を実施した。 Tried rocprofv3, --attach, and ROCBLAS_LAYER — all blocked. By elimination, adopted phase proxy (response JSON) and ran probes. 近似のみ確定・直接確認は不可Approximate only — direct confirmation blocked
current best reading は pure-MMVQ ではなく mixed-by-phase。decode 側は MMVQ-compatible、short/medium prompt 側はすでに MMQ-eligible と読むのが最も筋がいい。ただし dispatch-safe observer がなく、直接確認はできない。これが現在の frontier。 The current best reading is not pure MMVQ but mixed-by-phase. Decode looks MMVQ-compatible, while short/medium prompt windows already read as MMQ-eligible. However, without a dispatch-safe observer, direct confirmation is still impossible. This is the current frontier.
ここで分かること： Phase A〜E は期待通りに進んだ。ソースは読め、バイナリは存在し、逆アセンブリも成功した。 躓いたのは Phase F——「実際に動いた」を直接確認する段階で、全ての observer が機能しなかった。 ただし proxy からの narrowing 自体は進んでおり、`q4_k_m` short resident case も pure-MMVQ ではなく mixed-by-phase 読みまで狭まっている。 それでもなお「バイナリが存在すること」と「実際に使われること」の間にギャップが残り続けている。 What this tells us: Phases A–E proceeded as expected. Source was readable, binaries exist, disassembly succeeded. The stumbling point was Phase F — at the stage of directly confirming "it actually ran," all observers failed. But proxy-based narrowing did progress: even the `q4_k_m` short resident case has already narrowed from "pure MMVQ" to a mixed-by-phase reading. Even so, the gap between "the binary exists" and "it was actually used" still persists.
ただし practical ingest 側では前進がある。`IQ1_S` は requantize helper では狭いが、existing single-file GGUF の direct import は current host で通り、 `ollama create -> api/generate -> ollama ps` まで成立した。つまり今の `IQ1_S` の未確定は「入口が通るか」ではなく、「live path の主体がどの `IQ1_S` family か」に移っている。 There is practical progress on the ingest side, though. `IQ1_S` remains narrow through the requantize helper, but direct import of an existing single-file GGUF does work on the current host through `ollama create -> api/generate -> ollama ps`. So the remaining `IQ1_S` uncertainty is no longer "can it enter?" but rather "which `IQ1_S` family owns the live path?"
さらに `IQ1_S` は、source と binary だけでなく bounded phase proxy でも narrowing が始まっている。 `bundle_0030` には `IQ1_S` MMVQ (`1..8`) が、`bundle_0082` には `IQ1_S` MMQ (`8..128`) があり、 small direct-import case の live prompt window でも `7` は MMVQ-compatible、`13` は MMQ-eligible と読むのが自然。 ここでも `存在する`・`import できる`・`その family が live case を支配する` を別段階として扱うのが重要。 `IQ1_S` is now narrowing not only in source and binary space, but also through bounded phase proxy. `bundle_0030` carries `IQ1_S` MMVQ (`1..8`) while `bundle_0082` carries `IQ1_S` MMQ (`8..128`), and in the small direct-import live case the prompt windows read `7` as MMVQ-compatible and `13` as MMQ-eligible. Here too, it is important to treat `exists`, `can be imported`, and `owns the live case` as separate stages.
`Q5_K_M` も practical path が前進した。current Ollama helper path は依然 narrow だが、 existing single-file GGUF の direct import は current host で通り、resident case も `100% GPU` だった。 bounded phase proxy では short window が `7`、medium window が `13` で、current best reading は `Q5_K` MMVQ-compatible / MMQ-eligible である。 さらに `Q5_K` MMQ 側も `bundle_0098` として exact extracted hsaco まで取れた。 そのため medium 側の読みは以前より強くなったが、なお `object exists` と `that object owned the live window` は別段階なので、 `Q5_K_M` の medium 側 ownership には still unresolved を残す。 `Q5_K_M` has also moved forward on the practical path. The current Ollama helper path remains narrow, but direct import of an existing single-file GGUF now works on the current host, and the resident case stayed `100% GPU`. In bounded phase proxy, the short window was `7` and the medium window was `13`, so the current best reading is `Q5_K` MMVQ-compatible / MMQ-eligible. The `Q5_K` MMQ side has now also been isolated as exact extracted hsaco in `bundle_0098`. That strengthens the medium-window reading, but `object exists` and `that object owned the live window` are still different stages, so medium-window ownership remains unresolved.
この比較を 3 本で並べると役割がはっきりする。 `Q4_K_M` は main case study、`Q5_K_M` は最も近い near-neighbor、`IQ1_S` は最初の branch broadener である。 3 本とも、source skeleton / exact extracted MMVQ-MMQ anchor / practical ingest viability / bounded phase proxy までは揃った。 つまり現在の共通ボトルネックは live ownership であり、method 自体は 1 本の量子化に閉じていない。 Putting the three side by side makes their roles clearer. `Q4_K_M` is the main case study, `Q5_K_M` is the closest near-neighbor, and `IQ1_S` is the first branch broadener. All three now reach source skeleton, exact extracted MMVQ/MMQ anchors, practical ingest viability, and bounded phase proxy. In other words, the shared bottleneck is now live ownership, and the method is no longer tied to a single quantization family.

Observer 問題 — 「観測しようとしたら壊れた」は発見である The Observer Problem — "It Broke When We Tried to Observe" Is a Finding

観測ポイント： 各 observer を試した結果、何が起きたか。

Observation target: What happened when each observer was tried?

RDNA4 調査で最初に直面する壁は、観測ツールが軒並み機能しないことだ。 これは「調査が失敗した」のではなく、「RDNA4 の runtime model がツールチェーンの想定と異なる」という発見。 各手段を試した結果：

The first obstacle in the RDNA4 investigation is that observation tools fail across the board. This is not "the investigation failed" — it is the discovery that RDNA4's runtime model differs from what the toolchain assumes. Results from each attempt:

手段Approach 結果Result 何が起きたか / 何が分かったかWhat happened / What it reveals
rocprofv3 wrap 無効（CPU fallback）Invalid (CPU fallback) profiler の介在が Ollama の GPU bootstrap discovery を壊す。GPU 自体が使われなくなる。この状態のデータは observer 汚染として扱う。 Profiler interposition breaks Ollama's GPU bootstrap discovery. The GPU stops being used. Data in this state is treated as observer-contaminated.
rocprofv3 --attach ブロックBlocked ホストの ptrace_scope=1 が ptrace attach を拒否。OS のカーネルセキュリティ設定であり、変更はスコープ外。 Host ptrace_scope=1 rejects ptrace attach. OS kernel security setting — changing it is out of scope.
ROCBLAS_LAYER=9 部分のみPartial only rocblas_create_handle のみ記録。per-GEMM の dispatch 追跡には不足。Q4_K の custom kernel 経路は見えない。 Only rocblas_create_handle logged. Insufficient for per-GEMM dispatch tracking. Custom kernel paths for Q4_K are invisible.
phase proxy（response JSON）phase proxy (response JSON) 有効（近似）Valid (approximate) runtime を壊さない唯一の現実的な手段。他が全滅した消去法の結果として残った。 The only practical means that doesn't break the runtime. What remains by elimination after all others failed.

ここから分かること： ツールが壊れた理由を特定できることが重要。これらは「観測の失敗」ではなく「観測窓の構造的限界の確認」。
→ 次の問い： observer が壊れているとき、「見えない」ことは2種類ある。それをどう切り分けるか？

What this tells us: Being able to identify why tools fail is what matters. These are "confirmed structural limits of the observation window," not "observation failures."
→ Next question: When an observer is broken, "invisible" has two possible causes. How do we distinguish them?

「見えない」の切り分け：observer の問題か、実際に起きていないのか Distinguishing "Invisible": Broken Observer vs. Genuinely Not Happening

観測ポイント： 何かが観測されないとき、それはどちらの原因か。

Observation target: When something is not observed, which cause is it?

原因A — 観測窓の問題（observer が壊れている）： Cause A — Observation window limit (observer is broken): 起きていることが見えない。例：rocprofv3 が CPU fallback を引き起こすため、GPU 側の実行を追跡できない。 What is happening cannot be seen. Example: rocprofv3 causes CPU fallback, so GPU-side execution cannot be traced.
原因B — target 固有の特徴（実際に起きていない）： Cause B — Target-specific characteristic (genuinely not happening): 実際に起きていない、または起きる経路に入っていない。例：MIOpen が /proc/maps に現れないのは、Ollama の LLM 推論経路が MIOpen を使用しないから（bundle にも含まれていない）。 It is genuinely not happening, or the path leading to it is not entered. Example: MIOpen not appearing in /proc/maps can be read as Ollama's LLM inference path not using MIOpen (it is also absent from the bundle).

この切り分けが、本調査の中心的な知的作業。 dispatch-safe observer がない状態では、多くの「見えない」は原因A と B が区別できない。 それを区別できた事例・できなかった事例を正直に記録することが、この調査の価値になる。
→ 次の問い： ライブラリがロードされているとき、それは「使われている」ことを意味するのか？

This distinction is the central intellectual work of this investigation. Without a dispatch-safe observer, cause A and cause B cannot be distinguished for many "invisibles." Honestly recording which cases could be distinguished and which could not is where this investigation's value lies.
→ Next question: When a library is loaded, does that mean it is being used?

load ≠ dispatch — ライブラリが読み込まれることと演算が実行されることは別の問題 load ≠ dispatch — Library Loading and Operation Dispatch Are Separate Questions

観測ポイント： /proc/maps でライブラリが見えるとき、そのライブラリの演算は実行されているのか。

Observation target: When a library appears in /proc/maps, is that library's compute actually being dispatched?

Phase C の調査で、Ollama runner の /proc/maps に libhipblaslt.so.0.10.60303 が記録されていることを確認した。 これは事実：「hipBLASLt がロードされている」。 しかしここから「hipBLASLt の演算が LLM 推論ループで dispatch された」とは言えない。

Phase C confirmed that libhipblaslt.so.0.10.60303 appears in the Ollama runner's /proc/maps. This is the fact: "hipBLASLt is loaded." From this, however, it cannot be stated that "hipBLASLt operations were dispatched during the LLM inference loop."

なぜ区別が必要か。Q4_K_M 推論では ne11（バッチサイズ相当）の値に応じてカーネルが分岐する。 decode フェーズ（トークン生成）の多くは ne11 が小さく、MMVQ や MMQ の custom HIP kernel が主経路になる。 BLAS fallback（hipBLAS → rocBLAS → hipBLASLt）は ne11 が大きい prefill か、特定 shape の場合に限られる可能性がある。

Why does the distinction matter? In Q4_K_M inference, kernels branch based on ne11 (analogous to batch size). Most of the decode phase (token generation) likely has small ne11, where MMVQ and MMQ custom HIP kernels are the primary paths. The BLAS fallback (hipBLAS → rocBLAS → hipBLASLt) may only be reached during the large-ne11 prefill phase, or for specific shapes.

dispatch-safe observer がない現状では、「ロードされている」と「実行された」の間の溝を埋めることができない。 これは本調査の現時点での根本的な限界——そして phase proxy が登場する理由。 Without a dispatch-safe observer, the gap between "is loaded" and "was executed" cannot be bridged. This is the fundamental current limitation of this investigation — and the reason the phase proxy exists.

phase proxy — 消去法の末に残った唯一の観測手段 Phase Proxy — The Last Observation Method Standing

観測ポイント： runtime を一切壊さずに、dispatch 圏の近似を読める手段があるか。

Observation target: Is there any means to approximate the dispatch zone without touching the runtime at all?

wrap も attach も layer log も機能しない状況で、改めてソース側の dispatch（GPU へどの仕事を回すかの割り当て）分岐を確認した。 MMVQ（1トークンずつ返す場面で使われやすい専用カーネル経路）/ MMQ（最初にプロンプトをまとめて読む場面で使われやすい専用カーネル経路）の分岐は、 ne11（行列の列数——この調査では、だいたいプロンプトのトークン数に対応）で決まる。 Ollama の response JSON の prompt_eval_count（Ollama が返す「プロンプトを何トークン処理したか」の数値）を ne11 の近似指標として使えないか——それが phase proxy。

With wrap, attach, and layer log all non-functional, we revisited the source dispatch logic (dispatch = the assignment of which job goes to the GPU). The MMVQ (the dedicated kernel path favored in single-token generation) / MMQ (the dedicated kernel path favored when reading the full prompt upfront) branch is determined by ne11 (matrix column count — in this investigation, roughly corresponding to prompt token count). Could prompt_eval_count in Ollama's response JSON (the number "how many prompt tokens were processed" that Ollama returns) serve as a proxy for ne11? — that is the phase proxy.

実際に何をしているか（手の動き）：
① 短いプロンプトを Ollama に投げる（例：3トークン）
② 返ってきた response JSON の prompt_eval_count を読む → 3
③ これは ne11 = 3 ≤ 8 → MMVQ の担当圏に入っているはず と推理

④ 長いプロンプトを投げる（例：11トークン）
⑤ prompt_eval_count を読む → 11
⑥ これは 8 < ne11 ≤ 256 → MMQ の担当圏に入っているはず と推理

→ GPU の動作は壊れない。「プロンプトの長さを変えることで、どの担当圏に入るかを自分で制御できる」のがこの手法の強み。 Concretely, what you do (the hand movements):
① Send a short prompt to Ollama (e.g., 3 tokens)
② Read prompt_eval_count from the response JSON → 3
③ Since ne11 = 3 ≤ 8 → infer "it's in MMVQ territory"

④ Send a longer prompt (e.g., 11 tokens)
⑤ Read prompt_eval_count → 11
⑥ Since 8 < ne11 ≤ 256 → infer "it's in MMQ territory"

→ GPU operation is not disrupted. The key strength: "by varying prompt length, you can deliberately control which zone gets entered."

With wrap, attach, and layer log all non-functional, we revisited the source dispatch logic. The MMVQ/MMQ branch is determined by ne11 (≈ prompt token count). Could prompt_eval_count in Ollama's response JSON serve as a proxy for ne11? — that is the phase proxy.

probe ケースProbe Case prompt_eval_count 現時点で最も筋のいい読みStrongest current reading
threshold 未満Below threshold 7 / 8 MMVQ 圏（ne11 ≤ 8）MMVQ zone (ne11 ≤ 8)
threshold 超えAbove threshold 9 / 11 / 12 MMQ-eligible 圏（bundle_0096 候補）MMQ-eligible zone (bundle_0096 candidate)
boundary 候補Boundary candidate 290 MMQ/BLAS 境界付近（未確定）Near MMQ/BLAS boundary (unresolved)
decode-heavyDecode-heavy 3 eval_count=16 / 0.1479s（decode ベースライン）eval_count=16 / 0.1479s (decode baseline)
proxy の限界： prompt_eval_count は response JSON の値であり、kernel launch 時の実際の ne11 を直接読んでいない。 chunked prefill や KV cache の挙動によって同じトークン数でも ne11 が変わりうる。 上の読みは「最も筋のいい解釈」であり、本来は dispatch-safe observer による直接確認が必要。 Proxy limitations: prompt_eval_count is from the response JSON — not a direct reading of ne11 at kernel launch. Chunked prefill or KV cache behavior could cause ne11 to differ even for the same token count. The readings above are "strongest current interpretations" — direct confirmation by a dispatch-safe observer is the proper requirement.

この調査を読む際の注意 — バッジと callout の意味 A Note on Reading This Investigation — Badge and Callout Meanings

各ページで使われているバッジと callout は、情報の確度を色で示す。なぜなら、「未確定が多い」ことは品質の低さではなく、むしろ未確定を正直に記録することがこの調査における品質基準だからである。

Badges and callouts across pages indicate confidence level by color. "Many unresolved items" does not indicate low quality — honestly recording what is unresolved is the quality standard of this investigation.

表示Display 意味Meaning
確認済み / Confirmed 実機コマンド・ログ・ファイル存在・逆アセンブリ等で直接確認した事実Directly confirmed via live commands, logs, file presence, or disassembly
推論 / Inferred 確認した事実から論理的に導かれる解釈。直接確認はしていないInterpretation logically derived from confirmed facts. Not directly confirmed
未確定 / Unresolved 現時点では確定できない事項。observer 制約または未実施の観測によるCannot be determined at this stage, due to observer constraints or observations not yet performed
制約あり / Limited 部分的に確認できているが、tool 制限により完全ではないPartially confirmed, but incomplete due to tool limitations

次にどこを読むか Where to Go Next

このページで文脈を掴んだら、以下の順で読むと調査の流れを追いやすい。

Now that you have the context, reading in the following order makes the investigation flow easiest to follow.

本ページは本調査への入口として機能します。各データページは RX9070XT インデックス から参照できます。 This page serves as the entry point to this investigation. Individual data pages are accessible from the RX9070XT index.