このページで得られる理解：rocBLAS がどの GPU 世代を認識し、どの経路で kernel を選ぶのか——Processor enum による世代識別の仕組みと、gfx1201 での実機観測を記録する。
What you'll gain here: How rocBLAS identifies GPU generations and selects kernels — the Processor enum identification mechanism and live observations on gfx1201.
Processor enum に登録済みで、handle.cpp の getActiveArch() が実行時に識別する（ソース確認）。rocblas/library/ から読み込んで使用する。Kernels.so-000-gfx1201.hsaco が実機 bundle に存在し、56 ファイルを確認（Phase C）。ne11 閾値に依存するため、workload 次第で頻度が変わる。Processor enum; getActiveArch() in handle.cpp identifies them at runtime (source-confirmed).rocblas/library/ matching the identified generation.Kernels.so-000-gfx1201.hsaco exists in the live bundle with 56 gfx1201 files confirmed (Phase C).ne11 threshold — frequency varies by workload.rocBLAS は ROCm 数学スタックの基盤層です。hipBLAS・MIOpen・ユーザーアプリケーションから呼び出され、GEMM（行列積）を含む BLAS 演算を実行します。kernel は Tensile によって GPU アーキテクチャごとに事前生成され、rocblas/library/ 以下に配置されます。
rocBLAS is the foundation layer of the ROCm math stack. Called by hipBLAS, MIOpen, and user applications to execute BLAS operations including GEMM. Kernels are pre-generated per GPU architecture by Tensile and placed under rocblas/library/.
rocBLAS リポジトリ（retired 扱い）。現行の active 開発は rocm-libraries リポジトリに移管されています。
Primary source: rocBLAS repository (marked retired). Active development has moved to the rocm-libraries repository.
→ 次の問い：rocBLAS はどうやって GPU 世代を識別し、その世代に合った kernel を選ぶのか？
→ Next question: How does rocBLAS identify the GPU generation and select the matching kernel?
観測ポイント：rocBLAS が実行時に GPU 世代を識別するために使う enum と分岐ロジック。
文脈：どの世代に何の kernel が届くかを理解するには、識別機構を先に把握する必要がある。
ソース：rocBLAS リポジトリ / library/src/include/handle.hpp（line 76–101）および library/src/handle.cpp（line 81–158）
Observation point: The enum and branching logic rocBLAS uses to identify GPU generations at runtime.
Context: To understand which kernels reach which generation, the identification mechanism must be understood first.
Source: rocBLAS repository / library/src/include/handle.hpp (lines 76–101) and library/src/handle.cpp (lines 81–158)
| 世代Generation | GFX | Processor enum 値 |
getActiveArch() 分岐 |
世界観Worldview |
|---|---|---|---|---|
| GCN5 / MI25 | gfx900 |
Processor::gfx900 = 900 |
明示分岐あり（line 93–95）Explicit branch (line 93–95) | fallback が主戦場。ggml-hip に gfx900 custom kernel なしFallback as main stage. No gfx900 custom kernels in ggml-hip |
| RDNA4 / RX9070XT | gfx1201 |
Processor::gfx1201 = 1201 |
明示分岐あり（line 153–155）Explicit branch (line 153–155) | BLAS fallback 経路を調査中（dispatch 未確認）BLAS fallback path under investigation (dispatch unconfirmed) |
| CDNA4 / MI300X | gfx942 |
Processor::gfx942 = 942 |
明示分岐あり（line 109–111）Explicit branch (line 109–111) | 本命経路が通る世代——正規ルート確認が主Primary path expected — canonical route confirmation is the main work |
getActiveArch() 分岐が存在する。Processor enum の整数値（900 / 942 / 1201）が handle.arch に格納され、以降の kernel 選択に使用される。
All three generations have explicit branches in getActiveArch(). The Processor enum integer values (900 / 942 / 1201) are stored in handle.arch and used for subsequent kernel selection.
→ 次の問い：識別した世代に対応する kernel が実際に bundle に存在するか？
→ Next question: Does the bundle actually contain kernels matching the identified generation?
観測ポイント：gfx1201 向け rocBLAS kernel が実機 bundle に存在し、逆アセンブリで確認できるか。
Observation point: Whether gfx1201 rocBLAS kernels exist in the live bundle and can be confirmed via disassembly.
/usr/local/lib/ollama/rocm/rocblas/library/）に gfx1201 ターゲットのファイルが 56 本存在（Phase C 確認）。Kernels.so-000-gfx1201.hsaco（standalone ELF AMD GPU object）が含まれる。Cijk_S_GA（sgpr:14, vgpr:10）・Cijk_S_PostGSU3（sgpr:21, vgpr:11）等を確認。wavefront_size=32。/usr/local/lib/ollama/rocm/rocblas/library/ (Phase C confirmed).Kernels.so-000-gfx1201.hsaco (standalone ELF AMD GPU object) confirmed present.Cijk_S_GA (sgpr:14, vgpr:10) and Cijk_S_PostGSU3 (sgpr:21, vgpr:11), wavefront_size=32.| 項目Item | 確認値Confirmed Value |
|---|---|
| gfx1201 ファイル数gfx1201 file count | 56（Kernels.so + TensileLibrary variants） |
| wavefront_size | 32（Cijk_S_GA metadata より / from Cijk_S_GA metadata） |
| 主要 kernel familyPrimary kernel families | Cijk_S_GA・Cijk_S_PostGSU3・Cijk_S_PostGSU3_VW2 など |
| 命令カウント（Kernels.so scan）Instruction count (Kernels.so scan) | v_fma: 245, v_fmac: 84, global_load: 2948, global_store: 230 |
| v_mfma | inspect window 内では未検出Not detected in inspection window |
ROCBLAS_TENSILE_LIBPATH 差分を用いた lane 比較対照実験を実施し、AETS lane / system lane 間の throughput 差を観測。BLAS 経路が実際に使われていることの間接証拠。ROCBLAS_LAYER=9 で rocblas_create_handle,atomics_allowed のみ記録。per-GEMM attribution には不足（Phase B 観測）。handle.cpp の rocblas_initialize() が gfx1201 の handle 初期化時に呼ばれることをソース確認。ROCBLAS_TENSILE_LIBPATH differences observed throughput differences — indirect evidence that the BLAS path is actually used.ROCBLAS_LAYER=9 only recorded rocblas_create_handle,atomics_allowed — insufficient for per-GEMM attribution (Phase B observation).rocblas_initialize() is called during gfx1201 handle initialization in handle.cpp.| 示せることCan Show | 示せないことCannot Show |
|---|---|
| 3 世代の Processor enum 登録と getActiveArch() 分岐の存在（ソース確認）Processor enum registration and getActiveArch() branches for all 3 generations (source-confirmed) | dispatch-safe observer なしには、BLAS 経路への入射頻度を定量化できないWithout a dispatch-safe observer, the frequency of BLAS-path invocations cannot be quantified |
| gfx1201 向け Cijk_* kernel の bundle 内存在（56 ファイル、Phase C/E 確認）Cijk_* kernel existence in the gfx1201 bundle (56 files, Phase C/E confirmed) | catalog 読み込みと実際の dispatch は別——kernel が catalog にあっても dispatch されるとは限らないCatalog read and actual dispatch are separate — a kernel in the catalog is not necessarily dispatched |
| MI25 で BLAS 経路が実際に使われている間接証拠（ROCBLAS_TENSILE_LIBPATH 実験）Indirect evidence of BLAS path usage on MI25 (ROCBLAS_TENSILE_LIBPATH experiment) | gfx900 と gfx1201 で同一 shape に対する Cijk_* 選択がどう異なるかは未観測How Cijk_* selection differs between gfx900 and gfx1201 for the same shape is unobserved |
Cijk_* family が実際に dispatch されたかは未確定（dispatch-safe observer なし）。Cijk_* family の動作差は未観測。Cijk_* family was dispatched in the RX9070XT short live case is unresolved (no dispatch-safe observer).Cijk_* family are unobserved.