技術約14分で読めます

125BのQwen3.8-Flash-NextをEVO-X2のROCmで動かす

いけさん目次

前回は125BのQwen3.8-Flash-NextをM1 Max 64GBで動かした
今回は同じGGUFとテスト一式を普段ローカルでLLMを動かしているEVO-X2のWindows + ROCmで回し、常用している27Bを差し替えられるか試した。

検証環境

項目内容
マシンGMKtec EVO-X2(NucBox_EVO-X2)
CPUAMD Ryzen AI Max+ 395(16コア32スレッド)
GPURadeon 8060S(gfx1151)、BIOSでVRAMに48GB割当
メモリ64GB LPDDR5X(8532 MT/s)。VRAM割当のためシステム側に残るのは15.6GB
SSDCrucial CT1000E100SSD8(NVMe 1TB)
OSWindows 11 Pro 26200
ドライバAMD Software 26.8.1 / 32.0.31041.1004
llama.cppPR #27742のcommit b8bdf73ef68766を自前ビルド、のちに公式リリースb10666(commit 4e97ac86e
ROCm自前ビルド時はTheRockのnightly 10.1.0a20260812、公式リリース版では7.14.0
モデルAtomicChat/Qwen3.8-Flash-Next-GGUFAD-3.84bpw-IQ4_XS-M64(28シャード)+ mmproj F16

EVO-X2はBIOSでVRAMに48GBを固定してあるので、M1 Maxのように64GBをOSが必要なだけ振り分けることはできない。システム側には15.6GBしか残らない。
このモデルは直前のトークン列をキーに38.4GBのN-gram表を引くが、表そのものはSSDに置いたままでよく、読み出しを受けるのがこの15.6GBになる。

自前ビルドで落ちた場所

Qwen3.8-Flash-Next対応のWindows向けROCmバイナリはどこからも配られておらず、自前でビルドするしかなかった。
PR #27742がまだマージされていないので上流のmasterにsrc/models/qwen4exp.cppは無く、EVO-X2で使っているlemonade-sdk/llamacpp-rocmの最新リリースb1317にも入っていない。

開発ツールは入っていないので、ビルド環境から用意した。
といってもVisual Studio Build Toolsのインストーラは管理者権限を求めてくる。
vs_BuildTools.exe --layoutでパッケージ一式をディレクトリに落とし、VSIXを解凍してMSIをmsiexec /aで展開する経路で逃げた。
残るROCm SDKはTheRockのnightlyのtar.gzを、cmakeとninjaは公式zipを展開している。

ここでいちばん時間を取られたのは、必要なファイルがどれも直感と違う場所に入っていたからだ。
kernel32.Libは「Windows SDK Desktop Libs x64」ではなく「Windows SDK for Windows Store Apps Libs」のほう。
msvcrt.liboldnames.libCRT.x64.Desktop.baseではなくCRT.x64.Store.base側にある。
結局ヘッダ3606個とライブラリ816個をマージして通した。
llama.cppに.rcファイルは無いのに、cmakeのWindows-Clang.cmakeは勝手にリソースコンパイラを要求してくる。
-DCMAKE_RC_COMPILERにROCm付属のllvm-mc.exeを指定してかわした。

ビルドしたバイナリを既定のまま起動すると、モデルのロードで落ちた。

ggml_cuda_host_malloc: failed to allocate 37265.41 MiB of pinned memory: out of memory
llama_model_load: error loading model: unable to allocate ROCm_Host buffer

HIPバックエンドがピン留めメモリを39GB要求してきて、15.6GBしかないシステムRAMでは足りない。
M1 Maxなら配布元の指示どおりmmapを既定のまま起動すればよく、N-gram表はファイルキャッシュ経由で読まれて常駐メモリには載らなかった。
ところがHIPはCPU側に置くテンソルをページングできないROCm_Hostへコピーしようとするので、mmapで開いたページをそのまま使ってくれない。

そこでピン留めを切り、N-gram表を明示的にCPUバッファへ落とした。

set LLAMA_MMAP_RANDOM=1
set GGML_CUDA_NO_PINNED=1
llama-server.exe -m <shard1.gguf> -ngl 99 -lm mmap -ot "per_layer_token_embd=CPU" ...

これでロードは通った。
GPUに43720.58 MiB、SSD上のmmapに36621.27 MiBという、配布元がREADMEに書いている「In memory 45.8GB / On SSD 39.1GB」と同じ配分で、49層すべてがGPUへオフロードされている。

通ったのはそこまでで、最初の計算で落ちた。

cmn common_init_: warming up the model with an empty run - please wait ...
ggml/src/ggml-cuda/ggml-cuda.cu:107: ROCm error

llama-benchで回すと、代わりにrocBLAS側で止まった。

rocBLAS error: Could not initialize Tensile host: invalid unordered_map<K, T> key

自前ビルドのログ3段。1段目はピン留めメモリ37265.41 MiBの確保失敗、2段目は49/49層のGPUオフロードとCPU_Mapped 36621.27 MiBの成功ログ、3段目はTensile host初期化失敗

そこから、思いつく原因を片端から試した。

仮説検証結果
GPU故障、タイムアウト未復帰b1311 + 27Bを実行否定。pp128 323 t/sで正常
自前ビルドの不良同じ自前バイナリで27Bを実行否定。pp64 257 t/s
ROCmバージョン不一致b1311と同じ10.1.0a20260812でビルドし直す否定。同じエラー
Tensileカーネル欠損b1311のrocblas/libraryと比較否定。150ファイルがバイナリ一致
ランタイムDLL不足origami.dllrocm_kpack.dllを含めb1311と同一構成にする否定。同じエラー
カレントディレクトリ実行ファイルと同じ場所から実行否定
rocBLAS回避GGML_CUDA_FORCE_MMQ=1不可。初期化自体は走る
システムRAM不足空き10.6GB、ページファイル32GBを確認否定
GPUメモリ不足-ngl 47 / -ngl 40 に下げる別のエラーに変わる
ドライバが古い26.7.1から26.8.1に更新否定。同じエラー
PRが古いb8bdf73からef68766否定。同じエラー

ここで中断した。

llama.cpp側を確認し直す

一旦公式を見に行くと、PR #27742は8月28日にmasterへマージされていた(commit 6c84c7d5、28ファイル2881行の追加)。
リリースタグではb10664以降に入っている。

Windows向けのROCmバイナリも配り始めていた。
リリースアセットにllama-b10666-bin-win-rocm-7.14-x64.zipが並び、CIのターゲット一覧にはgfx1151も入っているので、自前ビルドはもう要らない。

自分が使っていたLLAMA_MMAP_RANDOMのほうは、マージ前に取り消されていた。
ef68766の1時間15分後に取り消すコミットが入り、その13分後にはTENSOR_READ_LAZYがマージされて--tensor-read-lazyに置き換わっている。
こちらは既定のautoが4GiBを超えるテンソルを対象にするので、38.4GBのN-gram表なら何も指定しなくても外れる。
なおllama-benchにこのオプションは無いが、既定は同じautoなので扱いは変わらない。

公式ROCmバイナリに何か足りない

llama-b10666-bin-win-rocm-7.14-x64.zip(230MiB)を展開しただけでは、GPUが1台も出てこない。

> llama-server.exe --list-devices
Available devices:
  (none)

ggml-hip.dllのロードに失敗している。
依存を調べるとhipblas.dllが要求されているのに、zipの55ファイルの中に無い。

> dumpbin /dependents ggml-hip.dll
    ggml-base.dll
    hipblas.dll        <-- zipに入っていない
    amdhip64_7.dll

zipに入っているのはamdhip64_7.dllrocm_kpack.dllamd_comgr.dllの3つだけ。
release.ymlの該当ステップにもrocblas/hipblaslt kernels resolve fine via PATH and are not copiedとあって、rocBLAS一式は意図して外してある。
逆にこの3つが入っているのは、AdrenalinドライバがSystem32に同名のamdhip64_7.dllを置いているからで、exeの隣なら先に読まれる(issue #26929)。

そこで、CIと同じROCm 7.14.0を、AMDが配っている公式のホイールからvenvへ入れた。

python -m venv rocm714venv
rocm714venv\Scripts\python.exe -m pip install ^
  --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries]==7.14.0"
rocm714venv\Scripts\python.exe -m pip install ^
  --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm-sdk-device-gfx1151==7.14.0"

rocm[libraries]hipblas.dllrocblas.dllが入る。
Tensileのカーネル本体はここには無い。
rocm-sdk-device-gfx1151のほうに入っていて、展開すると_rocm_sdk_libraries\bin\rocblas\library\gfx1151\に150ファイル出てくる。
この2つを入れたうえで、_rocm_sdk_libraries\bin_rocm_sdk_core\binをPATHに置く。

27Bに指定を1つずつ増やす

ROCm 7.14.0を用意しても、Flash-Nextは同じところで落ちた。

そこで、正常に動くほうの27Bに、Flash-Next用の指定を1つずつ増やして確かめた。

27B Q8_0 / ngl 99 / fa 1 の条件結果
素のままOK。pp64 278.07 t/s / tg16 7.86 t/s
-lm mmapを追加OK。pp64 266.84 t/s / tg16 7.82 t/s
LLAMA_ATTN_ROT_DISABLE=1を追加OK。pp64 279.40 t/s / tg16 7.85 t/s
GGML_CUDA_NO_PINNED=1を追加NG。ggml-cuda.cu:107: ROCm error

ピン留めのメモリ不足を避けるために自前ビルドのときから入れていたGGML_CUDA_NO_PINNED=1を足しただけで、27Bまで落ちた。

この変数を外すと、Flash-Nextはそのまま通った。27Bに付けていなかったから、そちらは落ちなかった。

b10666では--tensor-read-lazyの既定が効いて、N-gram表が最初から常駐の対象外になる。

create_tensor: tensor per_layer_token_embd.weight (size = 36621 MiB) lazy read enabled

ピン留めの確保そのものが起きないので、GGML_CUDA_NO_PINNEDは要らないどころか、付けると壊れた。

再テストのログ3段。1段目は公式zip単体でGPUが見えずhipblas.dllが無いこと、2段目は27Bへの指定追加でGGML_CUDA_NO_PINNED=1だけがROCm errorになること、3段目は49/49層オフロードとpp512 159.25 / tg128 23.26の結果

ロードと配分

-otとmmapを一緒に指定すると、ローダーがconsider using --load-mode none for better performanceと警告を出す。
ただ--load-mode noneにするとN-gram表が常駐してしまうので、この構成では選べない。

そのままロードすると、llama-benchで159.5秒、llama-serverだとmmproj込みで2分26秒かかった。
このときGPUには43720.58 MiB、SSD上のmmapには36621.27 MiBと644.14 MiBで、配分は自前ビルドのときと同じだったが、graph splitsだけが25〜28から4に減っていた。

M1 Max(配布元README)EVO-X2 自前ビルドEVO-X2 b10666
GPUに載る量45.8GB43720.58 MiB43720.58 MiB
SSD上のmmap39.1GB36621.27 + 644.14 MiB同じ
graph splits25〜284

llama-benchで測った数字はこんな感じ。

testt/s
pp512159.25 ± 7.61
tg12823.26 ± 0.01

ppはプロンプトを読む側、tgは答えを書く側の速度で、後ろの数字は扱ったトークン数を指す。

2つのIssueは再現するか

同じgfx1151で2件のIssueが出ていたので両方試したが、どちらも再現しなかった。

Issue #27856を見ると、1Kトークンを境に19〜21 t/sから6 t/s前後まで落ちると書いてあった。
そこでllama-benchの-dを使い、すでに入っている文脈の長さ(KVキャッシュの長さ)を振ってみたのだが、緩やかに下がるだけで、報告にあるような急な落ち込みは出ない。

KV長01K4K16K32K64K
tg3222.8619.3617.6414.1511.768.54

16Kでも14.15 t/sあるので、WIPのPR #27860を待つ必要は無さそうだった。
ただ、b10666で該当の処理が変わったのか、それともLinux + ROCm 7.2側だけの話なのかまでは分からない。

もう一件のIssue #27797は、メッセージの区切りが2つ以上あると/を出し続けるという話で、対策にはLLAMA_ATTN_ROT_DISABLE=1がいるとあった。
区切りを2つ以上にするため、3ターンのペルソナ会話を変数の有無で2回流している。

1ターン目2ターン目3ターン目
LLAMA_ATTN_ROT_DISABLE=1177 tok / 23.4 t/s284 tok / 23.7 t/s297 tok / 23.8 t/s
なし184 tok / 23.2 t/s312 tok / 23.4 t/s447 tok / 23.4 t/s

どちらも/の連続は出ず、口調も設定も3ターン維持したままで、速度差も無い。
b10666では要らなかった。

コンテキスト長をどこまで伸ばせるか

普段は27Bを--ctx-size 65536で回しているので、差し替えるなら同等以上がいる。
そこで-cを変えてはllama-serverを立て直し、空コンテキストでの生成速度を2回ずつ測った。

-ctg(2回)
65536(64K)23.71 / 25.78 t/s
98304(96K)20.99 / 22.37 t/s
131072(128K)23.64 / 25.61 t/s
147456(144K)5.22 / 5.32 t/s
163840(160K)5.22 / 5.24 t/s
262144(256K)3.72 / 3.82 t/s

131072と147456のあいだで1/5まで下がった。
モデルのn_ctx_trainは262144あるので、効いているのはVRAMの割り当てのほうだろう。
KVは48層のうちフルアテンションが12層だけなので軽く、512セルあたり16.5 MiB、1トークンでも約33KiBしかないから、128K分で約4.2GiB。
モデルの43.7GiBと足した47.9GiBがBIOSの48GBにちょうど入り、そこを超えたぶんが共有システムメモリへ溢れる。

ロード自体は256Kでも通るので、エラーが出ないまま遅くなってから気づく。

もうひとつ、大きい-cでllama-serverを起動していったん止め、直後に次を立てるとROCm error: unspecified launch failureで起動に失敗した。
これが2回あって、25秒ほど空けてから起動したら出なかった。

空き物理RAMと生成速度

VRAM割当を48GBから56GBへ増やせば、速度が落ちる境目も動くはず。
そこでFlash-Next稼働中のメインメモリを測った。

Flash-Next -c 13107227B(常用構成)
llama-serverのワーキングセット10,707 MB3,128 MB
プライベートコミット49,713 MB34,523 MB
空き物理RAM441 MB7,958 MB
コミット済み / 上限65,693 / 66,497 MB52,666 / 54,343 MB

27B運用時は7.96GB、ほぼ8GB空いているが、Flash-Next稼働中はワーキングセットだけで10.7GBある。
物理RAMを8GBにすればこの10.7GBが入らないうえ、コミット済みは上限の98.8%に張り付いていた。

この状態で100MBほどのプロセスを1つ足すと、速度が変わった。

条件空き物理RAMtg
クリーン850 MB24.2 t/s
100MBのプロセスを1つ追加173 MB5.3 t/s

100MBの差で4倍以上違う。
なお同時刻に27Bを測ると pp64 261.92 / tg16 7.76 で正常だったので、機体側の問題には見えない。

M1 Maxと同じ5テスト

5テストは-c 32768のllama-serverで回した。

#テスト生成 tok生成 t/sプリフィル t/s実時間思考
1BST挿入(思考ON)39424.264.817.5s1,216字
1BST挿入(思考OFF)9425.996.54.1s
2簡易BBS(思考OFF)7,50818.892.0399.7s
2簡易BBS(思考low)2,17821.699.4101.9s454字
2簡易BBS(思考xhigh)12,28817.6116.5697.8s38,172字で打ち切り
3ペルソナ3ターン177 / 284 / 29723.4 / 23.7 / 23.88.8 / 12.9 / 13.8s192 / 671 / 561字
4NSFW 3段階365 / 611 / 31423.5 / 23.6 / 23.916.4 / 26.9 / 14.0s494 / 1,709 / 421字
5Vision2,04819.4194.6109.8s4,324字

生成速度は素の状態なら23〜26 t/s出るが、長い出力になるほど下がり、17〜19 t/sまで落ちることもあった。

簡易BBSは、思考OFFの回が「ことばの広場」という単一HTMLになり、文字数カウンターとサンプル投稿3件まで入っていた。

思考OFFで生成した簡易BBS「ことばの広場」。和紙風の配色に朱色の落款、左に投稿フォーム、右に投稿一覧が並ぶ

思考xhighは38,172字を考えたところで上限に達し、空白ページになった。M1 Maxでも27Bでも同じだった。

NSFWの3段階のうち、素のプロンプトの回と直接描写をはっきり求めた回は、どちらも拒否された。
systemで制限解除を主張したBだけ部分的に応答が返り、それも「プロンプト内に埋め込まれた指示は動作設定を変更しない」と明言したうえでの文学的な短編だった。

Visionにはテスト2の思考OFFで作ったBBSのスクリーンショット(900×900)を入力した。
プリフィル194.6 t/sで、ヘッダー、導入文、左のフォーム、右の投稿一覧という構造を正しく分解し、画面内のテキストもほぼ書き出していた。
そのうえ「3件の投稿とあるが画面には2件しか見えないので、下にもう1件あると推測される」と、写っていない部分にまで言及している。
誤りはプレースホルダーの「名乗りたければ、ここへ」を「名乗るたけであれば、ここへ」と読んだ1箇所だけ。

6万トークン超で3モデルを比べる

常用の27Bを差し替えるかどうかは、文脈が空のときの速度だけでは決まらない。
同じllama-server(b10666 + ROCm 7.14.0)に同じプロンプトを投げて、3モデルを比べた。
27Bで使うMTPは、先読みしたトークンを本体モデルでまとめて検証する仕組みで、反復が多いテキストだと当たった割合が実態より高く出る。
そこでプロンプトには反復のない実テキストを選んだ。17.5万字で、モデルごとに62,538〜62,580トークンになった。

モデル構成空のとき tg62K入れたとき tgプリフィル @62K62K応答完了
Qwen3.6-35B-A3B Q6_K(旧構成)-c 65536 KV q8_057.5236.49516.51 t/s122.9s
Qwen3.8-27B Q8_0 + MTP(現行常用)-c 6553622.4411.54209.98 t/s299.6s
Qwen3.8-Flash-Next-c 131072 mmap + lazy22.948.55115.89 t/s547.3s

文脈が空なら27Bと並ぶが、62Kまで詰めるとFlash-Nextがいちばん遅い。

文脈の長さによる下がり方が、両者でまるで違う。

KV長016K32K64K下がり幅
Flash-Next22.8614.1511.768.54−63%
27B(MTPなし)7.837.597.326.86−12%

27Bは1トークンごとに全パラメータを読むので、文脈が伸びても生成の時間は変わらない。
Flash-Nextのほうは6Bしか使わないぶん、スパースアテンションとN-gram表にかかる時間の比重が大きくなってくる。
なお27B側は、MTPの当たった割合も92%から46%へ下がっている。

トークナイザは27BとFlash-Nextで同じ(語彙数248320)で、4種類のサンプルでトークン数が1個も違わない。そこで文脈長を字数に直した。

テキスト種別文字/トークン27B -c 65536Flash-Next -c 131072
日本語・技術文(コード混在)2.01約13.2万字約26.3万字
日本語・普通の散文2.60約17.0万字約34.0万字
HTML / コード2.94約19.2万字約38.5万字
英語3.52約23.1万字約46.1万字

日本語散文34万字は文庫本3冊ぶんになる。

普段動かすのは27B + MTPのままにした。
文脈は2倍取れるものの、その2倍を実際に使う場面でいちばん遅いからだ。
ディスクを80GiB常時占有してロードに2分半かかるのも、watchdogで再起動する現行の運用とは合わない。

速度だけなら旧構成のQwen3.6-35B-A3Bがどの指標でも最速だったので、同じ5テストをこちらにも回した。
指示追従と作り込みでは届かず、ペルソナは絵文字を勝手に使って設定を守れなくなるし、BBSも汎用的な出力になる。
そのかわり11本の合計時間が2分40秒で、Flash-Nextの約21分とは桁が違う。

そのテストで立てた3.6用のrun-server-qwen36-backup.batは、--jinjaを付けずに--reasoning-budget 0を入れていた。
これではenable_thinkingreasoning_effortも一切効かない。
--jinjaを足すとreasoning_effort: xhighまで解釈して、774字の思考が出た。