Metal関数ポインタの紹介

Metal関数ポインタの紹介

Metalは、低レベル・低オーバーヘッドのハードウェア加速グラフィックスフレームワークであり、Appに素晴らしい視覚効果をもたらすシェーダーAppプログラミングインタフェースです。関数ポインタを使用して、Metal Shading言語で書かれたシェーダーのプログラム制御性と拡張性を改善する方法を紹介します。Metalシェーダーの動的フロー制御にこの新しい機能を利用する方法について学びます。関数ポインタを使用して、レイトレーシングAppのカスタム交差機能を指定する方法について説明します。GPUパイプラインサイズとランタイムパフォーマンスのバランスをとるために、どのように関数ポインタが複数のコンパイルモデルを許可するかについて説明します。

リソース
関連ビデオ

WWDC21
- Metalのコンパイルワークフロー
- Metalのデバッグ、プロファイリング、アセット作成ツール
ダウンロード

こんにちは WWDCへようこそ “Metal関数ポインタの紹介” 私はリッチ･フォースター AppleのGPUソフトウェアエンジニアですこのセッションでは今年Metalに追加した新しい関数ポインタAPIについてお話しします基本の関数ポインタから始めましょう関数ポインタは呼び出すことができるコードを参照する力を与えますそれらはコードを記述できるようにすることでコードを拡張可能にしこれまでに見たことのないコードを呼び出すことができるようになります
典型的な例はコールバックで実行は関数ポインタで識別されるコードにジャンプしますこれによりプラグイン特殊化または通知の機能を提供できますしかしもっと多くのことができます今日はMetalで関数ポインタのサポートを公開する方法と visible関数を使用して目標を達成する方法についてお話します私たちがサポートする様々なコンパイルモデルが何を意味しそれらをいつ使用するかについて説明します次にvisible関数の表について説明し最後はパフォーマンスについて説明しますそれではMetalの関数ポインタから始めましょう
Metalに関数ポインタを追加すればそれを使用する様々な機会があることを知っていました明らかなケースはレイトレーシングで関数ポインタを使用してカスタムintersection関数を指定しますそしてこのセッションに関連するレイトレーシングのトークをご覧になることをお勧めします
レイトレーシングは関数ポインタの他の使用についてお話しする最高の場所であり今日はここから始めたいと思います従来のCornell Boxにレイトレーシングを適用するとシーンに光線を発射しそしてそれらが表面にぶつかった時交差をシェーディングする必要があります通常交差するサーフェスのマテリアルがあり次にライトに当たるまでトレーシングを継続しますその時点でライトの寄与を評価しますレイトレーシングのプロセスフローをもう一度見てみましょうまずカメラから始まりシーンに放出される光線を生成します次にシーン内のジオメトリとの交差についてそれらの光線をテストします
次に各交点の色を計算し画像を更新しますこのプロセスはシェーディングと呼ばれますシェーディングプロセスでは追加の光線を生成することもできますシーンとの交差についてそれらの光線をテストしますこのプロセスを何度も繰り返しシーンの周囲に光が跳ね返るのをシミュレートしますこのプレゼンテーションの内容を簡略化し単一の交差を実行するパストレーサーを検討することから始め次にその交差の結果をシェーディングしますこれはこれらの３つのステージの各コードを含む計算パイプラインを作成する― 単一のコンピュータカーネルになりますしかし私が注目したいのはシェーディングですこれは画像を出力する前の最後のステップであり関数ポインタを使用する様々な機会を提供します
レイトレーシングカーネル内ではシェーディングは最後の方で行います交差ができたら一致するマテリアルを見つけそのマテリアルのシェーディングを実行しますこの例ではマテリアルとライティングコードはすべてファイルの他の場所にある shade関数に含まれています shade関数をさらに詳しく見てみましょう
shade関数にはいくつかのステップがありますこれは単純なパストレーサーであるためサーフェス上の交差ポイントでライトから即座にライティング計算されます次にマテリアルを使用して交差ポイントにライティングを適用します
フローに関してはライティングとマテリアルを別々のステージと見なすことができますしかし複数のタイプのライトと複数のタイプのマテリアルがあります
ライティングを別々のタイプのライトに分割することができますこれにはライトタイプごとに異なるコードが必要になります
マテリアルのコードはライトのコードよりもさらに多様である可能性があります Metalの新しいvisible関数タイプを使用してライティングおよびmaterial関数の操作します
Visibleは頂点フラグメントカーネルなどの新しい関数修飾属性です関数定義でvisible属性を使用できます visible属性を使用する時は Metal APIから関数を操作することを宣言しています APIを使用してデバイスクエリでこの操作を実行できるかどうかを確認します visible関数を使用するとコードを参照可能なフレキシブルオブジェクトと見なすことができますこの場合エリアライトのコードを考えてみましょうそのコードはパイプラインを表すカーネル外に存在できるオブジェクトですコードは別のMetalファイルまたは別のMetalライブラリに含まれます
lighting関数をvisibleとして宣言するには関数の定義の前にvisible属性を追加しますこうすることで関数内のコードを表す Metal関数オブジェクトを作成できます
次のステップはMetalシェーダーコードをパイプラインに接続して呼び出すことができるようにすることですまずエリアライトコードを Metal関数オブジェクトでラップします次に新しいMetal関数オブジェクトをパイプラインに追加できます
area light関数のvisible属性を使用してそれをCPUのMetal関数オブジェクトにラップできます標準のメソッドを使用し visible関数の名前で Metal関数オブジェクトを作成します
次にlinkedFunctionsの一部として関数オブジェクトを計算パイプライン記述子に追加してパイプライン作成プロセスでパイプラインに追加できるようにしますすると後で関数ポインタで参照できるようなりますパイプラインのvisible関数の追加について話しているので利用可能なコンパイルモデルの選択について説明する必要があります
各ライティングおよびmaterial関数の Metal関数オブジェクトを一旦取得すればそれらをパイプラインに追加できます
パイプラインを構築する時これらの visible関数の各コピーを取得できますこれを“単一コンパイル”と呼びますパイプラインとそれが使用するすべてのvisible関数を表す― 単一オブジェクトがあるためです
先ほど見たlinkedFunctionsオブジェクトを使用してパイプラインから呼び出すことができる関数をリストします次に標準のパイプライン作成APIを使用して ComputePipelineStateを作成します
パイプラインを作成するとカーネルコードの特殊化されたバージョンとすべての関数の特殊化されたコピーの両方を含むオブジェクトが生成されますこの特殊化はリンク時最適化に似ておりカーネルコードとvisible関数をその使用方法に基づいて最適化できます
しなしながらこの最適化によりパイプラインの作成にかかる時間が長くなりパイプラインに追加する関数オブジェクトのコピーによりパイプラインオブジェクトが大きくなりますしかしパイプラインが呼び出すことができコピーの必要がない個別のオブジェクトとしてパイプライン外の関数オブジェクトを持つことができますこれは個別にコンパイルされたパイプラインの基礎です
各Metal関数オブジェクトを作成する時各関数をスタンドアロンのGPUバイナリー形式にコンパイルできるため関数コードを個別に保ち複数のパイプラインで再利用できます
関数をバイナリにコンパイルするには関数記述子を使用しますこれによりMetal関数オブジェクトの作成にオプションを追加できます関数を作成する場合このコードスニペットは nameパラメーターを持つ同じメソッドと同じ結果のMetal関数を作成しますただしfunctionDescriptorではさらに多くの構成オプションを指定できます
記述子の“オプション”プロパティを使用してバイナリのプリコンパイルを要求します関数を作成すると関数がバイナリにプリコンパイルされます
次にlinkedFunctionsオブジェクトのバイナリ関数の配列に新しい関数を提供しますこれは計算パイプラインの関数のコピーと特殊化を要求するのではなく関数のバイナリバージョンを参照していることを示していますご覧のようにプリコンパイルされた関数を特殊化したい関数と組み合わせることができますそして以前のように標準の呼び出しを使用してパイプラインを作成します
単一コンパイルと個別コンパイルの両方をご説明したのでそれぞれの利点を比較しましょう単一コンパイルでは標準のメソッドを使用して visible関数を作成します個別コンパイルではバイナリにプリコンパイルしてコンパイルバイナリを共有しパイプラインの作成中に関数のバイナリコンパイルを回避できるようにします
パイプライン記述子を構成する時は関数配列を使用しそれらを特殊化することを示し binaryFunctions配列を使用してバイナリバージョンを使用することを示します
特殊化の結果単一コンパイルの場合のパイプラインが大きくなります個別コンパイルがバイナリ関数への呼び出しを追加するパイプラインになりバイナリ関数が共有されます
関数の特殊化にはバイナリ関数の呼び出しを追加するよりも長いパイプライン作成時間が必要です先ほどお伝えしたようにこれはリンク時最適化に似ていますコンパイラはパイプライン全体を知ることができますが特殊化を適用するには追加のビルド時間が必要です
そのすべての関数の特殊化は単一コンパイルが実行時に高いパフォーマンスを発揮することを意味します個別にコンパイルされたパイプラインのフレキシビリティはバイナリ関数の呼び出しに実行時オーバーヘッドがあることを意味します先ほども触れましたがプリコンパイルされた関数と特殊化のための関数のセットを混在させることができますが完全に特殊化された単一コンパイルパイプラインが最高のパフォーマンスを提供します既存のパイプラインで個別にコンパイルされた関数を置き換えるためにバックグラウンドで新しい単一コンパイルパイプラインをいつでも作成することができます別のコンパイルパイプラインに戻るとパイプラインから呼び出すことができる関数の配列がありますしかしながら関数の固定集合がある場合常に別の関数が常に表示されます
この場合パイプラインに追加する新しいウッドマテリアルがありますインクリメンタルコンパイルは既存のパイプラインに新しい関数を追加するためのものです新しい関数の使用は通常動的環境特にストリーミングが新しいアセットを新しいシェーダーでロードするゲームで非常に一般的です
これをコードに組み込むとまず最初の計算パイプラインを作成する時に最初にインクリメンタルコンパイルをサポートするように選択する必要がありますデフォルトではインクリメンタルコンパイルはサポートされませんバイナリ関数を後で追加できるようにすることはたとえパイプラインの作成時に何も指定されていなくてもバイナリ関数の呼び出しが可能になることが期待されるからです
次に関数記述子を使用してプリコンパイルされたバイナリ関数としてウッドシェーダーの Metal関数オブジェクトを作成します
最後にComputePipelineStateに新しく追加されたメソッドを使用し追加のバイナリ関数を含む計算パイプラインを作成します
次にvisible関数テーブルについてお話します visible関数テーブルは visible関数への関数ポインタを GPUに渡す方法です
シェーディングのために検討した一連のライティング関数とmaterial関数に戻ると GPUで実行されているカーネルにこれらを提供する必要があります
関連する関数をグループ化してGPUに渡すため Metal APIを使用して visible関数テーブルを作成します visible関数テーブルはMetalシェーダーコードに渡すことができるオブジェクトです
始める前に使用する宣言をいくつか追加して関数ポインタの型を定義しコード例が長くなり過ぎないようにしましょう vsible関数テーブルはカーネルパラメータとして指定できバッファーバインディングポイントを使用しますそれらを引数バッファーに渡すこともできます
次にシェーダーにてインデックスで関数にアクセスします関数へのポインタを取得するかテーブルから直接呼び出すことができます
CPUでテーブルに必要な関数エントリの数に基づいて ComputePipelineStateから visible関数テーブルを割り当てます
次にパイプラインステートを作成する時に指定した関数へのハンドルをテーブルに入力します
最後にテーブルを使用する時は必要に応じてcomputeCommandEncoder またはargumentEncoderに設定します argumentEncoderを使用する場合は useResourceを呼び出しを忘れないでくださいインクリメンタルコンパイルされたパイプラインをできるだけ低コストで使用できるようにするために祖先パイプラインからvisible関数テーブルを再利用できるようにしていますテーブルに既に存在するハンドルは引き続き有効であり新しい関数ハンドルをテーブルに追加するだけですまったく新しいテーブルを作成して構築する必要はありません最後に古いパイプラインから新しく追加された関数ハンドルにアクセスしないように気をつけてください今日の最後になりましたがアプリケーションで関数ポインタを使用する際のパフォーマンスに関する考慮事項について説明しますパフォーマンスでは３つの主要な領域を取り上げます最適化のための関数グループから始めましょう前の図に戻ると特定の目的のための関数のグループがあることを確認しましたねただしパイプライン作成ではこれらの関係は認識されません
シェーダーでの使用に基づいて関数をグループ化できますライティングに使用する関数とマテリアルに使用する関数をどれにするかは分かっています
シェーディング関数で関数のグループを表すためにそれを更新して実行する呼び出しの関数グループを含めました Metalでは関数グループによって関数が使用される場所をコンパイラに指示できますシェーダーコードでは関数を呼び出す行に関数グループ属性を適用しその場所から呼び出し可能な関数グループに名前を付けます
次に計算パイプライン記述子を構成する時にグループごとに呼び出すことができるセッター関数を使用してディクショナリにグループを指定しますこれによりコンパイラは可能な限り多くの情報を入手できパイプラインを生成する際の最適化に役立ちます関数配列の関数に対してコンパイラが実行する特殊化に加えてコンパイラは各グループのセッター関数間の共通点を利用してそれを関数グループ属性でマークされた呼び出しサイトに適用できますこのすべての新しいフレキシビリティにより再帰関数を実装できるようになります
レイトレーシングプロセスフローを開始したところに戻るとマルチバウンスを考慮するこのレイトレーシングのモデルは反復的または再帰的な方法で実装できます
交差とシェーディングの関係に注目するとこのコールグラフをたどることができます
再帰パストレーサーを使用して交差後に交差ポイントをシェーディングしますこの場合新しいwood material関数を使用しますこのウッドマテリアルには反射コンポーネントが含まれているため別のレイとシーンと交差させ交差をシェーディングします
この場合もウッドマテリアルにおいて覆われた表面と交差します次にこれがシーンを反射して再び交差しますもう一度ウッドマテリアルをヒットします続行することもできますが通常パフォーマンスとスタックのオーバーフローを避けるために再帰的レイトレーサーの深度を制限するためここでストップします visible関数のチェーンを呼び出す計算カーネルの様々なスタックの使用をサポートするために計算パイプライン記述子で maxCallStackDepthプロパティを公開しカーネルから関数を呼び出す深さを指定できるようにしますデフォルトは１なので visible関数とintersection関数の一般的な使用例はそのまますぐ使用できますこの値はvisible関数呼び出しのチェーンに使用されることが予想されますがレイトレーシングの場合このコードはスタックの消費量を減らすために反復的に記述されていたかもしれません今日お話ししたい最後のパフォーマンスの考慮事項は関数ポインタを使用した場合のダイバージェンスのインパクトでしたただし関数ポインタについて説明する前に昨年の“Ray Tracing with Metal” プレゼンテーションでは後続のバウンスの一貫性が低下した時にレイトレーシングに固有のダイバージェンスを処理する高レベルメソッドについて説明しましたこれにはブロック Linear Layoutによるレイのコヒーレンスアクティブなレイのレイコンパクション GPU間でのロードバランシングのためのインターリーブタイリングが含まれますレイトレーシング最適化のアイデアの優れたレビューについてはプレゼンテーションを確認されることをお勧めしますただし関数ポインタの場合はスレッドレベルでのダイバージェンスを考慮する必要がありますスレッドグループをディスパッチする時 SIMDグループとして定義されたより小さなスレッドグループで実行されることがわかりますこの例では８つのSIMDグループがあります
SIMDグループ内のスレッドは同じ命令を同時に実行した時に最高のパフォーマンスを発揮します
ダイバージェンスは次に実行する命令が SIMDグループのスレッドで異なる場合に起こります通常シェーダーコードの分岐に到達した場合に発生しますこれは通常スレッドのサブセットに命令を実行させ他のスレッドは未使用のままにすることで処理されますすべてのスレッドが有用な命令をアクティブに実行している時に最高のパフォーマンスが得られます
関数ポインタはダイバージェンスにつながる可能性があるもう１つのケースです関数ポインタを介して関数を呼び出す時はセッター関数ポインタがどの程度分岐するかを考慮する必要があります最悪の場合SIMDグループ内では各スレッドが異なる関数を呼び出している可能性があり実行時間は必要な各関数を順番に実行する時間になります
ただしスレッドグループ全体を検討すると完全なSIMDグループを作成する為の十分な作業が必要になる場合があります
ダイバージェンスを処理するために関数呼び出しを並べ替えて SIMDグループにcoherenceを導入しますつまりあるスレッドの関数呼び出しはスレッドグループ内の別のスレッドによって実行される可能性があります
コードに関してはすべての関数パラメータスレッドのインデックススレッドグループメモリを呼び出す関数を書くことから始めます次にお気に入りのソートを使用して関数のインデックスを並べ替えます
次に並べ替えられた順序で関数を呼び出します関数の結果はスレッドグループのメモリに戻り各スレッドはそこから結果を読み取ることができますパラメータと結果のデバイスメモリを使用して同様のバージョンを実装することもできます複雑な関数を呼び出す場合これはダイバージェンスのオーバーヘッドを減らすはずです Metalの関数ポインタでできることはたくさんあります今日はvisible関数の作成と使用およびそれらをvisible関数テーブルに追加して計算カーネルから呼び出すことができる関数を提供する方法について説明しました関数を動的に追加できるようにパイプラインのサイズと作成時間と実行時のパフォーマンスと必要性のバランスを選択できる― 様々なコンパイルモデルについて説明しましたまたアプリケーションで関数ポインタを使用する場合のパフォーマンスに関する考慮事項についても説明しました私は関数ポインタの作成をとても楽しんでおり皆さんと共有できて本当に嬉しいです WWDCをお楽しみください

3:09 - Our simple Path Tracer in Metal Shading Language:

float3 shade(...);

[[kernel]] void rtKernel(...
                         device Material *materials,
                         constant Light &light)
{
    // ...

    device Material &material = materials[intersection.geometry_id];
    float3 result = shade(ray, triangleIntersectionData, material, light);

    // ...
}

3:30 - Our shading function

float3 shade(...)
{
    Lighting lighting = LightingFromLight(light, triangleIntersectionData);

    return CalculateLightingForMaterial(material, lighting, triangleIntersectionData);
}

4:59 - Declare a function as visible

[[visible]]
Lighting Area(Light light, TriangleIntersectionData triangleIntersectionData)
{
    Lighting result;
    
    // Clever math code ...
    
    return result;
}

5:30 - Single compilation pipeline on CPU

// Single compilation configuration
let linkedFunctions = MTLLinkedFunctions()
linkedFunctions.functions = [area, spot, sphere, hair, glass, skin]
computeDescriptor.linkedFunctions = linkedFunctions

// Pipeline creation
let pipeline = try device.makeComputePipelineState(descriptor: computeDescriptor,
                                                   options: [],
                                                   reflection: nil)

7:43 - Introducing MTLFunctionDescriptor

// Create by function descriptor:
let functionDescriptor = MTLFunctionDescriptor()
functionDescriptor.name = "Area"
// More configuration goes here
let areaBinaryFunction = try library.makeFunction(descriptor: functionDescriptor)

8:08 - Binary pre–compilation

// Create and compile by function descriptor:
let functionDescriptor = MTLFunctionDescriptor()
functionDescriptor.name = "Area"
functionDescriptor.options = MTLFunctionOptions.compileToBinary
let areaBinaryFunction = try library.makeFunction(descriptor: functionDescriptor)

8:20 - Binary functions

// Specify binary functions on compute pipeline descriptor
let linkedFunctions = MTLLinkedFunctions()
linkedFunctions.functions = [spot, sphere, hair, glass, skin]
linkedFunctions.binaryFunctions = [areaBinaryFunction]
computeDescriptor.linkedFunctions = linkedFunctions

// Pipeline creation
let pipeline = try device.makeComputePipelineState(descriptor: computeDescriptor,
                                                   options: [],
                                                   reflection: nil)

11:04 - Incremental compilation pipeline

// Create initial pipeline with option
computeDescriptor.supportAddingBinaryFunctions = true

// Create and compile by function descriptor
let functionDescriptor = MTLFunctionDescriptor()
functionDescriptor.name = "Wood"
functionDescriptor.options = MTLFunctionOptions.compileToBinary
let wood = try library.makeFunction(descriptor: functionDescriptor)

// Create new pipeline from existing pipeline
let newPipeline =
   try pipeline.makeComputePipelineStateWithAdditionalBinaryFunctions(functions: [wood])

12:22 - Visible function tables on the GPU

// Helper using declaration in Metal
using LightingFunction = Lighting(Light, TriangleIntersectionData);
using MaterialFunction = float3(Material, Lighting, TriangleIntersectionData);

// Specify tables as kernel parameters
visible_function_table<LightingFunction> lightingFunctions [[buffer(1)]],
visible_function_table<MaterialFunction> materialFunctions [[buffer(2)]],

// Access via index
LightingFunction *lightingFunction = lightingFunctions[light.index];
Lighting lighting = lightingFunction(light, triangleIntersection);
return materialFunctions[material.index](material, lighting, triangleIntersection);

12:49 - Visible function tables on the CPU

// Initialize descriptor
let vftDescriptor = MTLVisibleFunctionTableDescriptor()
vftDescriptor.functionCount = 3
let lightingFunctionTable = pipeline.makeVisibleFunctionTable(descriptor: vftDescriptor)!

// Find and set functions by handle
let functionHandle = pipeline.functionHandle(function: spot)!
lightingFunctionTable.setFunction(functionHandle, index:0)

// Find and set functions by handle
computeCommandEncoder.setVisibleFunctionTable(lightingFunctionTable, bufferIndex:1)
argumentEncoder.setVisibleFunctionTable(lightingFunctionTable, index:1)

14:23 - Function groups on GPU

// Add function groups to our shading function
float3 shade(...)
{
    LightingFunction *lightingFunction = lightingFunctions[light.index];
    [[function_groups("lighting")]] Lighting lighting = lightingFunction(light,
triangleIntersection);
    
    MaterialFunction *materialFunction = materialFunctions[material.index];
    [[function_groups("material")]] float3 result = materialFunction(material, lighting, triangleIntersection);
    return result;
}

14:49 - Function groups on CPU

// Function Group configuration
let linkedFunctions = MTLLinkedFunctions()
linkedFunctions.functions = [area, spot, sphere, hair, glass, skin]
linkedFunctions.groups = ["lighting" : [area, spot, sphere ],
                          "material" : [hair, glass, skin ] ]
computeDescriptor.linkedFunctions = linkedFunctions

特定のトピックをお探しの場合は、上にトピックを入力すると、関連するトピックにすばやく移動できます。

クエリの送信中にエラーが発生しました。インターネット接続を確認して、もう一度お試しください。

リソース

関連ビデオ

WWDC21