MH:W | GPU Particle - モンスターハンター:ワールドにおけるGPU Particleの実装

MH:W | GPU Particle モンスターハンター:ワールドにおけるGPU Particleの実装 株式会社カプコン 米山哲平 1

講演時との差異 • 一部不要なスライドの削除 – ノートはなにかのヒントになるように残しておきます • 一部アニメーションの削除 – Slide Share等でそのまま見えるように • コード例等を追加 – 講演時は目が泳ぐので非表示にしていました • GCC中/Twitterの質問の回答を追加 – アイテム周りの具体的な構造 – SRV/UAVの使い分けによる最適化について 2

講演内容 • GPU Particle システムの詳細 – コンピュートシェーダの処理フロー • コードを交えつつデータマップを中心に図解 3

対象者 • GPU Particle 及び GPGPU に興味のある方 – プロ・アマチュア・学生を問いません • 複雑な数式は無し • HLSL(Compute Shader)の基礎知識があると読みやすい • はじめての人にも実装のヒントになるはず！ 4

登壇者紹介 • 米山 哲平 • 株式会社カプコン – 技術研究開発部 技術開発室 – 2013年入社 • 「モンスターハンター:ワールド」 エフェクトモジュールのメインプログラマ – 自社開発エンジン「RE ENGINE」のエフェクトモ ジュールにも携わった – 入社以前は個人でゲームエンジンを開発し、Microsoft Imagine Cup 2013 日本代表として選出 5

アジェンダ • 前知識 – GPU Particleとは？ – MH:WのEffectとは？ • 実装の前に – 実装背景/要求/方針 • 実装詳細 – 全体概要 – データ構造の詳細 – 各GPU処理の詳細 • TIPS – 可読性に関する話 – デバッグツール – 高速化案 6

前知識① GPU PARTICLE 7

GPU Particle とは？ GPUを使って大量のエフェクトパーティクルを処理するシステム 8

GPUを使う利点 一般的なCPUは8スレッド前後 GPUでは1000スレッド以上で動作可能 (Stream Processor数換算) • 超並列 – 今やGPUは描画以外も計算できるように – ただし制限や癖がある 9

代表例1 • 導蟲で使用 – 光点の一粒一粒が個別に動作 – 導蟲単体で最大30,000パーティクル発生 10

代表例2 • GPU Particleを一番使用しているモンスター 11

代表例2 • 約100エミッター • 最大約50,000パーティクル – 1.00～1.50msで動作 12

前知識② EFFECT 13

MH:Wでのエフェクト 専用のエディターで作成 14

エフェクトの構成 • エフェクト エフェクト エミッター1 – エミッター1 エミッター2 • Velocityアイテム • Lifeアイテム – エミッター2 • Velocityアイテム • Scale Animation アイテム アイテム 15

アイテムとパーティクル Scale Animation Rotate Animation Velocity • パーティクルはアイテムに従って動かす – パーティクルはエミッター毎に複数発生 16

アイテムの具体的な構造 D&Dで追加 • 必要な機能のアイテムを追加 • アイテムごとにパラメータを設定 17

複数のエミッター 火の粉エミッター 煙エミッター 炎エミッター 18

要求/実装方針 実装の前に 19

実装背景 • シェーダーコードでシステム構築 – GPU上で動作するのでデバッグが難しい – クラス等も使用できないので拡張が難しい 20

実装背景 • 期間は（物量に対して）かなり短い – エフェクトシステムはフルスクラッチ • 必要な機能が揃ってない – GPU Particle, 雷表現, レーザー表現… – エフェクトアセットは並行して作成 21

実装背景 • CPUに余裕がない – Dispatchによる描画コマンドの増大を防ぎたい • 処理を気軽にGPUに投げたい – アーティストにどんどん利用してほしい 22

方針 • 可読性を第一に！ – 保守コストを抑える – 高い拡張性で先の新規実装を簡単に • エミッター毎にDispatchしない！ – CPUに空きがないので描画コマンドを少なくしたい – GPUの空き時間もできるだけ抑えたい • Interlock(Atomic)命令を0に！ – GPUの並列性を阻害したくない 23

実装詳細① 全体概要 24

大まかな流れ CPUで処理 CPUで処理 GPUで処理 エフェクト生成 フレーム開始 フレーム開始 エミッター生成 エミッター更新 全エミッター 一括更新 エミッター用 データ領域を確保 アイテム更新 パーティクル用 データ領域を確保 エミッター用 データ領域の更新 パーティクル 新規追加 全パーティクル 一括更新 全パーティクル 一括ソート エミッター 管理情報更新 エフェクトの描画 25

エフェクト生成時の流れ CPUで処理 エフェクト生成 Emitter Header Emitter Table エミッター生成 エミッター用 データ領域を確保 パーティクル用 データ領域を確保 0 0 3 1 Not Use 4 2 Not Use 5 3 4 Tableに登録 5 Header領域のインデックスに相当 ※ Emitter Table等のデータ構造については後述します 26

エフェクト生成時の流れ CPUで処理 エフェクト生成 Emitter Header Emitter Binary Particle Binary 0 エミッター生成 エミッター用 データ領域を確保 1 Not Use 2 Not Use 3 4 5 パーティクル用 データ領域を確保 Head Size Emitter Binary領域を 末尾から確保 27

エフェクト生成時の流れ CPUで処理 エフェクト生成 Emitter Binary Emitter Header Particle Binary 0 エミッター生成 エミッター用 データ領域を確保 1 Not Use 2 Not Use 3 4 5 パーティクル用 データ領域を確保 Head Size Particle Binary領域を 末尾から確保 28

CPUで毎フレーム行う処理 CPUで処理 フレーム開始 エミッター更新 アイテム更新 エミッター用 データ領域の更新 • 親の移動値の反映 等 ※パーティクルの座標ではない事に注意 29

CPUで毎フレーム行う処理 CPUで処理 フレーム開始 エミッター更新 アイテム更新 エミッター用 データ領域の更新 • アイテム単位の更新 – パーティクルを出す数の決定 – 外部パラメーターの反映 – タイムライン制御 …等 30

CPUで毎フレーム行う処理 CPUで処理 フレーム開始 Emitter Header Emitter Binary 0 エミッター更新 1 Not Use 2 Not Use 3 アイテム更新 5 エミッター用 データ領域の更新 Velocity アイテム 4 • • Head Size Headerから領域を参照 アイテムを順番に書き込む Scale Anim アイテム Rotate Anim アイテム 31

GPUで毎フレーム行う処理 GPUで処理 フレーム開始 全エミッター 一括更新 Emitter Table パーティクル 新規追加 0 全パーティクル 一括更新 4 全パーティクル 一括ソート 3 5 テーブルに登録されているエミッターを更新 エミッター 管理情報更新 エフェクトの描画 32

GPUで毎フレーム行う処理 GPUで処理 Particle Header フレーム開始 全エミッター 一括更新 前のフレームのパーティクル パーティクル 新規追加 全パーティクル 一括更新 全パーティクル 一括ソート エミッター 管理情報更新 Emt0 Emt3 Emt5 エフェクトの描画 新規パーティクルを末尾に追加 33

GPUで毎フレーム行う処理 GPUで処理 Particle Header フレーム開始 全エミッター 一括更新 パーティクル 新規追加 全パーティクル 一括更新 全パーティクル 一括ソート エミッター 管理情報更新 エフェクトの描画 Header上の全パーティクルを更新 34

GPUで毎フレーム行う処理 GPUで処理 フレーム開始 全エミッター 一括更新 Alive Particle Header Particle Header Dead パーティクル 新規追加 全パーティクル 一括更新 Sort 全パーティクル 一括ソート エミッター 管理情報更新 エフェクトの描画 生きてるパーティクル と 死んだパーティクル をソートで分離する 35

GPUで毎フレーム行う処理 GPUで処理 Particle Header フレーム開始 Emitter Range 全エミッター 一括更新 0 1 Not Use パーティクル 新規追加 2 Not Use 全パーティクル 一括更新 4 3 5 全パーティクル 一括ソート エミッター 管理情報更新 エフェクトの描画 各エミッター配下のパーティクルの分布を収集 Interlock命令を使用せずにパーティクル数をカウントするための重要な処理 36

GPUで毎フレーム行う処理 GPUで処理 フレーム開始 全エミッター 一括更新 パーティクル 新規追加 全パーティクル 一括更新 全パーティクル 一括ソート エミッター 管理情報更新 エフェクトの描画 描画 37

実装詳細② データ構造編 38

データ構造 • バッファは起動初期に大きなメモリ空間を確保 – エフェクトを再生するたびにメモリを確保し直さない • メモリの断片化防止 • そもそもメモリアロケーションはコストが高い 39

データ構造 • バッファの種類は大きく分けて2種類 – 毎フレームCPUからGPUへ書き込まれるバッファ • GPUからは読み取りのみ可能 – GPU内でのみ読み書きされるバッファ • CPUからは読み取りができない 40

CPUからGPUへ書き込まれるバッファ Emitter Table Emitter Headers Emitter Binary 41

GPU内でのみ読み書きされるバッファ 1/3 Particle Binary Particle Headers Particle Index List 42

GPU内でのみ読み書きされるバッファ Particle Headers Emitter Range 2/3 43

GPU内でのみ読み書きされるバッファ Dispatch Indirect Args(X/Y/Z) Particle Num Draw Indirect Args VertexNum / InstanceCount / VertexOffset / InstanceOffset Prev Particle Num 3/3 44

データの取り出し方 for (uint threadIdx = 0; threadIdx < EmitterNum; ++threadIdx) { uint emitterId = EmitterTable[threadIdx]; uint emtHead = EmitterHeaders[emitterId].EmtBinHead; void* emtPtr = &EmitterBinary[emtHead]; //emtPtrをキャストしてアクセス！ //auto* item = static_cast<CommonItem*>(emtPtr); //emtPtr += sizeof(CommonItem); } • スレッドごとにエミッターデータへアクセスする方法 • 簡易擬似コードでの書き方 – 実際にはGPUでアクセスするのでHLSLで書く必要がある – ここではイメージを掴んでもらう程度で・・・ 45

データの取り出し方 for (uint threadIdx = 0; threadIdx < TotalParticleNum; ++threadIdx) { auto& particle = ParticleHeaders[threadIdx]; auto& emitter = EmitterHeaders[particle.EmitterId]; uint ptHead = particle .Index * emitter.PtSize; void* emtPtr = &EmitterBinary[emitter.EmtBinHead]; void* ptPtr = &ParticleBinary[emitter.PtBinHead + ptHead]; /* emtPtrとptPtrでアクセス！ */ } • 「パーティクルのバイナリ領域」と「エミッターのバイナリ領域」 にアクセスする方法 46

メモ • データを脳内にマップするのが一番大変 – 慣れないうちは自分で理解しやすいマップ をエクセルとかで作るのが良い 開発時のメモ画像 47

まとめ • Headerを使ってBinaryにアクセスする • データを脳内でマップするのが一番大変 – 慣れない間はエクセル等で早見表を作ろう 48

実装詳細③ GPU処理フロー 49

GPU上の処理の流れ 初回初期化処理(2 pass) Clear System Clear Particles 毎フレーム実行される処理(10 pass) Begin Update Fill Unused Index Buffer Bitonicsort Spawn Particles Range Particles Initialize Particles Terminate Particles Update Particle Build Emitter Draw Args Build Primitive ※概要一覧表は付録に用意してあります 50

GPU上の処理の流れ 初回初期化処理(2 pass) Clear System Clear Particles 毎フレーム実行される処理(10 pass) Begin Update Fill Unused Index Buffer Bitonicsort Spawn Particles Range Particles Initialize Particles Terminate Particles Update Particle Build Emitter Draw Args Build Primitive 51

Clear System [compute shader] 1/12 [numthreads(1, 1, 1)] void ClearSystem() { IndirectArgs.Store(IA_PARTICLE_COUNTER, 0); IndirectArgs.Store(IA_PREV_PARTICLE_COUNTER, 0); } • アプリケーション起動後、初回のみ実行 • システム全体で共有される値を初期化 • Dispatch(1, 1, 1)で実行 52

Clear Particles [compute shader] 2/12 • アプリケーション起動後、初回のみ実行 • 後述するパーティクル毎のキーとなる情報を初期化 – INVALID_TAG(0xffffffff or -1) 53

Clear Particles
[compute shader] 2/12

[numthreads(PARTICLE_PER_THREAD, 1, 1)]
void ClearParticles(uint3 id : SV_DispatchThreadID)
{
if (TotalParticleMax <= id.x) return;
ParticleHeader[id.x].tag = INVALID_TAG;
ParticleHeader[id.x].depth = 0.0f;
}
•

Dispatch(ceil(TotalParticleMax / (float)PARTICLE_PER_THREAD), 1, 1) で実行
–
–

つまりパーティクルの数分スレッドを実行
参考としてPARTICLE_PER_THREADはMH:Wでは256を設定しています

54

GPU上の処理の流れ 初回初期化処理(2 pass) Clear System Clear Particles 毎フレーム実行される処理(10 pass) Begin Update Fill Unused Index Bitonicsort Spawn Particles Range Particles Initialize Particles Terminate Particles Update Particle Build Emitter Draw Args Build Primitive 55

GPU上の処理の流れ 初回初期化処理(2 pass) Clear System Clear Particles 毎フレーム実行される処理(10 pass) Begin Update Fill Unused Index Bitonicsort Spawn Particles Range Particles Initialize Particles Terminate Particles Update Particle Build Emitter Draw Args Build Primitive 56

Begin Update [compute shader] 3/12 Fill Unused Index Spawn Particles [numthreads(1, 1, 1)] void BeginUpdate() { uint prePtCount = IndirectArgs.Load(IA_PARTICLE_COUNTER); uint curPtCount = min(prePtCount + TotalSpawnCount, TotalParticleMax); IndirectArgs.Store(IA_PARTICLE_COUNTER, curPtCount); IndirectArgs.Store(IA_PREV_PARTICLE_COUNTER, prePtCount); } • • 「前のフレームの総パーティクル数」＋「発生”予定”のパーティクル数」 →「”仮”の総パーティクル数」 Dispatch(1, 1, 1)で実行 57

Begin Update Fill Unused Index [compute shader] 4/12 Spawn Particles Initialize Particles • 新しく割り当てられたParticle Index List領域 をINVALID_TAG(-1 or 0xffffffff)で埋める – 新規エミッターが発生した際に実行 – 割当領域は必ず末尾から取られる • そのため範囲は連続しているので一括して埋められる 58

Begin Update Fill Unused Index Spawn Particles [compute shader] 5/12 Initialize Particles Update Particle 13番目から4要素 • 新しく発生するパーティクル数(Spawn数)を決定する – 新規パーティクルはParticle Headersの末尾に割り当てる • ついでにEmitter Rangeも0で初期化 59

Fill Unused Index Spawn Particles Initialize Particles [compute shader] 6/12 Update Particle Bitonicsort 使用中Particle Index Listの数と等しい • Spawnしたパーティクルにインデックスを割り当てる – Particle Index Listから未使用インデックスを取得 – Dispatch(iEmitterCount, 1, 1) で実行 • 1スレッドで複数パーティクルを処理する 60

Fill Unused Index Spawn Particles Initialize Particles [compute shader] 6/12 Update Particle Bitonicsort 後ろからこの個数分割り当てる －1 の場合は現在のParticleNumの連番を割り当てる • Spawnしたパーティクルにインデックスを割り当てる – Particle Index Listから未使用インデックスを取得 – Dispatch(iEmitterCount, 1, 1) で実行 • 1スレッドで複数パーティクルを処理する 61

Fill Unused Index Spawn Particles Initialize Particles [compute shader] 6/12 Update Particle Bitonicsort for (uint i = threadIdx.x; i < emitterData.spawnNum; i += PARTICLE_PER_THREAD) { // TODO. } – 1スレッドで複数パーティクル分を処理する場合のコード例 • 256スレッドで処理する場合は0～255, 256～511, 512～767…という感じ – 最悪想定で1フレームに大量のパーティクルが出る可能性はある 62

Spawn Particles Initialize Particles Update Particle Scale Animation [compute shader] 7/12 Rotate Animation Bitonicsort Range Particles Velocity • パーティクルをアイテムに従って動かしていく – 1スレッドで1パーティクル処理する 63

Spawn Particles Initialize Particles Update Particle Bitonicsort [compute shader] 7/12 Range Particles • 更新はParticle Headerを参照する – 生存している最小限のパーティクルだけ効率的に処理される 64

Spawn Particles Initialize Particles Update Particle [compute shader] 7/12 Bitonicsort Range Particles 次のページで説明 • Uber Shader的に全アイテム実行 • GPUで動かすコードとしては一見やばそう・・・ 65

Update Instance Item 次のページで説明 • 殆どのアイテムで同じ書き方 – 読み出して、更新して、書き出す という一連の流れ – 実際にアイテムが設定されてない限り実行されない • stream****関数は後述 66

Update Item • アイテムごとの処理のみ • BufferのLoad/Storeとは完全に分離 67

Initialize Particles Update Particle Bitonic Sort [compute shader] 8/12 Range Particles Terminate Particles • Particle Headersを並び替え – Alive, EmitterID, Depthの順番で並び替える • Indexは見ない 68

Initialize Particles Update Particle Bitonic Sort Z-Sortなし 半透明描画 前後関係がおかしくて気持ち悪い Range Particles [compute shader] 8/12 Z-Sortあり 半透明描画 ボリューム感が有る Terminate Particles Z-Sortなし加算 違和感はない • 同時に深度ソートも行う – MH:Wでは半透明エフェクトが多いので深度ソートが必要 – 加算エフェクトだけに制限した場合は深度ソートが不要になる 69

Bitonic Sort [compute shader] 8/12 • Bitonicsortを選択した理由 – アルゴリズムを理解していたからというだけ – 全体をソートできれば何でも良い – Bitonicsortの注意点 • パーティクル数に応じてDispatch回数が増加 – 再生中の全エフェクトのMax Particle Numの合計から計算 » 実際の生存パーティクル数がCPU上から見れないため • ここだけ同期命令の一つである GroupMemoryBarrierWithGroupSyncを使用 – 並列度を下げるものではない 70

Sort Key Key1 Key2 Alive 1bit パーティクルの生存フラグ Emitter ID 13bit パーティクルの親エミッターの番号 Particle Index 18bit EmitterIDと合わせれば直接パーティクルにアクセスできるように Depth 32bit カメラからの距離(float) • Particle Headersは1要素64bit – Shared Memoryに一旦読み込む場合に丁度いいサイズ – AliveとEmitterIDをまとめてuint値として比較出来る • 比較ではParticle Indexを使用しない 71

比較関数 • Aliveは実際にはInvalidフラグな点に注意(trueの場合に無効) – 最上位ビットなのでKey1をuintとして見た場合に非常に大きな値になる • • Key1をキーとして昇順に並び替えると無効化されたパーティクルは後ろに移動する 結果的にINVALID_TAG(0xffffff)を使用すると勝手に末尾へ移動する 72

Update Particle Bitonicsort Range Particles [compute shader] 9/12 Terminate Particles Build Emitter Draw Args これを構築するプロセス • • ソートしたParticle Headersから各エミッター配下のパーティクルの境界を調べる これを利用すればカウントアップ等をする必要がなくなりInterlock命令が排除できる – (End – Head)で簡単にパーティクル数が求まる 73

Update Particle Bitonicsort Range Particles Particle Headers [compute shader] 9/12 Terminate Particles Build Emitter Draw Args • 各スレッドで2要素ずつ取り出して比較するだけ – 隣接している要素のEmitter IDが違うならデータの境界ということ 74

Update Particle Bitonicsort Range Particles Particle Headers [compute shader] 9/12 Terminate Particles Build Emitter Draw Args • それだけだと抜けが出るので一要素ずらして再比較 – 各スレッドで2+1要素を比較するだけで全パーティクルの分布が求まる 75

Bitonicsort Range Particles Terminate Particles [compute shader] 10/12 Build Emitter Draw Args Build Primitive ※ここの値は実際には不定 返却済み領域だけ正しくなる • 死んだパーティクルのIndexを返却する 76

Range Particles Terminate Particles Build Emitter Draw Args [compute shader] 11/12 Build Primitive End - Alive • Emitter Rangeから最終的なパーティクル数を求めEmitter Dataを更新 77

Range Particles Terminate Particles Build Emitter Draw Args [compute shader] 11/12 Build Primitive VertexNum / InstanceCount / VertexOffset / InstanceOffset • Draw Indirect用の引数を構築 – MH:WではEmitter毎にDraw Callを行っている – x6しているのはビルボードの場合の頂点数(正確にはインデックス数) 78

Terminate Particles Build Primitive Build Emitter Draw Args • • • • • • • [compute shader] 12/12 Color Size Rotation Scale Anim Rot Anim UV Sequence etc… Particle Build… Billboard • パーティクル情報から頂点を構築 – 1スレッドで1パーティクル分実行 – ビルボード及びリボン形状をサポート 79

GPU上の処理の流れ 初回初期化処理(2 pass) Clear System Clear Particles 毎フレーム実行される処理(10 pass) Begin Update Fill Unused Index Buffer Bitonicsort Spawn Particles Range Particles Initialize Particles Terminate Particles Update Particle Build Emitter Draw Args Build Primitive 80

まとめ • 各Dispatchは基本的に単純なことしかしない • Emitter RangeによってInterlockで実装しがちな カウントアップ処理を賄う 81

実装詳細④ データーの整頓 82

Emitterの追加 buffer buffer Emt0 Emt0 Emt1 Empty Empty Emt1 New! Emt2 New! Emt3 • Emitterが追加されるとデータ領域が末尾に割り当てられる 83

Emitterの削除 buffer Emt0 Delete! Emt1 buffer Emt0 Empty Emt2 Emt3 Emt2 Emt3 • 削除する際は領域を開放する – しかしバッファーは最初に確保した大きなメモリ空間だけなの で断片化してメモリが枯渇する 84

Bufferの詰め直し buffer0 Emt0 Delete! Emt1 Copy buffer1 Emt0 Emt2 Emt3 Copy Emt2 Emt3 Empty • エミッターが削除された場合はデータ領域を詰め直す必要がある – ガベージコレクション的なメモリ再配置処理 • このためにダブルバッファにする必要がある 85

Bufferの詰め直し対象 • Particle Index List – Copy Sub Resourceを何回か呼び出す必要がある • Particle Binary – こちらはUpdate Particle時の書き出し先に新しい領域を指定すれば良い • 簡単に済ませたいならCopy Sub Resourceでの実装でも良い 86

Emitter削除時の注意 Particle Headers Emitter Headers Alive EmittertID ParticletID EmtID EmtBinHead TRUE 0 0 0 0 TRUE 1 1 1 ??? TRUE 1 2 2 1088 TRUE 1 3 3 2348 TRUE 1 4 4 ??? TRUE 2 0 5 ??? TRUE 2 1 TRUE 3 0 不正な領域にアクセス！ • Emitterの削除要求が発生したら1フレーム待機して 必ず配下のパーティクルを削除するように！ – Particle Headersに配下のパーティクルが残っている可能性 87

実装結果 88

処理速度 • 約100エミッター / 最大50,000パーティクル – 1.00～1.50msで動作 – 実際に生存しているパーティクルは平均25,000パーティクル • ソート処理によって必要最小限の更新が行われる 89

処理速度 Graphics Pipe Compute Pipe • Async ComputeによってZ-Prepassの裏で動作 – 処理負荷はほぼ隠蔽される 90

GPU Particleの実装期間 • 約1ヶ月 – シェーダー及び管理システムを含む • エディターはCPU実装版と互換性をとりそのまま使用 • この時点で基本的なアイテム/機能まで実装 – 読みやすいコードを心がけたのでサクサク書けた • コードの可読性に関しては後述 91

TIPS 92

If分岐に関して • If分岐使いまくりだけど大丈夫？ – とくにUpdate Particle 93

If分岐に関して 実行マスク スレッド開始 If(threadIdx % 2) ↓↓↓↓↓↓↓↓↓↓ ↓ ↓ Process A Process B END ↓ ↓ ↓ ↓ ↓ ↓ ↓ ↓ 両方実行したときと 同じサイクル数がかかる ↓↓↓↓↓↓↓↓↓↓ GPUはif分岐に弱い 94

If分岐に関して • Particle Headerを参照しているため 隣接スレッドは基本同じエミッターの配下 – つまりアイテム構成も同じ • バラバラな分岐はあまり起こらない 95

コードの可読性 • できるだけコードを統一したい • 冗長で規則性のあるコードをまとめたい • コピペミスなどを事前に防ぎたい 96

コードの可読性 struct PtVelocity { float3 velocity; float gravityAccume; float gravity; }; • 構造体をByteAddressBufferで管理する場合で問題に – 冗長 – Load/Storeする場合はほぼ同じコードを書く必要がある – コピペミスしやすい 97

Stream Function struct PtVelocity { float3 velocity; float gravityAccume; float gravity; }; void streamPt_Velocity3D(uint rw, inout uint offset, inout PtVelocity value) { streamPt(rw, offset, value.velocity); // float3 streamPt(rw, offset, value.gravityAccume); // float streamPt(rw, offset, value.gravity); // float } • HLSL上であらゆる型の Load/Store を統一的 に記述するためのヘルパー関数 98

Stream Function void streamPt(uint rw, inout uint offset, inout float3 value) { if (rw == readOnly) stream_read(srv, offset, value); else if (rw == read) stream_read(uav, offset, value); else if (rw == write) stream_write(uav, offset, value); else stream_none(offset, value); } • 中身はread/writeによって分岐してるだけ 99

これ大丈夫なの？ void streamPt(uint rw, inout uint offset, inout float3 value) { if (rw == readOnly) stream_read(srv, offset, value); else if (rw == read) stream_read(uav, offset, value); else if (rw == write) stream_write(uav, offset, value); else stream_none(offset, value); } • SRV/UAV , Load/Store命令が混在してるが・・・ • rwに”定数”(read/readOnly/write/none)を指定すればOK – コンパイル時最適化で競合する命令は全部消える！ – If分岐命令も全部消える！ 100

Stream Function void stream_read(ByteAddressBuffer buffer, inout uint offset, out float value) { value = asfloat(buffer.Load (offset)); offset += 4 * 1; } void stream_read(ByteAddressBuffer buffer, inout uint offset, out float2 value) { value = asfloat(buffer.Load2(offset)); offset += 4 * 2; } void stream_read(ByteAddressBuffer buffer, inout uint offset, out float3 value) { value = asfloat(buffer.Load3(offset)); offset += 4 * 3; } void stream_read(ByteAddressBuffer buffer, inout uint offset, out float4 value) { value = asfloat(buffer.Load4(offset)); offset += 4 * 4; } • あとは基本形のオーバーロードを用意 101

Stream Function PtVelocity velocity; // 読み取り streamPt_Velocity3D(read, read_offset, velocity); // 更新 updateVelocity3D(velocity); // 書き込み streamPt_Velocity3D(write, write_offset, velocity); • read/write定数を切り替えるだけ – “定数”というところが大事 102

Stream Function STREAM_R_RW(streamPt, ParticleBinary, ParticleBinaryUAV); • プリプロセッサマクロにしてあるので任意の名前で 全基本型のオーバーロードメソッドが用意できる 103

Sizeofに応用 uint sizeof_Velocity3D() { uint size = 0; PtVelocity temp = (PtVelocity)0; streamPt_Velocity3D(none, size, temp); return size; } • Noneを指定すればノーコストで構造体のsizeofもできる – HLSLだと基本型しか出来ないけどこれで解決！ – noneだとoffsetを進めるだけなのでコンパイル時最適化で最終的な定数しか残らない！ – C++側とHLSL側で定義している構造体サイズが一致してるかのチェックも簡単に！ 104

コンパイラを信じよう！ • 分かりやすくシンプルなコードを – 一つのローカル変数を共有するのはほぼ無意味 • 細かい粒度で関数化しよう – HLSLでは殆どの場合インライン展開される – Load/Store命令は連続していればまとめられる • Stream Functionを信じよう！ 105

デバッグツール GPU Particle Profiler 106

デバッグツール • エミッター毎の使用領域を色分け 107

デバッグツール • 大量のパーティクルを出してるエミッターを調査 – 処理負荷調査に役立つ 108

デバッグツール • データーの不整合も監視 – システムのバグ調査に非常に役立つ 109

デバッグツール • できるだけ最初に用意して！ – バグ調査でかなりの時間を奪われる • 不整合が起こるとドライバハングにつながる – デバッグツールで予兆を監視しよう！ • 実装後すぐにバグを見つけることが出来た 110

高速化案 • 加算ブレンドエフェクトに制限する – メリット • ソート処理を無くす事が可能 – デメリット • 内部では全パーティクルを常に動作させる必要がある 111

高速化案 • エミッター単位でソート – メリット • パーティクル数が一定以下なら1Dispatchでソート可能 • 加算ブレンドエフェクトの場合はソートを省ける – デメリット • Dispatch数が増大 • パーティクル数に気をつける必要がある 112

高速化案 • Update Particleを複数に分ける – メリット • 並列実行性能を上げやすい – 細かい粒度で動作させられる – 使用レジスタを抑えやすい – デメリット • 細かすぎるとオーバーヘッドが増えて逆効果 113

高速化案 • アイテムの値をシェーダーに焼き込む – メリット • Emitter Binaryから読み込む必要がないので爆速 – 最適化もかかるので10倍位早くなるかも・・・ – デメリット • • • • Dispatch数が増大 シェーダーファイルが増大 アイテムのパラメータを動的に変更できない 焼き込みシステムを構築する必要がある 114

最後に • 簡単ではないが超難しいわけではない – システムの規模が大きいぶんやることは多い – データ構造を脳内でちゃんとマップできれば行けるはず！ • みんなもGPU Particle実装しよう！ – かっこいいエフェクトを作ろう！ – パーティクル管理のシステムは応用が効くはず！ 115

おしまい 116

今回の内容 • 前知識 – GPU Particleとは？ – MH:WのEffectとは？ • 実装の前に – 実装背景/要求/方針 • 実装詳細 – 全体概要 – データ構造の詳細 – 各GPU処理の詳細 Clear System Clear Particles Begin Update Fill Unused Index Buffer Bitonicsort Spawn Particles Range Particles Initialize Particles Terminate Particles Update Particle Build Emitter Draw Args Build Primitive • TIPS – コードの可読性について – デバッグツール – 高速化案 117

付録 118

ClearSystem 説明 システム変数を全て初期化します。 アプリケーション起動後、初回にのみ実行。 ClearParticles Particle Headerの全領域をInvalid Tagで埋める。 アプリケーション起動後、初回にのみ実行。 BeginUpdate 更新を始めるためにシステム変数を準備。毎フレーム実行。 FillUnusedIndexBuffer このフレームでエミッターが追加された場合に実行。 新しく割り当てられたParticle Index Listを-1で埋める。 SpawnParticles 有効な全エミッターをなめてエミッター毎にデータを構築。 InitializeParticles 新しく発生したパーティクルのヘッダ情報を初期化。 このフレームで生成されるパーティクルのヘッダ情報は Particle Headersの末尾(未使用領域の先頭)に追加される。 UpdateParticle Core Updateと呼んでいる部分。 設定されているアイテムによってパーティクルを動かします。 具体的にはVelocityアイテムで加速させたり、 Lifeアイテムによって半透明にしたりします。 さらにLifeが0になったり特定条件でパーティクルが死んだ場合、 ヘッダのAliveフラグをFalseにします。 Bitonicsort Particle Headersを並び替えます。 パーティクル数によって複数回Dispatchされます。 RangeParticles ソート済みのParticle Headersから各エミッター毎に、 最終的な生存パーティクル数、このフレームで死んだパーティクル数、 さらに全エミッター合計の生存パーティクル数を算出します。 もちろんInterlock無し。 TerminateParticles 死んだパーティクルのIndexを、 Particle Index Listに返却します。 BuildEmitterDrawArgs エミッター毎に実行されるDrawIndirectの引数を構築 BuildPrimitive パーティクルからビルボードとかリボンの頂点を作成 PHASE 119

GCC中の質問 Q.SRV/UAVの使い分けによる最適化について A.実際には使い分けています。 – 講演中は混乱しないように省いてました 120

GCC中の質問 Q.SRV/UAVの使い分けによる最適化について • STREAM FUNCTIONマクロでも対応済み – STREAM_R_RW(name, srv, uav) • SRV/UAVを併用して読み書き – STREAM_RW(name, uav) • UAVのみで読み書き – STREAM_R(name, uav) • SRVのみで読み取りのみ 121

GCC中の質問 Q.SRV/UAVの使い分けによる最適化について void streamPt(uint rw, inout uint offset, inout float3 value) { if (rw == readOnly) stream_read(srv, offset, value); else if (rw == read) stream_read(uav, offset, value); else if (rw == write) stream_write(uav, offset, value); else stream_none(offset, value); } • STREAM_R_RW(name, srv, uav)を使用した場合 – readOnlyでsrv読み出し – readでuav読み出し – writeでuav書き込み • Dispatch中のバッファの使用用途で使い分けてます 122

GCC中の質問 Q.一番好きなエフェクトは？ A.テオのスーパーノヴァ 123