WGPU
wgpuは素晴らしいぞ！
この先、AIの未来があるぞ
学習方法
CGを触っている人なら実はLLMの現実的なパイプラインとGPUの動きというのは同じになる
逆にCGを触っていないとGPU設計をAIに使うのは厳しすぎる(あとrust)
パイプランはこのように対応している
(大事なところ以外削っている)
CG用パイプライン renderPass->viewport->vertexBuffer(0,1,2)->indexBuffer->bindGroup(0,1,2)
LLM用パイプライン embedLayer->seq2seq->attention
GPUの構造
GPUは早いか？
A.クロック数はおそいGPUはだいたい2.5GHz(CPUは3.5~5GHzとか)
プログラミングについて
GPUってのはそもそも、普通のプログラミングじゃなくて、物流センターの操作が一番近い
アーキテクチャ
Unified Shader Unit(統合シェーダーユニット)
昔はfragmentやvertexで専用回路が組まれていたが、今は統合されている。
それがcompute shader
昔の名残＋使い勝手のために@vertexなどの命令は残っている
この@は属性（attribute）という
blender,hudini触ったことのある人はわかるだろう
VRAM(strage buffer)
校舎の倉庫みたいなもの、遠い
巨大、GPUの９割はこれ
Shared Memory(Workgroup memory)
プロセッサの近くにある
爆速、ただし、KB程度の保存しかできない
Register(普通の変数)
最速、１変数しか見れない
CG pipeline直通
vertex
| 定義 | 型 | 説明 |
|---|---|---|
| input | ||
| @builtin(vertex_index) | u32 | 頂点番号 |
| @builtin(instance_index) | u32 | インスタンスID（メッシュのID） |
| output | ||
| @builtin(position) | vec4f | GPU上でのxyzポジション（配列だが、構造体に近く,id.xみたいにアクセスする） |
fragment
| 定義 | 型 | 説明 |
|---|---|---|
| input | ||
| @builtin(vertex_index) | vec4f | ピクセル座標,画面上のx,z+z+w(遠近補正) |
| @builtin(front_facing) | bool | 表面かどうか |
| @builtin(sample_index) | u32 | マルチサンプルのindex,MSAA |
| output | ||
| @builtin(frag_depth) | f32 | 深度地の上書き,zbufferの値を自分で操作するときに使う |
compute Stage(Compute Shader?)
| 定義 | 型 | 説明 |
|---|---|---|
| input | ||
| @builtin(global_invocation_id) | vec3u | workgroupID,x,y,zで一意に決める |
| @builtin(local_invocation_id) | vec3u | ワークグループ内での相対ID,SharedMemへのIOに使う |
| @builtin(workgroup_id) | vec3u | ワークグループの指定 |
| @builtin(num_workgroups) | vec3u | すべてのワークグループ数 |
| @builtin(local_invocation_index) | u32 | グループ内通し番号,local_invocation_idの１次元圧縮 |
computeshaderについての説明
- global_invocation_id:クラス全体での通し番号
- num_workgroups:全クラス数
- workgroup_id:何組か
- local_invocation_id:クラス内での出席番号
attribute一覧
基本
@builtin(name):gpuが提供する変数をマッピングする
@group(n):bindgroupのグループ番号を指定
@binding(n):bindgroupのリソースのバインディング番号を指定
@workgroup_size(x, y, z):compute shaderのスレッドサイズを指定
描画周り
@vertex:vertex shaderのエントリーポイントであることを示す
@fragment:fragment shaderのエントリーポイントであることを示す
@compute:compute shaderのエントリーポイントであることを示す
@location(n) 頂点バッファ、fragmentなどの入出力の先
@interpolate(type, sampling):vertex shaderからfragmentへのデータ変換方法を指定
@invariant:vertex shaderの出力が入力に対して、同じ値を生成することを保証
アライメント
@align(n):アライメントをバイト単位で指定
@size(n):構造体メンバーのサイズをバイト単位で指定する
@must_use:関数の戻り値を無視するとエラーを出す
@id(n):override変数に対して、パイプライン生成時に値を上書きする識別子を指定
@diagnostic:特定のコードブロックに対して診断メッセージのフィルタリングルールを設定
strage_shaderたちの設定
@group(0) @binding(1) var<storage, read_write> rnn_w: array<f32>;
@group(0):０番目のデータ群を選択
@binding(1):１番目の層にあるデータ群
:strage(いわゆるVRAM)にread_writeでおく。ほかにはreadだけなのもありこっちのほうが早い
array:jsとかならなじみがあるだろう
構造的には
value = GPU_Mem[Group_ID(0..4)][Binding_ID(0..999)]
GroupBindingは４つまでしか使えない
Bindingは０～９９９の１０００個
さらに、WEBGPUは同時にworkgroupが動いて値をめちゃくちゃにすることを保証しない
- group0 全体での定数、グローバル情報
- group1 今描画しているモデル、素材
- group2 今描画しているモデルの素材などの情報
- group3 予備
bindingは999まで
つまり１度に認識できるPORTは４０００個までしか配列が入らない
（配列のなかに配列を入れることで制限は回避できるー＞GPUではindexを入れてＶＲＡＭ内を多重参照することができる）
極論をいえばgroupはbindingが枯渇するまで上げる必要はない
あとは割とrustっぽくかける
余談vertexshader
@vertex
fn vs(@builtin(vertex_index) in_v: u32) -> @builtin(position) vec4f {
let pos = positions[in_v];
return vec4f(pos.x, pos.y, pos.z, 1.0);
}
fragment
@fragment
fn fs() -> @location(0) vec4f {
return vec4f(1, 0, 0, 1);
}
compute
//workinggroup_sizeを指定する必要がある
@compute @workgroup_size(1) fn computeSomething{
@builtin(global_invocation_id) vec3u {
}
}
なぜcompute shaderの返り値がvec3uなのか、これはCG周りの名残りらしい＜－悪手！
HLSLなんかより百倍簡単だろう？