MetalのGPGPUによるPerlinノイズ

Jan 25, 2016 · old
metal gpgpu swift2

GPGPUでリアルタイムにパーリンノイズを生成しテクスチャに書き込んで表示するデモ。

パーリンノイズはGPU Gems2の改良パーリンノイズで、それを元にした数種類を選択可能。デモでは3次元ノイズを生成し、z値だけを時間で加算していきノイズを変化させている。

動作イメージ

ノイズ生成デモ

@m_ike__が投稿した動画 - 2016 1月 26 5:57午前 PST

小ネタ

ソース自体はGPUGemsのサンプルをほぼほぼ書き換えただけなので、ノイズ生成のアルゴリズムなどは参考リンクの方で。。。

ノイズ生成の流れ

初期化時に順列テーブルpermBufferを生成
初回のみなのでCPU側で生成
実行時毎にノイズの模様が変化する様にランダムに並び替えを行う
毎フレームごと、ComputeShaderでノイズを生成
シェーダは各種パラメータNoiseParameterや順列テーブルなどを受け取って、生成した結果をテクスチャに書き込んで返却する
2のテクスチャを板ポリに貼って描画する

ComputeShader用の設定

MTKViewの設定

ComputeShaderを利用する場合は、MTKViewのframebufferOnlyをfalseにしておくこと。
ただし、パフォーマンスに影響を与えるので、利用しない時は触らないこと。

waitUntilCompleted()

各種コマンドの実行の終了するまで待機するメソッド。
以下のようにcommandBuffer.commit()より後に置く。
描画時のような毎フレームごとの呼び出しの時は、別途 dispatch_semaphore_waitで処理を待機するので不要だが、初期化時の1回のみ実行するような場合は、これを使って待機させる。

/* ~~ （ここでGPGPUのコマンド処理） ~~ */
commandBuffer.commit()
commandBuffer.waitUntilCompleted()

スレッド

GPGPUは大量のスレッドを並列動作させて圧倒的な処理を行うのが特徴だが、モバイルのGPUのため、スレッドグループ毎の最大のスレッド数は512までになっている(公式)。

なお、スレッド数は32の倍数にするのが良い。nVIDIAの資料(URL忘れた…)によると64−192ぐらいが一番パフォーマンス的に良いらしいがA系チップに当てはまるかは不明。

今回は以下のように、スレッドグループは1つのみだが、グループのサイズをテクスチャのサイズと同じにしている。

// スレッドグループの事前設定
threadgroupSize = MTLSize(width: TexSize, height: TexSize, depth: 1)
threadgroupCount = MTLSize(width: 1, height: 1, depth: 1)

// コマンドバッファへスレッドグループの設定
computeEncoder.dispatchThreadgroups(threadgroupSize, threadsPerThreadgroup: threadgroupCount)

ここでの注意ポイントは、スレッドの次元に合わせてシェーダ側の引数も変更が必要なこと。今回はサイズが2次元(width * height * 1)なので、シェーダ側はuint2 id [[ thread_position_in_grid ]]と、 uint2で受け取ることになる(間違えるとシェーダのコンパイルが通らない親切?仕様)。1次元ならuintで受け取れば良い。

また、グループが複数だったり1次元のバッファで2次元のグループサイズを指定している様な場合は、threads_per_gridやthread_position_in_threadgroupを追加で受け取る必要がある。詳細は公式のAttribute Qualifiers for Kernel Function Inputの項目を参照。

同じスレッドグループ内であれば、メモリの共有ができたりするのだが、そのあたりも公式のサンプルMetal N-Bodyが参考になる。

デバッグ

XcodeのCapture GPU Frameがとても便利。

これを使うと、実行中の各バッファやテクスチャのイメージの様子をキャプチャしてくれる。バッファのバイナリも見れるので、例えば構造体のアライメントが崩れたりしているのを簡単に見つけることができる。

また、プロファイル機能も充実していてどこがボトルネックがわかりやすい。

感想

Metalの発表を聞いて一番ワクワクして触りたかったのが、このGPGPUだった。これまでMacでは統合開発環境でGPGPUが組めなかった（普通のシェーダだけならUnityがあった）ので、使い慣れたXcodeで触れるのは楽しみだった。

元々ソース自体はGemsのものや参考リンクのものがあったので、移植しただけ。実際にはパラメータやノイズの加工を工夫すれば、雲になったり地形になったりしていくが、パラメータの調整はどうも苦手というか時間がものすごくかかるので、とりあえず今回はここまでで。

並列化によってどれくらい高速化するのかとか、モバイルのGPGPUでどのくらいパフォーマンスが上がるのか（もちろんバッファの転送が無いだけでも効果は大きいだろうけど）とかが気になったのだけど、それもまた別の機会に・・・

にしても、シェーダ関係は環境毎に用語が微妙に変わるのは何とかしてほしい。。。

参考リンク

日本語の解説：パーリンノイズを理解する(POSTD)
本家：Ken Perlin’s homepage
GPUGems Chapter5：Implementing Improved Perlin Noise
GPUGems2 Chapter26：Implementing Improved Perlin Noise
FragmentShaderで実装や応用例：Perlin Noise on GPU

開発環境

Xcode 7.2
iOS 9.2
iPhone 6+

ソース

こちら (iOS9 A7以降搭載機種のみ)

M.Ike