Bring your Metal app to Apple silicon Macs

ハイライト

Apple Silicon Mac は Apple GPU の TBDR アーキテクチャを Mac にもたらします。Metal アプリは代わりに API 駆動の機能検出、明示的なロード/ストア アクション、位置不変性、およびスレッドグループ同期を使用して、Rosetta 互換パスのパフォーマンス コストを回避する必要があります。

主要内容

多くの Mac グラフィックス アプリケーションは、当初は Intel、AMD、または Nvidia GPU でのみ動作しました。これらはイミディエイト モード レンダラー (イミディエイト モード レンダラー) の動作に慣れています。深度、カラー、ステンシル バッファーはビデオ メモリ内で頻繁に読み書きされ、一部のドライバーでは多くの誤った使用法が直ちに公開されません。

Apple Silicon Mac は Apple が開発した GPU に置き換えられます。この GPU はタイルベースの遅延レンダリング (TBDR) を使用します。まずジオメトリを画面タイルに分割し、次に深度テスト、混合、および一部のカラー操作をオンチップ タイル メモリ内で完了します。適切に作成されたレンダラは、使用するメモリ帯域幅を大幅に削減し、消費電力とパフォーマンスの両方にメリットをもたらします。

移行パスは 2 つのステップに分かれています。古い x86 アプリケーションは最初に Rosetta を通じて実行でき、システムは画像を可能な限り正確にするためにいくつかの Metal 互換性の回避策を開きます。次に、開発者はネイティブ Apple Silicon バージョンを再コンパイルし、未定義の動作に依存する Metal コードを修正する必要があります。互換性の回避策にはパフォーマンスが犠牲になるため、ネイティブ バージョンでは長期的な解決策として使用できません。

このセッションの価値は、「再コンパイル」で停止しないことです。これは 4 種類の実際の問題を直接示しています。GPU 名による判断不能、不正なロード/ストア アクション、位置不変性の欠如、不正なスレッドグループ メモリ同期の仮定、同じレンダー パスでの現在の深度/ステンシル アタッチメントのサンプリングです。これらの問題は、TBDR アーキテクチャ上で画像エラーや帯域幅の無駄に変わります。

詳細

Metal API を使用して GPU 機能を検出する

(14:23) Apple は、アプリが Metal GPU 機能を直接クエリすることを明示的に要求しています。機能を推測するためにオペレーティング システム名や GPU 名を使用しないでください。 Apple Silicon Mac は、Mac 2 ファミリと Apple GPU ファミリの両方をサポートします。 「macOSにはApple GPU機能はない」という過去の判断は無効となった。

{
    self.appleGPUFeatures = metalDevice.supportsFamily(.apple5)

    self.simdgroupSize = computePipeline.threadExecutionWidth

    self.isLowPower = metalDevice.isLowPower
}

キーポイント:

• supportsFamily(.apple5)Apple GPU ファミリのサポートに直接クエリして、プログラム可能なブレンディング、タイル シェーダー、ローカル イメージ ブロックなどの Apple GPU 機能が利用可能かどうかを確認します。
• threadExecutionWidth計算パイプラインから SIMD グループ幅を取得します。これは、後続のスレッドグループ同期ロジックに必要です。
• isLowPower統合 GPU スタイルと個別 GPU スタイルを区別するために使用されるコード パス。セッションが具体的に指定されており、Apple GPU ではそれが返されますfalse、ディスクリート GPU のパフォーマンス特性の観点から見る必要があります。

ロード/ストアアクションを正しく選択してください

(15:23) ロード/ストア アクションは、レンダー パスの開始時と終了時に、色、深度、ステンシル アタッチメントがオンチップ タイル メモリにどのように出入りするかを制御します。 TBDR GPU は、システム メモリから古いコンテンツを読み込むかどうか、およびタイル メモリをシステム メモリに書き戻すかどうかを決定するため、この選択に敏感です。

判断方法は非常に簡単です。レンダー パスの開始時に、前のコンテンツを保持する必要がある場合にのみ、ロード アクション = ロードを選択します。それ以外の場合は、不要な負荷を避けてください。レンダー パスの最後に、コンテンツが後のパスで消費される場合のみ、ストア アクション = ストアを選択します。それ以外の場合は保管を避けてください。

キーポイント:

• セッションの例は、前のパスでテクスチャをレンダリングし、その後、最後のパスで古いコンテンツを描画し続けることです。使用する場合DontCareロード アクションでは、Apple GPU はシステム メモリ内のテクスチャをタイル メモリにアップロードしません。
• 後続の描画でフレームバッファ全体がカバーされない場合、初期化されていないタイル メモリによりグラフィックス エラーが残る可能性があります。古いコンテンツを保存する必要がある場合は、load を選択する必要があります。
• ストア アクションの判断もリソース フローから行われます。添付ファイルが後続のレンダー パスによって消費される場合にのみ、タイル メモリはシステム メモリにフラッシュ バックされます。
• その後の読み取りを必要としない添付ファイルはストアを回避するため、追加のメモリ トラフィックを削減できます。

パス間で位置を正確に再利用する必要がある場合は、位置不変性を明示的にオンにします。

(20:58) マルチパス アルゴリズムの一般的な方法は、最初のパスで深さを記述し、equal深度テストは、ジオメトリの同じバッチを再度描画します。 2 つの頂点シェーダーが同じ位置計算関数を呼び出した場合でも、コンパイラーの最適化により、位置出力にわずかな違いが生じる場合があります。 Apple GPU は、デフォルトでは位置の不変性を保証しません。

// Renderer.swift
let options = MTLCompileOptions()
options.preserveInvariance = true

library = try device.makeLibrary(source: sourceString,
                                 options: options)


// vertex.metal
struct VertexOut
{
    float4 pos [[position, invariant]];
    float data;
};

キーポイント:

• MTLCompileOptions()Metal コンパイル オプションを作成します。
• options.preserveInvariance = trueコンパイラは、位置の不変性を維持する必要があります。
• makeLibrary(source:options:)このオプションを指定したコンパイル構成を使用してシェーダー ライブラリをビルドします。
• [[position, invariant]]頂点の出力位置にマークが付けられ、この位置がシェーダー間で一貫している必要があることをコンパイラーに伝えます。
• この機能にはパフォーマンスのコストがかかるため、同じ深度値に依存するシェーダーにのみ使用する必要があります。使用してequal深さ比較パス。

スレッドグループのメモリは SIMD グループのサイズに応じて同期する必要があります

(24:25) コンピューティング シェーダーのスレッドグループ メモリは、同じスレッドグループ内のスレッドによって共有されます。 GPU が異なれば SIMD グループ サイズも異なります。 Apple GPU は 32 です。アプリケーションが SIMD グループ内で 64 スレッドが自然に実行されると想定している場合、SIMD グループ間の同期が失われる可能性があります。

// Correct synchronization

// launched with threadgroup size = 64
kernel void kernelMain(uint tid [[ thread_index_in_threadgroup ]],
                       uint simd_size [[ threads_per_simdgroup ]],
                       device uint * res [[ buffer(0) ]])
{
    threadgroup uint buf[64];

    buf[tid] = initBuffer(tid);

    if (simd_size == 64u)
        simdgroup_barrier(mem_flags::mem_threadgroup);
    else
        threadgroup_barrier(mem_flags::mem_threadgroup);

    uint index = (tid < 32) ? tid + 32 : tid - 32;
    res[tid] = buf[tid] + buf[index];
}

キーポイント:

• thread_index_in_threadgroupスレッドグループ内の現在のスレッドのインデックスを与えます。
• threads_per_simdgroupシェーダー内の現在の GPU の SIMD グループ サイズを取得します。
• threadgroup uint buf[64]64 スレッドによって共有されるスレッドグループ メモリです。
• buf[tid] = initBuffer(tid)共有メモリに書き込むと、後で他のスレッドがそれを読み取ります。
• simdgroup_barrier同じ SIMD グループ内のスレッドのみを同期します。
• threadgroup_barrierスレッドグループ全体を同期します。スレッドグループに複数の SIMD グループが含まれる状況に適しています。
• res[tid] = buf[tid] + buf[index]別のスレッドによって書き込まれた値を読み取るため、正確性を保つためには先行するバリアが必要です。

同じレンダー パスで現在の深度/ステンシル アタッチメントをサンプリングしないでください。

(26:11) セッションは最終的に深度/ステンシル バインディング エラーを指摘しました。同じテクスチャが現在のレンダー パスのアタッチメントとして使用されている場合、同時にサンプリング テクスチャとして使用することはできません。これにより、同じ基礎となるテクスチャに対して同時読み取りと書き込みが行われます。

修正は、サンプリングされたテクスチャ バインディングから現在の深度/ステンシル アタッチメントを削除することです。現在のアタッチメントの内容を読み取る必要がある場合は、レンダー パスの前に 2 番目のコピーを作成し、シェーダーにこのコピーを読み取らせます。

キーポイント:

• セッションによって説明されたエラーは、現在の深度アタッチメントがフラグメント シェーダーによって同時にサンプリングされ、同じ基礎となるテクスチャの同時読み取りと書き込みが発生するというものです。
• 不透明なジオメトリの場合、フラグメント シェーダが実行される前に深度が蓄積されます。レンダー パスの最後に、Apple GPU はオンチップの深度/ステンシル メモリをシステム メモリにフラッシュして戻します。現時点では、読み取りと書き込みを同時に行うと、正確性の問題が発生します。
• この競合状態は、最後の描画だけでなく、レンダー パス内のどの描画でも発生する可能性があります。
• この問題を回避するためにテクスチャ バリアやメモリ バリアを使用しないでください。セッションでは、TBDR アーキテクチャではそのような回避策は高価であると明示的に述べられています。
• アルゴリズムが現在のアタッチメントを実際にサンプリングする必要がある場合は、最初に 2 番目のコピーを作成してから、シェーダーにこのコピーを読み取らせます。

重要ポイント

1. Apple GPU 機能レイヤーをレンダラーに追加します

• 内容: 起動時の集中ログsupportsFamily、threadExecutionWidthそしてisLowPower。
• 実行する価値がある理由: 同じ Mac コード セットが Intel/AMD/Nvidia と Apple GPU に同時に直面する必要があり、名前判定はすぐに無効になります。
• 開始方法: カプセル化RendererCapabilities、プログラマブル ブレンディング、タイル シェーダー、simdgroup サイズ、GPU パフォーマンス カテゴリをそれに入れると、レンダリング パスはこのオブジェクトのみを読み取ります。

2. レンダーパスのロード/ストア監査を実行する

• 対処方法: 色、深さ、ステンシルの添付のロード/ストア動作を 1 つずつ確認します。
• 実行する価値がある理由: TBDR アーキテクチャは添付ファイルをタイル メモリに配置しますが、これは間違っています.dontCare画面エラーが発生するため、冗長です.load/.store帯域幅が無駄になります。
• 開始方法: パスに従ってリソース フロー テーブルを描画し、後続のパスでどの添付ファイルが消費されるかをマークします。これらの添付ファイルのみが使用されます.store。

3. 複数のパス深度が等しいプロセスのマーク位置不変性

• やるべきこと: 最初のパスで深さを記述し、それを 2 番目のパスで使用する方法を確認します。equalレンダリングプロセスの比較。
• 実行する価値がある理由: 異なる頂点シェーダーの最適化の違いにより、深度値がまったく同じにならず、その結果、一部のピクセルが誤って破棄されます。
• 開始方法: 関連ライブラリに対してのみオープンしますpreserveInvariance、正確な一致が必要な頂点出力に追加します。[[position, invariant]]。

4. 計算シェーダーの SIMD グループ バージョン戦略を確立する

• 内容: スレッドグループ メモリに依存するカーネル向けに、SIMD グループ サイズによる選択の実装を準備します。
• 実行する価値がある理由: Apple GPU の SIMD グループは 32 です。古いコードが 64 を想定している場合、それは失われます。threadgroup_barrier。
• 開始方法: 最初に使用しますthreads_per_simdgroup同期を修正し、高周波カーネルを 32 とその他の幅の 2 つのバージョンに分割します。走行中に押すthreadExecutionWidth選ぶ。

5. 深度/ステンシル サンプリングをデュアル テクスチャ プロセスに変更します。

• 何をすべきか: 現在のパスは深度/ステンシルを書き込み、シェーダーは前の深度/ステンシルのコピーのみを読み取ります。
• 実行する価値がある理由: 同じレンダー パス内で現在の添付ファイルをサンプリングすることは未定義の動作です。Apple の互換性のあるスナップショット メカニズムは古い SDK パスでのみ機能し、パフォーマンスが低下します。
• 開始方法: 深度を読み取る必要があるエフェクトのパスの前に、サンプル可能なコピーを作成します。現在の深度/ステンシル アタッチメントをサンプリングされたテクスチャとして同時にバインドすることは禁止されています。

関連セッション

• Apple シリコン Mac 向けの Metal パフォーマンスの最適化 - このセッションの直接のフォローアップであり、Metal レンダラーに Apple GPU の帯域幅とタイル メモリを活用させる方法について引き続き説明します。
• Harness Apple GPUs with Metal — Apple GPU の機能から始まり、プログラム可能なブレンディングやタイル シェーダーなどの機能に適したレンダリング シナリオを追加します。
• GPU カウンターを使用して Metal アプリとゲームを最適化する — GPU カウンターを使用して、ロード/ストア、バリア、帯域幅の最適化によって実際にボトルネックが軽減されるかどうかを検証します。
• Discover ray tracing with Metal — 2020 年の Metal グラフィックス パイプラインのより高度なレンダリング機能について学び続けるのに最適です。