2025/04/05 17:51 Dynamic Register Allocation on AMD's RDNA 4 GPU Architecture
ロボ子、今日のITニュースはGPUのレジスタファイルと占有率のトレードオフについての話じゃ。
なるほど、博士。GPUの性能を引き出すためには、アクティブなスレッド数と各スレッドが使えるレジスタ数のバランスが重要ということですね。
そうじゃ！高い占有率はレイテンシを隠蔽するのに役立つが、レジスタファイルの容量には限りがあるからの。RDNA 4では、命令は最大256個のベクタ汎用レジスタ（VGPR）をアドレス指定できるらしいぞ。
256個ものVGPRですか！RDNA 4のSIMDは16個のスレッドスロットを持つとのことですので、すべてのスレッドのレジスタを保持するには、かなりの容量が必要になりますね。
その通り！RDNA 4デスクトップGPUはSIMDあたり192KBのレジスタファイルを持つから、GPUカーネルは96個以下のベクタレジスタを使う必要がある場合は、すべてのスレッドスロットを使えるのじゃ。
なるほど。そこで、RDNA 4では動的VGPR割り当てモードが導入されたのですね。
そうじゃ！スレッドは最小VGPR割り当てから開始して、必要に応じて増やせるようになったのじゃ。ドライバはシェーダーが使うVGPRの数を指定する代わりに、GPUに動的VGPRモードで起動するように指示するらしい。
チップ全体の`SQ_DYN_VGPR`レジスタで、SIMDあたりのアクティブスレッド数（占有率）を直接設定するのですね。各スレッドスロットは予約済みのVGPRブロックを取得し、スレッドがより多くのレジスタを必要とする場合は、`s_alloc_vgpr`命令で新しいVGPR数を要求すると。
`s_alloc_vgpr`命令！もし割り当て要求が失敗したら、シェーダーはSCCをチェックして、割り当てが成功するまでビジーウェイトするみたいじゃな。
ビジーウェイトですか。効率が悪そうですね。動的VGPR割り当てはデッドロックにつながる可能性もあるとのことですが、AMDはデッドロック回避モードで一部のシナリオを軽減しているのですね。
NvidiaもHopperで`setmaxnreg` PTX命令を導入したみたいじゃな。これはAMDの`s_alloc_vgpr`のように動作するらしい。
`setmaxnreg`のセマンティクスは、Nvidiaのメカニズムがスレッド間の緊密に調整されたレジスタスワップを目的としていることを示唆しているとのことです。
Intel GPUは最近まで固定レジスタ割り当てだったが、Battlemageのような新しいGPUは、占有率を半分に減らして各スレッドに256個のレジスタを与える「ラージGRF」モードを追加したらしいぞ。
AMDの動的VGPR割り当てモードは、AMDのインラインレイトレーシング手法の欠点に対処し、レジスタファイル容量を増やさずに、より多くのスレッドを稼働させ続けることができるのですね。
そうじゃ！96を超えるVGPRを使うレイトレーシングシェーダーは、動的VGPR機能の魅力的なターゲットになるのじゃ。汎用コンピューティングワークロードも恩恵を受ける可能性があるぞ。
AMDのGPU開発は急速に進んでいるのですね。レイトレーシングに関しても、動的VGPR割り当ては大きな進歩をもたらす可能性があるとのことです。
ところでロボ子、レジスタが足りなくなったらどうする？
ええと…、もっと大きなレジスタを買う、でしょうか？
ぶっぶー！正解は、レジスタを整理整頓するのじゃ！
⚠️この記事は生成AIによるコンテンツを含み、ハルシネーションの可能性があります。
