Metalを使用したレイトレーシング

Metalを使用したレイトレーシング

レイトレーシング（MetalグラフィックスフレームワークとShading Languageの中核部分）によって、Appやゲームでフォトリアリスティックな3Dシーンを実現します。MetalレイトレーシングAPIおよびレイトレーシングのためのShading Language Extensionの基本を紹介し、それらをグラフィックスAppおよびゲームで使用する方法および最適なパフォーマンスを得るためにカーネルを管理し、シングルコンピューティングカーネルに組み込む方法について学びます。

リソース
関連ビデオ

WWDC23
- Metalレイトレーシングのガイド
WWDC21
- Metalのデバッグ、プロファイリング、アセット作成ツール
ダウンロード

こんにちは WWDCへようこそ “Metalを使用したレイトレーシングの検出” GPUソフトウェアエンジニアのショーン･ジェームスです Metalに今年追加された新しいレイトレーシングAPIについて説明します新しいAPIでは任意の計算カーネル内から直接交差テストを行えますまた交差関数を記述して交差プロセスをカスタマイズ可能ですレイトレーシングアプリケーションは光線の軌跡をたどります音響や物理シミュレーションなど多くの分野に適用される中でフォトリアリスティックレンダリングは人気の１つですレンダリングではレイトレーシングで個々の光線をモデル化し反射ソフトシャドウ間接照明などの効果をシミュレートできますまずレイトレーシングについて話し新しいMetalレイトレーシングAPIをご紹介します
初めにカメラからシーンに放たれる光線を生成次に光線とシーン内のジオメトリとの交差をテストします各交点は表面で跳ね返る光を表し跳ね返る光の量と方向でオブジェクトの見え方が決まりますそして各交点の色を計算し画像を更新これはシェーディングと呼ばれさらに跳ねてシーンに入る追加の光線も生成できますこれらの交差もテストしこのプロセスを何度も繰り返してシーン内で跳ね返る光をシミュレートします過去２年間で Metal Performance Shadersフレームワークの活用方法をご紹介しましたまず光線の初期セットを生成する計算カーネルを起動し Metalバッファに書き込みます次にMPSRayIntersectorクラスで光線をシーンと交差させインターセクタが結果を別のバッファに書き込みますそして最終カーネルを起動し交差の読み取りとシェーディングを行います追加の光線は光線バッファに書き戻され計算カーネル２つをループで起動すると跳ね返りをシミュレートできます有効な手法ですがコードを３つの計算カーネルに分けた上光線と交差をメモリに渡す必要もあります新しいMetalレイトレーシングAPIではこのインターセクタオブジェクトをシェーディング言語から直接使えます
その結果３つのカーネルを１つに結合できますカーネル間で光線データを渡さないので光線バッファと交差バッファも排除できます光線や交差を渡すためのメモリの読み書きも不要ですこれは柔軟なプログラミングモデルで計算カーネルを何度も起動せず外側のループをシェーディング言語の単純なループで表せますでは基本的なレイトレーシングカーネルの話です各ピクセル１スレッドの２次元計算カーネルを起動しますまずrayを生成しますこれは所定のピクセルの初期光線を生成する単なる計算です次にIntersectorオブジェクトを作成このオブジェクトはシーン内のジオメトリと光線との交差を検出するものでここに多様なプロパティを設定しその動作をカスタマイズ可能です例えばアンビエントオクルージョンやシャドウなどで高速化を図れます近くの交差を追求せず最初の交差を受け入れるのですそしてrayをIntersectorに渡し交差結果を回収最後に交差結果を使ってシェーディングを実行します交差ステップは多くの作業を伴うので複雑なシェーディングコードとの結合では占有率が低い計算カーネルになる可能性がありますカーネル間の最適な分担にはアプリケーションのプロファイルが必要です以上が基本的な交差テストの要件です交差関数の別の引数である加速構造はレイトレーシングプロセスを高速化するためのデータ構造です空間を再帰的に分割するため特定の光線と交差し得ない三角形を素早く削除できます MPSと同様に Metalはこのデータ構造を築きますユーザーはジオメトリを提供するだけですその仕組みは？初めに構築したい加速構造を記述する記述子オブジェクトを作成次に加速構造を格納するメモリを割り当てて最後に加速構造を構築します
まずは記述子です Metalはプリミティブ加速構造とインスタンス加速構造をサポートします今回はプリミティブ加速構造を作成しましょう三角形などのプリミティブを含むものですプリミティブ加速構造は個々のジオメトリで構成されます
Metalも２種類のジオメトリをサポートしますが今回は三角形の例です三角形ジオメトリの各部には独自の頂点バッファインデックスバッファ三角形の数などが含まれます PrimitiveAccelerationStructureDescriptorを作成したら次にgeometryDescriptorを作成しますここでは単一のTriangleGeometryDescriptorですそしてvertexBufferを接続し triangleCountを指定三角形のデータは加速構造に埋め込まれるため加速構造を築いたら vertexBufferの維持は不要ですあとはgeometryDescriptorsを accelerationStructureDescriptorに追加する記述子を作成したのでメモリを割り当てられますがこれは大きなメモリ割り当てになり得ます新しいMetalレイトレーシングAPIに加えた改善の１つは Metalで加速構造が割り当てられるタイミングと場所の完全な制御ですつまりアプリケーションの都合でメモリの再利用や割り当てが可能ですまずMetal deviceのメソッドを呼び出しますこれは割り当てに必要なsizesを含む accelerationStructureを返すものですそしてMetal deviceから AccelerationStructureを割り当てて最後にscratchBufferも割り当てますこれは加速構造の構築中に一時的なストレージとしてMetalに使用され完了したら破棄できますこのバッファ内のデータにはアクセスしないのでパフォーマンスを上げるためプライベートリソースストレージモードを使用します加速構造にメモリを割り当てましたがまだ構築されていませんそのため最後に加速構造を構築します
ここでもいくつかの改善を行いました加速構造の構築は時間を要すことがあるため MetalでどのGPU commandQueueや commandBufferで構築するのかを含め構築のタイミングを完全に制御できるようになりましたまずcommandBufferを作成し新しいAccelerationStructure CommandEncodersの１つを作成します他のMetalエンコーダと同様このオブジェクトは GPU上での作業スケジュールに使用されますこの場合実際のbuildコマンドをエンコードし accelerationStructure scratchBufferを提供します
最後にエンコーディングを終了し commandBufferをコミットしますこれでGPUは加速構造の構築を開始できます加速構造の構築はCPU同期なしで GPUタイムライン上で実行されるようになりましたこのコマンドバッファの後にGPUで交差作業をスケジュールしても安全です加速構造の使用法は既に見たように Intersectorに渡すだけですご覧のように加速構造は通常のMetalバッファバインディングポイントにバインドしますこれらのバインディングポイントに加速構造をバインドするため対応するメソッドが計算コマンドエンコーダと引数エンコーダにあります新しいMetalレイトレーシングAPIの基本的な構成要素は以上です今日取り上げなかった高度なトピックがいくつかあるのでドキュメントとサンプルコードを確認することをお勧めします概念の多くは同様なので前の２つのレイトレーシングに関する講演もご欄ください例えばMPSのようにMetal加速構造はプリミティブ加速構造の２レベルのインスタンス化をサポートしていますこれによりメモリ使用量を減らせます Metalはインスタンス化と組み合わせて動的ジオメトリに対応可能な既存の加速構造の再フィットもサポートしています最後に Metal加速構造はコンパクションもサポートしていますコンパクションを使用すると加速構造の構築時相当量のメモリを回収できます次に交差関数を使用した交差プロセスのカスタマイズ方法を説明しますまずはこれまで取り上げたすべてのデモを見てみましょう
これは高度なコンテンツチームによって作成されたアプリケーションであり複雑なシーンに適用される新しいMetalレイトレーシングAPIのデモですこれはMac Proで記録されたシーンのインタラクティブプレビューです新しいAMD Radeon Pro W5700X GPUを使用していますインタラクティブプレビューはピクセルごとに１つのサンプルのみを使用するためノイズは増えますがカメラの動きが止まると収束し始めます
これは編集中にアーティストがシーンを見る方法です高品質のオフラインレンダリングの場合各フレームで十分な時間をとってノイズのない画像に収束するようにします最終結果がどのようになるかを把握するために１ピクセルあたり400サンプルを使用してシーンのフライスルーをレンダリングしました
Megascansライブラリの使用について Quixel社に感謝しますこちらはMegascansアセットから組み立てられ 3200万を超える三角形を含んでいます
このアプリケーションは前に説明したアウトラインと非常によく似た働きをしますパストレースと呼ばれるレンダリングアルゴリズムを使用してソフトシャドウ被写界深度間接照明などの高度な効果をシミュレートします
これでインタラクティブアプリケーションと完全なオフラインレンダラーの両方で新しいMetalレイトレーシングAPIを確認できました次に新しいMetal APIの最新機能の１つである交差機能の話に移りますこれまでの弊社のプログラミングモデルでは光線がインターセクタに入り交差が出てきますがインターセクタは常にブラックボックスでした実際にはインターセクタの動作をカスタマイズしたい場合があります
例えばアルファテストなどラスタライゼーションの一般的な手法でテクスチャで透明にするべき三角形の部分をマスクしますこれにより実際の三角形の数を増やさず多くの幾何学的詳細を追加できますこの例ではアルファテストをオンにすると葉の細部まで表現できます
これをラスタライゼーションで実装するのは非常に簡単でフラグメントシェーダーからフラグメントを破棄するだけですただしレイトレーシングを使うと効率的な実装の難易度が上がります光線をキャストするたびインターセクタは最も近い交差を返しますただしこの場合交点が実際には透明なので無視する必要があります
従って最初の交点から始まる新しい光線をキャストする必要があります最後に不透明な表面にぶつけて停止しシェーディングに移行できますが各光線はインターセクタまで新たに移動する必要がありますこれではツリー構造のルートから加速構造トラバーサルを何度も再開することになるためとても非効率ですより効率的に実装するには三角形の交差関数を使用します
この仕組みを理解するにはインターセクタが内部でどう機能するか知る必要があります
まず光線とジオメトリ一部の交差が見つかるまで加速構造をトラバースし交差が前の最も近い交差よりも近いかどうかを確認します近ければその交差を新しい最短の交差にして加速構造のトラバースに戻ります
最終的に最も近い交差を終了して返します交差関数はこのプロセスに直接プラグインしますこれらの関数は交差が見つかる度に呼び出され交差機能は交差を受け入れるか拒否するかを選択できます受け入れた場合その交差は通常どおり新しい最短の交差になりますただし交差の拒否もできその場合は加速構造のトラバースに戻りその交差は破棄しますアルファテストの例ではアルファテストロジックを交差関数に入れることができますその場合加速構造を１回のトラバースでできてはるかに効率的です基本的な交差関数は次のようになりますまず関数を宣言し“triangle”が交差関数であることを示し “triangle data”は交点の重心座標にアクセスすることを示しますこの関数は交差を受け入れるか拒否するかを示すbool値を返しますインターセクタからprimitiveIndex geometryIndex交点の重心座標などの情報を受け取り使用されない他のデータも一緒に受け取りますアルファテクスチャにアクセスするために使う独自のリソースをバインドすることもできますまずこのプリミティブのアルファテクスチャを調べます次に重心座標を使用してテクスチャ座標を補間しますがこれはシェーディング手順とまったく同じです次に実際のテクスチャルックアップを実行し最後に交点が透明かどうかに応じてtrueまたは falseを返しますただし交差関数は三角形のジオメトリに限定されません Metalでは境界ボックスの交差関数を記述することもできますこれらの関数はカスタムプリミティブをモデル化するために使用できる指定した境界ボックスと光線が交差する時に呼び出されますこのシーンでは球は境界ボックスの交差機能を使用して実装されています光線は三角形でモデル化するのではなく球に対して数学的にテストされますこれは非常に単純な例ですが境界ボックスの交差関数を使用すると曲線などのより複雑なサーフェスをモデル化することもできますこの例では境界ボックスの交差機能を使用して髪の毛をレンダリングしました各髪束はカスタムの３次ベジェ曲線プリミティブでありこれを１万個使って髪に全体的な形を作っています
各曲線プリミティブを交差させるとさらに労力を要しますが三角形の代わりにカーブを使えばアーチファクトのない完全に滑らかな髪の毛をレンダリングできますまた各曲線は多数の三角形の代わりに少数の制御点しか必要としないため使用するメモリを減らせますこのメモリの効率的な使用は毛量が多いキャラクターの場合に重要です球プリミティブで作られたシーンを想定します
まず各球を囲む軸に沿った境界ボックスを提供し次に三角形の場合と同様に Metalはボックス上で加速構造を構築します
その後光線を提供するとMetalは境界ボックスと光線の間の交差を検索します三角形の場合と同様にMetalでは潜在的な交差を見つける度に境界ボックス交差関数を呼び出しますこの加速構造を構築するには小さな変更をいくつか加えるだけで済みます
前と同様にPrimitiveAccelerationStructure Descriptorを作成しますただし今回は三角形のジオメトリ記述子を作成する代わりに BoundingBoxGeometryDescriptorを作成します
この記述子はboundingBoxBufferとはboundingBoxCountの接続場所を提供します
三角形交差関数と同様に境界ボックス交差関数はtrueを返して交差の受け入れを示し falseを返して交差の拒否を示しますただしこれらは光線の原点からカスタムプリミティブ上の任意の交点までの最短距離を計算する役割も果たします Metalはこの距離を使用してシーン内のさまざまな三角形と境界ボックスプリミティブ全体に渡り最も近い交差を計算します境界ボックスの交差関数の記述を示すよう宣言が少し変更されます複数の値を返す必要があるためユーザー定義の構造体を返しますこの例では交差を受け入れの可否を示す bool値を含む構造体を作成しました交差距離を含むfloatも含まれます三角形交差関数と同様に光線のoriginや directionなどの光線データを受け取りますまたprimitiveIndexなどの交差に関する情報も受け取りますさらに独自のリソースもバインドできこのリソースを使用して球の配列を渡しますまず境界ボックスで囲まれた球と光線が交差するかどうかを数学的に確認し球と光線が交差しない場合直ちにfalseを返すことができます次に三角交差関数の場合と異なり交差関数が許容範囲内にあるかどうかを確認するもう１つの作業を行います許容範囲内であれば交差を受け入れ交差距離を返すことができますしかし追加のデータを返したい場合はどうでしょうか？例えばシェーディングのために交点で表面法線を返したい場合には光線のpayloadを使用しますこれは交差関数から交差プロセスを開始した計算カーネルまで遡って渡される小さなデータですペイロードタイプへの参照を交差関数の引数として宣言します次にそれに書き込みをしてペイロードを変更できますペイロードへの変更は交差の受け入れに関わらず表示されることに注意してくださいそのため多くの場合交差の受け入れ時のみペイロードを変更する必要がありますペイロード値を取得するには intersectorへの呼び出しを変更して参照によってペイロードを渡すだけです intersectorが戻ると交差関数からのペイロードへの変更はすべて表示されますこれまで計算カーネルからのレイインターセクタの使用法と交差関数の作成方法について説明しました次にこの２つを組み合わせる方法について説明しますインターセクタが実際に交差関数を呼び出せるようにするにはそれらを計算パイプライン状態でリンクする必要があります
まずMetal libraryから InertsectionFunctionをロードしますがこれは他のMetal関数と同様に行われます
次にMTLLinkedFunctions オブジェクトを使用して交差関数をPipelineDescriptorにアタッチします
最後に通常どおりPipelineの状態をコンパイルしますコンパイラは計算カーネルと交差関数のコードをリンクし交差関数を計算カーネル内のレイインターセクタから呼び出せるようにしますこれですべての交差関数を計算パイプライン状態にリンクしたのでこれらの交差関数を個々のジオメトリにマッピングする必要があります
これには交差関数テーブルと呼ばれるものを使用しますこのテーブルにはパイプライン状態の実行可能コードへのポインターが含まれています加速構造ジオメトリの各部分が独自の頂点バッファを持てるのと同様にそれぞれが独自の交差関数を持つこともできますまたインスタンス加速構造を使用している場合は各インスタンスに独自の交差関数のセットを含めることもできます AccelerationStructureGeometryと InstanceDescriptorの両方に個別の intersectionFunctionTableOffsetがあります光線がジオメトリの一部に当たるとこれらの２つの値が追加されて呼び出す交差関数テーブルエントリが決まります交差関数テーブルを作成するにはまず記述子を作成しますこの例ではすべての交差関数を格納できる記述子を作成します次にfunctionTableを割り当てますこのテーブルは特定の計算パイプライン状態の実行可能コードをポイントするためそのパイプライン状態からも作成されますこのテーブルはこのパイプライン状態または派生したパイプライン状態でのみ使用できます次にcomputePipelineから各intersectionFunctionのHandleを取得しますこのHandleはパイプライン状態にリンクされている関数の実行可能コードのアドレスを表します次にこのHandleを functionTableに挿入しますこれはテーブルの交差機能で使われるリソースをバインドする場所でもあります
例えばsphereBufferを Index:0にバインドします交差関数テーブルの使用法は加速構造とよく似ています
この場合も交差関数テーブルを intesectorに渡すだけです intersection function tableはMetalバッファバインディングポイントにもバインドしますまた交差関数テーブルを計算コマンドエンコーダと引数エンコーダにバインドする際に対応する APIメソッドもあります以上が交差関数を設定するために必要な作業です交差関数は交差プロセスをより詳細に制御できる強力なツールですまとめると最初に新しい MetalレイトレーシングAPIと任意の計算カーネル内から直接利用できるようになったレイインターセクタオブジェクトについて説明しましたこれははるかに柔軟なプログラミングモデルでありアプリケーションに最適な方法で作業を分割できます
また新しいAPIを使用し加速構造を割り当て構築するタイミングと場所をより詳細に制御する方法も説明しました
最後に交差関数を作成して交差プロセスをカスタマイズする方法をお話ししました交差関数についてお話できることはまだたくさんあります他のアプリケーションで使用できるよう基礎となるメカニズムも用意しました
例えば光線はいつでもどの面にも当たる可能性があるためシェーディングステップでは必要に応じて任意のマテリアルのコードを実行できる必要がありますこのためこのセッションに付随する関数ポインターの講演をご覧になりこれを可能にするもう１つの強力なツールについて学ぶことをお勧めします本日は新しいレイトレーシングAPIの基本を説明したのみでしたこのセッションに付属するサンプルコードとドキュメントを確認することもお勧めします引き続きWWDCをお楽しみください

[[kernel]]
void rtKernel(primitive_acceleration_structure accelerationStructure [[buffer(0)]],
              /* ... */)
{
    // Generate ray
    ray r = generateCameraRay(tid);

    // Create an intersector
    intersector<triangle_data> intersector;

    // Intersect with scene
    intersection_result<triangle_data> intersection;

    intersection = intersector.intersect(r, accelerationStructure);

    // shading...
}

4:48 - Create an acceleration structure descriptor

let accelerationStructureDescriptor = MTLPrimitiveAccelerationStructureDescriptor()

// Create geometry descriptor(s)
let geometryDescriptor = MTLAccelerationStructureTriangleGeometryDescriptor()

geometryDescriptor.vertexBuffer = vertexBuffer
geometryDescriptor.triangleCount = triangleCount

accelerationStructureDescriptor.geometryDescriptors = [ geometryDescriptor ]

5:46 - Allocate acceleration storage

// Query for acceleration structure sizes
let sizes = device.accelerationStructureSizes(descriptor: accelerationStructureDescriptor)

// Allocate acceleration structure
let accelerationStructure =
    device.makeAccelerationStructure(size: sizes.accelerationStructureSize)!

// Allocate scratch buffer
let scratchBuffer = device.makeBuffer(length: sizes.buildScratchBufferSize,
                                      options: .storageModePrivate)!

6:24 - Build acceleration structure

// Create command buffer/encoder
let commandBuffer = commandQueue.makeCommandBuffer()!
let commandEncoder = commandBuffer.makeAccelerationStructureCommandEncoder()!

// Encode acceleration structure build
commandEncoder.build(accelerationStructure: accelerationStructure,
                     descriptor: accelerationStructureDescriptor,
                     scratchBuffer: scratchBuffer,
                     scratchBufferOffset: 0)

// Commit command buffer
commandEncoder.endEncoding()
commandBuffer.commit()

7:15 - Pass acceleration structure to ray intersector

[[kernel]]
void rtKernel(primitive_acceleration_structure accelerationStructure [[buffer(0)]],
              /* ... */)
{
    // generate ray, create intersector...

    intersection = intersector.intersect(r, accelerationStructure);

    // shading...
}

7:25 - Bind acceleration structure with compute command encoder
```
computeEncoder.setAccelerationStructure(accelerationStructure, bufferIndex: 0)
```

12:16 - Triangle intersection functions

[[intersection(triangle, triangle_data)]]
bool alphaTestIntersectionFunction(uint primitiveIndex        [[primitive_id]],
                                   uint geometryIndex         [[geometry_id]],
                                   float2 barycentricCoords   [[barycentric_coord]],
                                   device Material *materials [[buffer(0)]])
{
    texture2d<float> alphaTexture = materials[geometryIndex].alphaTexture;

    float2 UV = interpolateUVs(materials[geometryIndex].UVs,
        primitiveIndex, barycentricCoords);

    float alpha = alphaTexture.sample(sampler, UV).x;

    return alpha > 0.5f;
}

14:38 - Creating a bounding box acceleration structure

// Create a primitive acceleration structure descriptor
let accelerationStructureDescriptor = MTLPrimitiveAccelerationStructureDescriptor()

// Create one or more bounding box geometry descriptors:
let geometryDescriptor = MTLAccelerationStructureBoundingBoxGeometryDescriptor()

geometryDescriptor.boundingBoxBuffer = boundingBoxBuffer
geometryDescriptor.boundingBoxCount = boundingBoxCount

accelerationStructureDescriptor.geometryDescriptors = [ geometryDescriptor ]

15:29 - Bounding Box Result

struct BoundingBoxResult {
    bool accept [[accept_intersection]];
    float distance [[distance]];
};

15:38 - Bounding box intersection functions

[[intersection(bounding_box)]]
BoundingBoxResult sphereIntersectionFunction(float3 origin            [[origin]],
                                             float3 direction         [[direction]],
                                             float minDistance        [[min_distance]],
                                             float maxDistance        [[max_distance]],
                                             uint primitiveIndex      [[primitive_id]],
                                             device Sphere *spheres   [[buffer(0)]])
{
    float distance;

    if (!intersectRaySphere(origin, direction, spheres[primitiveIndex], &distance))
        return { false, 0.0f };

    if (distance < minDistance || distance > maxDistance)
        return { false, 0.0f };

    return { true, distance };
}

16:20 - Ray payload

[[intersection(bounding_box)]]
BoundingBoxResult sphereIntersectionFunction(/* ... */,
                                             ray_data float3 & normal [[payload]])
{
    // ...

    if (distance < minDistance || distance > maxDistance)
        return { false, 0.0f };

    float3 intersectionPoint = origin + direction * distance;
    normal = normalize(intersectionPoint - spheres[primitiveIndex].origin);

    return { true, distance };
}

16:48 - Ray payload 2

[[kernel]]
void rtKernel(/* ... */)
{
    // generate ray, create intersector...

  float3 normal;

    intersection = intersector.intersect(r, accelerationStructure, functionTable, normal);

    // shading...
}

17:18 - Linking intersection functions

// Load functions from Metal library
let sphereIntersectionFunction = library.makeFunction(name: “sphereIntersectionFunction”)!
// other functions...

// Attach functions to ray tracing compute pipeline descriptor
let linkedFunctions = MTLLinkedFunctions()

linkedFunctions.functions = [ sphereIntersectionFunction, alphaTestFunction, ... ]

computePipelineDescriptor.linkedFunctions = linkedFunctions

// Compile and link ray tracing compute pipeline
let computePipeline = try device.makeComputePipeline(descriptor: computePipelineDescriptor,
                                                     options: [],
                                                     reflection: nil)

18:17 - Intersection function table offsets

class MTLAccelerationStructureGeometryDescriptor : NSObject {

    var intersectionFunctionTableOffset: Int

// ...

}

struct MTLAccelerationStructureInstanceDescriptor {
    var intersectionFunctionTableOffset: UInt32
    // ...
};

18:35 - Creating an intersection function table

// Allocate intersection function table
let descriptor = MTLIntersectionFunctionTableDescriptor()

descriptor.functionCount = intersectionFunctions.count

let functionTable = computePipeline.makeIntersectionFunctionTable(descriptor: descriptor)

for i in 0 ..< intersectionFunctions.count {
    // Get a handle to the linked intersection function in the pipeline state
    let functionHandle = computePipeline.functionHandle(function: intersectionFunctions[i])

    // Insert the function handle into the table
    functionTable.setFunction(functionHandle, index: i)
}

// Bind intersection function resources
functionTable.setBuffer(sphereBuffer, offset: 0, index: 0)

19:26 - Pass intersection function table to ray intersector

[[kernel]]
void rtKernel(primitive_acceleration_structure accelerationStructure   [[buffer(0)]],
              intersection_function_table<triangle_data> functionTable [[buffer(1)]],
              /* ... */)
{
    // generate ray, create intersector...

    intersection = intersector.intersect(r, accelerationStructure, functionTable);

    // shading...
}

19:33 - Bind intersection function table with compute command encoder
```
encoder.setIntersectionFunctionTable(functionTable, bufferIndex: 1)
```

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

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

リソース

関連ビデオ

WWDC23

WWDC21