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

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

amdgpu 21.40 で RDNA2 で HIP(CUDA っぽいの)が動くようになりました(experimental) - Qiita

Mesa 21.3 RADV で AMD GPU で Vukan レイトレするメモ - Qiita

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

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

RadeonOpenCompute/ROCm: ROCm - Open Source Platform for HPC and Ultrascale GPU Computing

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

κ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

OpenCL

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

κ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

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

Vulkan TutorialのRust版をやってみた〜目次〜 - Qiita

κeen

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

Rustで始めるVulkan Raytracing

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

Ubuntu – jammy の mesa-vulkan-drivers パッケージに関する詳細

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

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

Jammy Jellyfish Release Schedule - Release - Ubuntu Community Hub

κeen

GPU Gemsが無料公開されてるのを知ったのでVulkan入門が済んだらこれ読みながらGSLSをVulkanに移植すると練習になりそう。

κ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

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

κeen

学習用に：
GLSLからSPIR-Vへの変換器
https://github.com/KhronosGroup/glslang

Porting WebGL shaders to WebGPU - Ashley's blog

GLSLからWGSLへの変換ができるツールがあれば上手くいかないときにそれ参照すればいいのかなって思ったけど見当らなかった

κeen

ところでWgpuのチュートリアルみつけた。初手はこれかな。

Introduction | Learn Wgpu

κeen

因みに今読んでるのはこれ
[連載]やってみれば超簡単！ WebGL と GLSL で始める、はじめてのシェーダコーディング（１） - Qiita

κeen

nagaでGLSL→WGSL変換できそう

https://github.com/gfx-rs/naga

κeen

wgpuが（多分naga経由で）GLSLも受け取れそうなのでドライバ側はwgpu一本槍で勉強できそう

ShaderSource in wgpu - Rust

κeen

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

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

この手順でいこう。

κeen

ちょうどbevy0.6でwgpuに全乗りになった。WGSLサポートとかも入る。
https://bevyengine.org/news/bevy-0-6/
ライフゲームのサンプルなど

κeen

mesa22.0でRADVが速くなったらしい。2月にstableになるって言ってるからもしかしたらUbuntu 22.04に入るかも

κ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

openclのカーネルに直接spirv使えるみたいだからrust-gpu経由でRustでカーネル書けるかも。libclcがどうなるかは知らないけど。

κeen

AMDGPUの追加のリソースみつけた。

アセンブリ
AMD ISA Documentation - GPUOpen

AMD公式のトレーニング
ROCm™ Learning Center - AMD

κ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

あー、あれだ。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

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

κeen

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

wgpuのチュートリアルをやってシェーダ言語を動かせるようになる ← イマココ
- cf https://zenn.dev/link/comments/6f90ac2d599f49
GLSLで書かれたシェーダをWGSLに移植しながらWGSLの仕様を一通り理解する
- 題材
- 仕様
  - https://www.khronos.org/registry/OpenGL/specs/gl/GLSLangSpec.1.40.pdf
  - https://gpuweb.github.io/gpuweb/wgsl/
自分で何かWGSLを使って書いてみる
- マンデルブロ集合とかの模様
- 魚眼レンズカメラ
- グレースケール変換やエッジ検出とかの古典的CV
- GEMMとかのGPGPU的使い方
(未定)VulkanチュートリアルをやってVulkanのAPIを触ってみる
- wgpuのAPIを極める方向にいくのかも

κeen

3でやること候補

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

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

κeen

shadertoyというサイトを知ったので↑に追記しておく。

https://www.shadertoy.com

あとGPUSoundという分野があるというのもshadertoyで知った。

GPUSound入門 - Qiita

κeen

この実例がコード短くて良かった

思考過程が書かれててイメージしやすかった
GLSLで音楽（はじめに） - Qiita

κeen

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

κeen

GLSLのシェーダーステージをよく分かってない。どのステージで何ができるとかが違うっぽいけど分からない。多分ロードマップの2番で勉強することになると信じてる。

κeen

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

κ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

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

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

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

κeen

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

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

コードはここなんだけど、読み解いた挙動と合致しないのがあってよく分からない。もしかしたらデプロイされてるのとGitHubのコードで差異があるかも

κeen

featmouseブランチだとそれっぽいので、多分これの開発途中だったものがデプロイされてるのかな。

マウスの扱いはまだ不要だからmasterブランチの方の移植頑張ろう。

κeen

サンプルが少し古くて（GL 2.0）、wgpu（GL 3.0）だとちょいちょい動かないところがある。

attribute → in
gl_FragColor → out 変数を定義して使う

3D-Graphics/GLSL(shader言語) - 神嶋教授の独白

OpenGL(GLSL)のvarying,attribute,in,outについて - げぇむぷろぐらみんぐ

κeen

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

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

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

Uniformのドキュメントにany typeとある。

しかし実際はエラーになる

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

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

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

κeen

まずGLSL Editor開いたときのデフォルトのシェーダをWGSLに移植した。

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

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

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

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

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

κeen

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

κeen

あんまり関係ないけどOSSのゲームエンジンがVulkanサポートしたらしい。

Godot 4.0 Alpha 1 Released With Vulkan Renderer & Other Shiny Features - Phoronix

godbotはRustバインディングがあるらしいということだけ知ってる。

κeen

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

κeen

Vulkan 1.3が出てる。まだ内容みても何が変わったかは分からないけど理解できるようになるといいな

Vulkan 1.3 and Roadmap 2022 - The Khronos Group Inc

κeen

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

κeen

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

wgsl-analyzer/wgsl-analyzer: a language server implementation for the WGSL shading language

κeen

wgpuのチュートリアルに載ってる文法とWGSLの公式にある文法が違うんだけどって思ったらwgpuの方が古いか何かのよう。master版だとWGSLの公式の文法に書き換えられてる。

κeen

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

κeen

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

κeen

上げといた。まあないよりマシでしょクオリティで。
https://github.com/KeenS/wgsl-mode.el

κeen

VulkanとOpenCLの違いについてここの回答が簡潔で分かりやすかった。

κeen

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

κeen

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

κ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

なんかもやっとしたのでポインタ使って 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

このチュートリアルを

wgpuで動くように整えてGLSLとそれをWGSLに翻訳したもの両方で全部動かしてみた。

κeen

このあとGPU Gemsに進む予定。2月中には終わらないだろうから適当なところで区切りつける。

κeen

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

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

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

κeen

今AMDのドキュメントにあったOpenCLプログラミングガイドを読んでる。ちょっとGPU Gemsと入れ違いになるかも

κeen

ここ。mesaのopenclだとclcのコンパイルに失敗したのでrocm使ってる。
https://github.com/KeenS/opencl-learner

κeen

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

κeen

シェーダーステージについてはここが分かりやすかった。

WGPUだとひとまとめになってるけど、モデルの座標のローカル→ワールド変換とワールド→カメラ変換で分割されるのね。あとテッセレーションとかステンシルとかも分かりやすかった。

κeen

今ならこの記事の内容が分かる

κeen

AMDのOpenCL解説があるならNVIDIAはどうだろうと思ったらサンプルが結構いっぱいあった。有り難い。
https://developer.nvidia.com/opencl

κeen

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

κeen

OpenCL Cのコンパイルとリンク

κeen

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

κeen

rustのopenclバインディングも試してみた

κeen

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

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

κeen

RDNA2のパフォーマンスガイドみつけた
https://gpuopen.com/performance/

κeen

https://zenn.dev/link/comments/479134467e527f の問題、masterのnagaで修正されてた。該当PRはこれ↓

バグじゃなくて最新の仕様に追い付いてなかっただけみたい。

κ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

↑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

jammyのmesaドライバが21.3になってる。よかった。

κeen

mesa22がリリースされた。Ubuntu 22.04には入らなそうかな。dynamic_renderingとか試したいんだけどねー

κeen

jammyのmesaが22.0になってる！！

κeen

このページ探してた。OpenCLのリファレンスマニュアル。

OpenCL 2.2 Reference Pages

κeen

CL Cにprintfあったんかい

κeen

GPUを使ったエイプリルフール記事錬成したからようやくGPU Gemsに戻れる。チャプターによるけどコードがあまりなくて理論だけだから苦戦するかも。計算ももちろんのこと、可視化の部分も自分で書かないといけない。

κeen

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

κeen

GCC 12.1でOpenACCのAMD GPUサポートが改善されたっぽい

κ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ドライバもまともに使えるようになってる。

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