Metal #21 [付録] コンピュートシェーダー

• Lingheng Tao
• 2024年6月8日
• 読了時間 3 分

Unity Shader Content Table Metal #22 [Appendix] SwiftUI Shader CG Paper #2 KinectFusion: Real-time 3D Reconstruction and Interaction Using a Moving Depth Camera

#GameEngine#ComputerGraphics#TechnicalArt

このノートは Metal のコンピュートシェーダーに関する部分についてです。

グラフィックスカードで計算する

並列計算

グラフィックスカードの利点は並列計算（parallelism）。スレッドセーフなループは CPU から GPU に移して計算できる。

Metal 公式ドキュメントの例：

Plain Textvoid add_arrays(const float *inA, const float *inB, 
			  float* result, int length) 
{
	for (int i = 0; i < length; i++) 
	{
		result[i] = inA[i] + inB[i];
	}
}

この例では、配列の各要素が独立して読み取られるため、順序が計算結果に影響しない。この種のループは並列計算に適しており、ループの各イテレーションを別の演算ユニットに割り当てられる。

Metal で上記を GPU コードに変換すると：

Plain Textkernel void add_arrays(device const float *inA, device const float *inB, device float *result, uint i [[thread_position_in_grid]])
{
	result[i] = inA[i] + inB[i];
}

主な変更点：

• ループ構造がなくなった。add_arrays は1回のイテレーションのみを担当し、具体的なイテレーションは異なる計算ユニットに割り当てられる。
• kernel キーワード：kernel は公開 GPU 関数をマーク。public GPU 関数のみがアプリから見える。kernel は同時に計算関数をマークし、独立スレッドで並列計算可能な関数であることを示す。
• device キーワード：device はこれらのポインタがデバイスアドレス空間に格納されることをマーク。
• [[thread_position_in_grid]]：C++ の属性構文。Metal が各スレッドに異なるインデックス値 i を計算し、この関数の実引数として渡すことを宣言。

GPU デバイスの取得

Metal#1 の初期化で紹介した概念。グラフィックスカードを指す抽象が必要。Metal では MTLDevice が GPU の抽象。

Plain Textid<MTLDevice> device = MTLCreateSystemDefaultDevice();

macOS は複数 GPU をサポートするが、Metal はそのうち1つをデフォルトデバイスとしてこの関数で返す。

Metal 関数の参照を取得

Initializer はまずこの関数をロードし、GPU 上で実行できるように準備する。アプリをコンパイルするとき、Xcode はこの関数（例：add_arrays）を Metal のデフォルト関数ライブラリに追加し、アプリに埋め込む。MTLLibrary と MTLFunction オブジェクトで、Metal の関数ライブラリとその中の関数への参照を取得できる。