Zenn
Open81

なんかAMD GPUを使っていきたい

AMD
#
GPU
#
wgpu
#
opencl
#
wgsl
κeenκeen

大きな目標はないけどなんか触りたい。

参考になりそうなリソース:

amdgpu 21.40 で RDNA2 で HIP(CUDA っぽいの)が動くようになりました(experimental) - Qiita
https://qiita.com/syoyo/items/5e6f4cc056a94060329e

Mesa 21.3 RADV で AMD GPU で Vukan レイトレするメモ - Qiita
https://qiita.com/syoyo/items/9b162222cf7de9d6dd12

EmbarkStudios/rust-gpu: 🐉 Making Rust a first-class language and ecosystem for GPU shaders 🚧
https://github.com/EmbarkStudios/rust-gpu

gfx-rs/wgpu: Safe and portable GPU abstraction in Rust, implementing WebGPU API.
https://github.com/gfx-rs/wgpu

ROCm-Developer-Tools/HIP: HIP: C++ Heterogeneous-Compute Interface for Portability
https://github.com/ROCm-Developer-Tools/HIP

RadeonOpenCompute/ROCm: ROCm - Open Source Platform for HPC and Ultrascale GPU Computing
https://github.com/RadeonOpenCompute/ROCm

RadeonOpenCompute/ROCm-docker: Dockerfiles for the various software layers defined in the Radeon Open Compute Platform
https://github.com/RadeonOpenCompute/ROCm-docker

κeenκeen

HIP

普通にグラフィックやるのかGPGPUで計算するのか→
GPGPUの方が分かりやすいしGPGPUで→
HIP使うかー→

ROCmの環境整えるの面倒そうなのでROCm-dockerを使ったが、コンパイルしたバイナリからGPUデバイスが見えなくて動かない。恐らくrocm-dkmsが必要だがUbuntu 21.10だと色々な理由で動かないので難しそう。
（色々な理由=パッケージのMakefileにバグがあって適切にGCCのバージョンを判定できない、Ubuntu 20.04向けのパッケージなので依存のバージョンが古くて21.10では入らない、など）。

来春の22.04がLTSでAMDも22.04向けのドライバを出してくると思うのでそこまで待ち。

κeenκeen

OpenCL

多分HIPと同じ理由で動かないんじゃないかと思ってるけど、まだ試せてない。Radeon SoftwareじゃないOpenCLドライバで動きそうな気もする。要検証。

κeenκeen

グラフィック

あんまりよく分かってないけどVulkan使えばいいのかな？
Vulkanのドライバも色々あり、よく分からない。

Ubuntuパッケージは（AMDのドライバをインストールしていれば）3種類ある。

  • mesa-vulkan-drivers: RADVと言われるもの
  • vulkan-amdgpu: AMDVLKと言われるもの
  • vulkan-amdgpu-pro: amdgpu-proと言われるもの

RADV使ってると警告が出るが、↑のsyoyoさんの記事によるとRADVが一番おすすめらしい。よくわからない。

Rust的にはVulkan周りを触る選択肢は色々ある。
wgpuの他、ashみたいな素のVulkanバインディングがある。

MaikKlein/ash: Vulkan bindings for Rust
https://github.com/MaikKlein/ash

また、↑のrust-gpuではRustでGPUカーネルを書ける。VulkanとSPIR-Vの関係がよく分かってなかったけどSPIR-VがGPUカーネルのための中間言語、VulkanがGPUドライバのためのAPIってことでいいのかな？

両者Rustで扱えるので、まずはwgpuあたりで素のGPUを体験したあとにrust-gpuを使ってカーネルもRustで書くというのをやってみるとよさそう

κeenκeen

手のつけどころとしてはこの辺な気がする。
ただし、ちょいと古い部分もあるし自分でC++→Rustへの翻訳力をつけといた方がいいと思うので元のC++版をみてやって、つまったらRust版を参照する形にする方針で。

Vulkan TutorialのRust版をやってみた 〜目次〜 - Qiita
https://qiita.com/kbone/items/1f8bc5d571dcc3ddaac3

κeenκeen

リソースに記述忘れ。自分のやりたいこととしてはこのWeb Bookが同じ方向性っぽい。ただし自分のアクセスできるドライバにVK_KHR_ray_tracingがないのでこれそのままは不可能。

Rustで始めるVulkan Raytracing
https://zenn.dev/hatoo/books/52bcb9e9f7c87d

AMDVLKはまだray_tracingのサポートがないよう。↑のsyoyoさんの記事でRADVはmesa 21.3ではray_tracingのサポートが入るとしているが、次のUbuntu 22.04でもまだmesa 21.2.2のようなのでうーん。

Ubuntu – jammy の mesa-vulkan-drivers パッケージに関する詳細
https://packages.ubuntu.com/jammy/mesa-vulkan-drivers

因みにDebianはsidでは21.3.3になってるのでパッケージは存在しそう。
Debian -- sid の mesa-vulkan-drivers パッケージに関する詳細
https://packages.debian.org/sid/mesa-vulkan-drivers

Ubuntu 22.04のリリーススケジュール的には2/24がFeature Freezeで多分そこまでにUbuntuに取り込まれるかどうかって感じなんですかね。

Jammy Jellyfish Release Schedule - Release - Ubuntu Community Hub
https://discourse.ubuntu.com/t/jammy-jellyfish-release-schedule/23906

κeenκeen

そういやシェーダ言語のまとめやってなかった。

比較的高級

HLSL: 主にMSのDirectX向け。ゲームがWindowsターゲットなことが多いので割とよく見る。
GSLS: OpenGLで使うシェーダ言語。色々機能が拡張されて低レベルなこともできるらしい。WebGLでも使える。

比較的低級

WGSL: WebGPUで使うシェーダ言語。
SPIR-V: Vulkanで使うシェーダ言語。WGSLとほぼ1:1対応が取れるらしい。SPIR-Vはバイナリフォーマットでもあるし、そのテキストフォーマットでもあるっぽい。syntaxとか眺めとく。

今回の目的としてはGPUプログラミングなので比較的低級なSPIR-Vを中心にやっていきたいが、汎用性とかGPU Gemsみたいに教材がGLSLで書かれてるのを考えるとGLSLも並列して学ぶことになりそう。とはいえ（狭義）言語仕様はそこまで複雑ではなくてGPUでどう書くかとかAPIがうんぬんとかが多いと思うのでそこまで困らない予定。あとrust-gpuでRustで書くのもあるので3言語で同じGPUコード実装しながら進めることになるのかな？

κeenκeen

wgpuが（元々WebGPUのバックエンドとして開発されたのもあって）WGSLをネイティブサポートしてるからそっちの方が扱いやすいかも

κeenκeen

今やりたいこと:
学習の効率を上げるためにWGSL or GLSLを受け取って描画するプログラムを作りたい

  1. WGSLを実行するプログラムを書く
  2. GLSLを実行するプログラムを書く
    • 多分頂点を先に用意しておいてフラグメントシェーダを書くことになるのかな？
  3. ファイルを受け取って拡張子を見てWGSL/GLSLを実行するプログラムを書く
  4. ↑をファイルが変更されたらリロードするようにする

この手順でいこう。

κeenκeen

ふと気が向いてopenclのサンプルコードを動かしたら動いた。

https://github.com/kenba/opencl3

ガチャガチャやったので何が決定打なのかは分からないが、mesaのopenclをアンインストールしてrocmのopenclランタイムを入れたら動いたような気がする。

$ sudo apt remove mesa-opencl-icd:amd64
$ sudo apt autoremove
$ sudo apt install rocm-opencl-runtime
$ cargo run --example basic
results front: 301
results back: 1300
kernel execution duration (ns): 2080

因みに clinfo ではプラットフォームが1つ（Platform Vendorが Advanced Micro Devices, Inc.のみでMesaがない）で、Compiler AvailableがYes。

$ clinfo
Number of platforms                               1
  Platform Name                                   AMD Accelerated Parallel Processing
  Platform Vendor                                 Advanced Micro Devices, Inc.
  Platform Version                                OpenCL 2.2 AMD-APP (3361.0)
  Platform Profile                                FULL_PROFILE
  Platform Extensions                             cl_khr_icd cl_amd_event_callback
  Platform Extensions function suffix             AMD
  Platform Host timer resolution                  1ns
...
  Compiler Available                              Yes
...

AMDのOpenCL実装はバグが多くて不評らしいが、動かないよりは動いた方がいいのでよかった。

κeenκeen

ROCm dockerでHIPが使えないといったの、嘘だった。sudoつけたら動いた。

rocm-user@c6750dec623a:/opt/rocm/hip/samples/0_Intro/square$ sudo make
/opt/rocm/hip/bin/hipify-perl square.cu > square.cpp
/opt/rocm/hip/bin/hipcc  square.cpp -o square.out
rocm-user@c6750dec623a:/opt/rocm/hip/samples/0_Intro/square$ ./square.out
error: 'hipErrorInvalidDevice'(101) at square.cpp:61
rocm-user@c6750dec623a:/opt/rocm/hip/samples/0_Intro/square$ sudo ./square.out
info: running on device
info: allocate host mem (  7.63 MB)
info: allocate device mem (  7.63 MB)
info: copy Host2Device
info: launch 'vector_square' kernel
info: copy Device2Host
info: check result
PASSED!

状況を説明しておくと、rocm-dkmsが導入するのは恐らく/dev/kfdなど。今自分が使ってるLinuxのバージョンにはこれ相当の機能が取り込まれているので原理的には動くはず。

何故動かないかというとdkmsと取り込まれたLinuxの機能で /dev/kfd のグループが違うからっぽい。

Linuxではrenderグループに入ってる。

$ ls -l /dev/kfd
crw-rw---- 1 root render 506, 0  1月  8 00:39 /dev/kfd

一方でdockerを起動するときはvideoをマッピングするしdocker内でもvideoを用意してある（renderはない）

docker run -it \
     --device=/dev/kfd \
     --device=/dev/dri \
     --security-opt seccomp=unconfined \
    # ↓ここでvideoを追加してる
     --group-add video \
     rocm/rocm-terminal

前回試したときはこれに気付いていて手動でvideoを追加してたつもりだったが、失敗していたようだ。 ふと思い直して rocminfo を叩いたら権限エラーが出たのでsudoをつけてみたら動いたことから気付けた。

κeenκeen

あー、あれだ。renderのgidを506と思って操作してたけど109だった。 --group-add 109 したら動く。

$ cat /etc/group | grep render | cut -d: -f3
109

1000未満のgroup idは予約済みだから多分システム毎に変わりはしないはず。

--group-add 109 で済むことに気付く前にrender入れてdockerをビルドするようにしてみたのも供養。

$ git clone -b rocm4.5.2_update https://github.com/RadeonOpenCompute/ROCm-docker
$ cd ROCm-docker/rocm-terminal
$ git apply <<'EOF'
diff --git a/rocm-terminal/Dockerfile b/rocm-terminal/Dockerfile
index e9fe270..2d2ff0c 100644
--- a/rocm-terminal/Dockerfile
+++ b/rocm-terminal/Dockerfile
@@ -45,7 +45,8 @@ COPY sudo-nopasswd /etc/sudoers.d/sudo-nopasswd
 # This is meant to be used as an interactive developer container
 # Create user rocm-user as member of sudo group
 # Append /opt/rocm/bin to the system PATH variable
-RUN useradd --create-home -G sudo,video --shell /bin/bash rocm-user
+RUN groupadd -g 109 render && \
+    useradd --create-home -G sudo,video,render --shell /bin/bash rocm-user
 #    sed --in-place=.rocm-backup 's|^\(PATH=.*\)"$|\1:/opt/rocm/bin"|' /etc/environment

 USER rocm-user
EOF
$ docker build -t rocm/rocm-terminal .
κeenκeen

wgpuチュートリアルのBeginnerまで終わった。長かった。

κeenκeen

全体的なロードマップ。先にシェーダに馴染んでからホスト側のAPIを攻めたい

  1. wgpuのチュートリアルをやってシェーダ言語を動かせるようになる ← イマココ
  2. GLSLで書かれたシェーダをWGSLに移植しながらWGSLの仕様を一通り理解する
  3. 自分で何かWGSLを使って書いてみる
    • マンデルブロ集合とかの模様
    • 魚眼レンズカメラ
    • グレースケール変換やエッジ検出とかの古典的CV
    • GEMMとかのGPGPU的使い方
  4. (未定)VulkanチュートリアルをやってVulkanのAPIを触ってみる
    • wgpuのAPIを極める方向にいくのかも
κeenκeen

3でやること候補

  • マンデルブロ集合とかの模様を描く
  • 魚眼レンズカメラ
  • グレースケール変換やエッジ検出とかの古典的CV
  • GEMMとかのGPGPU的使い方

このあとOpenCLかHIPでGPGPUを触ってみるつもりなので比較できるように広い範囲の使い方を試しときたい

κeenκeen

OpenCLとかのGPGPUだけでなくグラフィック系の用途にも興味あるのはシェーダ言語覚えとくと一発芸の幅が広がるかなという思惑もあるので、サウンドも扱えるのは有り難い。こういうのはプロシージャルオーディオとか呼ばれる分野なのかな？

κeenκeen

amdgpuのドライバの21.40.2が出た。なぜかlibffi6に依存するようになって、Ubuntu 21.10にはそのパッケージは存在しないのでエラーになる。
というか20.04にもlibffi6は存在しないんだけど、これ大丈夫なのかな

κeenκeen

間違ってbionic用を使ってる疑惑がでたのでまっさらにして入れ直した

$ sudo amdgpu-install --uninstall
$ sudo apt remove amdgpu-install
$ sudo apt-add-repository --remove 'deb https://repo.radeon.com/amdgpu/21.40.2/ubuntu bionic main' 
$ sudo apt-add-repository --remove 'deb [arch=amd64] https://repo.radeon.com/rocm/apt/4.5.2 ubuntu main'

$ wget https://repo.radeon.com/amdgpu-install/21.40.2/ubuntu/focal/amdgpu-install_21.40.2.40502-1_all.deb
$ sudo apt install ./amdgpu-install_21.40.2.40502-1_all.deb
$ sudo amdgpu-install --usecase=graphics --no-dkms
κeenκeen

rust-gpuのexampleにwgpuでcompute shaderを走らせるところがある

https://github.com/EmbarkStudios/rust-gpu/blob/main/examples/runners/wgpu/src/compute.rs

compute shader 101というリポジトリで多少のWGSLでcompute shaderやるサンプルコードと説明がある。GLSLが一番成熟していて、VulkanのAPIを全部使える。WGSLはWebGPUの標準ということで将来は主流になりそうだが、まだ機能が欠けている部分がある。また、どちらのシェーダ言語も低級なのでもうちょっと高級な仕組みも考えられていて、中の人の注目はrust-gpuと。
https://github.com/googlefonts/compute-shader-101

あとなんかWebGPUのCompute Shaderの入門書いてる人もいる。
https://web.dev/gpu-compute/

κeenκeen

今、↓の連載をやるためにGLSLエディタのドライバ部分をwgpuに移植しようとしている。

[連載]やってみれば超簡単！ WebGL と GLSL で始める、はじめてのシェーダコーディング（１） - Qiita
https://qiita.com/doxas/items/b8221e92a2bfdc6fc211

コードはここなんだけど、読み解いた挙動と合致しないのがあってよく分からない。もしかしたらデプロイされてるのとGitHubのコードで差異があるかも
https://github.com/doxas/js4kintro/blob/master/editor/script.js

κeenκeen

先に進まないけどnagaのバグかもしれない

これは期待としてはバリデーションが通ってほしい

layout(set=0, binding=0) uniform float time; // time
void main() {}

Uniformのドキュメントにany typeとある。
https://www.khronos.org/opengl/wiki/Uniform_(GLSL)
しかし実際はエラーになる

cargo run simple.frag
    Finished dev [unoptimized + debuginfo] target(s) in 0.01s
     Running `target/debug/naga simple.frag`
error:
  ┌─ simple.frag:1:40
  │
1 │ layout(set=0, binding=0) uniform float time; // time
  │                                        ^^^^ naga::GlobalVariable [1]

Global variable [1] 'time' is invalid:
        Type isn't compatible with the storage class

明日覚えたたら起票してみる。

κeenκeen

ひとまず手元ではこれで凌いだ。

layout(set=0, binding=0) uniform Time { float t; } time; // time
κeenκeen

まずGLSL Editor開いたときのデフォルトのシェーダをWGSLに移植した。
https://github.com/KeenS/wgpu-wgsl-learner/commit/89cc85b1b9b230cebc5c65d8e9d11cc102fa64b5

1 mod がWGSLになかったので移植した

fn mod(d: f32, m: f32) -> f32 {
    return d - floor(d/m);
}

2 gl_FragCoord は main の引数で受け取り、gl_FragColor は main の返り値にする形になった

fn main([[builtin(position)]] FragCoord : vec4<f32>) -> [[location(0)]] vec4<f32> {

3 インクリメント演算子がなかった（多分）ので i = i + 1 を使った

グローバル変数をやめて引数と返り値を使ってるところはGLSLが古くなってるからWGSLで刷新したいといってるモチベーションの1つなのかな。

κeenκeen

ロードマップの2番に入ってる

  1. wgpuのチュートリアルをやってシェーダ言語を動かせるようになる
  2. GLSLで書かれたシェーダをWGSLに移植しながらWGSLの仕様を一通り理解する ← イマココ
  3. 自分で何かWGSLを使って書いてみる
  4. (未定)VulkanチュートリアルをやってVulkanのAPIを触ってみる

1番のできればやりたいといっていたシェーダの変更を検出して自動で再ロードするやつ、シェーダのコンパイルエラーを考えるとちゃんとエラーハンドリングが必要だったり、winitのイベントループとの協調考えないといけなかったりで少し大変。

前者は真面目にコンパイラ呼ぶことになりそう。後者は EventLoop::with_user_event() で任意のイベントを送れるようにして、そこに ShaderUpdated みたいなイベントを飛ばせるようにして、別スレッドでコンパイルしたシェーダを飛ばすことになるのかな。

κeenκeen

modの移植間違ってね？ m を掛けないといけない

κeenκeen

GLSLとWGSLの差分がないか確認するのに2つシェーダのソースを喰わせて左右に並べて表示する機能があると便利

κeenκeen

EmacsにWGSLのサポートがないのでwgsl-mode作らなきゃ

κeenκeen

wgsl-analyzerがあったから使ってみたらメモリ喰い続けて死亡した。0.3.0での現象で、0.2.0では起きないっぽい。

wgsl-analyzer/wgsl-analyzer: a language server implementation for the WGSL shading language
https://github.com/wgsl-analyzer/wgsl-analyzer

κeenκeen

master版使ってgot事なき。まだ全然移植してない段階でよかった。

κeenκeen

wgsl-mode、まあまあ様になった。理想的にはハイライトもインデントもwgsl-analyzerがやってくれるとうれしいんだけど、その場しのぎとしてこれくらい実装しとけばいいかな。型名のハイライトあたりがまだまだだけど、これやるには真面目な構文解析が必要だからあきらめよう。

κeenκeen

Vulkanだけだとできないことがありそうなのが分かったので3月あたりからOpenCLに切り替えてみる。

κeenκeen

WGSLに mod はないけど modf があることに気付いた

κeenκeen

glslだと引数を書き換える関数があるけど

void circle(vec2 p, vec2 offset, float size, vec3 color, inout vec3 i){
    float l = length(p - offset);
    if(l < size){
        i = color;
    }
}

wgslだとベースの値を受け取って関数から返すようにした方がいいんだろうか

fn circle(p: vec2<f32>, offset: vec2<f32>, size: f32, color: vec3<f32>, baseColor: vec3<f32>) -> vec3<f32> {
   let l = length(p - offset);
   if(l < size) {
      return color;
   } else {
      return baseColor;
   }
}
κeenκeen

なんかもやっとしたのでポインタ使って inout っぽいことしてみた。

fn sunrise(p: vec2<f32>, i: ptr<function, vec3<f32>>) {
   let t = time.t;
   let f  = atan2(p.y, p.x) + t;
   let fs = sin(f * 10.0);
   *i = mix(lightColor, backColor, fs);
}

呼出側ではこう呼ぶ。

   sunrise(p, &destColor);

ref と ptr があってややこしい。

κeenκeen

このscrapを分割して記事にするとしたら3つくらいになりそう？

  • 導入編: GPUを叩くAPIやその実装をいくつか紹介
  • グラフィック編: wgpuやそのチュートリアルの紹介、GLSLとWGSL、vulkanやWebGPUなど
  • GPGPU編: OpenCL、Cuda、HIPなど

どれも他人に教えられるほどは分かってないので自分が参考にした資料を紹介する感じかな。あとドライバのインストール方法とか。

κeenκeen

opencl、エラー情報の取り出しが面倒でつらい。確かにこれはC++使いたくなる。

κeenκeen

シェーダーステージについてはここが分かりやすかった。
https://yttm-work.jp/shader/shader_0002.html
WGPUだとひとまとめになってるけど、モデルの座標のローカル→ワールド変換とワールド→カメラ変換で分割されるのね。あとテッセレーションとかステンシルとかも分かりやすかった。

κeenκeen

薄々気付いてたけど入門終わったあとにすることがないんだよな。グラフィックはbevy使ってゲームでも作る？GPGPUはゼロからはじめるディープラーニングとかやってみて自分でコンピュートライブラリ作ってみようかな。

κeenκeen

amdのOpenCLのドキュメント読んだ。ちょっとWGSLのcompute shaderでopenclのコード書き直してみたくなった。

κeenκeen

WGSLのcompute shaderでOpenCLのコード書き直すのやってみてる。

https://github.com/KeenS/wgpu-compute-sample

まだ今の段階では本質的な表現力の違いとかは分からないけど、wgpuは記述が煩雑でコードが長くなる。そろそろWGSLのストレージクラスとかちゃんと勉強した方がよさそう。

κeenκeen

openclをWebGPUのcompute shaderに移植するやつやってみた。おおむねそのまま移植できるけど、作法が違ったり機能がなかったり仕様としてはあるけどwgpuで実装されてなかったりした。

  • OpenCLで前のイベントを待って次のイベントを実行するやつができない。
  • OpenCLと違ってlocal work sizeがGPU側（シェーダ言語）で設定することになってる
    • local work sizeに合わせてローカルバッファとかを設定したくなるのでまあ分からんでもない
  • local work sizeをホスト側から与えるようにもできるための仕様があるが、wgpuでは未実装
    • override変数 + workgroup_sizeにmodule constantを渡せるやつ
  • workgroupがいくつあるか取れないっぽい…？
    • num_workgroupsがそれかと思ったら全体のジョブ数だった
    • num_workgroups / local work size で計算できるので致命的ではない
    • wgpuのバグかもしれないので要調査。
κeenκeen

↑nagaでSPIR-Vにコンパイルしてspirv-disでディスアセンブルしたけどNumWorkgroupsだった。そういう仕様っぽい。

               OpDecorate %global_id BuiltIn GlobalInvocationId
               OpDecorate %local_id BuiltIn LocalInvocationId
               OpDecorate %workgroup_id BuiltIn WorkgroupId
              ; ここ
               OpDecorate %global_size BuiltIn NumWorkgroups
κeenκeen

GPU Gems読んでるけどつらい。ゲームとか作ってる訳じゃないからモチベーションが沸かないのと、GPUを使うための体系的な知識じゃなくて「〜の表現をするには〜という近似があって〜」のようなネタ帳的な内容だからコツコツ積み上げるのもつらい。
流し読みして興味ありそうな部分だけちゃんと読むのがいいかな。

κeenκeen

Ubuntu 22.04に上げてから色々あってGPU触ってなかったけど、ちょっと最近の話題上げるか。

22.04に上げてしばらくの間はamdが配ってるGPUドライバが22.04サポートしてなかったけど、最近(ROCm 5.3以降)は22.04サポートしてるっぽいから多分インストールできる。

https://docs.amd.com/category/ROCm™ v5.x

ただしどうしても依存のトラブルは絶えないのでdockerで使えるなら使った方がいい。blenderとかその他アプリケーションから使いたい場合は気合でインストールすることになる。

もう1つの話題としてはmesa 22.3から新しいOpenCLドライバ(rusticl)が追加されており、ビルドで有効になってればmesaのOpenCLドライバもまともに使えるようになってる。

https://docs.mesa3d.org/relnotes/22.3.0.html

OpenCL 3.0(1.2のスーパーセット。2.0とは互換性がない。…っぽい)準拠。GIMPやdarktableが動いたという報告もあり、恐らくこれが取り込まれるUbuntu 23.04あたりで希望が出てくるかも。もしかしたら22.10にもアップグレード降ってくるかも？