BIOSでの手動割り当てなしでRyzen APUでStable Diffusion WebUIを実行する方法

# 画像生成AI Stable Diffusion
画像生成AIのStable Diffusionの統合環境である[Stable Diffusion WebUI(AUTOMATIC1111版)](https://github.com/AUTOMATIC1111/stable-diffusion-webui)を、AMDのLinux用GPGPU実行環境[ROCm](https://github.com/RadeonOpenCompute/ROCm)を使ってRyzen 5600GのGPUで動かす方法の解説です。

以前の記事ではBIOSで手動でVRAMを割り当てていました。

https://qiita.com/asfdrwe/items/1ca2d07777415f399a3b

ノートパソコンではBIOSでの手動のVRAM割り当てができない場合が多いようなので、あまりスマートなやり方ではないのですが、VRAMから割り当てるhipMallocをメインメモリから割り当てるhipHostMallocに動的に置き換える[force-host-alloction-APU](https://github.com/segurac/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)](https://dri.freedesktop.org/wiki/GART/)に15727M割り当てられています。

GTTは標準ではCPUのメモリサイズから自動で計算されて割り当てられますが、カーネルモジュールamdgpuの`gttsize`オプションにより手動で指定して割り当てることもできます。必要なら[amdgpuの設定](https://dri.freedesktop.org/docs/drm/gpu/amdgpu.html)を参考にgrub等で設定してください。

※ [実験した結果GTTの容量は無関係のようです](https://qiita.com/asfdrwe/items/758fa4b9dc3aa15ba7e1)(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の命令セットは同じ](https://llvm.org/docs/AMDGPU/AMDGPUAsmGFX900.html))。[force-host-alloction-APUのREADME.md](https://github.com/segurac/force-host-alloction-APU/blob/main/README.md#usage-reports-by-another-user)よりROCm6.0では`HSA_ENABLE_SDMA=0`にしないと不安定になるようです。

```
$ export HSA_OVERRIDE_GFX_VERSION=9.0.0
$ export HSA_ENABLE_SDMA=0
```

ROCmでpytorchが正常に動作するか[test-rocm.py](https://gist.github.com/damico/484f7b0a148a0c5f707054cf9c0a0533)をダウンロードして実行します。

```
$ 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](https://github.com/segurac/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](https://github.com/ROCm/MIOpen/issues/2734)に-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を割り当てた場合と同等の速度です。


## 参考ページ
- [HIPでのメモリ管理](https://rocblas.readthedocs.io/en/rocm-6.0.0/conceptual/gpu-memory.html)