hipBLASLt

このページで得られる理解:hipBLASLt は LLM 推論で直接呼ばれているのか——ロードは確認済み、呼び出しは未確定。その理由と証拠の限界を記録する。

What you'll gain here: Is hipBLASLt directly called during LLM inference? Loading is confirmed; actual invocation is unresolved. This page records why and where the evidence runs out.

RX9070XT: /proc/maps でロード確認済み(Phase C) RX9070XT: Load confirmed in /proc/maps (Phase C) LLM 推論中の直接呼び出し: 未確定 Direct call during LLM inference: Unresolved

このページの結論Page Conclusion

ROCm スタックにおける位置づけPosition in the ROCm Stack

hipBLASLt は rocBLAS と並列に位置する特化 GEMM ライブラリです。backend に TensileLite(Tensile の軽量版)を持ち、activation fusion・sparse GEMM などの拡張機能を提供します。ggml-hip の現行ソースでは直接の hipblasLt* 呼び出しサイトは主たる dispatch 経路に現れていませんが、rocBLAS 内部依存により常にロードされます。

hipBLASLt is a specialized GEMM library positioned alongside rocBLAS. Its backend is TensileLite (a lightweight Tensile variant), providing extended features like activation fusion and sparse GEMM. In the current ggml-hip source, direct hipblasLt* call sites do not appear in the primary dispatch path, but it is always loaded due to the internal rocBLAS dependency.

→ 次の問い:「ロードされている」≠「呼ばれている」。では hipBLASLt は実際に LLM 推論で使われているのか?

→ Next question: "Loaded" ≠ "called." So is hipBLASLt actually used in LLM inference?

世代横断:TensileLite アーキテクチャ対応Cross-Generation: TensileLite Architecture Support

観測ポイント:hipBLASLt の TensileLite バックエンドがどの世代をサポートするか。
ソース:hipBLASLt リポジトリ / tensilelite/Tensile/Common/Architectures.py(line 32–63)

Observation point: Which generations are supported by hipBLASLt's TensileLite backend.
Source: hipBLASLt repository / tensilelite/Tensile/Common/Architectures.py (lines 32–63)

世代Generation GFX architectureMap コード名architectureMap codename xnack バリアントxnack variants 世界観Worldview
GCN5 / MI25 gfx900 "vega10"(line 36) なしNone 登録済み。ggml-hip custom kernel なし → BLAS 経路が主Registered. No ggml-hip custom kernels → BLAS path is primary
RDNA4 / RX9070XT gfx1201 "gfx1201"(line 63) なしNone ロード確認済み、直接呼び出しは未確定Load confirmed; direct call unresolved
CDNA4 / MI300X gfx942 "aquavanjaram"(line 46) gfx942:xnack+ / gfx942:xnack-(line 70) 本命経路。hipBLASLt の直接利用が最も見込まれる世代Primary path. Most likely generation where hipBLASLt is directly used
gfx900 / gfx942 / gfx1201 すべてが TensileLite の architectureMap に登録されており、Tensile 本体と同等のカバレッジ。ただし「登録済み」は「LLM 推論で直接使われる」を意味しない。 gfx900 / gfx942 / gfx1201 are all registered in TensileLite's architectureMap, matching Tensile proper in coverage. However, "registered" does not mean "directly used in LLM inference."

観測済み事実(RX9070XT / gfx1201)Observed Facts (RX9070XT / gfx1201)

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

示せることCan Show 示せないことCannot Show
実行時ロードの確認(/proc/maps, readelf NEEDED)Runtime load confirmed (/proc/maps, readelf NEEDED) LLM 推論中に実際に call された証拠がないNo evidence of actual calls during LLM inference
TensileLite Architectures.py による全 3 世代の登録確認All 3 generations registered in TensileLite Architectures.py dispatch-safe observer なしには、hipBLASLt が rocBLAS 内部で呼ばれているかを確認できないWithout a dispatch-safe observer, cannot confirm if hipBLASLt is called inside rocBLAS
ggml-hip ソース上に直接 call site がないことAbsence of direct call sites in ggml-hip source 「呼ばれていない」の証明にはならない——indirect 経路の可能性が残るThis does not prove it is never called — indirect paths remain possible

未確定事項Open Questions

次の観測点Next Observation Points

掲載情報は観測記録に基づきます。未確定事項は確定次第更新します。 Content is grounded in observation logs. Open questions are updated as findings solidify.