大規模言語モデルの利用とAPU

大規模言語モデル（Large Language Models, LLM）は自然言語処理(Natural Language Processing, NLP)で使われる深層学習モデルです。LLMを利用することでChatGPTなどの対話型生成AIを実現できます。

LLMの利用にはメモリを多く使用しますが、通常の外付けGPUではVRAMの増設は困難です。Ryzen APUはCPUとGPUでメインメモリを共有するUnified Memory Architecture(UMA)で、メインメモリは容易に増設可能です。大きなメインメモリを容易に持てるAPUならではのLLMの活用が可能かもしれないので、Fedora 40を実行するRyzen 5600G上で、Fugaku-LLMをROCm対応llama.cppとkoboldcppで動かし、ベンチマークを取りました。

結論を先に言うと、Ryzen 5600Gは、メインメモリの活用することで大きなLLMをGPUで利用が可能ですが、速度が遅くあまり有効ではありません。より高速なGPUとメモリを持つAPUでは意味が出てくるかもしれませんが、現時点ではAPUはLLMの利用に向いているとは言えないと思います。

動作環境

種類	内容
CPU	Ryzen 5600G
マザーボード	ASRock B450M-HDV
メモリ	CFD W4U3200CS-16G (16GB×2=32GB)
OS	Fedora 40

ROCm対応llama.cpp

メタ社のLLaMAに対応するllama.cppのROCm対応版をインストールし、ビルドします。HIPでのUMAに対応しているので、VRAMの割り当てやLD_PRELOAD等は不要です。

準備

rocmやhip、hipblasの開発パッケージが必要なので、インストールしていなければdnfでインストールしてください。

$ sudo dnf install rocminfo rocm-smi rocm-device-libs clang hipcc lld lld-devel compiler-rt hipblas-devel

また、llama.cppのissueより、Fedora 40でVEGA世代(gfx900)のhipblas関係をビルドするためには特別な設定が必要なので、もしrocm-rpm-macros-modulesが入っていない場合は入れてください。

$ sudo dnf install rocm-rpm-macros-modules

rocminfoで正常にROCm環境が動いているか調べてください。

$ rocminof
ROCk module is loaded
=====================    
HSA System Attributes    
=====================    
Runtime Version:         1.1
System Timestamp Freq.:  1000.000000MHz
Sig. Max Wait Duration:  18446744073709551615 (0xFFFFFFFFFFFFFFFF) (timestamp count)
Machine Model:           LARGE                              
System Endianness:       LITTLE                             
Mwaitx:                  DISABLED
DMAbuf Support:          YES

==========               
HSA Agents               
==========               
*******                  
Agent 1                  
*******                  
...
*******                  
Agent 2                  
*******                  
  Name:                    gfx90c                             
  Uuid:                    GPU-XX                             
  Marketing Name:          AMD Radeon Graphics                
  Vendor Name:             AMD                                
  Feature:                 KERNEL_DISPATCH                    
  Profile:                 BASE_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        128(0x80)                          
  Queue Min Size:          64(0x40)                           
  Queue Max Size:          131072(0x20000)                    
  Queue Type:              MULTI                              
  Node:                    1                                  
  Device Type:             GPU                                
...

おかしければ/dev/kfdのパーミッション等をチェックしてください。ROCmの設定についてはStable Diffusion WebUIの記事を参照してください。

Ryzen 5600G用の環境変数の設定

5600GのGPUのgfx90cはvega10(Vega 56/64)のISAであるgfx900と互換なので、環境変数HSA_OVERRIDE_GFX_VERSIONをgfx900にします。ROCm 6.0ではHSA_ENABLE_SDMAを無効にしないとうまく動かないので、これも設定します。

$ export HSA_OVERRIDE_GFX_VERSION=9.0.0
$ export HSA_ENABLE_SDMA=0

llama.cppの取得とビルド

gitでファイルを取得します。

$ git clone https://github.com/ggerganov/llama.cpp
$ cd llama.cpp

llama.cppのissueに従い、moduleコマンドでrocm/gfx9を読み込みます。

cmakeにROCmのCUDA互換環境であるHIPを有効にするようLLAMA_HIPBLASをONにし、APUのUMAを有効にするようにLLAMA_HIP_UMAをONにし、使うGPUがgfx900なのでAMDGPU_TARGETS=gfx900を指定し、build以下でビルドします。Fedora 40ではさらにCXXFLAGSに--rocm-device-lib-path=/usr/lib/clang/17/amdgcn/bitcodeを指定しないと正常にコードがビルドされないようです。

$ module load rocm/gfx9
$ mkdir build
$ cd build
$ CC=/usr/bin/hipcc CXX=/usr/bin/hipcc cmake .. -DLLAMA_HIPBLAS=ON -DLLAMA_HIP_UMA=ON -DAMDGPU_TARGETS=gfx900 -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_FLAGS="--rocm-device-lib-path=/usr/lib/clang/17/amdgcn/bitcode"
$ make -j

問題なければbin/以下にmainができているはずです。

ちなみにmodule load rocm/gfx9の中身は/usr/share/modulefiles/rocm/gfx9です。Fedora 40で以外ではこの作業は不要だと思います。

#%Module 1.0
#
#  ROCm module for use with 'environment-modules' package:
# 

conflict rocm

setenv        ROCM_BIN      /usr/lib64/rocm/gfx9/bin
setenv        ROCM_LIB      /usr/lib64/rocm/gfx9/lib
setenv        ROCM_GPUS     "gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack+;gfx90a:xnack-;gfx940;gfx941;gfx942"
prepend-path  PATH          /usr/lib64/rocm/gfx9/bin
prepend-path  LD_LIBRARY_PATH /usr/lib64/rocm/gfx9/lib

LLAMA_HIP_UMA=ONの意味

cmakeがMakefileに対して-DGGML_HIP_UMAを指定しています。これにより、ggml-cuda.cuのggml_cuda_device_mallocに関して、cudaMallocをhipMallocManagedとhipMemAdviseに置き換えるよう、122行目から記述されています。

static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device) {
    ggml_cuda_set_device(device);
#if defined(GGML_USE_HIPBLAS) && defined(GGML_HIP_UMA)
    auto res = hipMallocManaged(ptr, size);
    if (res == hipSuccess) {
        // if error we "need" to know why...
        CUDA_CHECK(hipMemAdvise(*ptr, size, hipMemAdviseSetCoarseGrain, device));
    }
    return res;
#else
    return cudaMalloc(ptr, size);
#endif
}

この変更の意味はROCmの文書のメモリ割り当てを参照してください。

Stable Diffusion WebUIの記事のhipMalloc(cudaMalloc)をhipHostMallocに置き換える手法や後述するkoboldcppの手法とは少し違いますが、メインメモリから取ってくるように変更していることは同じです。

Fugaku-LLMのダウンロードとllama.cppの実行

こちらからFugaku-LLM-13B-instruct-0325b-q5_k_m.ggufとFugaku-LLM-13B-instruct-0325b.ggufをダウンロードしてllama.cppのbuild以下に置いてください。huggingfaceのアカウント取得とログインとFugaku-LLMの利用規約への同意が必要です。

他にも量子化したものがこちらにあるので、こちらでもかまいません。

LLMの量子化に関してはこの辺を参照してください。

https://sc-bakushu.hatenablog.com/entry/2024/02/26/062547
https://sc-bakushu.hatenablog.com/entry/2024/04/20/050213

Fugaku-LLM-13B-instruct-0325b-q5_k_m.ggufをllama.cppで実行します。『スーパーコンピュータ富岳について教えてください。』という文章を指定して回答してもらいます。bin/mainが実行ファイルす。-n'で回答トークン数(文字数)、-mでLLMファイルの指定、-pで取り合わせ内容の文章の指定ができます。また、llama.cppはどれくらいGPUにオフロードするか指定でき、Fugaku-LLM-13B-instruct-0325b-q5_k_m.ggufでは0から41まで指定できるのですが、とりあえず-ngl 10000`で全部確実にGPUにオフロードできます。

$ bin/main -m Fugaku-LLM-13B-instruct-0325b-q5_k_m.gguf -n 128 -p "スーパーコンピュータ富岳について教えてください。" -ngl 10000

...
llm_load_print_meta: model ftype      = Q5_K - Medium
llm_load_print_meta: model params     = 13.42 B
llm_load_print_meta: model size       = 10.42 GiB (6.67 BPW) 
...

llm_load_tensors: ggml ctx size =    0.46 MiB
llm_load_tensors: offloading 40 repeating layers to GPU
llm_load_tensors: offloading non-repeating layers to GPU
llm_load_tensors: offloaded 41/41 layers to GPU
llm_load_tensors:      ROCm0 buffer size = 10472.08 MiB
llm_load_tensors:        CPU buffer size =   202.50 MiB
...

 スーパーコンピュータ富岳について教えてください。
A. 富岳は、理化学研究所と富士通が共同開発したスーパーコンピュータです。「京」の後継機として、2020年6月に本格稼働を開始しました。富岳は、ピーク性能で毎秒43.5テラビットの計算性能を誇り、国内のスーパーコンピュータの中で最も高性能なマシンです。また、「京」が1ラックあたり1.536kWだったのに対し、富岳は1ラックあたり最大720kWと大容量の電力を供給できるため
llama_print_timings:        load time =   19096.74 ms
llama_print_timings:      sample time =       3.61 ms /   128 runs   (    0.03 ms per token, 35417.82 tokens per second)
llama_print_timings: prompt eval time =    2931.48 ms /    13 tokens (  225.50 ms per token,     4.43 tokens per second)
llama_print_timings:        eval time =   35159.27 ms /   127 runs   (  276.84 ms per token,     3.61 tokens per second)
llama_print_timings:       total time =   38111.22 ms /   140 tokens
Log end

こんな感じで返答してくれるはずです。

なおFugaku-LLM-13B-instruct-0325b-q5_k_m.ggufの大きさは10.42 GiBなので、Geforce 3060 12GBなら全部GPUにオフロードできますが、8GBの外付けGPUでは全部GPUにオフロードすることはできないはずです。

Fugaku-LLM-13B-instruct-0325b.ggufの方も実行してみましょう。

$ bin/main -m Fugaku-LLM-13B-instruct-0325b.gguf -n 128 -p "スーパーコンピュータ富岳について教えてください。" -ngl 10000
...
llm_load_print_meta: model type       = ?B
llm_load_print_meta: model ftype      = F16
llm_load_print_meta: model params     = 13.42 B
llm_load_print_meta: model size       = 25.01 GiB (16.00 BPW) 
...
llm_load_tensors: ggml ctx size =    0.46 MiB
llm_load_tensors: offloading 40 repeating layers to GPU
llm_load_tensors: offloading non-repeating layers to GPU
llm_load_tensors: offloaded 41/41 layers to GPU
llm_load_tensors:      ROCm0 buffer size = 25100.07 MiB
llm_load_tensors:        CPU buffer size =   506.25 MiB
...

 スーパーコンピュータ富岳について教えてください。
A: Fujitsu's HPC Cluster、略して「富岳」は、スーパーコンピュータのファミリーであり、ベクトルプロセッサ、メモリ、およびインターコネクトをベースとしたシステムです。理化学研究所と富士通によって開発され、2020年6月に本格稼動した。
Fujitsu's HPC Cluster "FuGaku "は、ベクトル処理ユニット（VPU）、メモリ、およびインターコネクトをベースとしたスーパーコンピュータのファミリーである。理化学研究所と富士通によって開発され、2020年6月に本格稼動
llama_print_timings:        load time =  506844.15 ms
llama_print_timings:      sample time =       3.91 ms /   128 runs   (    0.03 ms per token, 32719.84 tokens per second)
llama_print_timings: prompt eval time =    1881.28 ms /    13 tokens (  144.71 ms per token,     6.91 tokens per second)
llama_print_timings:        eval time =  112732.85 ms /   127 runs   (  887.66 ms per token,     1.13 tokens per second)
llama_print_timings:       total time =  114636.73 ms /   140 tokens
Log end

遅いですが、25.01 GiBのモデル全部をGPUにオフロードして実行できます。

32GBのメインメモリがあればAPUなら25.01 GiBのモデルであっても全部GPUへのオフロードが可能です。Geforce 3060 12GBでは全部GPUにオフロードすることはできません。

ROCm対応koboldcpp

koboldcppはGUIの設定画面とWebインターフェイスで対話的にLLMを利用できるソフトです。llamacppを内蔵しています(最新バージョンではないようですが)。LLMを活用するならばllama.cppを直接使うより使いやすいと思います。

環境設定

python3.8なので、python3.8をインストールして、venvで環境を作り、有効にします。どこでもいいのですが~/以下で作業します。

$ cd ~
$ sudo dnf install python3.8
$ python3.8 -m venv koboldcpp
$ cd koboldcpp
$ . bin/activate

ROCm対応koboldcppの取得とビルド

ROCm対応版koboldcppからファイルをダウンロードします。KoboldCPP-v1.66.1.yr1-ROCmで検証しているので、リリースページから落として展開してください。

$ wget https://github.com/YellowRoseCx/koboldcpp-rocm/archive/refs/tags/v1.66.1.yr1-ROCm.tar.gz
$ tar xvfz v1.66.1.yr1-ROCm.tar.gz
$ cd koboldcpp-rocm-1.66.1.yr1-ROCm/

pipで必要なpythonモジュールをインストールします。

$ pip install -r requirements.txt

MakefileがFedora 40では32bitバイナリ用の/usr/libをLDFLAGSに設定しているので修正します。このパッチをkoboldcpp.patchとして~/koboldcpp/koboldcpp-rocm-1.66.1.yr1-ROCmにダウンロードし、次のコマンドで当ててください。

$ patch -p1 < koboldcpp.patch

LLAMA_HIPBLAS=1を指定してビルドしてください。

$ make LLAMA_HIPBLAS=1 -j

後述するベンチマークでは比較用にVULKAN、CLBLAST(OpenCL)、OPENBLAS(CPU)も有効して次のようにビルドしています。

$ make LLAMA_VULKAN=1 LLAMA_OPENBLAS=1 LLAMA_CLBLAST=1 LLAMA_HIPBLAS=1 -j

force-host-alloction-APUのビルド

そのままでは標準ではVRAMに512MBしか割り当てられていないので、hipMallocをhipHostMallocに置き換えるforce-host-alloction-APUを利用してHIP利用時にメインメモリからメモリを確保するようにします。gitで取得してビルドします。

$ git clone https://github.com/segurac/force-host-alloction-APU
$ cd force-host-alloction-APU/
$ hipcc forcegttalloc.c -o libforcegttalloc.so  -shared -fPIC

正しくビルドできたらlibforcegttalloc.soが生成されます。

元のディレクトリに戻ってください。

$ cd ..

libforcegttalloc.soをLD_PRELOADに指定してkoboldcpp.pyを実行します。

$ LD_PRELOAD=./force-host-alloction-APU/libforcegttalloc.so python3 koboldcpp.py

こんなランチャーが表示されるはずです。

Model FilesからModelにFugaku-LLM-13B-instruct-0325b-q5_k_m.ggufを指定し、Quick LaunchのPresetsをUse hipBLAS(ROCm)にし、GPU Layersを10000にして、下のLaunchを押してください。ブラウザが起動し
127.0.0.1:5001が開かれ、koboldcppの対話画面が表示されるはずです。

上のScenariosからNew Instructを選び、下の入力欄に聞きたい内容を入力してEnterを押してください。

次の文章を入力したときの実行例です。
-『スーパーコンピュータ富岳について教えてください。』
-『もっと詳しく教えてください。』

koboldcppのベンチマーク

koboldcppを再起動して、ランチャーのHardwareからベンチマークを動かすことができます。Fugaku-LLM-13B-instruct-0325b-q5_k_m.ggufに関してUse hipBLAS(ROCm)GPU Layersを0にしてベンチマークを動かした結果が次です。

======
Timestamp: 2024-05-31 08:20:22.163267+00:00
Backend: koboldcpp_hipblas.so
Layers: 0
Model: Fugaku-LLM-13B-instruct-0325b-q5_k_m
MaxCtx: 2048
GenAmount: 100
-----
ProcessingTime: 116.30s
ProcessingSpeed: 16.75T/s
GenerationTime: 29.99s
GenerationSpeed: 3.34T/s
TotalTime: 146.29s
Output: 
-----
===

ProcessingTimeが入力文を処理するのにかかった時間(秒)で、ProcessingSpeedが入力文を処理する速度（トークン/秒)です。ProcessingSpeedが大きいほど高性能です。GenerationTimeが回答文を生成するのにかかった時間(秒)で、GenerationSpeedが回答文を生成する速度（トークン/秒)です。GenerationSpeedが大きいほど高性能です。

Fedora 40+Ryzen 5600G+32GBでのhipBLAS(GPU)、Vulkan(GPU)、OpenBLAS(CPU)に加え、比較用に別マシンのWindows 11+Ryzen 5600+32GB+Geforce 3060でのcuBLAS(GPU)、CLBlast(GPU)、OpenBLASのベンチマーク結果をまとめてみました。Windowsの方はkoboldcppのビルド済みバイナリを利用しています。

なお、オフロードの上限にする数値がバラバラなのは単なるミスです。

Fugaku-LLM-13B-instruct-0325b-q5_k_m.gguf

Ryzen 5600G

	ProcessingTime(秒)	ProcessingSpeed（トークン/秒)	GenerationTime(秒)	GenerationSpeed（トークン/秒)
hipBLAS(オフロード2000)	112.22	17.36	47.89	2.09
hipBLAS(オフロード2000、UseQuantMat未指定)	120.66	16.15	45.48	2.20
hipBLAS(オフロード10000)	108.4	17.96	44.50	2.25
hipBLAS(オフロード20)	114.27	17.05	32.53	3.07
hipBLAS(オフロード0)	116.30	16.75	29.99	3.34
OpenBLAS	439.92	4.43	30.25	3.31
CLBLAST	動かない
Vulkan	85.72	22.73	30.24	3.31

VulkanがProcessingSpeedもGenerationSpeedも速いです。hipBLASはGPUオフロードするとProcessingSpeedは上がるけどGenerationSpeedが下がるという傾向があります。OpenBLASはProcessingSpeedが大幅に遅いので何かビルドミスしているかもしれません。いまいちすっきりしない結果です。

Geforce 3060 + Ryzen 5600

	ProcessingTime(秒)	ProcessingSpeed（トークン/秒)	GenerationTime(秒)	GenerationSpeed（トークン/秒)
cuBLAS(オフロード20000)	20.33	95.84	15.33	6.52
cuBLAS(オフロード0)	10.27	189.66	32.94	3.04
OpenBLAS	137.49	14.17	33.41	2.99
CLBLAST(オフロード20000)	29.79	65.40	21.83	4.58
CLBLAST(オフロード22)	30.00	64.94	28.33	3.53
Vulkan	動かない

cudaを利用するcuBLASが速いですが、Ryzen 5600Gと同様にGPUオフロードするとProcessingSpeedは上がるけどGenerationSpeedが下がるという傾向がありました。OpenCLを利用するCLBLASはcuBLASより明らかに遅いです。

Fugaku-LLM-13B-instruct-0325b.gguf

Ryzen 5600G

	ProcessingTime(秒)	ProcessingSpeed（トークン/秒)	GenerationTime(秒)	GenerationSpeed（トークン/秒)
hipBLAS(オフロード200000)	137.06	14.21	132.54	0.75
hipBLAS(オフロード22)	152.47	12.78	1214.47	0.08
hipBLAS(オフロード0)	117.67	16.55	219.96	0.45
OpenBLAS	437.75	4.45	63.56	1.57
CLBLAST	動かない
Vulkan	動かない

hipBLASは全部オフロードするとGenerationSpeedが速く、全部しないとProcessingSpeedが速い、中途半端にオフロードするとGenerationSpeedが大幅に遅くなりました。OpenBLASが一番GenerationSpeedが速いです。

よくわからない結果です。

Geforce 3060 + Ryzen 5600

	ProcessingTime(秒)	ProcessingSpeed（トークン/秒)	GenerationTime(秒)	GenerationSpeed（トークン/秒)
cuBLAS(オフロード20000)	動かない
cuBLAS(オフロード9)	9.50	205.03	62.85	1.59
OpenBLAS	142.34	13.69	69.26	1.44

cuBLASでオフロード可能な範囲でGPUオフロードした方が速いです。

ベンチマークの考察

Geforce 3060は基本的にはcuBLASを使って可能な限りGPUオフロードした方が速くなります。

Ryzen 5600Gは謎です。中途半端な性能のGPUなせいか、ROCmに正式に対応していないからかわかりませんが、どれもなんともいいがたい結果です。Linux上でのVulkanには正式に対応しているはずので、下手にHIPを利用するよりVulkanを利用するほうがよいのかもしれません。

ただしVulkanの場合おそらくVRAMの容量512M+GTTの容量15725M=16GBまでしか利用できない制限があると思います。VRAMとGTTの容量は

$ sudo dmesg | grep amdgpu
...
[    4.072847] [drm] amdgpu: 512M of VRAM memory ready
[    4.072849] [drm] amdgpu: 15725M of GTT memory ready.
...

などとすれば確認できます(2024/6/2追記)。

まとめ

以上Ryzen 5600GでGPUを利用したLLMの利用について一通り調べました。

Ryzen 5600GではLLMをあまり有効活用できるとはいえないですが、より高性能なAPUが出てきた際に今回調べた内容が役立つかもしれないので、参考にどうぞ。

関連記事

https://qiita.com/asfdrwe/items/758fa4b9dc3aa15ba7e1
hipでのメモリ関係の調査
https://github.com/XiongjieDai/GPU-Benchmarks-on-LLM-Inference
NVIDIAとApple SiliconでのLLMの実行速度の比較

更新履歴

2024/6/2
Vulkanに関する補足と関連記事の追加
2024/6/1
公開