画像生成AI Stable Diffusion
画像生成AIのStable Diffusionの統合環境であるStable Diffusion WebUI(AUTOMATIC1111版)を、AMDのLinux用GPGPU実行環境ROCmを使ってRyzen 5600GのGPUで動かす方法の解説です。
以前の記事ではBIOSで手動でVRAMを割り当てていました。
ノートパソコンではBIOSでの手動のVRAM割り当てができない場合が多いようなので、あまりスマートなやり方ではないのですが、VRAMから割り当てるhipMallocをメインメモリから割り当てるhipHostMallocに動的に置き換えるforce-host-alloction-APUを利用することで、BIOSの変更なしにRyzen APUのGPUでのStable Diffusion WebUIを実現しました。
動作確認した機器構成
| 種類 | 内容 |
|---|---|
| CPU | Ryzen 5600G |
| マザーボード | ASRock B450M-HDV (BIOS 4.70) |
| メモリ | CFD W4U3200CS-16G (16GB×2=32GB) |
| OS | Fedora 40 (Native実行、wsl2では動かない) |
準備
Linuxを起動しVRAMとGTTの割り当てを確認
Linuxを起動してください。
dmesgコマンドでカーネルログを確認してください。
$ sudo dmesg | grep amdgpu
...
[ 4.585655] [drm] amdgpu: 512M of VRAM memory ready
[ 4.585656] [drm] amdgpu: 15727M of GTT memory ready.
...
GPUのVRAMに割り当てられたメモリに512M、GTT(Graphics Translation Table)に15727M割り当てられています。
GTTは標準ではCPUのメモリサイズから自動で計算されて割り当てられますが、カーネルモジュールamdgpuのgttsizeオプションにより手動で指定して割り当てることもできます。必要ならamdgpuの設定を参考にgrub等で設定してください。
※ 実験した結果GTTの容量は無関係のようです(2024/4/26追記)
必要なパッケージのインストール
dnfコマンドでrocminfo等のrocm関係のパッケージをインストールしてください。
$ sudo dnf install rocminfo rocm-smi rocm-device-libs
Stable Diffusion WebUIはPython 3.10.6推奨で、Fedora40の標準はPython 3.12なので、python3.10パッケージをインストールします。
$ sudo dnf install python3.10
gitも使うので入っていなければインストールしてください。
$ sudo dnf install git
/dev/kfdのパーミッション設定
ROCmでは/dev/kfd経由でGPUにアクセスします。一般ユーザからアクセス可能にするためにusermodコマンドでROCmを使用するユーザ(今回はasfdrwe)をrenderグループに入れてください。
$ sudo usermod -aG render $USER
group、group-ファイルを見てrenderグループに入っているか確認してください。
$ grep $USER /etc/group*
/etc/group:wheel:x:10:asfdrwe
/etc/group:video:x:39:asfdrwe
/etc/group:render:x:105:asfdrwe
/etc/group:asfdrwe:x:1000:
/etc/group-:wheel:x:10:asfdrwe
/etc/group-:video:x:39:asfdrwe
/etc/group-:asfdrwe:x:1000:
rocminfoでのROCmの動作確認
rocminfoコマンドでROCmの実行環境の確認ができます。正常に動作するか確認してください。
/dev/kfdのパーミッションが不適切だと実行に失敗します。失敗するならls -l /dev/kfdでrenderグループが読み書きできるか確認してsudo chmod g+rw /dev/kfd等でパーミッションを設定してください。
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
*******
Name: AMD Ryzen 5 5600G with Radeon Graphics
Uuid: CPU-XX
Marketing Name: AMD Ryzen 5 5600G with Radeon Graphics
Vendor Name: CPU
Feature: None specified
Profile: FULL_PROFILE
Float Round Mode: NEAR
Max Queue Number: 0(0x0)
Queue Min Size: 0(0x0)
Queue Max Size: 0(0x0)
Queue Type: MULTI
Node: 0
Device Type: CPU
Cache Info:
L1: 32768(0x8000) KB
Chip ID: 0(0x0)
ASIC Revision: 0(0x0)
Cacheline Size: 64(0x40)
Max Clock Freq. (MHz): 4464
BDFID: 0
Internal Node ID: 0
Compute Unit: 12
SIMDs per CU: 0
Shader Engines: 0
Shader Arrs. per Eng.: 0
WatchPts on Addr. Ranges:1
Features: None
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: FINE GRAINED
Size: 32210200(0x1eb7d18) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
Pool 2
Segment: GLOBAL; FLAGS: KERNARG, FINE GRAINED
Size: 32210200(0x1eb7d18) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
Pool 3
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 32210200(0x1eb7d18) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
ISA Info:
*******
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
Cache Info:
L1: 16(0x10) KB
L2: 1024(0x400) KB
Chip ID: 5688(0x1638)
ASIC Revision: 0(0x0)
Cacheline Size: 64(0x40)
Max Clock Freq. (MHz): 1900
BDFID: 2304
Internal Node ID: 1
Compute Unit: 7
SIMDs per CU: 4
Shader Engines: 1
Shader Arrs. per Eng.: 1
WatchPts on Addr. Ranges:4
Features: KERNEL_DISPATCH
Fast F16 Operation: TRUE
Wavefront Size: 64(0x40)
Workgroup Max Size: 1024(0x400)
Workgroup Max Size per Dimension:
x 1024(0x400)
y 1024(0x400)
z 1024(0x400)
Max Waves Per CU: 40(0x28)
Max Work-item Per CU: 2560(0xa00)
Grid Max Size: 4294967295(0xffffffff)
Grid Max Size per Dimension:
x 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
Max fbarriers/Workgrp: 32
Packet Processor uCode:: 471
SDMA engine uCode:: 40
IOMMU Support:: None
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 524288(0x80000) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: FALSE
Pool 2
Segment: GLOBAL; FLAGS:
Size: 524288(0x80000) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: FALSE
Pool 3
Segment: GROUP
Size: 64(0x40) KB
Allocatable: FALSE
Alloc Granule: 0KB
Alloc Alignment: 0KB
Accessible by all: FALSE
ISA Info:
ISA 1
Name: amdgcn-amd-amdhsa--gfx90c:xnack-
Machine Models: HSA_MACHINE_MODEL_LARGE
Profiles: HSA_PROFILE_BASE
Default Rounding Mode: NEAR
Default Rounding Mode: NEAR
Fast f16: TRUE
Workgroup Max Size: 1024(0x400)
Workgroup Max Size per Dimension:
x 1024(0x400)
y 1024(0x400)
z 1024(0x400)
Grid Max Size: 4294967295(0xffffffff)
Grid Max Size per Dimension:
x 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
FBarrier Max Size: 32
*** Done ***
Stable Diffusion WebUIのインストール
venvでPythonの仮想環境を構築
以下~/で作業します。作業場所はどこでも問題ないはずです。
まず~/sdにpython3.10でvenvで仮想環境を構築します。
$ cd ~
$ python3.10 -m venv sd
ROCm対応pytorchをインストール
~/sdに移動し仮想環境を有効にします。
$ cd sd
$ . bin/activate
先にpipやwheelのアップデートをしておきます。
$ python3 -m pip install --upgrade pip wheel
必要なパッケージをインストールします。ROCm対応pytorchを入れます。
$ pip3 install --upgrade diffusers transformers scipy ftfy
$ pip3 install --pre torch torchvision torchaudio --index-url https://download.pytorch.org/whl/nightly/rocm6.0
ROCm対応pytorchの動作確認
環境変数HSA_OVERRIDE_GFX_VERSIONとHSA_ENABLE_SDMAを設定してください。gfx900(vega10)として動作させるためにHSA_OVERRIDE_GFX_VERSIONを9.0.0に設定しています(LLVMによるとgfx90cとgfx900の命令セットは同じ)。force-host-alloction-APUのREADME.mdよりROCm6.0ではHSA_ENABLE_SDMA=0にしないと不安定になるようです。
$ export HSA_OVERRIDE_GFX_VERSION=9.0.0
$ export HSA_ENABLE_SDMA=0
ROCmでpytorchが正常に動作するかtest-rocm.pyをダウンロードして実行します。
$ wget https://gist.githubusercontent.com/damico/484f7b0a148a0c5f707054cf9c0a0533/raw/43c317bfbde626d9112d44462d815613194988e0/test-rocm.py
$ python3 test-rocm.py
正常に動作している場合次のように表示されるはずです。
Checking ROCM support...
GOOD: ROCM devices found: 2
Checking PyTorch...
GOOD: PyTorch is working fine.
Checking user groups...
GOOD: The user asfdrwe is in RENDER and VIDEO groups.
GOOD: PyTorch ROCM support found.
Testing PyTorch ROCM support...
Everything fine! You can run PyTorch code inside of:
---> AMD Ryzen 5 5600G with Radeon Graphics
---> gfx900
BADが出た場合どこかおかしいので一旦PCを再起動して作業をやりなおしてください。
force-host-alloction-APUのビルド
ビルドに必要なパッケージをインストールします(2024/4/24追記)。
$ sudo dnf install clang hipcc lld lld-devel compiler-rt
force-host-alloction-APUを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が生成されるはずです。
force-host-alloction-APUのテストとMIOpenのビルド
この状態でforce-host-alloction-APUのテストツールを実行しても、Fedora 40では次のようなMIOpenのエラーが出るはずです。
$ LD_PRELOAD=./libforcegttalloc.so python test/diffusion.py 'astronaut sitting on the moon' output.jpg cuda
...
MIOpen(HIP): Error [Compile] 'hiprtcCompileProgram(prog.get(), c_options.size(), c_options.data())' naive_conv.cpp: HIPRTC_ERROR_COMPILATION (6)
MIOpen(HIP): Error [BuildHip] HIPRTC status = HIPRTC_ERROR_COMPILATION (6), source file: naive_conv.cpp
MIOpen(HIP): Warning [BuildHip] error: unknown argument: '-fno-offload-uniform-block'
1 error generated when compiling for gfx900.
MIOpen Error: /long_pathname_so_that_rpms_can_package_the_debug_info/src/extlibs/MLOpen/src/hipoc/hipoc_program.cpp:304: Code object build failed. Source: naive_conv.cpp
...
MIOpenのIssueに-fno-offload-uniform-blockに関するワークアラウンドがあり、Fedora Rawhideに取り込まれているようなので、Fedora Rawhideのmiopen-6.0.2-4.fc41.src.rpmを取得しビルドします。
$ wget https://dl.fedoraproject.org/pub/fedora/linux/development/rawhide/Everything/source/tree/Packages/m/miopen-6.0.2-4.fc41.src.rpm
(バージョンアップしていてmiopen-6.0.2-4.fc41.src.rpmがなくなっている場合は、 https://dl.fedoraproject.org/pub/fedora/linux/development/rawhide/Everything/source/tree/Packages/m/ からmiopen-6.0.*-*.fc41.src.rpmを探してください)
miopenのビルドに必要なものが足らないはずなので、必要なパッケージをインストールしてください。
$ sudo dnf install rpm-build bzip2-devel boost-devel cmake eigen3-devel fplus-devel frugally-deep-devel half-devel libzstd-devel ninja-build nlohmann_json-devel rocblas-devel rocm-cmake rocm-comgr-devel rocm-hip-devel rocm-runtime
rpmbuildコマンドでビルドします。
$ rpmbuild --rebuild miopen-6.0.2-4.fc41.src.rpm
正しくビルドできれば標準では~/rpmbuild/RPMS/x86_64/以下に
miopenのバイナリパッケージが作られるはずです。
作成したmiopenのバイナリパッケージをインストールします。
$ sudo dnf install ~/rpmbuild/RPMS/x86_64/miopen-6.0.2-4.fc40.x86_64.rpm
そのままではこのパッケージのlibMIOpen.so.1.0がうまく読み込まれないようなので、libforcegttalloc.soに加えてこのファイルもLD_PRELOADするように指定してforce-host-alloction-APUのテストツールを実行し、正しく動くか確認します。
$ LD_PRELOAD=./libforcegttalloc.so:/usr/lib64/libMIOpen.so.1.0 python3 test/diffusion.py 'astronaut sitting on the moon' output.jpg cuda
正しく動けば月面で座っている宇宙飛行士の画像(output.jpg)が生成されるはずです。
Stable Diffusion WebUIをインストール
上のディレクトリに移動し、gitコマンドでStable Diffusion WebUIをインストールし、stable-diffusion-webuiに移動します。
$ cd ..
$ git clone https://github.com/AUTOMATIC1111/stable-diffusion-webui
$ cd stable-diffusion-webui
Stable Diffusion WebUIの実行
環境変数COMMANDLINE_ARGSに--no-half --no-half-vaeに指定し、LD_PRELOADにlibforcegttalloc.soとlibMIOpen.so.1.0を指定して、Stable Diffusion WebUIをwebui.shで実行してください。
$ export COMMANDLINE_ARGS="--no-half --no-half-vae"
$ LD_PRELOAD=../force-host-alloction-APU/libforcegttalloc.so:/usr/lib64/libMIOpen.so.1.0 ./webui.sh
初回はStable Diffusion WebUIに必要なパッケージのインストールがされるので時間がかかります。次回以降はもっと早く起動します。正常に実行できている場合 http://127.0.0.1:7860 にアクセスするよう表示され、ブラウザが自動的に起動して http://127.0.0.1:7860 を開きます。
...
Running on local URL: http://127.0.0.1:7860
To create a public link, set `share=True` in `launch()`.
...
SettingのUser InterfaceのLive previewsのShow live previews of the created imageを外し、それ以外デフォルトで実行して59.4秒で画像が生成されました。export COMMANDLINE_ARGS="--use-cpu all --no-half --no-half-vae --skip-torch-cuda-test"を指定してCPUのみで生成した場合は1分43.9秒でした。BIOSを変更してVRAMを割り当てた場合と同等の速度です。