Radeon VIIとUbuntu 26.04でROCmを動かし、Flash-NextをOpenCodeにつないだ
5年ほど前に組んだデスクトップがある。当時、GPUにはすでに型落ちだったRadeon VIIを5万円くらいで入手して載せた。Stable Diffusionで遊んでいた時期もあったが、飽きてからはすっかり出番がなくなっていた。
最近はローカルLLMが盛り上がっている。とはいえ、パソコンもGPUもクソ高い。新しいGPUを買う前に、あまり期待せず、久しぶりにこのデスクトップに火を入れて環境を組み直してみることにした。
気になったのは、Radeon VIIのROCm公式サポートが切れていることだ。かなり苦戦する覚悟だったが、実際にはChatGPTのGPT-6.1 Terra（medium）に相談しながら、環境の調査から構築、動作確認まで進められた。多少の修正はあったものの、さくさく進む。最終的に、Qwen3.8-Flash-Nextが動き、OpenCodeからローカルモデルへツール付きでアクセスできるようになった。古いデスクトップでvibe codingを試せる環境ができてしまい、だいぶびっくりしている。
この記事は、そのRadeon VIIでROCmを動かし、ローカルLLMをOpenCodeにつなぐまでの記録だ。
結果は、Q2_0量子化版のFlash-Nextで短めの会話の生成が約35 tokens/s。別途、Qwen3.8-27BのQ3量子化版もllama.cppで約21〜22 tokens/s出ている。両者は量子化も実行エンジンも違うため、モデル同士の純粋な比較ではない。
この記事は2026年11月3日時点の実機メモだ。gfx906のROCmはAMDの通常の動作保証対象とは別扱いで、以下は一つの環境で動いた記録だ。アップデート後も同じように動くとは限らない。
環境と構成
項目
今回の環境
GPU
Radeon VII、VRAM 16GB、gfx906
CPU
Ryzen 5 3600、6コア12スレッド、AVX2
RAM
公称64GB、OSから約60.7GiB
OS
Ubuntu 26.04.1 LTS
カーネル
7.0.0-42-generic、標準のamdgpu/KFD
ROCm
TheRock nightly 10.3.0a20261102、gfx906向け
HIPのバージョン表示
7.18.26434
推論
Strata v0.2.10、gfx906用ソースビルド
モデル
ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF、Q2_0
コーディングUI
OpenCode v1.19.5
Flash-Nextは本体125B、1トークンで有効になるのは約6BのMoEモデルで、別に51BのN-gram参照テーブルがある。Q2_0はパラメーターを削った小型モデルではなく、重みを量子化したものだ。Qwenの説明と量子化版の配布元を参照してほしい。
構成は大きく3層だ。
Ubuntu標準のamdgpu/KFDがGPUを扱う。
gfx906向けROCmのユーザー空間ライブラリとHIPが計算を担当する。
StrataがモデルをGPU・RAM・ストレージに配置し、ローカルAPIを公開する。
コンテナは使わず、ROCmやビルド成果物はホームディレクトリ下に分離した。システムのドライバやAPTリポジトリは変更していない。
まずamdgpuとKFDを確認する
画面が映っているだけではGPU計算が動くとは限らない。まずカーネルドライバとデバイスノードを確認した。
uname -a
lspci -nnk | grep -A 4 -i 'vga|display|3d'
ls -l /dev/kfd /dev/dri
getfacl /dev/kfd /dev/dri/renderD128
このPCでは1002:66afのRadeon VIIをamdgpuが使用し、/dev/kfdとrenderD128にはログインユーザー向けのACLが付いていた。アクセス権の設定方法は環境によって異なる。/dev/dri/renderD128も固定とは限らないので、自分の環境で確認しよう。
gfx906向けROCmをユーザー領域へ
TheRockの対応表ではgfx906はビルド対象だが、Sanity Tested・Release Readyにはチェックがない。配布されていることと、実機で動作検証されていることは別だ。
今回はAMDのnightlyから、同じ日付のSDK・ライブラリ・開発ツール・gfx906向けデバイスコードを仮想環境に揃えた。uvは別途用意した。以下の版は実験時点の固定値で、nightlyの保持期間や配布状態によっては取得できなくなる可能性がある。
uv venv ~/rocm-gfx906/venv
uv pip install --python ~/rocm-gfx906/venv/bin/python \
--index-url nightly.repo.amd.com \
'rocm[libraries,devel,device-gfx906]==10.3.0a20261102'
source ~/rocm-gfx906/venv/bin/activate
rocm-sdk init
export ROCM_PATH="$(rocm-sdk path --root)"
export HIP_PATH="$ROCM_PATH"
export PATH="$ROCM_PATH/bin:$PATH"
export LD_LIBRARY_PATH="$ROCM_PATH/lib:$ROCM_PATH/lib64${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}"
rocminfo
これは配置先を分かりやすく書き直した再現例だ。実機では~/.local/share/rocm-gfx906/venvを使った。Ubuntu標準のPythonではvenvのensurepipがなかったため、uvで作成した。
rocminfoではRadeon VIIがgfx906:sramecc+:xnack-として見えた。さらにHIPで4,096要素のベクトル加算、rocBLASで2×2と512×512の行列積を検証した。SDK付属テストも27項目が通過した。
HIP compiled vector addition: PASS (4096 elements)
rocBLAS SGEMM 512x512 identity product: PASS
Ran 27 tests ... OK
amdflangは最初、libquadmath.so.0不足で起動しなかった。今回はlibquadmath0をapt downloadしてユーザー領域に展開し、ライブラリパスを通して解決した。同様にC++開発ヘッダーとGCCの開発ライブラリもユーザー領域へ展開した。一般的な環境では、必要な開発パッケージをディストリビューションの方法で揃える方が簡単だ。無関係なUbuntu版のROCmパッケージを混ぜないようにした。
最初にllama.cppでGPU推論を確認
Strataに進む前に、llama.cppをHIP有効・gfx906指定でビルドした。使用したllama.cppのコミットはf9448cddc48c129a4abc7ff9c2a84c2b4c101c61だ。ビルド方法の基本はllama.cppの公式ドキュメントにある。
llama-cli --list-devicesにROCm0: AMD Radeon VIIが表示され、Qwen2.5-0.5BのGGUFで短い生成に成功。次にQwen3.8-27BのUD-Q3_K_XLも試した。モデルファイルは約13.2GBだ。
llama-server -m /path/to/Qwen3.8-27B-UD-Q3_K_XL.gguf \
-ngl 99 -c 4096 -np 1 --host 127.0.0.1 --port 8082
日本語の短い質問に正常回答し、会話では約21〜22 tokens/s。後にログを詳細化して、66/66層がGPUに配置されていることも確認した。別の8K設定では、約7,390トークンの合成入力を約240 tokens/sで処理した。いずれも今回のハードウェアと設定での観測だ。
Strataをgfx906向けにビルドする
StrataのAMD向けドキュメントには、gfx906をwave64互換レイヤーでビルドする手順がある。配布済みのRX 6000/7000/9000向けバイナリをRadeon VIIに流用するのではなく、専用のSTRATA_HIP_GFX906経路を使う。
Strata v0.2.10のソースと、プロジェクトが指定するllama.cppコミット3cf03257f219afbe7334045ff7c6a06ac68c627dを用意した。システムにC/C++開発環境、CMake、Ninja、ROCmの開発ツールが必要だ。
STRATA_SRCとLLAMA_SRCは、それぞれ取得・展開したソースディレクトリ
cmake -S "$STRATA_SRC" -B "$STRATA_SRC/build-906" -G Ninja \
-DSTRATA_HIP_GFX906=ON \
-DCMAKE_HIP_ARCHITECTURES=gfx906 \
-DSTRATA_GGML_DIR="$LLAMA_SRC" \
-DCMAKE_C_COMPILER="$ROCM_PATH/bin/amdclang" \
-DCMAKE_CXX_COMPILER="$ROCM_PATH/bin/amdclang++" \
-DCMAKE_PREFIX_PATH="$ROCM_PATH" \
-DCMAKE_BUILD_TYPE=Release
cmake --build "$STRATA_SRC/build-906" --target strata strata-device -j 6
実機ではシステムにGCC開発環境がなかったため、APTからダウンロードして展開したlibstdc++-15-devとlibgcc-15-devを指す--gcc-toolchainをC/C++/HIPの各フラグに追加した。上は開発環境が通常の場所にある場合の基本形だ。
ビルド時に互換レイヤーの定義を追加した
v0.2.10をそのままビルドすると、cudaDeviceGetStreamPriorityRangeとcudaStreamCreateWithPriorityが未定義になった。gfx906互換ヘッダーinclude/strata/platform/hip_compat/strata_hip.hに、対応するHIP APIへのマッピングを追加して再ビルドした。
#define cudaStreamCreateWithFlags hipStreamCreateWithFlags
+#define cudaStreamCreateWithPriority hipStreamCreateWithPriority
+#define cudaDeviceGetStreamPriorityRange hipDeviceGetStreamPriorityRange
これは今回の版とビルドで必要だったローカル修正だ。将来の版では修正済みかもしれない。ビルド後、strata-device --selftest、gfx906向けGDN演算テスト、専門家演算の比較テストで失敗がないことを確認した。これらは全機能や長時間の安定性を保証するものではない。
Flash-Nextを準備して起動
モデルはISTA-DASLabのGGUFのQ2_0を選んだ。37.6GBの本体と28.8GBの参照テーブル、合わせて約66.4GBだ。今回使用したリビジョンはed59f92082b1e93c0e96d60a8b11aab089b52f09。両ファイルを配布元のSHA-256と照合した。
Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00001-of-00002.gguf
SHA256 69820c02ec7d0b45ef2ebb19d6620299db749fe2aded7f39f93c6b88b199b720
Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00002-of-00002.gguf
SHA256 316b46f3a2dbd68c900f43136ab9449f9dcc3725dfd8c794847c204bc161e113
転送が途中で停滞して再始動した際、最初の不完全ファイルを引き継がず、新しいファイルとして再取得した。完成後に古い不完全ファイル約17.6GBを削除した。大きなモデルでは、途中再開の動作と空き容量を確認したいところだ。
取得後、Strataに付属するtools/iq_pack.pyでインデックス、dense重み、トークナイザーを準備する。STRATA_GGUF_PYは上で用意した固定コミットのllama.cppにあるgguf-pyを指す。
export STRATA_GGUF_PY="$LLAMA_SRC/gguf-py"
python "$STRATA_SRC/tools/iq_pack.py" \
--gguf /path/to/Q2_0/Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00001-of-00002.gguf \
--out /path/to/pack-q2_0
サーバー設定には、ビルドしたstrata、pack、2つのGGUF、付属のdata/expert-profile.binを指定する。主要なエンジン引数は次のとおりだ。
{
"exe": "/path/to/build-906/strata",
"cwd": "/path/to/Strata",
"tokenizer": "/path/to/pack-q2_0/tokenizer",
"model_name": "Qwen3.8-Flash-Next-Q2_0-gfx906",
"backend": "hip",
"gpu": 0,
"parallel": 1,
"drafting": "off",
"args": [
"--pack", "/path/to/pack-q2_0",
"--native", "/path/to/Q2_0/Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00001-of-00002.gguf",
"--ple-gguf", "/path/to/Q2_0/Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00002-of-00002.gguf",
"--expert-profile", "/path/to/Strata/data/expert-profile.bin",
"--expert-cache", "auto",
"--prefill", "auto",
"--max-context", "262144",
"--kv", "int8",
"--kv-resident", "32768",
"--ple-io", "direct",
"--spec", "2"
]
}
パスは自分の環境に置き換える。Pythonの依存パッケージは別の仮想環境にrequirements.txtから導入し、ROCmのライブラリパスも通した状態で、ソースルートからサーバーを起動した。
cd "$STRATA_SRC"
python -m serve.server --engine strata --config /path/to/config.json \
--host 127.0.0.1 --port 8081
最初は--specを付け忘れ、ネイティブpackには--spec T (T >= 2)が必要だとエラーになった。--spec 2を追加すると読み込みが進み、日本語での応答に成功した。ここでの--spec 2はエンジンの指定で、MTP補助モデルを導入したという意味ではない。今回draftingはoffだ。
サーバーはローカルのみにbindしている。稼働確認は以下で行える。
curl -s 127.0.0.1
curl -s 127.0.0.1
8Kから32K、256Kへ
まず8,192トークンで動作確認し、その後32,768へ増やした。約30,417トークンの合成入力の中央に置いた合言葉を正しく回収できた。入力処理は約71.8秒、約424 tokens/sだった。
次に--max-context 262144 --kv int8 --kv-resident 32768で起動。約248,458トークンの入力でも、中央の合言葉を正しく回収できた。入力処理は約640秒、約388 tokens/sだった。
試行
結果
8K設定、短い会話
約35 tokens/sの生成を観測
32K設定、約30K入力
合言葉を正しく回収、入力約424 tokens/s
256K設定、約248K入力
合言葉を正しく回収、入力約388 tokens/s
長文は同じ日本語の文章を繰り返した合成データで、検索課題も一つだ。複雑な複数文書の理解、長時間稼働、異なるコンテンツで同じ精度・速度が出ることまでは確認していない。短い出力6トークンの生成速度を、長文回答の安定速度として扱うこともしない。
256K起動時は専門家キャッシュが8,235個、空きVRAMは約790MiBだった。GPUがデスクトップ表示も担当しているので余裕は小さめだ。--expert-cache autoやVRAM予約量の調整が必要になる場合がある。Strataのメモリ設定を確認してほしい。
また、今回のStrataの既定設定ではSTRATA_ROUTE_TAIL_SKIP=7が有効で、キャッシュにない専門家の一部を省略する。ログにも回答が少し変わり得ると表示される。7,390トークンの合成入力で1回ずつ試したところ、生成は有効時36.7 tokens/s、無効時33.7 tokens/sで、回答文も少し変わった。各1回の短い応答なので、この差を一般化したり品質差を判断したりはできない。
OpenCodeにつなぐ
StrataのOpenAI互換APIを使う。OpenCode v1.19.5をユーザー領域に導入し、~/.config/opencode/opencode.jsonにローカルプロバイダーを追加した。設定形式はOpenCodeのプロバイダー設定を参照した。
{
"$schema": "opencode.ai",
"model": "strata/Qwen3.8-Flash-Next-Q2_0-gfx906",
"provider": {
"strata": {
"npm": "@ai-sdk/openai-compatible",
"name": "Strata (local Radeon VII)",
"options": {
"baseURL": "127.0.0.1",
"apiKey": "local",
"timeout": 1200000
},
"models": {
"Qwen3.8-Flash-Next-Q2_0-gfx906": {
"name": "Qwen3.8 Flash Next Q2_0 (local)",
"tool_call": true,
"limit": { "context": 262144, "output": 8192 }
}
}
}
}
}
モデルIDは/v1/modelsが返すIDと一致させる。既存設定がある場合は丸ごと置き換えず、必要なプロバイダーを統合してほしい。apiKeyはこのローカルの無認証サーバーに渡すダミー値だ。
opencode models strata
opencode
簡単な接続確認として、OpenCodeにテストファイルを読むよう依頼したところ、readツールを呼び出し、ファイル内の合言葉を正しく返した。コードの変更や複雑なエージェント作業の品質までは、この時点では検証していない。
振り返り
Radeon VIIの16GB VRAMだけにモデル全体を載せるのではなく、GPUの専門家キャッシュ、RAM、SSDを組み合わせることで、Flash-Nextの推論までたどり着けた。gfx906は標準的な導入経路から外れているため、GPU認識だけで成功とせず、HIP、行列演算、エンジン、実モデル、ツール呼び出しの順に確かめるのが役立った。
一方、今回はnightly ROCm、ソースビルド、互換ヘッダーへの小さな修正を使っている。Qwen3.8-27BはStrataではなくllama.cppで動かしており、Strataで同じ27Bが動くという話ではない。まずは手元で動いた版と設定を記録し、更新するときは一つずつ比較していこうと思う。
この記事について
実はこの記事も、実際に環境を構築したChatGPTのセッションで書いた。会話の中で調べたこと、実行したコマンド、ログに残った測定結果をもとにChatGPTが下書きした。OpenCodeではファイル読み取りツールの呼び出しまで確認したが、複雑なコード修正の成功率はまだ測っていない。これから、この環境でどこまでできるか試していきたい。
ROCm 최신은 Core SDK 10.1.0이다.
릴리스가 ROCm/TheRock 태그로 옮겨갔고(therock-10.1, 10/6 KST), ROCm/ROCm GitHub 릴리스 페이지는 7.14.0에서 멈춰 있다. 거기만 보면 버전을 잘못 알기 쉽다.
github.com/ROCm/TheRock/releases
13 views
FeatherAttn: Attention backward kernel on Strix Halo, 4x fast as FlashAttn
woct0rdho/feather-attn
On RDNA3/3.5 GPUs, FlashAttn wires to Triton kernels in AITER (AMD's kernel collection). The fwd kernel is fast enough, but it seems no one has optimized the bwd kernel. I've made a bwd kernel with a better designed pipeline. When tested on Strix Halo, it's 4x as fast for head dim 128, and 1.3x as fast for head dim 64.
This is especially relevant because most recent models use head dim 128, including Anima, Krea 2, MOSS, H3 (while SDXL uses head dim 64). I've tested that it makes Anima LoRA training 40% faster.
This kernel does not use int8/fp8 quantization or sparsity. It's still an open question how to make use of them on RDNA3/3.5 .
By the way, besides attn bwd kernel, another bottleneck of LoRA training on Strix Halo is matmul bwd kernel. You may see the discussion in TheRock#5314 and install my forked rocm-libraries for faster GEMM NT.
8 points · 3 comments
FeatherAttn: Attention backward kernel on Strix Halo, 4x fast as FlashAttn
woct0rdho/feather-attn
On RDNA3/3.5 GPUs, FlashAttn wires to Triton kernels in AITER (AMD's kernel collection). The fwd kernel is fast enough, but it seems no one has optimized the bwd kernel. I've made a bwd kernel with a better designed pipeline. When tested on Strix Halo, it's 4x as fast for head dim 128, and 1.3x as fast for head dim 64.
This is especially relevant because most recent models use head dim 128, including Anima, Krea 2, MOSS, H3 (while SDXL uses head dim 64). I've tested that it makes Anima LoRA training 40% faster.
This kernel does not use int8/fp8 quantization or sparsity. It's still an open question how to make use of them on RDNA3/3.5 .
By the way, besides attn bwd kernel, another bottleneck of LoRA training on Strix Halo is matmul bwd kernel. You may see the discussion in TheRock#5314 and install my forked rocm-libraries for faster GEMM NT.
20 points · 4 comments
Strata検証②：EVO-X2（Ryzen AI Max+ 395）で256K コンテクストが安定動作！Qwen3.8-Flash-Next × OpenCode でAIエージェント開発
Strata検証②：EVO-X2（Ryzen AI Max+ 395）で256K コンテクストが安定動作！ Qwen3.8-Flash-Next × OpenCode でAIエージェント開発
はじめに
前回の記事「EVO-X2（Ryzen AI Max+ 395）＋ Vulkan で Qwen3.8-Flash-Next が動いた！ でもコンテクスト長を大きくすると不安定！残念！」では、llama.cpp（Vulkan）で Qwen3.8-Flash-Next を 256K コンテクストで動かしたところ、10 万トークン前後の生成中に EVO-X2 がリセットしました。そのため、Flash-Next の常用は DGX Spark に移しました。
その直後、Qwen3.8-Flash-Next 専用の推論エンジン Strata が v0.1.40 で Ryzen AI Max+ 395（Strix Halo）に正式対応しました。そこで、同じ EVO-X2、同じ BIOS 設定とメモリの配分のまま、エンジンだけを Strata に替えて 256K を試しました。
結論
1. Strata なら、EVO-X2 で 256K（262,144 トークン）のコンテクストが動いた。 速度などの測定も含めて約 10.6 時間動かし続け（うち GPU 使用率 95% 以上が約 8.2 時間）、リセットは 0 回でした。前回リセットが起きた「10 万トークン前後で生成中」の状況も通過し、最大コンテクストは約 23.7 万トークンでした。
2. ハーネスの OpenCode と組み合わせて、256K のまま実際のエージェントとして使えた。 OpenCode V1（1.18.35）と V2（2.0.24）の両方で、ツールの呼び出し、思考、コンテクストの要約（compaction）、2 本同時の実行（Strata の "parallel": 2）まで確かめました。
3. 速度も上がった。 1 件ずつの出力速度は、MTP（モデルに付いている投機デコード用の層）ありで 43.6 tok/s、MTP なしで 28.5 tok/s です（前回の llama.cpp は 20.4 tok/s）。長いプロンプトの読み込みは、前回の約 3〜4.5 倍の速さでした。
4. 日本語の品質は変わらない。 JamC-QA（300 問）で 66.0%（前回の llama.cpp は 66.67%）。
5. ただし、前回のリセットの原因は分かっていない。 「Strata にしたから直った」とは言えません。くわしくは「前回のリセットについて」の節に書きました。
構成
!図 0 構成図
| 項目 | 内容 |
|---|---|
| 本体 | GMKtec EVO-X2（Ryzen AI Max+ 395、メモリ 128GB） |
| GPU | Radeon 8060S Graphics（Ryzen AI Max+ 395 の内蔵 GPU） |
| メモリの配分 | 前回と同じ。BIOS で VRAM に 96GB、残りがシステム RAM（OS から使えるのは約 30GB） |
| BIOS | 1.09（2025-09-13）。GMKtec の公開ページに新しい BIOS は見当たらなかった |
| OS／カーネル | Ubuntu 26.04 LTS / Linux 7.0.0-38（前回のリセットは 7.0.0-34 のとき） |
| 推論エンジン | Strata v0.1.40.2 |
| ROCm | AMD TheRock 7.14.1（Ryzen AI Max+ 395 向け）。ユーザー領域に展開し、システムの ROCm には触れていない |
| モデル | Qwen3.8-Flash-Next（unsloth の GGUF、UD-IQ4_XS）。前回と同じファイル |
| クライアント | OpenCode とベンチマークは、EVO-X2 とは別の PC から LAN 越しにつないだ |
推論をする機械の上では推論だけを動かし、OpenCode などのハーネスは別の PC に置いています。推論機が固まったときに、記録やクライアントを道連れにしないためです。監視の記録を別の PC にも送っているのは、前回のリセットで直前の記録ごと消えてしまった反省からです。
Strata とは
Strata は Qwen3.8-Flash-Next 専用の推論エンジンです。MoE（Mixture of Experts）モデルのエキスパートのうち、よく使うものを GPU に、それ以外を CPU のメモリに置くことで、12GB 程度のゲーミング GPU でも 125B のモデルを動かせることで話題になりました。OpenAI 互換と Anthropic 互換の API を出すので、OpenCode などからそのままつなげます。
v0.1.40 で Ryzen AI Max+ 395 に対応し、ROCm（HIP）で動くようになりました。前回の llama.cpp は Vulkan で動かしていたので、GPU を動かす仕組みそのものが違うことになります。
ちなみに、Strata の開発者のマシンも Ryzen AI Max+ 395・128GB・Ubuntu 26.04 で、EVO-X2 とよく似ています。ただしメモリの配分は逆で、開発者は GPU のメモリ割り当てを小さくし、システム RAM 側を広げる構成です。
動かすときのコツ 1：EVO-X2 の配分では、既定のままだと黙って落ちる
Strata の既定（--expert-cache auto）は、全エキスパート（約 59.5GB）をシステム RAM 側に確保しようとします。当ラボの EVO-X2 は BIOS で VRAM に 96GB を割り当てているため、システム RAM は約 30GB しかなく、確保できずにエラー表示なしでエンジンが終了しました。ログの最後に残っていたのは次の 1 行です。
ArenaExpertSource: the arena could not be reserved (59523276800 B)
次の 2 つを明示すると解決しました。
--mmap-experts --expert-cache 24576
--mmap-experts は GGUF をファイルのまま対応づけて読む指定、--expert-cache 24576 は全エキスパート分の枠を取る指定です。これで全 24,576 エキスパート（約 55.4GiB）が、BIOS で割り当てた 96GB の VRAM に載りました。起動は約 29 秒、読み込み後の VRAM の空きは約 27.7GiB（256K・KV キャッシュ int8・MTP あり・1 本）、システム RAM の使用は約 5.9GB でした。
前回の llama.cpp では、PLE（約 29GB の n-gram の表）を SSD に置く工夫をしましたが、今回はモデル全体が VRAM に収まっています。「BIOS で VRAM を大きく取る」配分だからこそ取れるやり方です。
実際の手順
実際の手順です（v0.1.40.2、Strata の公式ドキュメント docs/STRIX_HALO.md に沿ったもの）。ROCm は AMD の TheRock ビルド（7.14.1、Ryzen AI Max+ 395 向け）を自分のフォルダに展開して使います。root は要らず、システムに入っている ROCm には触れません。
# 1. ROCm（TheRock 7.14.1、Ryzen AI Max+ 395 (gfx1151) 向け）をユーザー領域に展開
mkdir -p ~/rocm && cd ~/rocm
curl -LO https://repo.amd.com/rocm/tarball-multi-arch/therock-dist-linux-gfx1151-7.14.1.tar.gz
mkdir install && tar -xzf therock-dist-linux-gfx1151-7.14.1.tar.gz -C install
export SDK=$HOME/rocm/install
export ROCM_PATH=$SDK HIP_PATH=$SDK HIP_PLATFORM=amd
export PATH=$SDK/bin:$SDK/lib/llvm/bin:$PATH
export LD_LIBRARY_PATH=$SDK/lib:$SDK/lib/rocm_sysdeps/lib:$SDK/lib/llvm/lib:${LD_LIBRARY_PATH:-}
# 2. Strata をビルド（公式手順どおり。EVO-X2 では約 4 分）
git clone https://github.com/Niko1221/Strata.git && cd Strata && git checkout v0.1.40.2
cmake -S . -B build-halo -G Ninja -DCMAKE_BUILD_TYPE=Release \
-DCMAKE_HIP_ARCHITECTURES=gfx1151 \
-DCMAKE_HIP_COMPILER=$SDK/lib/llvm/bin/clang++ -DCMAKE_HIP_COMPILER_ROCM_ROOT=$SDK \
"-DCMAKE_PREFIX_PATH=$SDK;$SDK/lib/rocm_sysdeps;$SDK/lib/llvm" \
"-DCMAKE_HIP_FLAGS=--rocm-path=$SDK --rocm-device-lib-path=$SDK/lib/llvm/amdgcn/bitcode"
cmake --build build-halo --target strata -j 8
# 3. GGUF（unsloth UD-IQ4_XS）を Strata の形式に変換（約 30 秒）
.venv/bin/python tools/iq_pack.py \
--gguf <モデル置き場>/Qwen3.8-Flash-Next-UD-IQ4_XS-00001-of-00003.gguf \
--out packs/ud-iq4_xs --compat-bf16
--compat-bf16 は必須です。この GGUF が Q8_0 で持っている小さな射影を、Strata は BF16 で読むため、変換時に丸めが入ります（この丸めが、llama.cpp と出力が完全には一致しない理由です）。MTP のドラフトは、公式ドキュメント（docs/ORCA.md の準備の節）のとおりに作って mtp/rt に置きます。
サーバの設定は strata-ud-iq4_xs.json のような JSON に書きます。EVO-X2 で使ったものは次のとおりです。
{
"exe": "build-halo/strata",
"args": [
"--pack", "packs/ud-iq4_xs",
"--native", "<モデル置き場>/Qwen3.8-Flash-Next-UD-IQ4_XS-00001-of-00003.gguf",
"--expert-profile", "data/expert-profile.bin",
"--expert-cache", "24576", "--mmap-experts",
"--prefill", "16384", "--max-context", "262144", "--kv", "int8",
"--spec", "4", "--spec-min-p", "0.5",
"--mtp", "mtp/rt", "--lookup-chain", "3", "--mtp-q4", "all"
],
"cwd": ".",
"tokenizer": "packs/ud-iq4_xs/tokenizer",
"model_name": "qwen3.8-flash-next-ud-iq4_xs",
"log": "strata-ud-iq4_xs.log",
"fit_max_tokens": true
}
- --expert-cache 24576 --mmap-experts が、上で書いた EVO-X2 向けの指定です（公式の例は "--expert-cache", "auto"）
- --prefill 16384 にしているのは、後述の hipBLASLt の調整表がこの大きさのプロンプトまで含んでいるためです
- 2 本同時に処理するときは "parallel": 2 を足します（1 本あたり VRAM を約 3.65GiB 多く使い、読み込み後の空きは約 20.2GiB でした）
起動のときは、公式ドキュメントの「速い構成」の環境変数と、hipBLASLt の調整表を渡します。調整表は、対応したカードなら setup が渡してくれますが、Ryzen AI Max+ 395 では自分で指定します（指定しないと、プロンプトの読み込みが遅くなります）。
export STRATA_HIPBLASLT_TUNING=$PWD/tools/hip/gfx1151-hipblaslt-100401.txt
export STRATA_PF_FUSED=1 STRATA_PF_GEMM=1 STRATA_HC_UPMIX=1 STRATA_PA_FAST=1
export STRATA_HIP_WMMA=1 STRATA_SELECT_WMMA=1 STRATA_HC_Q8=1
export STRATA_PF_SWITCH_MIN_T=4096
.venv/bin/python -m serve.server --engine strata --config strata-ud-iq4_xs.json --host 0.0.0.0 --port 8095
起動ログに prefill gemm: hipBLASLt tuning enabled (… gfx1151, version 100401) と出れば、調整表が効いています。環境変数の 7 つは、既定より速い代わりに丸め方が変わるもので、公式は KL（出力分布の差）を測ったうえで「速い構成」として案内しています。Ryzen AI Max+ 395 向けの既定の 18 個のスイッチは、エンジンが自動で有効にします。
動かすときのコツ 2：OpenCode と組み合わせるなら "fit_max_tokens": true
Strata は、「プロンプト＋max_tokens」がコンテクスト長を超えるリクエストを HTTP 400 で断ります。OpenCode（V1）は毎回 max_tokens として約 32,000 を送るので、会話が約 23 万トークンに達したところで、compaction が始まる前に止まってしまいました。
Strata の設定に "fit_max_tokens": true（残りの余地に合わせて max_tokens を縮める）を入れると解消しました。
256K の安定性
5 つの試験（合計 約 6.8 時間）を行いました。速度などの測定も含めると、約 10.6 時間動かし続けています。どの試験も、前回リセットが起きた 8.8 万〜10.4 万トークンの帯を通り越して、20 万トークン前後まで到達しています。
!図 1 各試験で到達したコンテクスト
| 試験 | やり方 | 時間 | 最大コンテクスト | 結果 |
|---|---|---|---|---|
| API で会話を育てる | 毎ターン資料を足し、思考あり（high）で長い回答を生成 | 55 分（31 ターン、生成 約 13.9 万トークン） | 約 23.1 万（最終 23.7 万） | 異常なし。前回リセットした 8.8 万〜10.4 万トークンを、GPU 使用率 99〜100% で生成しながら約 32 分かけて通過 |
| OpenCode V1 で実使用 | 同じセッションに依頼を重ねる（調査、レビュー、編集、要約） | 約 1.5 時間 | 約 23.5 万（compaction で 9.2 万へ） | 異常なし |
| OpenCode V2 で実使用 | 同じ依頼 | 84 分 | 約 19 万（compaction 3 回） | 異常なし |
| 2 本同時（parallel 2、V1 と V2） | 別々のセッションを同時に | 約 97 分 | 各 約 18 万 | 異常なし |
| 2 本同時（parallel 2、API で最大まで） | 会話を 2 本同時に育てる | 82 分 | 21.7 万と 23.2 万を同時に保持 | 異常なし。毎ターン全文を読み直す厳しい条件で、21.4 万トークンの読み込みに 7.6 分かかる場面もあった |
- 通算: 経過 約 10.6 時間のうち、GPU 使用率 95% 以上が約 8.2 時間。リセット 0 回、サーバの停止 0 回
- 監視: 1 秒ごとに GPU の状態・温度・電力・メモリを記録し、別の PC にも同時に送りました（前回、リセットで記録ごと消えた反省から）。記録の欠けはありませんでした
- 温度と電力: GPU 温度は最大 94℃、電力は最大 約 117W
コンテクストが長くなっても、出力速度は大きく落ちませんでした（API で会話を育てる試験、思考あり、MTP あり）。10 万トークンを超えたところで約 50 tok/s から約 43 tok/s に下がりますが、その先は 40 tok/s 台を保っています。
!図 4 コンテクスト長と出力速度
| コンテクスト | 出力速度（平均） |
|---|---:|
| 5 万未満 | 約 50 tok/s |
| 5〜10 万 | 約 52 tok/s |
| 10〜15 万 | 約 43 tok/s |
| 15〜20 万 | 約 42 tok/s |
| 20 万以上 | 約 43 tok/s |
2 本同時は、OpenCode を 2 つ並べて使う（たとえば調査用と実装用）個人利用を想定したものです。Strata の設定で "parallel": 2 にして測りました。1 本あたりの出力速度は、1 本だけのときの 45〜50 tok/s から 15〜21 tok/s に下がり、合計は 1 本のときとほぼ同じでした。2 本同時に使うと速くなるわけではなく、待たずに済むようになる、という使い方です。
OpenCode V1 と V2
| 項目 | V1（1.18.35） | V2（2.0.24） |
|---|---|---|
| 256K での実使用 | ○（最大 約 23.5 万） | ○（最大 約 19 万、compaction 3 回） |
| ツールの呼び出し | read / grep / glob / edit / write / bash が動作 | read / grep / glob / edit / write / shell / execute が動作。形式の崩れなし |
| 設定ファイル | opencode.json | opencode.jsonc（キーの名前や構造が V1 と別物） |
- compaction は時間がかかる。 要約のリクエストは通常と先頭が違うため、コンテクスト全体を初めから読み直します。23.5 万トークンで 221 秒、15.8 万トークンで約 2 分かかり、その間エージェントは無音になります。OpenCode のタイムアウトは 900 秒にしました
- ふだんはキャッシュがよく効く。 実使用では、入力の 92〜96% がキャッシュから読まれていました（2 本同時でも効いていました）
- V2 はモデルを固定しておくこと。 V2 は -m でモデルを指定しないと、外部の既定のモデルを選びます。ローカルの LLM だけで使うなら、設定ファイルでモデルを固定するか、-m を必ず付けてください
実際の開発に使ってみて
この構成は、その後ラボの実際の開発の推論先として使っています（Strata はその間に版を上げています）。たとえば、自分のデータを整理するエージェントの開発です。ブラウザを人と同じ速さで操作して、自分のデータ（画像とその情報）を集め、一覧で見られる Web 画面まで、OpenCode に作らせました。
データを外に出さずに、調査・操作・コーディングまでをローカルの LLM で回せる。外に出せないデータを扱う開発で、これが現実的な選択肢になってきたと感じています。
速度
1 件ずつ、思考なしで測りました。MTP の有無で 2 本に分けているのは、前回の llama.cpp と DGX Spark が MTP なしで測ったものだからです。
!図 2 出力速度
| 構成 | 出力速度（3 回の中央値） |
|---|---:|
| EVO-X2 ＋ Strata（MTP あり） | 43.6 tok/s |
| EVO-X2 ＋ Strata（MTP なし） | 28.5 tok/s |
| 前回：EVO-X2 ＋ llama.cpp（Vulkan、MTP なし） | 20.4 tok/s |
| 参考：DGX Spark ＋ SGLang（NVFP4、MTP なし） | 27.0 tok/s |
MTP なしどうしで比べても、Strata は llama.cpp の約 1.4 倍でした。
長いプロンプトを初めて読み込む（キャッシュなし）時間です。エージェントとして使うときは、資料やコードをまとめて読ませる場面が多いので、ここが速いかどうかが体感に効きます。
!図 3 長いプロンプトの初回読み込み時間
| 長さ | Strata（MTP あり） | Strata（MTP なし） | 前回：llama.cpp | 参考：DGX Spark |
|---:|---:|---:|---:|---:|
| 約 2,500 トークン | 3.1 秒 | 3.1 秒 | 10.5 秒 | 2.9 秒 |
| 約 22,800 トークン | 19.9 秒 | 19.7 秒 | 約 65 秒（推定） | 14.3 秒 |
| 約 63,300 トークン | 53.1 秒 | 53.5 秒 | 約 218 秒（推定） | 40.6 秒 |
- Strata の読み込みは、llama.cpp の約 3〜4.5 倍の速さで、DGX Spark の 8 割弱でした
- 同じプロンプトの 2 回目（キャッシュあり）は 0.05〜0.36 秒です
- llama.cpp の約 22,800 と約 63,300 トークンの値は、一部キャッシュが効いていた実測値をキャッシュなしに換算した推定値です（前回の記事と同じ値）
日本語の品質
| 構成 | JamC-QA（300 問） |
|---|---:|
| Strata（MTP あり） | 66.0% |
| Strata（MTP なし） | 66.0% |
| 前回：llama.cpp（Vulkan） | 66.67% |
300 問で 2 問の差です。Strata の形式に変換するときの丸め（後述）などの範囲とみています。
前回のリセットについて
「エンジンを替えたらリセットしなくなった」のは事実ですが、原因は分かっていません。前回のリセットの記録を読み直したので、分かったことを書いておきます。
記録から分かったこと
- 前回、予期しない停止は 2 回ありました。どちらも llama.cpp（Vulkan）が、約 10.4 万トークンと約 8.85 万トークンのコンテクストで生成している最中でした
- 2 回とも、GPU のエラー（計算キューのタイムアウト、Device lost、GPU のリセット、ページフォルトなど）はカーネルのログに残っていませんでした
- 2 回目は 46 秒後に自分で再起動し、次の起動のカーネルのログに、前回のリセット理由として「internal CPU shutdown event occurred」と記録されていました
最後の 1 行は、Linux 6.16 から入った機能です。AMD の CPU が持っている「前回のリセット理由」のレジスタを、起動時に読んで表示してくれます。
x86/amd: Previous system reset reason [0x00800800]: internal CPU shutdown event occurred
つまり、GPU のエラーではなく、CPU 側の事象としてシステム全体がリセットされたことになります。
「Vulkan だから落ちた」とは言えない理由
llama.cpp の Vulkan で、長いコンテクストのときに GPU の計算キューがタイムアウトする、という報告は複数あります（#21724、discussion #27154）。最初はこれを疑いました。ただし、これらの報告の多くはプロセスの異常終了で済んでいて、今回のようなシステム全体のリセットとは症状が違います。GPU のエラーも記録されていませんでした。
また、次の点からも「Strata だから直った」とは言えません。
- 電力は Strata のほうが高い。 前回のリセットの直前は約 80W、Strata の試験では最大 約 117W でした。「負荷が軽いから落ちなかった」という説明は成り立ちません
- カーネルも変わっている。 前回のリセットは 7.0.0-34 のときで、今回の試験は 7.0.0-38 です
- 回数が少ない。 偶然を否定できるほどの回数ではありません
Ryzen AI Max+ 395 では、似た報告が多い
Ryzen AI Max+ 395 を積んだ機械では、高負荷の推論中にシステムごと落ちる、という報告がいくつもの機種から出ています。記録される理由は報告ごとに違い、「ACPI の電源状態の遷移」や「訂正できないエラーによるデータファブリックの sync flood」などがあります（ROG Flow Z13 の報告、Beelink GTR9 Pro の報告、GPU と NPU を同時に使ったときの報告）。機種、ファームウェア、負荷のかけ方によって出たり出なかったりするようです。
同じ症状に当たった人は、次の起動でカーネルのログの「Previous system reset reason」を見てください。
journalctl -k -b | grep "Previous system reset reason"
GPU 側の問題か、本体側の問題かを切り分ける手がかりになります。
測定の条件と注意
- MTP の扱いが構成で違う。 前回の llama.cpp と DGX Spark は MTP なしです。公平に比べるなら、Strata の MTP なし（28.5 tok/s）を使ってください
- Strata の「MTP なし」は、投機デコードを完全には切っていない。 今回の形式のモデルでは最低限の投機デコードが必須で、簡易なドラフト（受理率 約 49%）だけが残る構成です
- モデルの変換。 Strata の形式に変換するとき、一部の値を BF16 に丸めています（公式の注記あり）。そのため llama.cpp と出力が完全には一致しません
- 思考の止め方。 前回は llama.cpp の --reasoning off、今回はリクエストの reasoning_effort: none です
- 速度の測り方。 前回と同じベンチマークのクライアント（ShareGPT の 200 プロンプト）を使いましたが、Strata に対してはクライアント側の一部の値が正しく出ませんでした（TTFT が数ミリ秒になる、MTP なしで出力を少なく数えるなど）。そのため、出力速度は Strata のサーバのログから計算しました
- 回数。 出力速度は 3 回の中央値、長いプロンプトは各 1 回です
- 版。 数字はすべて v0.1.40.2 での測定です。Strata はその後も版を上げていて、本記事の公開時点では 0.2 系が出ています
未解決の点
1. 一晩以上の連続運転は、この版ではしていない。 1 回の連続運転は最長 約 97 分、通算の高負荷は約 8.2 時間です。その後の 0.2.4.1 では、OpenCode V2 で 6 時間を超える連続運転をして異常がないことを確かめました。これは次の記事で書きます
2. 出力が 1 回だけ壊れた。 OpenCode V1 で約 16.4 万トークンのとき、最後の応答が壊れた短い文で止まりました。約 780 ステップ中この 1 回だけで、原因（モデル、量子化、エンジン、MTP）は切り分けていません
3. GPU 温度が最大 94℃。 長時間の高負荷では冷却に注意が必要です
まとめ
前回「残念！」で終わった EVO-X2 の 256K コンテクストは、Strata で動きました。速度も上がり、OpenCode と組み合わせて実際のエージェントとして使えています。
一方で、前回のリセットの原因はまだ分かっていません。Ryzen AI Max+ 395 で長時間の高負荷をかける人は、監視の記録を別の機械にも送り、リセットしたら「Previous system reset reason」を確かめることをおすすめします。
英語版のリポジトリ（起動の設定、監視・測定のスクリプト、生データ）: nabe2030/strata-evo-x2-256k
Strata への報告: #4032（Linux の Ryzen AI Max+ 395 での実測）、#4033（VRAM を大きく取った構成で --expert-cache auto が落ちる件）、#4034（ドキュメントの提案）
おわりに — 当ラボについて
筆者は NPO法人 IoTメディアラボラトリーで、IoT 研究とあわせて AI 分野の検証を担当しています。東京大学・本郷キャンパス 大学院工学系研究科の同名のラボで行っていた IoT 研究が独立し、研究機関として活動を続けている組織です。
私がやっていることは、本記事のようなローカル LLM の実機検証が中心です。新しいモデルが出るたびに手元のハードで動かし、速度・日本語知識・コーディング・マルチモーダルの 4 軸で測り、動かなければ原因を突き止めるところまでやります。ベンダーの公称値ではなく、実際に自分の機械で何が起きるかを知りたいからです。
そのぶん、計算資源はいくらあっても足りません。もっと GPU が欲しい（笑）。機材の提供、検証のご依頼、共同研究など、お声がけいただけると喜びます。
WIN 5 GPU + NPU crashes: AMD fixes AGESA, need new BIOS?
I have a GPD WIN 5 395 and I heard the system hangs during simultaneous GPU/NPU inference.
This report describes pinning the fabric clock at maximum as a workaround.
AMD says AGESA from 1.0.0.2 to 1.0.0.2b resolved their reproduction and a new firmware fix must be delivered through the manufacturer’s BIOS.
Does anyone know:
- Which AGESA version the current WIN 5 BIOS contains?
- Whether GPD plans a BIOS update incorporating this fix?
- The best official contact for GPD employee who could get this report to the BIOS engineering team?
6 points
MI50 ROCm 10.2 TheRock vs Vulkan Mesa 26.2 llama.cpp benchmarks
I prefer running llama.cpp Vulkan prebuilt binary. I just download the latest version and ready to roll. I finally took the hours necessary to get TheRock latest tarball version of ROCm 10.2 running on dual AMD Radeon Instinct MI50 gfx906 (32gb combined VRAM).
Same models benched in previous post. A mix of Dense and MoE models and quants that better utilize available VRAM.
- llama_bench_Swift-Qwen3.8-27B-Uncensored-MTP.Q6_K.gguf
- llama_bench_Swift-1.5-Qwen3.8-27B-Q6_K.gguf
- llama_bench_Gemma-4-MoonGem-31B.i1-Q6_K.gguf
- llama_bench_gemma-4-31B-it-UD-Q6_K_XL.gguf
- llama_bench_Nemotron-3.5-30B-A3B-Antislop-FTPO.i1-Q5_K_M.gguf
- llama_bench_Laguna-XS-2.1-APEX-I-Balanced.gguf
- llama_bench_Agents-A1-Q4_K_M.gguf
- llama_bench_Qwen3.6-35B-A3B-UD-Q5_K_XL.gguf
- llama_bench_Qwen3.6-35B-A3B-Uncensored-HauhauCS-Aggressive-Q6_K_P.gguf
Here is the backend performance comparison contrasting the native ROCm (v10.2 for gfx906) runtime against your optimized Vulkan (MESA_PPA_26.2) baseline. The data highlights a massive architectural split: ROCm significantly accelerates token generation across the board but suffers high variance and regressions in MoE pre-fills.
Both GPUs are power limited to 145 watts. Build versions used:
llama-b11382 used for Vulkan (prebuilt ubuntu binary)
llama-b11401 used for ROCm (compiled with proper flags)
Architectural Performance Breakdown: Vulkan vs. ROCm
|Model|Size|Params|Test|Vulkan Baseline (t/s)|ROCm 1st Run (t/s)|Performance Delta (%)|
|:-|:-|:-|:-|:-|:-|:-|
|qwen35 27B Q6_K|20.88 GiB|27.32 B|pp512 tg128|149.30 ± 0.15 18.35 ± 0.01|176.86 ± 17.28 20.21 ± 0.65|+18.46% +10.14%|
|qwen35 27B Q6_K|22.21 GiB|27.32 B|pp512 tg128|163.48 ± 0.18 18.52 ± 0.03|183.49 ± 15.75 20.04 ± 0.63|+12.24% +8.21%|
|gemma4 31B Q6_K|23.46 GiB|30.70 B|pp512 tg128|119.98 ± 0.12 15.72 ± 0.02|172.93 ± 0.89 17.26 ± 0.13|+44.13% +9.80%|
|gemma4 31B Q6_K|25.62 GiB|30.70 B|pp512 tg128|133.67 ± 0.24 16.38 ± 0.02|178.59 ± 2.31 17.50 ± 0.09|+33.61% +6.84%|
|nemotron_h_moe 31B|25.18 GiB|32.91 B|pp512 tg128|834.87 ± 1.72 63.82 ± 0.07|736.98 ± 134.77 103.86 ± 0.50|-11.73% +62.74%|
|laguna 30B.A3B Q5_K|22.64 GiB|33.44 B|pp512 tg128|723.34 ± 2.99 58.98 ± 0.28|856.14 ± 29.71 76.81 ± 0.20|+18.36% +30.23%|
|qwen35moe 35B Q4_K|19.70 GiB|34.66 B|pp512 tg128|957.62 ± 3.88 51.50 ± 0.07|834.58 ± 102.56 69.70 ± 0.22|-12.85% +35.34%|
|qwen35moe 35B Q5_K|24.76 GiB|34.66 B|pp512 tg128|909.20 ± 5.71 53.18 ± 0.07|763.93 ± 106.19 68.47 ± 0.34|-15.98% +28.75%|
|qwen35moe 35B Q6_K|28.53 GiB|34.66 B|pp512 tg128|763.05 ± 66.01 52.81 ± 0.04|754.82 ± 70.48 67.29 ± 0.21|-1.08% +27.42%|
Core Insight Strategy & Bottlenecks
1. Token Generation (tg128) Dominance: ROCm dominates pure text generation. The native AMD matrix kernels unleash your MI50 computation potential, unlocking a massive +62.74% boost for Nemotron but taking a small hit on pre-fill -11.73%.
2. Dense Model Pre-fills (pp512): Dense architectures scale cleanly under ROCm. Gemma 4 sees a +33% to +44% processing throughput spike over the Vulkan RADV driver driver bounds. MoE models take a hit with a -15.89% difference with llama_bench_Qwen3.6-35B-A3B-UD-Q5_K_XL.gguf being the happiest with Vulkan backend.
3. The MoE Prompt Processing Delinquency: Notice the massive standard deviations under ROCm for MoE pre-fills (e.g., Qwen3.5MoE Q4_K has an instability block of ± 102.56). ROCm suffers from severe scheduling thrashing when building prompt streams across multiple active experts.
Here is the structured layout with pp512 and tg128 separated into individual columns for a clean side-by-side comparison between the two backend architectures.
Backend Comparison Table (Vulkan vs. ROCm)
|Model|Size|Params|Vulkan pp512 (t/s)|ROCm pp512 (t/s)|Vulkan tg128 (t/s)|ROCm tg128 (t/s)|
|:-|:-|:-|:-|:-|:-|:-|
|qwen35 27B Q6_K|20.88 GiB|27.32 B|149.30 ± 0.15|176.86 ± 17.28|18.35 ± 0.01|20.21 ± 0.65|
|qwen35 27B Q6_K|22.21 GiB|27.32 B|163.48 ± 0.18|183.49 ± 15.75|18.52 ± 0.03|20.04 ± 0.63|
|gemma4 31B Q6_K|23.46 GiB|30.70 B|119.98 ± 0.12|172.93 ± 0.89|15.72 ± 0.02|17.26 ± 0.13|
|gemma4 31B Q6_K|25.62 GiB|30.70 B|133.67 ± 0.24|178.59 ± 2.31|16.38 ± 0.02|17.50 ± 0.09|
|nemotron_h_moe 31B|25.18 GiB|32.91 B|834.87 ± 1.72|736.98 ± 134.77|63.82 ± 0.07|103.86 ± 0.50|
|laguna 30B.A3B Q5_K|22.64 GiB|33.44 B|723.34 ± 2.99|856.14 ± 29.71|58.98 ± 0.28|76.81 ± 0.20|
|qwen35moe 35B Q4_K|19.70 GiB|34.66 B|957.62 ± 3.88|834.58 ± 102.56|51.50 ± 0.07|69.70 ± 0.22|
|qwen35moe 35B Q5_K|24.76 GiB|34.66 B|909.20 ± 5.71|763.93 ± 106.19|53.18 ± 0.07|68.47 ± 0.34|
|qwen35moe 35B Q6_K|28.53 GiB|34.66 B|763.05 ± 66.01|754.82 ± 70.48|52.81 ± 0.04|67.29 ± 0.21|
So yes it's worth the hassle of jumping through hoops to get ROCm working on MI50 setups. At least I have 2 backends working. Next up I'll try some RPC.
preview.redd.it
8 points · 1 comment
Mentions
Radeon VIIとUbuntu 26.04でROCmを動かし、Flash-NextをOpenCodeにつないだ 5年ほど前に組んだデスクトップがある。当時、GPUにはすでに型落ちだったRadeon VIIを5万円くらいで入手して載せた。Stable Diffusionで遊んでいた時期もあったが、飽きてからはすっかり出番がなくなっていた。 最近はローカルLLMが盛り上がっている。とはいえ、パソコンもGPUもクソ高い。新しいGPUを買う前に、あまり期待せず、久しぶりにこのデスクトップに火を入れて環境を組み直してみることにした。 気になったのは、Radeon VIIのROCm公式サポートが切れていることだ。かなり苦戦する覚悟だったが、実際にはChatGPTのGPT-6.1 Terra（medium）に相談しながら、環境の調査から構築、動作確認まで進められた。多少の修正はあったものの、さくさく進む。最終的に、Qwen3.8-Flash-Nextが動き、OpenCodeからローカルモデルへツール付きでアクセスできるようになった。古いデスクトップでvibe codingを試せる環境ができてしまい、だいぶびっくりしている。 この記事は、そのRadeon VIIでROCmを動かし、ローカルLLMをOpenCodeにつなぐまでの記録だ。 結果は、Q2_0量子化版のFlash-Nextで短めの会話の生成が約35 tokens/s。別途、Qwen3.8-27BのQ3量子化版もllama.cppで約21〜22 tokens/s出ている。両者は量子化も実行エンジンも違うため、モデル同士の純粋な比較ではない。 この記事は2026年11月3日時点の実機メモだ。gfx906のROCmはAMDの通常の動作保証対象とは別扱いで、以下は一つの環境で動いた記録だ。アップデート後も同じように動くとは限らない。 環境と構成 項目 今回の環境 GPU Radeon VII、VRAM 16GB、gfx906 CPU Ryzen 5 3600、6コア12スレッド、AVX2 RAM 公称64GB、OSから約60.7GiB OS Ubuntu 26.04.1 LTS カーネル 7.0.0-42-generic、標準のamdgpu/KFD ROCm TheRock nightly 10.3.0a20261102、gfx906向け HIPのバージョン表示 7.18.26434 推論 Strata v0.2.10、gfx906用ソースビルド モデル ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF、Q2_0 コーディングUI OpenCode v1.19.5 Flash-Nextは本体125B、1トークンで有効になるのは約6BのMoEモデルで、別に51BのN-gram参照テーブルがある。Q2_0はパラメーターを削った小型モデルではなく、重みを量子化したものだ。Qwenの説明と量子化版の配布元を参照してほしい。 構成は大きく3層だ。 Ubuntu標準のamdgpu/KFDがGPUを扱う。 gfx906向けROCmのユーザー空間ライブラリとHIPが計算を担当する。 StrataがモデルをGPU・RAM・ストレージに配置し、ローカルAPIを公開する。 コンテナは使わず、ROCmやビルド成果物はホームディレクトリ下に分離した。システムのドライバやAPTリポジトリは変更していない。 まずamdgpuとKFDを確認する 画面が映っているだけではGPU計算が動くとは限らない。まずカーネルドライバとデバイスノードを確認した。 uname -a lspci -nnk | grep -A 4 -i 'vga|display|3d' ls -l /dev/kfd /dev/dri getfacl /dev/kfd /dev/dri/renderD128 このPCでは1002:66afのRadeon VIIをamdgpuが使用し、/dev/kfdとrenderD128にはログインユーザー向けのACLが付いていた。アクセス権の設定方法は環境によって異なる。/dev/dri/renderD128も固定とは限らないので、自分の環境で確認しよう。 gfx906向けROCmをユーザー領域へ TheRockの対応表ではgfx906はビルド対象だが、Sanity Tested・Release Readyにはチェックがない。配布されていることと、実機で動作検証されていることは別だ。 今回はAMDのnightlyから、同じ日付のSDK・ライブラリ・開発ツール・gfx906向けデバイスコードを仮想環境に揃えた。uvは別途用意した。以下の版は実験時点の固定値で、nightlyの保持期間や配布状態によっては取得できなくなる可能性がある。 uv venv ~/rocm-gfx906/venv uv pip install --python ~/rocm-gfx906/venv/bin/python \ --index-url nightly.repo.amd.com \ 'rocm[libraries,devel,device-gfx906]==10.3.0a20261102' source ~/rocm-gfx906/venv/bin/activate rocm-sdk init export ROCM_PATH="$(rocm-sdk path --root)" export HIP_PATH="$ROCM_PATH" export PATH="$ROCM_PATH/bin:$PATH" export LD_LIBRARY_PATH="$ROCM_PATH/lib:$ROCM_PATH/lib64${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" rocminfo これは配置先を分かりやすく書き直した再現例だ。実機では~/.local/share/rocm-gfx906/venvを使った。Ubuntu標準のPythonではvenvのensurepipがなかったため、uvで作成した。 rocminfoではRadeon VIIがgfx906:sramecc+:xnack-として見えた。さらにHIPで4,096要素のベクトル加算、rocBLASで2×2と512×512の行列積を検証した。SDK付属テストも27項目が通過した。 HIP compiled vector addition: PASS (4096 elements) rocBLAS SGEMM 512x512 identity product: PASS Ran 27 tests ... OK amdflangは最初、libquadmath.so.0不足で起動しなかった。今回はlibquadmath0をapt downloadしてユーザー領域に展開し、ライブラリパスを通して解決した。同様にC++開発ヘッダーとGCCの開発ライブラリもユーザー領域へ展開した。一般的な環境では、必要な開発パッケージをディストリビューションの方法で揃える方が簡単だ。無関係なUbuntu版のROCmパッケージを混ぜないようにした。 最初にllama.cppでGPU推論を確認 Strataに進む前に、llama.cppをHIP有効・gfx906指定でビルドした。使用したllama.cppのコミットはf9448cddc48c129a4abc7ff9c2a84c2b4c101c61だ。ビルド方法の基本はllama.cppの公式ドキュメントにある。 llama-cli --list-devicesにROCm0: AMD Radeon VIIが表示され、Qwen2.5-0.5BのGGUFで短い生成に成功。次にQwen3.8-27BのUD-Q3_K_XLも試した。モデルファイルは約13.2GBだ。 llama-server -m /path/to/Qwen3.8-27B-UD-Q3_K_XL.gguf \ -ngl 99 -c 4096 -np 1 --host 127.0.0.1 --port 8082 日本語の短い質問に正常回答し、会話では約21〜22 tokens/s。後にログを詳細化して、66/66層がGPUに配置されていることも確認した。別の8K設定では、約7,390トークンの合成入力を約240 tokens/sで処理した。いずれも今回のハードウェアと設定での観測だ。 Strataをgfx906向けにビルドする StrataのAMD向けドキュメントには、gfx906をwave64互換レイヤーでビルドする手順がある。配布済みのRX 6000/7000/9000向けバイナリをRadeon VIIに流用するのではなく、専用のSTRATA_HIP_GFX906経路を使う。 Strata v0.2.10のソースと、プロジェクトが指定するllama.cppコミット3cf03257f219afbe7334045ff7c6a06ac68c627dを用意した。システムにC/C++開発環境、CMake、Ninja、ROCmの開発ツールが必要だ。 STRATA_SRCとLLAMA_SRCは、それぞれ取得・展開したソースディレクトリ cmake -S "$STRATA_SRC" -B "$STRATA_SRC/build-906" -G Ninja \ -DSTRATA_HIP_GFX906=ON \ -DCMAKE_HIP_ARCHITECTURES=gfx906 \ -DSTRATA_GGML_DIR="$LLAMA_SRC" \ -DCMAKE_C_COMPILER="$ROCM_PATH/bin/amdclang" \ -DCMAKE_CXX_COMPILER="$ROCM_PATH/bin/amdclang++" \ -DCMAKE_PREFIX_PATH="$ROCM_PATH" \ -DCMAKE_BUILD_TYPE=Release cmake --build "$STRATA_SRC/build-906" --target strata strata-device -j 6 実機ではシステムにGCC開発環境がなかったため、APTからダウンロードして展開したlibstdc++-15-devとlibgcc-15-devを指す--gcc-toolchainをC/C++/HIPの各フラグに追加した。上は開発環境が通常の場所にある場合の基本形だ。 ビルド時に互換レイヤーの定義を追加した v0.2.10をそのままビルドすると、cudaDeviceGetStreamPriorityRangeとcudaStreamCreateWithPriorityが未定義になった。gfx906互換ヘッダーinclude/strata/platform/hip_compat/strata_hip.hに、対応するHIP APIへのマッピングを追加して再ビルドした。 #define cudaStreamCreateWithFlags hipStreamCreateWithFlags +#define cudaStreamCreateWithPriority hipStreamCreateWithPriority +#define cudaDeviceGetStreamPriorityRange hipDeviceGetStreamPriorityRange これは今回の版とビルドで必要だったローカル修正だ。将来の版では修正済みかもしれない。ビルド後、strata-device --selftest、gfx906向けGDN演算テスト、専門家演算の比較テストで失敗がないことを確認した。これらは全機能や長時間の安定性を保証するものではない。 Flash-Nextを準備して起動 モデルはISTA-DASLabのGGUFのQ2_0を選んだ。37.6GBの本体と28.8GBの参照テーブル、合わせて約66.4GBだ。今回使用したリビジョンはed59f92082b1e93c0e96d60a8b11aab089b52f09。両ファイルを配布元のSHA-256と照合した。 Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00001-of-00002.gguf SHA256 69820c02ec7d0b45ef2ebb19d6620299db749fe2aded7f39f93c6b88b199b720 Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00002-of-00002.gguf SHA256 316b46f3a2dbd68c900f43136ab9449f9dcc3725dfd8c794847c204bc161e113 転送が途中で停滞して再始動した際、最初の不完全ファイルを引き継がず、新しいファイルとして再取得した。完成後に古い不完全ファイル約17.6GBを削除した。大きなモデルでは、途中再開の動作と空き容量を確認したいところだ。 取得後、Strataに付属するtools/iq_pack.pyでインデックス、dense重み、トークナイザーを準備する。STRATA_GGUF_PYは上で用意した固定コミットのllama.cppにあるgguf-pyを指す。 export STRATA_GGUF_PY="$LLAMA_SRC/gguf-py" python "$STRATA_SRC/tools/iq_pack.py" \ --gguf /path/to/Q2_0/Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00001-of-00002.gguf \ --out /path/to/pack-q2_0 サーバー設定には、ビルドしたstrata、pack、2つのGGUF、付属のdata/expert-profile.binを指定する。主要なエンジン引数は次のとおりだ。 { "exe": "/path/to/build-906/strata", "cwd": "/path/to/Strata", "tokenizer": "/path/to/pack-q2_0/tokenizer", "model_name": "Qwen3.8-Flash-Next-Q2_0-gfx906", "backend": "hip", "gpu": 0, "parallel": 1, "drafting": "off", "args": [ "--pack", "/path/to/pack-q2_0", "--native", "/path/to/Q2_0/Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00001-of-00002.gguf", "--ple-gguf", "/path/to/Q2_0/Qwen3.8-Flash-Next-GSQ-RCO-Q2_0-00002-of-00002.gguf", "--expert-profile", "/path/to/Strata/data/expert-profile.bin", "--expert-cache", "auto", "--prefill", "auto", "--max-context", "262144", "--kv", "int8", "--kv-resident", "32768", "--ple-io", "direct", "--spec", "2" ] } パスは自分の環境に置き換える。Pythonの依存パッケージは別の仮想環境にrequirements.txtから導入し、ROCmのライブラリパスも通した状態で、ソースルートからサーバーを起動した。 cd "$STRATA_SRC" python -m serve.server --engine strata --config /path/to/config.json \ --host 127.0.0.1 --port 8081 最初は--specを付け忘れ、ネイティブpackには--spec T (T >= 2)が必要だとエラーになった。--spec 2を追加すると読み込みが進み、日本語での応答に成功した。ここでの--spec 2はエンジンの指定で、MTP補助モデルを導入したという意味ではない。今回draftingはoffだ。 サーバーはローカルのみにbindしている。稼働確認は以下で行える。 curl -s 127.0.0.1 curl -s 127.0.0.1 8Kから32K、256Kへ まず8,192トークンで動作確認し、その後32,768へ増やした。約30,417トークンの合成入力の中央に置いた合言葉を正しく回収できた。入力処理は約71.8秒、約424 tokens/sだった。 次に--max-context 262144 --kv int8 --kv-resident 32768で起動。約248,458トークンの入力でも、中央の合言葉を正しく回収できた。入力処理は約640秒、約388 tokens/sだった。 試行 結果 8K設定、短い会話 約35 tokens/sの生成を観測 32K設定、約30K入力 合言葉を正しく回収、入力約424 tokens/s 256K設定、約248K入力 合言葉を正しく回収、入力約388 tokens/s 長文は同じ日本語の文章を繰り返した合成データで、検索課題も一つだ。複雑な複数文書の理解、長時間稼働、異なるコンテンツで同じ精度・速度が出ることまでは確認していない。短い出力6トークンの生成速度を、長文回答の安定速度として扱うこともしない。 256K起動時は専門家キャッシュが8,235個、空きVRAMは約790MiBだった。GPUがデスクトップ表示も担当しているので余裕は小さめだ。--expert-cache autoやVRAM予約量の調整が必要になる場合がある。Strataのメモリ設定を確認してほしい。 また、今回のStrataの既定設定ではSTRATA_ROUTE_TAIL_SKIP=7が有効で、キャッシュにない専門家の一部を省略する。ログにも回答が少し変わり得ると表示される。7,390トークンの合成入力で1回ずつ試したところ、生成は有効時36.7 tokens/s、無効時33.7 tokens/sで、回答文も少し変わった。各1回の短い応答なので、この差を一般化したり品質差を判断したりはできない。 OpenCodeにつなぐ StrataのOpenAI互換APIを使う。OpenCode v1.19.5をユーザー領域に導入し、~/.config/opencode/opencode.jsonにローカルプロバイダーを追加した。設定形式はOpenCodeのプロバイダー設定を参照した。 { "$schema": "opencode.ai", "model": "strata/Qwen3.8-Flash-Next-Q2_0-gfx906", "provider": { "strata": { "npm": "@ai-sdk/openai-compatible", "name": "Strata (local Radeon VII)", "options": { "baseURL": "127.0.0.1", "apiKey": "local", "timeout": 1200000 }, "models": { "Qwen3.8-Flash-Next-Q2_0-gfx906": { "name": "Qwen3.8 Flash Next Q2_0 (local)", "tool_call": true, "limit": { "context": 262144, "output": 8192 } } } } } } モデルIDは/v1/modelsが返すIDと一致させる。既存設定がある場合は丸ごと置き換えず、必要なプロバイダーを統合してほしい。apiKeyはこのローカルの無認証サーバーに渡すダミー値だ。 opencode models strata opencode 簡単な接続確認として、OpenCodeにテストファイルを読むよう依頼したところ、readツールを呼び出し、ファイル内の合言葉を正しく返した。コードの変更や複雑なエージェント作業の品質までは、この時点では検証していない。 振り返り Radeon VIIの16GB VRAMだけにモデル全体を載せるのではなく、GPUの専門家キャッシュ、RAM、SSDを組み合わせることで、Flash-Nextの推論までたどり着けた。gfx906は標準的な導入経路から外れているため、GPU認識だけで成功とせず、HIP、行列演算、エンジン、実モデル、ツール呼び出しの順に確かめるのが役立った。 一方、今回はnightly ROCm、ソースビルド、互換ヘッダーへの小さな修正を使っている。Qwen3.8-27BはStrataではなくllama.cppで動かしており、Strataで同じ27Bが動くという話ではない。まずは手元で動いた版と設定を記録し、更新するときは一つずつ比較していこうと思う。 この記事について 実はこの記事も、実際に環境を構築したChatGPTのセッションで書いた。会話の中で調べたこと、実行したコマンド、ログに残った測定結果をもとにChatGPTが下書きした。OpenCodeではファイル読み取りツールの呼び出しまで確認したが、複雑なコード修正の成功率はまだ測っていない。これから、この環境でどこまでできるか試していきたい。
ROCm 최신은 Core SDK 10.1.0이다. 릴리스가 ROCm/TheRock 태그로 옮겨갔고(therock-10.1, 10/6 KST), ROCm/ROCm GitHub 릴리스 페이지는 7.14.0에서 멈춰 있다. 거기만 보면 버전을 잘못 알기 쉽다. github.com/ROCm/TheRock/releases
FeatherAttn: Attention backward kernel on Strix Halo, 4x fast as FlashAttn woct0rdho/feather-attn On RDNA3/3.5 GPUs, FlashAttn wires to Triton kernels in AITER (AMD's kernel collection). The fwd kernel is fast enough, but it seems no one has optimized the bwd kernel. I've made a bwd kernel with a better designed pipeline. When tested on Strix Halo, it's 4x as fast for head dim 128, and 1.3x as fast for head dim 64. This is especially relevant because most recent models use head dim 128, including Anima, Krea 2, MOSS, H3 (while SDXL uses head dim 64). I've tested that it makes Anima LoRA training 40% faster. This kernel does not use int8/fp8 quantization or sparsity. It's still an open question how to make use of them on RDNA3/3.5 . By the way, besides attn bwd kernel, another bottleneck of LoRA training on Strix Halo is matmul bwd kernel. You may see the discussion in TheRock#5314 and install my forked rocm-libraries for faster GEMM NT.
FeatherAttn: Attention backward kernel on Strix Halo, 4x fast as FlashAttn woct0rdho/feather-attn On RDNA3/3.5 GPUs, FlashAttn wires to Triton kernels in AITER (AMD's kernel collection). The fwd kernel is fast enough, but it seems no one has optimized the bwd kernel. I've made a bwd kernel with a better designed pipeline. When tested on Strix Halo, it's 4x as fast for head dim 128, and 1.3x as fast for head dim 64. This is especially relevant because most recent models use head dim 128, including Anima, Krea 2, MOSS, H3 (while SDXL uses head dim 64). I've tested that it makes Anima LoRA training 40% faster. This kernel does not use int8/fp8 quantization or sparsity. It's still an open question how to make use of them on RDNA3/3.5 . By the way, besides attn bwd kernel, another bottleneck of LoRA training on Strix Halo is matmul bwd kernel. You may see the discussion in TheRock#5314 and install my forked rocm-libraries for faster GEMM NT.
Strata検証②：EVO-X2（Ryzen AI Max+ 395）で256K コンテクストが安定動作！Qwen3.8-Flash-Next × OpenCode でAIエージェント開発 Strata検証②：EVO-X2（Ryzen AI Max+ 395）で256K コンテクストが安定動作！ Qwen3.8-Flash-Next × OpenCode でAIエージェント開発 はじめに 前回の記事「EVO-X2（Ryzen AI Max+ 395）＋ Vulkan で Qwen3.8-Flash-Next が動いた！ でもコンテクスト長を大きくすると不安定！残念！」では、llama.cpp（Vulkan）で Qwen3.8-Flash-Next を 256K コンテクストで動かしたところ、10 万トークン前後の生成中に EVO-X2 がリセットしました。そのため、Flash-Next の常用は DGX Spark に移しました。 その直後、Qwen3.8-Flash-Next 専用の推論エンジン Strata が v0.1.40 で Ryzen AI Max+ 395（Strix Halo）に正式対応しました。そこで、同じ EVO-X2、同じ BIOS 設定とメモリの配分のまま、エンジンだけを Strata に替えて 256K を試しました。 結論 1. Strata なら、EVO-X2 で 256K（262,144 トークン）のコンテクストが動いた。 速度などの測定も含めて約 10.6 時間動かし続け（うち GPU 使用率 95% 以上が約 8.2 時間）、リセットは 0 回でした。前回リセットが起きた「10 万トークン前後で生成中」の状況も通過し、最大コンテクストは約 23.7 万トークンでした。 2. ハーネスの OpenCode と組み合わせて、256K のまま実際のエージェントとして使えた。 OpenCode V1（1.18.35）と V2（2.0.24）の両方で、ツールの呼び出し、思考、コンテクストの要約（compaction）、2 本同時の実行（Strata の "parallel": 2）まで確かめました。 3. 速度も上がった。 1 件ずつの出力速度は、MTP（モデルに付いている投機デコード用の層）ありで 43.6 tok/s、MTP なしで 28.5 tok/s です（前回の llama.cpp は 20.4 tok/s）。長いプロンプトの読み込みは、前回の約 3〜4.5 倍の速さでした。 4. 日本語の品質は変わらない。 JamC-QA（300 問）で 66.0%（前回の llama.cpp は 66.67%）。 5. ただし、前回のリセットの原因は分かっていない。 「Strata にしたから直った」とは言えません。くわしくは「前回のリセットについて」の節に書きました。 構成 !図 0 構成図 | 項目 | 内容 | |---|---| | 本体 | GMKtec EVO-X2（Ryzen AI Max+ 395、メモリ 128GB） | | GPU | Radeon 8060S Graphics（Ryzen AI Max+ 395 の内蔵 GPU） | | メモリの配分 | 前回と同じ。BIOS で VRAM に 96GB、残りがシステム RAM（OS から使えるのは約 30GB） | | BIOS | 1.09（2025-09-13）。GMKtec の公開ページに新しい BIOS は見当たらなかった | | OS／カーネル | Ubuntu 26.04 LTS / Linux 7.0.0-38（前回のリセットは 7.0.0-34 のとき） | | 推論エンジン | Strata v0.1.40.2 | | ROCm | AMD TheRock 7.14.1（Ryzen AI Max+ 395 向け）。ユーザー領域に展開し、システムの ROCm には触れていない | | モデル | Qwen3.8-Flash-Next（unsloth の GGUF、UD-IQ4_XS）。前回と同じファイル | | クライアント | OpenCode とベンチマークは、EVO-X2 とは別の PC から LAN 越しにつないだ | 推論をする機械の上では推論だけを動かし、OpenCode などのハーネスは別の PC に置いています。推論機が固まったときに、記録やクライアントを道連れにしないためです。監視の記録を別の PC にも送っているのは、前回のリセットで直前の記録ごと消えてしまった反省からです。 Strata とは Strata は Qwen3.8-Flash-Next 専用の推論エンジンです。MoE（Mixture of Experts）モデルのエキスパートのうち、よく使うものを GPU に、それ以外を CPU のメモリに置くことで、12GB 程度のゲーミング GPU でも 125B のモデルを動かせることで話題になりました。OpenAI 互換と Anthropic 互換の API を出すので、OpenCode などからそのままつなげます。 v0.1.40 で Ryzen AI Max+ 395 に対応し、ROCm（HIP）で動くようになりました。前回の llama.cpp は Vulkan で動かしていたので、GPU を動かす仕組みそのものが違うことになります。 ちなみに、Strata の開発者のマシンも Ryzen AI Max+ 395・128GB・Ubuntu 26.04 で、EVO-X2 とよく似ています。ただしメモリの配分は逆で、開発者は GPU のメモリ割り当てを小さくし、システム RAM 側を広げる構成です。 動かすときのコツ 1：EVO-X2 の配分では、既定のままだと黙って落ちる Strata の既定（--expert-cache auto）は、全エキスパート（約 59.5GB）をシステム RAM 側に確保しようとします。当ラボの EVO-X2 は BIOS で VRAM に 96GB を割り当てているため、システム RAM は約 30GB しかなく、確保できずにエラー表示なしでエンジンが終了しました。ログの最後に残っていたのは次の 1 行です。 ArenaExpertSource: the arena could not be reserved (59523276800 B) 次の 2 つを明示すると解決しました。 --mmap-experts --expert-cache 24576 --mmap-experts は GGUF をファイルのまま対応づけて読む指定、--expert-cache 24576 は全エキスパート分の枠を取る指定です。これで全 24,576 エキスパート（約 55.4GiB）が、BIOS で割り当てた 96GB の VRAM に載りました。起動は約 29 秒、読み込み後の VRAM の空きは約 27.7GiB（256K・KV キャッシュ int8・MTP あり・1 本）、システム RAM の使用は約 5.9GB でした。 前回の llama.cpp では、PLE（約 29GB の n-gram の表）を SSD に置く工夫をしましたが、今回はモデル全体が VRAM に収まっています。「BIOS で VRAM を大きく取る」配分だからこそ取れるやり方です。 実際の手順 実際の手順です（v0.1.40.2、Strata の公式ドキュメント docs/STRIX_HALO.md に沿ったもの）。ROCm は AMD の TheRock ビルド（7.14.1、Ryzen AI Max+ 395 向け）を自分のフォルダに展開して使います。root は要らず、システムに入っている ROCm には触れません。 # 1. ROCm（TheRock 7.14.1、Ryzen AI Max+ 395 (gfx1151) 向け）をユーザー領域に展開 mkdir -p ~/rocm && cd ~/rocm curl -LO https://repo.amd.com/rocm/tarball-multi-arch/therock-dist-linux-gfx1151-7.14.1.tar.gz mkdir install && tar -xzf therock-dist-linux-gfx1151-7.14.1.tar.gz -C install export SDK=$HOME/rocm/install export ROCM_PATH=$SDK HIP_PATH=$SDK HIP_PLATFORM=amd export PATH=$SDK/bin:$SDK/lib/llvm/bin:$PATH export LD_LIBRARY_PATH=$SDK/lib:$SDK/lib/rocm_sysdeps/lib:$SDK/lib/llvm/lib:${LD_LIBRARY_PATH:-} # 2. Strata をビルド（公式手順どおり。EVO-X2 では約 4 分） git clone https://github.com/Niko1221/Strata.git && cd Strata && git checkout v0.1.40.2 cmake -S . -B build-halo -G Ninja -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_HIP_ARCHITECTURES=gfx1151 \ -DCMAKE_HIP_COMPILER=$SDK/lib/llvm/bin/clang++ -DCMAKE_HIP_COMPILER_ROCM_ROOT=$SDK \ "-DCMAKE_PREFIX_PATH=$SDK;$SDK/lib/rocm_sysdeps;$SDK/lib/llvm" \ "-DCMAKE_HIP_FLAGS=--rocm-path=$SDK --rocm-device-lib-path=$SDK/lib/llvm/amdgcn/bitcode" cmake --build build-halo --target strata -j 8 # 3. GGUF（unsloth UD-IQ4_XS）を Strata の形式に変換（約 30 秒） .venv/bin/python tools/iq_pack.py \ --gguf <モデル置き場>/Qwen3.8-Flash-Next-UD-IQ4_XS-00001-of-00003.gguf \ --out packs/ud-iq4_xs --compat-bf16 --compat-bf16 は必須です。この GGUF が Q8_0 で持っている小さな射影を、Strata は BF16 で読むため、変換時に丸めが入ります（この丸めが、llama.cpp と出力が完全には一致しない理由です）。MTP のドラフトは、公式ドキュメント（docs/ORCA.md の準備の節）のとおりに作って mtp/rt に置きます。 サーバの設定は strata-ud-iq4_xs.json のような JSON に書きます。EVO-X2 で使ったものは次のとおりです。 { "exe": "build-halo/strata", "args": [ "--pack", "packs/ud-iq4_xs", "--native", "<モデル置き場>/Qwen3.8-Flash-Next-UD-IQ4_XS-00001-of-00003.gguf", "--expert-profile", "data/expert-profile.bin", "--expert-cache", "24576", "--mmap-experts", "--prefill", "16384", "--max-context", "262144", "--kv", "int8", "--spec", "4", "--spec-min-p", "0.5", "--mtp", "mtp/rt", "--lookup-chain", "3", "--mtp-q4", "all" ], "cwd": ".", "tokenizer": "packs/ud-iq4_xs/tokenizer", "model_name": "qwen3.8-flash-next-ud-iq4_xs", "log": "strata-ud-iq4_xs.log", "fit_max_tokens": true } - --expert-cache 24576 --mmap-experts が、上で書いた EVO-X2 向けの指定です（公式の例は "--expert-cache", "auto"） - --prefill 16384 にしているのは、後述の hipBLASLt の調整表がこの大きさのプロンプトまで含んでいるためです - 2 本同時に処理するときは "parallel": 2 を足します（1 本あたり VRAM を約 3.65GiB 多く使い、読み込み後の空きは約 20.2GiB でした） 起動のときは、公式ドキュメントの「速い構成」の環境変数と、hipBLASLt の調整表を渡します。調整表は、対応したカードなら setup が渡してくれますが、Ryzen AI Max+ 395 では自分で指定します（指定しないと、プロンプトの読み込みが遅くなります）。 export STRATA_HIPBLASLT_TUNING=$PWD/tools/hip/gfx1151-hipblaslt-100401.txt export STRATA_PF_FUSED=1 STRATA_PF_GEMM=1 STRATA_HC_UPMIX=1 STRATA_PA_FAST=1 export STRATA_HIP_WMMA=1 STRATA_SELECT_WMMA=1 STRATA_HC_Q8=1 export STRATA_PF_SWITCH_MIN_T=4096 .venv/bin/python -m serve.server --engine strata --config strata-ud-iq4_xs.json --host 0.0.0.0 --port 8095 起動ログに prefill gemm: hipBLASLt tuning enabled (… gfx1151, version 100401) と出れば、調整表が効いています。環境変数の 7 つは、既定より速い代わりに丸め方が変わるもので、公式は KL（出力分布の差）を測ったうえで「速い構成」として案内しています。Ryzen AI Max+ 395 向けの既定の 18 個のスイッチは、エンジンが自動で有効にします。 動かすときのコツ 2：OpenCode と組み合わせるなら "fit_max_tokens": true Strata は、「プロンプト＋max_tokens」がコンテクスト長を超えるリクエストを HTTP 400 で断ります。OpenCode（V1）は毎回 max_tokens として約 32,000 を送るので、会話が約 23 万トークンに達したところで、compaction が始まる前に止まってしまいました。 Strata の設定に "fit_max_tokens": true（残りの余地に合わせて max_tokens を縮める）を入れると解消しました。 256K の安定性 5 つの試験（合計 約 6.8 時間）を行いました。速度などの測定も含めると、約 10.6 時間動かし続けています。どの試験も、前回リセットが起きた 8.8 万〜10.4 万トークンの帯を通り越して、20 万トークン前後まで到達しています。 !図 1 各試験で到達したコンテクスト | 試験 | やり方 | 時間 | 最大コンテクスト | 結果 | |---|---|---|---|---| | API で会話を育てる | 毎ターン資料を足し、思考あり（high）で長い回答を生成 | 55 分（31 ターン、生成 約 13.9 万トークン） | 約 23.1 万（最終 23.7 万） | 異常なし。前回リセットした 8.8 万〜10.4 万トークンを、GPU 使用率 99〜100% で生成しながら約 32 分かけて通過 | | OpenCode V1 で実使用 | 同じセッションに依頼を重ねる（調査、レビュー、編集、要約） | 約 1.5 時間 | 約 23.5 万（compaction で 9.2 万へ） | 異常なし | | OpenCode V2 で実使用 | 同じ依頼 | 84 分 | 約 19 万（compaction 3 回） | 異常なし | | 2 本同時（parallel 2、V1 と V2） | 別々のセッションを同時に | 約 97 分 | 各 約 18 万 | 異常なし | | 2 本同時（parallel 2、API で最大まで） | 会話を 2 本同時に育てる | 82 分 | 21.7 万と 23.2 万を同時に保持 | 異常なし。毎ターン全文を読み直す厳しい条件で、21.4 万トークンの読み込みに 7.6 分かかる場面もあった | - 通算: 経過 約 10.6 時間のうち、GPU 使用率 95% 以上が約 8.2 時間。リセット 0 回、サーバの停止 0 回 - 監視: 1 秒ごとに GPU の状態・温度・電力・メモリを記録し、別の PC にも同時に送りました（前回、リセットで記録ごと消えた反省から）。記録の欠けはありませんでした - 温度と電力: GPU 温度は最大 94℃、電力は最大 約 117W コンテクストが長くなっても、出力速度は大きく落ちませんでした（API で会話を育てる試験、思考あり、MTP あり）。10 万トークンを超えたところで約 50 tok/s から約 43 tok/s に下がりますが、その先は 40 tok/s 台を保っています。 !図 4 コンテクスト長と出力速度 | コンテクスト | 出力速度（平均） | |---|---:| | 5 万未満 | 約 50 tok/s | | 5〜10 万 | 約 52 tok/s | | 10〜15 万 | 約 43 tok/s | | 15〜20 万 | 約 42 tok/s | | 20 万以上 | 約 43 tok/s | 2 本同時は、OpenCode を 2 つ並べて使う（たとえば調査用と実装用）個人利用を想定したものです。Strata の設定で "parallel": 2 にして測りました。1 本あたりの出力速度は、1 本だけのときの 45〜50 tok/s から 15〜21 tok/s に下がり、合計は 1 本のときとほぼ同じでした。2 本同時に使うと速くなるわけではなく、待たずに済むようになる、という使い方です。 OpenCode V1 と V2 | 項目 | V1（1.18.35） | V2（2.0.24） | |---|---|---| | 256K での実使用 | ○（最大 約 23.5 万） | ○（最大 約 19 万、compaction 3 回） | | ツールの呼び出し | read / grep / glob / edit / write / bash が動作 | read / grep / glob / edit / write / shell / execute が動作。形式の崩れなし | | 設定ファイル | opencode.json | opencode.jsonc（キーの名前や構造が V1 と別物） | - compaction は時間がかかる。 要約のリクエストは通常と先頭が違うため、コンテクスト全体を初めから読み直します。23.5 万トークンで 221 秒、15.8 万トークンで約 2 分かかり、その間エージェントは無音になります。OpenCode のタイムアウトは 900 秒にしました - ふだんはキャッシュがよく効く。 実使用では、入力の 92〜96% がキャッシュから読まれていました（2 本同時でも効いていました） - V2 はモデルを固定しておくこと。 V2 は -m でモデルを指定しないと、外部の既定のモデルを選びます。ローカルの LLM だけで使うなら、設定ファイルでモデルを固定するか、-m を必ず付けてください 実際の開発に使ってみて この構成は、その後ラボの実際の開発の推論先として使っています（Strata はその間に版を上げています）。たとえば、自分のデータを整理するエージェントの開発です。ブラウザを人と同じ速さで操作して、自分のデータ（画像とその情報）を集め、一覧で見られる Web 画面まで、OpenCode に作らせました。 データを外に出さずに、調査・操作・コーディングまでをローカルの LLM で回せる。外に出せないデータを扱う開発で、これが現実的な選択肢になってきたと感じています。 速度 1 件ずつ、思考なしで測りました。MTP の有無で 2 本に分けているのは、前回の llama.cpp と DGX Spark が MTP なしで測ったものだからです。 !図 2 出力速度 | 構成 | 出力速度（3 回の中央値） | |---|---:| | EVO-X2 ＋ Strata（MTP あり） | 43.6 tok/s | | EVO-X2 ＋ Strata（MTP なし） | 28.5 tok/s | | 前回：EVO-X2 ＋ llama.cpp（Vulkan、MTP なし） | 20.4 tok/s | | 参考：DGX Spark ＋ SGLang（NVFP4、MTP なし） | 27.0 tok/s | MTP なしどうしで比べても、Strata は llama.cpp の約 1.4 倍でした。 長いプロンプトを初めて読み込む（キャッシュなし）時間です。エージェントとして使うときは、資料やコードをまとめて読ませる場面が多いので、ここが速いかどうかが体感に効きます。 !図 3 長いプロンプトの初回読み込み時間 | 長さ | Strata（MTP あり） | Strata（MTP なし） | 前回：llama.cpp | 参考：DGX Spark | |---:|---:|---:|---:|---:| | 約 2,500 トークン | 3.1 秒 | 3.1 秒 | 10.5 秒 | 2.9 秒 | | 約 22,800 トークン | 19.9 秒 | 19.7 秒 | 約 65 秒（推定） | 14.3 秒 | | 約 63,300 トークン | 53.1 秒 | 53.5 秒 | 約 218 秒（推定） | 40.6 秒 | - Strata の読み込みは、llama.cpp の約 3〜4.5 倍の速さで、DGX Spark の 8 割弱でした - 同じプロンプトの 2 回目（キャッシュあり）は 0.05〜0.36 秒です - llama.cpp の約 22,800 と約 63,300 トークンの値は、一部キャッシュが効いていた実測値をキャッシュなしに換算した推定値です（前回の記事と同じ値） 日本語の品質 | 構成 | JamC-QA（300 問） | |---|---:| | Strata（MTP あり） | 66.0% | | Strata（MTP なし） | 66.0% | | 前回：llama.cpp（Vulkan） | 66.67% | 300 問で 2 問の差です。Strata の形式に変換するときの丸め（後述）などの範囲とみています。 前回のリセットについて 「エンジンを替えたらリセットしなくなった」のは事実ですが、原因は分かっていません。前回のリセットの記録を読み直したので、分かったことを書いておきます。 記録から分かったこと - 前回、予期しない停止は 2 回ありました。どちらも llama.cpp（Vulkan）が、約 10.4 万トークンと約 8.85 万トークンのコンテクストで生成している最中でした - 2 回とも、GPU のエラー（計算キューのタイムアウト、Device lost、GPU のリセット、ページフォルトなど）はカーネルのログに残っていませんでした - 2 回目は 46 秒後に自分で再起動し、次の起動のカーネルのログに、前回のリセット理由として「internal CPU shutdown event occurred」と記録されていました 最後の 1 行は、Linux 6.16 から入った機能です。AMD の CPU が持っている「前回のリセット理由」のレジスタを、起動時に読んで表示してくれます。 x86/amd: Previous system reset reason [0x00800800]: internal CPU shutdown event occurred つまり、GPU のエラーではなく、CPU 側の事象としてシステム全体がリセットされたことになります。 「Vulkan だから落ちた」とは言えない理由 llama.cpp の Vulkan で、長いコンテクストのときに GPU の計算キューがタイムアウトする、という報告は複数あります（#21724、discussion #27154）。最初はこれを疑いました。ただし、これらの報告の多くはプロセスの異常終了で済んでいて、今回のようなシステム全体のリセットとは症状が違います。GPU のエラーも記録されていませんでした。 また、次の点からも「Strata だから直った」とは言えません。 - 電力は Strata のほうが高い。 前回のリセットの直前は約 80W、Strata の試験では最大 約 117W でした。「負荷が軽いから落ちなかった」という説明は成り立ちません - カーネルも変わっている。 前回のリセットは 7.0.0-34 のときで、今回の試験は 7.0.0-38 です - 回数が少ない。 偶然を否定できるほどの回数ではありません Ryzen AI Max+ 395 では、似た報告が多い Ryzen AI Max+ 395 を積んだ機械では、高負荷の推論中にシステムごと落ちる、という報告がいくつもの機種から出ています。記録される理由は報告ごとに違い、「ACPI の電源状態の遷移」や「訂正できないエラーによるデータファブリックの sync flood」などがあります（ROG Flow Z13 の報告、Beelink GTR9 Pro の報告、GPU と NPU を同時に使ったときの報告）。機種、ファームウェア、負荷のかけ方によって出たり出なかったりするようです。 同じ症状に当たった人は、次の起動でカーネルのログの「Previous system reset reason」を見てください。 journalctl -k -b | grep "Previous system reset reason" GPU 側の問題か、本体側の問題かを切り分ける手がかりになります。 測定の条件と注意 - MTP の扱いが構成で違う。 前回の llama.cpp と DGX Spark は MTP なしです。公平に比べるなら、Strata の MTP なし（28.5 tok/s）を使ってください - Strata の「MTP なし」は、投機デコードを完全には切っていない。 今回の形式のモデルでは最低限の投機デコードが必須で、簡易なドラフト（受理率 約 49%）だけが残る構成です - モデルの変換。 Strata の形式に変換するとき、一部の値を BF16 に丸めています（公式の注記あり）。そのため llama.cpp と出力が完全には一致しません - 思考の止め方。 前回は llama.cpp の --reasoning off、今回はリクエストの reasoning_effort: none です - 速度の測り方。 前回と同じベンチマークのクライアント（ShareGPT の 200 プロンプト）を使いましたが、Strata に対してはクライアント側の一部の値が正しく出ませんでした（TTFT が数ミリ秒になる、MTP なしで出力を少なく数えるなど）。そのため、出力速度は Strata のサーバのログから計算しました - 回数。 出力速度は 3 回の中央値、長いプロンプトは各 1 回です - 版。 数字はすべて v0.1.40.2 での測定です。Strata はその後も版を上げていて、本記事の公開時点では 0.2 系が出ています 未解決の点 1. 一晩以上の連続運転は、この版ではしていない。 1 回の連続運転は最長 約 97 分、通算の高負荷は約 8.2 時間です。その後の 0.2.4.1 では、OpenCode V2 で 6 時間を超える連続運転をして異常がないことを確かめました。これは次の記事で書きます 2. 出力が 1 回だけ壊れた。 OpenCode V1 で約 16.4 万トークンのとき、最後の応答が壊れた短い文で止まりました。約 780 ステップ中この 1 回だけで、原因（モデル、量子化、エンジン、MTP）は切り分けていません 3. GPU 温度が最大 94℃。 長時間の高負荷では冷却に注意が必要です まとめ 前回「残念！」で終わった EVO-X2 の 256K コンテクストは、Strata で動きました。速度も上がり、OpenCode と組み合わせて実際のエージェントとして使えています。 一方で、前回のリセットの原因はまだ分かっていません。Ryzen AI Max+ 395 で長時間の高負荷をかける人は、監視の記録を別の機械にも送り、リセットしたら「Previous system reset reason」を確かめることをおすすめします。 英語版のリポジトリ（起動の設定、監視・測定のスクリプト、生データ）: nabe2030/strata-evo-x2-256k Strata への報告: #4032（Linux の Ryzen AI Max+ 395 での実測）、#4033（VRAM を大きく取った構成で --expert-cache auto が落ちる件）、#4034（ドキュメントの提案） おわりに — 当ラボについて 筆者は NPO法人 IoTメディアラボラトリーで、IoT 研究とあわせて AI 分野の検証を担当しています。東京大学・本郷キャンパス 大学院工学系研究科の同名のラボで行っていた IoT 研究が独立し、研究機関として活動を続けている組織です。 私がやっていることは、本記事のようなローカル LLM の実機検証が中心です。新しいモデルが出るたびに手元のハードで動かし、速度・日本語知識・コーディング・マルチモーダルの 4 軸で測り、動かなければ原因を突き止めるところまでやります。ベンダーの公称値ではなく、実際に自分の機械で何が起きるかを知りたいからです。 そのぶん、計算資源はいくらあっても足りません。もっと GPU が欲しい（笑）。機材の提供、検証のご依頼、共同研究など、お声がけいただけると喜びます。
WIN 5 GPU + NPU crashes: AMD fixes AGESA, need new BIOS? I have a GPD WIN 5 395 and I heard the system hangs during simultaneous GPU/NPU inference. This report describes pinning the fabric clock at maximum as a workaround. AMD says AGESA from 1.0.0.2 to 1.0.0.2b resolved their reproduction and a new firmware fix must be delivered through the manufacturer’s BIOS. Does anyone know: - Which AGESA version the current WIN 5 BIOS contains? - Whether GPD plans a BIOS update incorporating this fix? - The best official contact for GPD employee who could get this report to the BIOS engineering team?
MI50 ROCm 10.2 TheRock vs Vulkan Mesa 26.2 llama.cpp benchmarks I prefer running llama.cpp Vulkan prebuilt binary. I just download the latest version and ready to roll. I finally took the hours necessary to get TheRock latest tarball version of ROCm 10.2 running on dual AMD Radeon Instinct MI50 gfx906 (32gb combined VRAM). Same models benched in previous post. A mix of Dense and MoE models and quants that better utilize available VRAM. - llama_bench_Swift-Qwen3.8-27B-Uncensored-MTP.Q6_K.gguf - llama_bench_Swift-1.5-Qwen3.8-27B-Q6_K.gguf - llama_bench_Gemma-4-MoonGem-31B.i1-Q6_K.gguf - llama_bench_gemma-4-31B-it-UD-Q6_K_XL.gguf - llama_bench_Nemotron-3.5-30B-A3B-Antislop-FTPO.i1-Q5_K_M.gguf - llama_bench_Laguna-XS-2.1-APEX-I-Balanced.gguf - llama_bench_Agents-A1-Q4_K_M.gguf - llama_bench_Qwen3.6-35B-A3B-UD-Q5_K_XL.gguf - llama_bench_Qwen3.6-35B-A3B-Uncensored-HauhauCS-Aggressive-Q6_K_P.gguf Here is the backend performance comparison contrasting the native ROCm (v10.2 for gfx906) runtime against your optimized Vulkan (MESA_PPA_26.2) baseline. The data highlights a massive architectural split: ROCm significantly accelerates token generation across the board but suffers high variance and regressions in MoE pre-fills. Both GPUs are power limited to 145 watts. Build versions used: llama-b11382 used for Vulkan (prebuilt ubuntu binary) llama-b11401 used for ROCm (compiled with proper flags) Architectural Performance Breakdown: Vulkan vs. ROCm |Model|Size|Params|Test|Vulkan Baseline (t/s)|ROCm 1st Run (t/s)|Performance Delta (%)| |:-|:-|:-|:-|:-|:-|:-| |qwen35 27B Q6_K|20.88 GiB|27.32 B|pp512 tg128|149.30 ± 0.15 18.35 ± 0.01|176.86 ± 17.28 20.21 ± 0.65|+18.46% +10.14%| |qwen35 27B Q6_K|22.21 GiB|27.32 B|pp512 tg128|163.48 ± 0.18 18.52 ± 0.03|183.49 ± 15.75 20.04 ± 0.63|+12.24% +8.21%| |gemma4 31B Q6_K|23.46 GiB|30.70 B|pp512 tg128|119.98 ± 0.12 15.72 ± 0.02|172.93 ± 0.89 17.26 ± 0.13|+44.13% +9.80%| |gemma4 31B Q6_K|25.62 GiB|30.70 B|pp512 tg128|133.67 ± 0.24 16.38 ± 0.02|178.59 ± 2.31 17.50 ± 0.09|+33.61% +6.84%| |nemotron_h_moe 31B|25.18 GiB|32.91 B|pp512 tg128|834.87 ± 1.72 63.82 ± 0.07|736.98 ± 134.77 103.86 ± 0.50|-11.73% +62.74%| |laguna 30B.A3B Q5_K|22.64 GiB|33.44 B|pp512 tg128|723.34 ± 2.99 58.98 ± 0.28|856.14 ± 29.71 76.81 ± 0.20|+18.36% +30.23%| |qwen35moe 35B Q4_K|19.70 GiB|34.66 B|pp512 tg128|957.62 ± 3.88 51.50 ± 0.07|834.58 ± 102.56 69.70 ± 0.22|-12.85% +35.34%| |qwen35moe 35B Q5_K|24.76 GiB|34.66 B|pp512 tg128|909.20 ± 5.71 53.18 ± 0.07|763.93 ± 106.19 68.47 ± 0.34|-15.98% +28.75%| |qwen35moe 35B Q6_K|28.53 GiB|34.66 B|pp512 tg128|763.05 ± 66.01 52.81 ± 0.04|754.82 ± 70.48 67.29 ± 0.21|-1.08% +27.42%| Core Insight Strategy & Bottlenecks 1. Token Generation (tg128) Dominance: ROCm dominates pure text generation. The native AMD matrix kernels unleash your MI50 computation potential, unlocking a massive +62.74% boost for Nemotron but taking a small hit on pre-fill -11.73%. 2. Dense Model Pre-fills (pp512): Dense architectures scale cleanly under ROCm. Gemma 4 sees a +33% to +44% processing throughput spike over the Vulkan RADV driver driver bounds. MoE models take a hit with a -15.89% difference with llama_bench_Qwen3.6-35B-A3B-UD-Q5_K_XL.gguf being the happiest with Vulkan backend. 3. The MoE Prompt Processing Delinquency: Notice the massive standard deviations under ROCm for MoE pre-fills (e.g., Qwen3.5MoE Q4_K has an instability block of ± 102.56). ROCm suffers from severe scheduling thrashing when building prompt streams across multiple active experts. Here is the structured layout with pp512 and tg128 separated into individual columns for a clean side-by-side comparison between the two backend architectures. Backend Comparison Table (Vulkan vs. ROCm) |Model|Size|Params|Vulkan pp512 (t/s)|ROCm pp512 (t/s)|Vulkan tg128 (t/s)|ROCm tg128 (t/s)| |:-|:-|:-|:-|:-|:-|:-| |qwen35 27B Q6_K|20.88 GiB|27.32 B|149.30 ± 0.15|176.86 ± 17.28|18.35 ± 0.01|20.21 ± 0.65| |qwen35 27B Q6_K|22.21 GiB|27.32 B|163.48 ± 0.18|183.49 ± 15.75|18.52 ± 0.03|20.04 ± 0.63| |gemma4 31B Q6_K|23.46 GiB|30.70 B|119.98 ± 0.12|172.93 ± 0.89|15.72 ± 0.02|17.26 ± 0.13| |gemma4 31B Q6_K|25.62 GiB|30.70 B|133.67 ± 0.24|178.59 ± 2.31|16.38 ± 0.02|17.50 ± 0.09| |nemotron_h_moe 31B|25.18 GiB|32.91 B|834.87 ± 1.72|736.98 ± 134.77|63.82 ± 0.07|103.86 ± 0.50| |laguna 30B.A3B Q5_K|22.64 GiB|33.44 B|723.34 ± 2.99|856.14 ± 29.71|58.98 ± 0.28|76.81 ± 0.20| |qwen35moe 35B Q4_K|19.70 GiB|34.66 B|957.62 ± 3.88|834.58 ± 102.56|51.50 ± 0.07|69.70 ± 0.22| |qwen35moe 35B Q5_K|24.76 GiB|34.66 B|909.20 ± 5.71|763.93 ± 106.19|53.18 ± 0.07|68.47 ± 0.34| |qwen35moe 35B Q6_K|28.53 GiB|34.66 B|763.05 ± 66.01|754.82 ± 70.48|52.81 ± 0.04|67.29 ± 0.21| So yes it's worth the hassle of jumping through hoops to get ROCm working on MI50 setups. At least I have 2 backends working. Next up I'll try some RPC. preview.redd.it