M3とA17 ProでのGPUの進化

こんにちは Jedd Haberstroです AppleのGPU, Graphics, and Displays Software Groupのエンジニアです

A17 ProおよびM3ファミリーのチップの 新しいApple family 9 GPU アーキテクチャについて説明します これはiPhone 15 Proと新しいMacの 中核的なテクノロジーです

ユーザーから高い評価を集める Apple製品の 豊かなユーザー体験は GPUが支えています 新しいiPhone 15 Proで 外出先でのゲームプレイを提供したり

新しいiMacでアプリの滑らかな UIアニメーションを実現したり

新しいMacBook Proで 機械学習を利用した 高度なビデオ処理や 画像処理を実行したりなど これらのアプリ機能を実現する上で

GPUは重要な役割を果たします

Metal APIを使用すると AppleのGPUの演算能力を 活用しながらアプリ全体で 多種多様な一連のMetalシェーディング 言語プログラムを実行できます これらのシェーダプログラムは 数行のコードを実行する 小さくシンプルなシェーダから 数十万行のコードを実行する 大規模で複雑なシェーダや フレームワーク ライブラリまで 多岐にわたります これらのシェーダに共通しているのは 大規模なデータ並列処理です この並列実行により アプリのパフォーマンスが大幅に 向上する可能性があります このような並列処理は 3Dレンダリングされたシーンの 頂点ごとや画面のピクセルごとなど 各入力に対して Metalシェーダプログラムを何度も 並列実行することで実現します

最終的に GPUがこれらのシェーダを 並列実行します

各のGPUの中心となるのが シェーダコアです 各シェーダコアは数千スレッドを 並行実行できます パフォーマンスをさらに拡大するために GPUには並列動作できる シェーダコアが多数あり アプリで数万のスレッドを 並列実行できます

既存のiPhoneやiPad Macに搭載されているGPUは すでに優れたパフォーマンスを 備えています

またアプリのデベロッパがGPUの 可能性を最大限に活用できる 強力なデベロッパツール スイートも用意されています しかし 新しいApple family 9 GPUでは いくつかの素晴らしい新機能進化により パフォーマンスが かつてないレベルまで向上しています

1つ目は 既存のアプリのパフォーマンスと 電力効率を改善するまったく新しい シェーダコアアーキテクチャです 提供する体験に関して メリットがすぐに得られ これから構築する次世代アプリの 厳しい課題にも対応できます

ハードウェアアクセラレーテッド レイトレーシングは MetalのRay Tracing APIを 使用しているアプリに 明確なメリットをもたらし またレイトレーシングの使用機会を拡大して 優れたパフォーマンスと 多彩なレンダリング効果を実現します

ハードウェアアクセラレーテッド メッシュシェーディングでは アプリでこれまで以上に高度な ジオメトリ処理パイプラインを 構築できます

これらを詳しく説明する前に 新しいApple family 9 GPUにより 現状のアプリで 変更を加えることなく実現される 驚きのパフォーマンスを見てみましょう これはLarian Studiosの 「Baldur's Gate 3」です 上はM3を搭載した新しい MacBook Proで 下はM2を搭載した MacBook Proで実行しています いずれもレンダリングには1,800pの 超高品質ビデオ設定を使用しています M3搭載のMacでは より高いスレッド占有率で ゲームのMetalシェーダを実行する 次世代シェーダコアの機能により 大幅にパフォーマンスが向上します これはBlenderでCycles Path Tracerを使用して 理髪店のシーンをレンダリングしたもので M3搭載のMac上でMetalレイトレーシング を活用しています 両方のレンダリングを同時に開始しましたが ハードウェアアクセラレーテッド レイトレーシングと次世代シェーダコアの 効果で M3搭載のMacのレンダリングは 大幅に速くなります

これは「Toy Story 4」の アンティークショップのUSDを リアルタイムで視覚化し PixarのHydra Stormで レンダリングしています Hydra StormはM3搭載のMacのMetal メッシュシェーディングを使用しており ハードウェアアクセラレーテッド メッシュシェーディングと 組み合わせることでこれまでよりも 高速に実行されます

これらの機能を詳しく見てみましょう まず次世代シェーダコアです

Apple family 9 GPUは複数の 要素で構成されています 計算や頂点のコマンドプロセッサは Metalコマンドバッファを解析します ラスタライザは実行のために フラグメントシェーダを ディスパッチします

またキャッシュの階層があり 例えば GPUラストレベルキャッシュは すべてのGPUメモリ トラフィックを処理します

ただ どのGPUも 中心はシェーダコアです これらが アプリのMetalシェーダを 実行する構成要素です

シェーダコアはテクスチャユニットや 新しいレイトレーシングユニットと ペアになっており テクスチャユニットは テクスチャリソースをサンプリング して書き込むことができ レイトレーシングユニットは レイ交差リクエストを高速化します

シェーダコアはさらに その構成部分に細分化できます

各シェーダコアには一連の 実行パイプラインがあり 様々な種類の命令を実行します 例えば FP32 FP16 整数 演算などは Metalデータ型のfloatやhalf intなどを使用する シェーダ内の変数の演算に対応します またメモリパイプラインは テクスチャとバッファの 読み取りと書き込み操作用です

これら実行パイプラインを すべて使用状態に保つには 通常は複数のSIMDグループの 命令を実行する必要があります そのため シェーダコアで実行中の SIMDグループを追跡するプールと どのSIMDグループの命令を 次に実行するかを選択する スケジューラがあります

一般に シェーダプログラムが使用する 様々な種類のデータを格納する オンチップメモリもいくつかあります レジスタは変数の値を保存します スレッドグループとタイルメモリは 計算スレッドグループ全体で 共有されるデータや タイル全体で共有されるカラー アタッチメントデータを保存します

キャッシュはスタックと バッファへのアクセスの パフォーマンスを向上させます

シェーダコアの機能と その構成要素についての 理解を踏まえて Apple family 9 GPUシェーダ コアの魅力的な新機能を 3つ説明します これらの新機能より アプリに変更を加えることなく シェーダのパフォーマンスを向上できます 一方 この新しいシェーダコアの機能を より深く理解することで そのメリットをさらに活用できます

最初の変更点は 動的シェーダコアメモリです これによりアプリのスレッド占有率が上昇し 多くの場合 パフォーマンスが向上します

2番目の変更点は 柔軟なオンチップメモリです これによりシェーダがバッファやスタック スレッドグループおよびタイルメモリに アクセスする際の 効率が向上します

最後の変更点はシェーダコアの 高性能ALUパイプラインです 並列実行能力が向上します これにより浮動小数点演算や整数演算の 組み合わせを実行するアプリの パフォーマンスが向上します

これらの新機能を詳しく見ていく前に シェーダコアが実行パイプラインを 使用状態に保つ方法とスレッド占有率の 重要性について説明します

MetalシェーダがALUパイプラインを 使用して いくつかの数学演算を実行した後 バッファを読み取り その結果が 直後に使用されるとします バッファにアクセスするには デバイスメモリにまでアクセスしなければ ならない場合があり 大きな遅延が生じる可能性があります この間 SIMDグループは ほかの操作を実行できず ALUパイプラインは使用されません

これを緩和するために シェーダコアは 別のSIMDグループの 命令を実行でき このSIMDグループには独自の ALU命令が含まれることもあります これにより ALUが使用されない 時間を短縮し SIMDグループを 並列で実行できるため パフォーマンスが向上します

シェーダコアで実行する 追加のSIMDグループがある場合は これを何度も繰り返すことができ ALUやその他の実行パイプラインで 実行する命令が 不足することはありません

シェーダコア上で 同時に実行するSIMDグループの数を スレッド占有率と呼びます

しかし シェーダコアで同時に実行する SIMDグループの数は何で決まる のでしょうか

この疑問に答えるために例を見てみます これはプロトタイプのレイ トレーシング計算カーネルです 光線をアクセラレーション構造と交差させ

交差結果を調べ

交差するプリミティブの マテリアルに基づいて 異なるシェーディング関数を実行します この例ではガラスとレザーの 両方のマテリアルのシェーディングを サポートしています

コードの各行はプログラムの 変数を格納するために ある程度のレジスタを使用します プログラムの途中ではコードの 動作に応じて 使用するレジスタの量が 増減します この例では shadeGlass関数の実装がプログラムの ほかの部分よりもレジスタを多く使用します

Apple family 9 GPUより前は オンチップレジスタファイルから レジスタが割り当てられるまで SIMDグループはシェーダコアで 実行を開始できませんでした 割り当て量はプログラム内の 任意の時点での最大レジスタ 使用量に一致します SIMDグループは継続期間にわたって その割り当てられたレジスタを 確保し続けます ただし プログラムの大部分で それらのレジスタのほとんどが 使用されないこともあります したがって 最大レジスタ使用量に応じて 例えば シェーダコアで一度に4つの SIMDグループしか実行 できない場合があります それより多いと より多くのオンチップレジスタファイル メモリが必要になるためです ただし Apple family 9 GPUの 新しい動的シェーダコアメモリ機能により レジスタの最大使用量が 実行可能なSIMDグループの数に 影響を与えることがなくなりました オンチップレジスタメモリは プログラムの各部分が実際に 使用する量に応じて シェーダの存続期間中に 割り当てと割り当て解除が 動的に行われます これによりSIMDグループは オンチップレジスタファイルを より効率的に使用でき より多くのスペースを解放して 使用可能にします 同時に実行できる SIMDグループが増えるため アプリのスレッド占有率に加え パフォーマンスにも大きな影響を 与える可能性があります

先ほど述べたように SIMDグループの継続期間に わたってレジスタは 動的に割り当てと 割り当て解除が行われます これは部分的には レジスタファイルがレジスタの 永続的なストレージではなく キャッシュになっているためです つまり チップ上に格納できる限り より多くのレジスタを使用できます

柔軟なオンチップメモリ機能では スレッドグループやタイルメモリなど シェーダコアの残りのメモリタイプにも この処理を拡張し キャッシュとして実現しています

レジスタ スレッドグループ タイル スタックおよび バッファデータがすべてチップ上に キャッシュされるようになり オンチップメモリを これらすべてのメモリタイプを 処理する少数の大きなキャッシュとして 再設計できます この柔軟性は各メモリタイプを あまり使用しないシェーダに役立ちます 以前は 例えば計算カーネルが スレッドグループメモリを 使用しない場合 対応する オンチップストレージは まったく使用されませんでした 新しいアーキテクチャでは オンチップストレージが シェーダによって使用されるメモリ タイプに動的に割り当てられます これまでよりも多くのオンチップ ストレージを使用できるため パフォーマンスが向上します

例えば レジスタの使用量が 多いシェーダの場合 占有率が高くなる可能性があります

バッファデータの大きなワーキングセットに 繰り返しアクセスするシェーダの場合 キャッシュヒット率が向上し バッファアクセスの遅延が減って パフォーマンスが向上します また 関数ポインタ 可視の関数テーブル 動的にリンクされたシェーダライブラリなど 非インライン関数を頻繁に 使用するアプリの場合 関数パラメータを渡す オンチップスタック領域が増え 関数呼び出しが高速化することを 意味します

しかし アプリがオンチップ ストレージよりも多くのメモリを 使用する場合はどうなるでしょうか そのままでは データが 次のキャッシュレベルまたは メインメモリにまでスピルします 幸い シェーダコアはシェーダの動作を 動的に監視し 占有率レベルを調整して スピルの発生を防ぎます これによりデータが チップ上に保持されます そして最終的に実行パイプラインが 使用状態になります

ただしこれは シェーダの占有率が シェーダの動的なレジスタ使用率だけでなく スレッドグループや タイル スタック バッファメモリへの シェーダのアクセス状況よって 影響を受けるということです

これらの新しいハードウェア機能により 多くのアプリの占有率が改善されます つまりデベロッパが占有率を 最適化する必要性が これまでよりも大きく減少します ただし Apple family 9 GPUで占有率を さらに最適化する必要がある場合 当社が開発したプロファイリング ツールのスイートが役に立ちます 占有率を診断して最適化する 方法の詳細については これらの解説をご覧ください

Apple family 9 GPUの シェーダコアの機能として 最後に説明するのは 高性能ALUパイプラインです

Apple GPUシェーダコアには FP16命令など 様々な命令タイプに対応する 個別のALUパイプラインがあります Apple GPUはFP16演算を実行する ために高度に最適化されています そのため できる限りFP16データ型を 使用することをお勧めします

FP16演算命令はピーク スループットで実行されます

同等のFP32よりも使用する レジスタが少なくて済みます バッファがFP16にネイティブに データを保存する場合は メモリ帯域幅が減少します また 数学演算のソース変数または デスティネーション変数が 現状FP16でない場合 コストをかけずに変換できます

ただし アプリでFP32や整数など ほかの数学演算を実行する場合 Apple family 9 GPUシェーダコアでは 3つすべてのデータ型の命令を これまで以上のレベルで並列実行できます これにより 以前のApple GPUと比較して 最大2倍のALUパフォーマンスを 実現できます この特別な並列処理の効果を得るには 命令を複数のSIMDグループから 実行する必要があります つまり 占有率を増やすことで ALUパイプラインの使用率が向上します 例を考えてみましょう 2つのSIMDグループが 同時に実行されており 両方ともALU命令を実行するとします 従来 これらのSIMDグループは 1つずつ実行することが 求められる場合がありました

ただし ここに示すように 異なる時点で 実行するFP32命令と FP16命令がある場合 それらの実行を重複させることができ 並列性とパフォーマンスが向上します

次世代シェーダコアの新機能をまとめます シェーダの継続期間にレジスタの 割り当てと割り当て解除が動的に行われ シェーダのスレッド占有率が向上します

大きなオンチップキャッシュでレジスタ スレッドグループ タイル スタック およびバッファメモリを処理し これらのメモリタイプへの アクセス性能が向上します

シェーダコアは占有率を動的に調整し チップ上のデータと実行パイプラインを 使用状態に維持します

さらに FP16 FP32 整数演算を これまで以上のレベルで並列実行でき ALUのパフォーマンスが向上します

次に ハードウェアアクセラレーテッド レイトレーシングを見てみましょう

Metalレイトレーシングを使用すると アプリでApple GPUによる 大規模な並列処理を活用し 光線をシーンジオメトリと 交差させることができます Metalレイトレーシングに慣れていない場合や 詳細を知りたい場合は 「Metalレイトレーシングのガイド」や 「Metalレイトレーシングによる アプリの強化」をご覧ください

Metal Ray Tracing APIの 中心となるのは インターセクタオブジェクトであり アクセラレーション構造に含まれる プリミティブと 光線の交点を決定します このオブジェクトはレイトレーシング アプリのGPU関数 別名シェーダによって何度も 呼び出される場合があり アプリのパフォーマンスの面で重要です

先ほど このraytracingKernelの レジスタ使用量を確認するときに このようなGPU関数を示しました ここではインターセクタ オブジェクトを作成し オブジェクトのintersectメソッドを 呼び出して交差部分を見つけています

交点を決定するために インターセクタでは いくつかの重要なステージを実行します 最初にアクセラレーション 構造をトラバースして 候補になるプリミティブを見つけます 次に アプリが提供する交差関数を呼び出し 光線がプリミティブと交差するかを 判断します

交差する場合 その交点は前の交点と比較され 最も近い交点が見つかるまで これが繰り返されます

その後 最も近い交点が 呼び出し元のGPU関数に返され アプリ固有の処理が実行されます

Apple family 9 GPUで 新たに導入された インターセクタオブジェクトの実装は ハードウェアアクセラレーションを 使用しており この重要な操作のパフォーマンスが 大幅に向上します

ハードウェアアクセラレーションを 使用するインターセクタは GPU関数と連動しては実行されません そこで 両者間での光線と レイペイロードの やり取りを容易にするために データがオンチップメモリに 読み書きされます これは新しいXcodeの RTスクラッチパフォーマンス カウンタで確認できます

インターセクタの役割と機能について 説明したので 例を使用して 1回のインターセクタループで パフォーマンス特性を見てみましょう

アプリが2つのSIMD グループを実行しており それぞれが4つの光線と アクセラレーション構造を 交差させる場合を考えます

この例では アクセラレーション構造に 典型的なコーネルボックスが含まれ その中にボックスオブジェクトと 球オブジェクトが1つずつあります

光線 アクセラレーション構造 交差関数テーブルを渡して intersectメソッドを呼び出すことで 光線がシーンにキャストされます 各SIMDグループにはボックスと 交差する2つの光線と 球と交差する2つの光線があります この例では ボックスを不透明な三角形 プリミティブとして定義するために MTLAccelerationStructure

TriangleGeometryDescriptorを使用し opaqueプロパティをYESに設定しています したがって インターセクタは Metalに組み込みの 交差関数を使用して交差を計算できます

ただし 球はインターセクタが 呼び出す必要がある カスタム境界ボックス交差関数を使用して 手続き的に定義されます

カスタムBoundingBoxIntersection関数は bounding_boxパラメータを含む intersection属性を使用して宣言しています

前述のように intersectメソッドは アクセラレーション構造に対して光線を テストする各スレッドで呼び出されます この例を念頭に 従来の実装で 各交差のトラバースの呼び出しと 交差テストがどのように実行されるか 見てみましょう

一般的な方法では 光線をテストするプリミティブを見つける 所要時間はトラバースによって異なります これにより いわゆる実行分岐が発生し SIMDグループ内の各スレッドは 次のステージに進む前に そのグループの 最長トラバースの完了まで待機します

またこのように 交差関数を実行するときも 同様のオーバーヘッドが生じます 実行分岐により 各タイプの交差関数が1つずつ実行され 並列性がさらに低下します 全体的にすべてのステージを通じて 各スレッドはランタイムの 大部分をアイドル状態に費やし SIMDグループ内のほかのスレッドが 完了するのを待ちます これがパフォーマンスの 大きなボトルネックとなります

従来の実装のこのようなイメージを念頭に ハードウェアアクセラレーテッド レイトレーシングが こうした非効率性をどのように 最適化するか説明します

最初の大きな改善は ハードウェアインターセクタが固定機能の ハードウェアを使用して各トラバースを 完全に独立して実行できることです これが可能な理由の一部は 光線をGPU関数と連動して 処理するのではなく ハードウェアインターセクタに 送っていることにあります これによりトラバース時間が大幅に短縮され 従来のトラバースの実行分岐による オーバーヘッドが排除されます

一方で交差関数は Metalシェーディング言語コードであるため SIMDグループにグループ化して シェーダコア上で実行する必要があります ただし ハードウェアインターセクタは 各光線を独立して処理するため 別々のSIMDグループから発生した光線の 交差関数呼び出しを自由に グループ化できます

これが並べ替えステージの役割です 時間的に近い範囲にある光線が このステージに到達すると 交差関数呼び出しが 一貫したSIMDグループにグループ化され 従来の実装に存在する 実行分岐のオーバーヘッドが減少するか 完全に排除されます

ハードウェアアクセラレーテッド レイトレーシングがアプリの レイインターセクタ呼び出しの パフォーマンスを どのように向上させるかを説明しました 次に アプリの能力を引き出すための ベストプラクティスをいくつか説明します

最初の提案は できる限り インターセクタオブジェクトAPIを 使用することです Metalではレイトレーシングの実行に 交差クエリAPIも使用できます ただしこのAPIでは 読み書きすべきレイトレーシング スクラッチメモリの量が増え 並べ替えステージが無効になります

また カスタム交差関数を 作成するときは 多数の異なる論理交差ルーチンを 実行できる1つのウーバー関数を 作成するのは避けることをお勧めします 代わりに 論理交差ルーチンごとに 1つのMetal交差関数を作成してください 並べ替えステージのメリットが増します

また インターセクタ オブジェクトとやり取りする レイペイロード構造体のサイズを 最小限に抑えることも重要です これによりシェーダの遅延が減少し スレッド占有率が増加する可能性があります

レイトレーシングアプリの最適化方法の 詳細とガイダンスについては これらの解説をご覧ください

まとめです Apple family 9 GPUでは 固定機能のトラバースブロックと 交差関数の並べ替えステージを伴った 新しいハードウェア アクセラレーションにより レイトレーシングのパフォーマンスが 大幅に向上します

この新しいハードウェアでは すべてのMetalレイトレーシング アプリのパフォーマンスが向上しますが アプリで最大限のメリットを得るには できる限り 交差クエリAPIではなく 交差APIを使用することをお勧めします

最後に紹介するApple family 9 GPUの新機能は ハードウェアアクセラレーテッド メッシュシェーディングです

メッシュシェーディングは レンダリングパイプラインで柔軟に行える GPUを使用したジオメトリ処理ステージであり 従来の頂点シェーダ段階を2つの計算的な シェーダで置き換えます

オブジェクトシェーダは 最初のステージで実行され メッシュオブジェクトの全体など アプリ固有の入力の 粗粒度処理を実行する ために使用できます 各オブジェクトスレッドグループでは メッシュグループを生成して 後続のより細粒度の高い処理を 実行することもできます メッシュシェーダは第2のステージを構成します 通常 メッシュスレッドグループは 親オブジェクトの構成部分 つまり メッシュレットとも呼ばれる 部分を処理します

メッシュスレッドグループの出力は 従来のグラフィックスパイプラインの 残りの部分で処理される 頂点とプリミティブのリストを カプセル化した Metalメッシュオブジェクトです

メッシュシェーディングには 多数の用途があり 例えば 細粒度ジオメトリカリングや

手続き型ジオメトリ生成 圧縮形式のようなカスタムの アプリ固有ジオメトリ表現 さらに ほかのグラフィックスAPIからの ジオメトリシェーダや テッセレーションシェーダの ポーティングに使用されます

Metalのメッシュシェーディングに なじみのない方は これら2つの解説を チェックしてください

Apple family 9 GPUの ハードウェアアクセラレーテッド メッシュシェーディングにおける 最も注目の改善点は 既存のメッシュシェーディングコードでの パフォーマンスの大幅な向上です

Apple family 9 GPUは オブジェクトとメッシュスレッドグループを より効率的にスケジューリングし 中間メッシュレットデータを チップ上に保持できます そのため メモリトラフィックが減少します

新しい ハードウェアでは Metal APIの 機能強化もいくつか追加されています 1つ目は間接コマンドバッファに対する 描画メッシュコマンドの エンコーディングのサポートです これにより GPUを使用する レンダリングパイプラインで 従来の頂点シェーダに加えて メッシュシェーディングを処理できます

APIの2つ目の機能強化として メッシュグリッドあたりの スレッドグループの最大数が 1,024から100万以上に拡大されています

ここで最適なメッシュシェーディング パフォーマンスを実現する いくつかのベストプラクティスを 確認しましょう

メッシュスレッドグループによって 出力されるmetal::meshオブジェクトには いくつかのテンプレートパラメータがあり そのサイズをできる限り 小さく保つことが重要です

そのためには メッシュの頂点データ型と プリミティブデータ型の場合 例えば これらのデータ型を ほかの無関係な頂点関数または メッシュ関数と共有する目的で 未使用の属性を用意することがありますが この未使用の属性を削除することが 考えられます メッシュ型では 出力できる プリミティブと頂点の最大数も指定する 必要があります これらはアプリのジオメトリや パイプライン アセットが実際に必要とするよりも 大きく設定しないでください これらのサイズを考慮することで メモリトラフィックが減少し 占有率が増加する可能性があります

メッシュシェーダでプリミティブごとの カリングを実行する場合 後続のカリングステージにハードウェアで カリングする目的のみで頂点位置を メッシュオブジェクトに書き込む ことはお勧めしません 代わりに ハードウェアの残りの ジオメトリ処理ステージで 処理時間を大幅に節約できるように このようなプリミティブの書き込みを 完全に省略することをお勧めします

Apple family 9 GPUについて 説明した内容をまとめます

この次世代シェーダコアでは レジスタストレージを動的に割り当て 複数のメモリタイプ間でオンチップ メモリを共有することにより オンチップメモリの使用率が増加し スレッド占有率とパフォーマンスが 向上します

ハードウェアアクセラレーテッド レイトレーシングでは Metal Ray Tracing APIを 使用することで アプリのパフォーマンスが大幅に向上し 新しい精彩な視覚効果を実現します そして最後に ハードウェア アクセラレーションにより メッシュシェーディングの パフォーマンスが大幅に向上し より多くのアプリで ジオメトリ処理パイプラインを カスタマイズできます

ご視聴ありがとうございました