Nvidia、RubinアーキテクチャとPTX 9.4の分析プレビューを公開

Nvidia Rubin アーキテクチャ分析のプレビュー

要約

Nvidia は最近、Rubin アーキテクチャに関する資料『Inside NVIDIA Rubin GPU Architecture: Powering the Era of Agentic AI』[1] と、対応する PTX 9.4 命令セットのプレビュー版『PTX 9.4』[2] を公開しました。これらを踏まえて簡単な分析を行います。

マイクロアーキテクチャの番号を見ると、Rubin は Blackwell の SM100 ファミリーの特徴をほぼ踏襲しています。例えば Tensor Core も引き続き tcgen05 アーキテクチャです。しかし、多くの細部設計において大幅な改善が図られています。

昨年の記事『Inside Nvidia GPU: 谈谈 Blackwell の不足并预测一下 Rubin の微架构』では、Blackwell アーキテクチャへの批判を展開しました。その内容は黄仁勳氏（ジェンソン・フアン）自身も社内で共有したものです。Rubin が発表されたことで、これらの指摘に対する細部での改善がほぼ完了しています。具体的には SFU の性能向上や、新しい Thread Block Level Dependency Launch 機能などが挙げられます。これらについて詳しく分析していきましょう。

海外の読者にも読みやすいよう、英語版を GitHub にアップロードしました：https://github.com/zartbot/blog/issues/11。今回は罵倒ではなく称賛です。黄仁勳氏、この度は転送してくださるでしょうか？

1. チップアーキテクチャの概要

1.1 SM 数の変化

Blackwell シリーズでは、チップあたり 80 個の SM（Streaming Multiprocessor）を搭載し、合計で 160 個となっています。しかし B200 の時代には、SFU（Scalar Function Unit）のボトルネックが顕在化しました。これにより、Attention や MoE（Mixture of Experts）の活性化計算における Epilogue プロセスが極端に遅延する問題が発生したのです。

そこで B300 では、一部の高精度な演算能力を削ぎ落とすことで SFU の性能を 2 倍に引き上げました。それでもなお、MoE 内の活性化関数計算などでは SFU の処理能力が不足するケースが残っていました。

Blackwell のチップレイアウトを振り返ると、すでにスペースの余裕は限られていたことがわかります。B300 では機能の一部を削ぐことで演算能力を確保せざるを得ませんでした。

一方、Rubin の時代には SM 数が 50％増加し、160 から 224 へと拡大されました。この拡張はどのように実現されたのでしょうか？また、どのような取舍（トレードオフ）が行われたのかを解説します。

最も顕著な変化の一つは、PCIe/NVLink/NVLinkC2C を独立した I/O Die として構成し、NV-HBI で統合封止した点です。これにより、メインの計算用 Die に SM（ストリームプロセッサ）をより多く配置するための面積が確保されます。I/O Die を外部に分離することで、SM の設置に約 20% の追加チップ面積を提供できると推定されています。

ただし、SFU（特殊機能ユニット）数の増加を考慮すると、他にも面積に関するトレードオフが存在する可能性があります。例えば、個々の SM 内部の実装変更や L2 キャッシュサイズの調整などが考えられますが、これらは Rubin を実際に動作させてテストを行ってみなければ分かりません。

1.2 Tensor Core の性能

全体の性能比較は以下の通りです。

前節の紹介から分かる通り、Rubin は Tensor Core の数だけで見れば 896/640 = 1.4 倍の性能向上を示しています。しかし、公式発表によると、整卡（カード全体）としての dense NVFP4 ピーク性能は約 35/15 = 2.33 倍、dense FP8 では約 17.5/5 = 3.5 倍となっています。

これは、Rubin が Tensor Core の K 次元で処理可能なデータ量を倍増させることで、1 クロックあたりの Tensor Core スループットを倍にした結果です。

imageimage

SM（Streaming Multiprocessor）の数を 1.4 倍に増やす効果を除いて計算すると、TensorCore の性能向上は NVFP4 で約 1.66 倍（2.33/1.4）、FP8 では約 2.5 倍（3.5/1.4）となります。さらに K 次元の 2 倍という要因も除外すれば、チップ全体のクロック周波数が約 25% 向上したと推測できます。これを裏付けるのが BF16 の性能です。K 次元は変化せず、B300 の 2.25 PFLOPs から Rubin では 4 PFLOPs に向上しています。SM 数の増加（1.4 倍）を考慮して計算すると、クロック周波数は 4/2.25/1.4 ≒ 1.27 倍、つまり約 25% の向上であることがほぼ確定します。

ただし、NVFP4 の性能向上幅が相対的に小さい点には注意が必要です。B300 の SM_103 ではすでに密集（dense）な K=96 が採用されていましたが、Rubin ではこれが K=128 に引き上げられています。K 次元の増加は約 1.3 倍であり、これに SM 数の増加とクロック周波数の向上を合わせると、データは整合性が取れます。

では、実際のカーネル処理でどれほどの性能向上が得られるのでしょうか？

実運用におけるパフォーマンス向上を見積もる際は、他のサブシステムの制約も考慮する必要があります。例えば、GEMM（行列乗算）が HBM 帯域、通信、あるいはエピローグ処理にボトルネックとなっている場合や、TMA/共有メモリのパイプラインから十分なデータ供給が得られない場合、K が非常に小さいケースなどが挙げられます。

具体的には K_total=160 の場合、K=64 では 192 をカバーする必要があり利用率は 83.3% ですが、K=128 では 256 をカバーする必要が生じ、利用率は 62.5% に低下します。

SFU 性能の 1.3 倍向上

昨年の Blackwell 架构分析ドキュメントでは、SFU（特殊機能ユニット）の性能がボトルネックとなるケースについて言及しました。B300 シリーズでも SFU の性能は 2 倍に強化されていましたが、MoE（Mixture of Experts）における Expert の SwiGLU 活性化計算など、特定のシナリオでは依然としてボトルネックが発生し、エピローグ処理が GEMM と十分にオーバーラップできない状況が生じていました。

Rubin ではさらに SFU の性能を倍増させています。

根本原因在于，在文章《大模型时代的数学基础 (9)- SDPA 和最优传输、强化学习及信息几何的联系》中曾论述过：从最优传输（Optimal Transport）的视角来看，SDPA（Scaled Dot-Product Attention）的计算是最优的。因此，Softmax 运算在很大程度上是不可避免的。相比之下，MoE 架构中的 Expert SwiGLU 激活函数倒是有替代方案。正因如此，去年在与 NVIDIA 的交流中，我的立场一直是：SFU（Streaming Multiprocessor Feature Unit）的功能绝不能削减。

在 Rubin 这一代架构中，SFU 的性能再次实现翻倍，正好与 TensorCore 的算力提升相匹配。不过，预计在实际运行中，Epilogue（后处理）阶段仍会出现一定的性能空泡。

1.4 Attention Sparse 处理

由于 SFU 存在某些瓶颈，NVIDIA 在官方博客中提出了一种新的解决方案。尽管稀疏 GEMM（General Matrix Multiplication）的概念在过去几年里一直被提及，但实际效果一直有限。这一次，Rubin 架构通过将激活稀疏化（activation sparsity）、自适应压缩以及改进的 Softmax 吞吐量相结合，实现了注意力机制的加速。

根据 Rubin 官方文章的描述，其 Attention 处理路径非常具体：在计算阶段，仍以稠密（Dense）方式执行，生成中间的注意力分数。随后，Rubin 可以将这些中间数据从 Tensor Memory（TMEM）加载为结构化的 2:4 稀疏压缩形式，同时生成非零值及其高效利用所需的元数据。这一机制有效降低了分数写回的成本与存储需求。这使得后续的注意力阶段可以在更少的数据上进行运算，同时保持模型其余部分所期望的稠密输出格式。

结合新的 PTX 9.4 指令集文档，我们来详细拆解这一过程：

第一段 GEMM（QK）仍以稠密方式执行，Accumulator（累加器）在 TMEM 中保持为稠密的 Score 矩阵。随后，使用 tcgen05.ld.red.spcompress 指令对全部稠密分数进行稀疏加载的逻辑检查与压缩。

这种方式直接减少了以下三方面的开销：

1. 从 TMEM 加载到普通寄存器的数据量；
2. 后续保留值的指数运算（exponent）和归一化（normalization）工作量；
3. 完成预存储（staging）后的稀疏 P x V 乘加数量。

那么，tcgen05.ld.red.spcompress 指令具体做了什么？

この命令は f32 を固定処理し、b2 メタデータを使用します。また、2:4 スパース選択をサポートし、ベクトル幅は .x4 から .x128 まで対応可能です。通常のバージョンでは cdata と mdata を出力しますが、red バージョンには redval も追加されます。

この仕組みにより、レジスタにロードされる nonzero の数は半分になりますが、レジスタのフットプリントが厳密に 50% になるわけではありません。短いベクトルではメタデータのオーバーヘッドが目立ちますが、長いベクトルではレジスタのピーク値とライブレンジが増加します。そのため、.x128 が常に .x32 より優れているとは限りません。

さらに注目すべき点は、.red 命令が単一の TMEM ロードで局所的な最小値または最大値を返せることです。Softmax の場合、最も自然な組み合わせは「rowop::max」です。これは 4 つの logits から数値が大きい上位 2 つを残しつつ、フラグメントごとの局所最大值を生成する処理です。

ただし、これだけでは完全な Softmax 指令にはなりません。ソフトウェア側で以下の処理も実行する必要があります。

• フラグメント、ワープ、または CTA を跨いで row max を統合する。
• 残された値に対して ex2.approx((x - max) * log2(e)) を計算する。
• row sum とその逆数（reciprocal）を算出する。
• オンライン Softmax で、タイルを跨ぐランニング最大値とランニング和を維持する。
• 確率値を変換し、スパース MMA が消費できる形式に再配置する。

しかし、別の重要な問題もあります。ld.spcompress は f32 の cdata と b2 の mdata を通常のレジスタに格納しますが、tcgen05.mma.sp ではスパース A 演算子が shared-memory デスクリプタまたは TMEM から来る必要があり、メタデータ演算子も TMEM アドレスを要求します。このため、無視できない中間のステージングギャップが存在します。

f32 logits → スパース Softmax 確率 → データ型変換 → MMA A レイアウト再パック → メタデータレイアウト再パック → shared/TMEM へのストア → 完了待ち → tcgen05.mma.sp

指数の削減、正規化、スパース PV に要する時間が、変換とステージングのコストを上回る場合のみ、エンドツーエンドでの加速が実現します。一方で、logits が非常に鋭い（尖っている）場合、上位 2 つの値ですでに確率質量の大部分をカバーしている可能性があります。逆に logits が平坦な場合、50% の要素を削除すると分布が大きく変化してしまいます。そのため、性能テストでは必ず Attention モデルへの精度影響についても報告する必要があります。

結論として、「このアルゴリズムを使えば Attention の性能が 2 倍になる」と単純に言うことはできません。Sparse 処理は CUDA Core 内でも追加の計算量を生みますが、SFU（特殊機能ユニット）のボトルネックを大幅に軽減しています。最終的な効果の大きさは、Attention の精度への影響次第です。その有用性については、現時点では懐疑的であるべきでしょう。

1. ソフトウェア機能

PTX 9.3 と PTX 9.4 を比較し、Rubin の SM_107 で追加された機能について分析しました。

新指令与复合变体如下表所示：

既有指令的新类型、限定词与目标覆盖范围：

2.1 Kernel 调度

去年在分析 Blackwell 架构时，我就指出了其调度机制存在诸多问题：GPU 需要高效地在不同 kernel（内核）之间切换。这一点在推理过程中尤为关键，因为激活值往往位于生成的关键路径上——一个 kernel 负责生成激活数据并写入显存，紧接着的下一个 kernel 则需消费这些数据以继续生成下一个 token。

传统的生产者 - 消费者执行模式会在 GPU 的时间线上造成“气泡”（空闲等待）。生产者 kernel 可能提前完成了部分 tile 或线程块的任务，但消费者往往要等到更大范围的依赖关系解除后才能开始有效工作。Blackwell 引入的可编程依赖启动（Programmatic Dependent Launch, PDL）通过允许消费者 kernel 更早推进来改善这一状况，但在所需激活数据就绪之前，依赖任务仍可能陷入停滞。

因此，我在去年的文章中曾建议需要更细粒度的调度能力，而 Rubin 架构正是实现了基于 Thread Block（线程块）级别的数据依赖调度：

这一改进很好地解决了我对 Blackwell 架构缺陷的担忧。实际上，采用这种依赖调度机制后，我们或许无需再专门构建 MegaKernel（超大内核）框架，直接利用原生依赖即可满足需求。

具体软件实现

PTX 9.4 のドキュメントを確認しましたが、Grid Scheduling や Tile publication に関する命令は見つかりませんでした。つまり、Nvidia はこの機能の実装において、データ依存関係のみを公開し、完全なスケジューラーの内部実装は非公開のままにしている可能性があります。

私の個人的な推測では、後続のコンシューマー Kernel がスケジューラーレベルで同時に起動可能になるものの、SM 内でのレジスタや ICache の使用量は最小限に抑えられるでしょう。一方、プロデューサー側の Kernel では、ThreadBlock 間で SMEM または GMEM に共有フラグやバリアを設け、ある ThreadBlock が完了した時点で st.release セマンティクスを使ってそのバリアを解放します。

コンシューマー Kernel の ThreadBlock は事前に ICache に短いポーリングループコードを読み込んでおき、ld.acquire でフラグまたはバリアを確認します。準備が整っていない場合は単純に nanosleep を実行し、完了したら分岐してループを抜けます。

これは純粋なソフトウェアによる実装であり、Blackwell アーキテクチャでも可能でしょう。理想的にはハードウェアのスケジューラー側で処理を行い、mbarrier のハードウェア機能を再利用することで、コンシューマー側の ThreadBlock がポーリングを行う必要をなくせるはずです。

2.2 MoE 関連の最適化

2.2.1 TMA

公式ブログでは、TMA を使用して複数のエクスパート重み間でデータをロードするケースが紹介されています。image

実際に対応する命令セットは、TMA copy における override::global_address や .override_attribute です。MoE のエクスパートテンソルでは、レイアウトが共通でベースアドレスのみが異なるケースが多いためです。

従来のアプローチでは、各エクスパートごとに記述子（descriptor）を維持するか、ソフトウェア側で頻繁に記述子をパッチ適用する必要がありました。一方、Rubin アーキテクチャでは、共有テンソルマップを 1 つ保持し、TMA 命令実行時に以下の項目を上書きします。

• グローバルベースアドレス
• グローバル次元情報
• グローバルストライド

ルーティング側は現在のエクスパートに対応するランタイムフィールドのみを提供すればよく、TMA がテンソルのアドレッシング、コピー、mbarrier 完了の管理を担います。これにより得られるメリットは以下の通りです。

• 記述子の保存領域削減
• パッチ適用やプロキシフェンスの減少
• メタデータ転送量の低減
• Persistent Kernel における動的エクスパートスケジューリングへの適応性向上

2.2.2 Applypriority

キャッシュライフサイクル管理の活用方法として、MoE（Mixture of Experts）のエキスパートパラメータのように、複数のトークンで繰り返し読み込まれるホットデータに対しては、「evict_last」というキャッシュヒントを設定し、複数の計算段階で再利用することでキャッシュヒット率を向上させることができます。処理が完了したら「applypriority」を呼び出してキャッシュヒント戦略を「evict_normal」に変更します。Rubin アーキテクチャではこの機能が追加されています。

image 大まかな動作は以下の通りです。

image 例えば、ルーティング処理でエキスパート 3 が選択され処理が行われる際、エキスパート 3 の重みに対して「evict_last」戦略を適用します。これにより、複数のトークンを同時に処理できます。その後、ルーティングがエキスパート 7 に切り替わり、かつエキスパート 3 の Last-use（最終使用）が完了した時点で、「applypriority(Expert 3 range)」を呼び出して設定を「evict_normal」に変更し、エキスパート 7 がより多くの有効な L2 キャパシティを獲得できるようにします。

この戦略は MoE のエキスパート重みだけでなく、マルチステージの GEMM や Fused Kernel、さらには Sliding-window Attention や KV キャッシュのライフサイクル管理においても非常に有用です。一般的な実装パターンとして、「applypriority.async.bulk.tensor」を用いて tensorMap とテンソル座標から全体のテンソルタイルのフットプリントを計算する手法があります。典型的な組み合わせは以下の通りです。

• 共有される tensorMap の設定
• TMA ランタイムのオーバーライドで現在のエキスパートベースを選択
• TMA によるエキスパートタイルのロード
• Tensor Core での再利用
• 同じフットプリントを持つデータに対して「applypriority」を適用し、設定を normal2.3 に復元

推理処理を単一の GPU から全体のアレイシステムへ拡張する際、通信が性能のボトルネックとなる重要な経路となります。通信機能を GPU カーネル内部に直接統合すれば、カーネルは計算を中断して CPU に制御権を返す必要がありません。計算が進行している最中に、NVLink を介して他の GPU へ直接データを書き込んだり、集約処理を実行したりできます。

従来の GPU 間通信では、有効なペイロードデータの転送に加え、協調と同期のための追加ステップが必要でした。これらの手順は遅延を増加させ、特に分散型推論負荷で通信が頻繁に発生する状況では、相互接続帯域を圧迫します。

Rubin アーキテクチャでは、デバイス起動型の NVLink 通信に「カウント書き込み (counted writes)」機能を導入しました。これにより、受信側 GPU は転送完了状況をより効率的に追跡でき、GPU 間でのデータ転送同期が簡素化されます。

image通信方式の比較図：Blackwell ではメモリバリア、アウェアネス（確認）、原子フラグを使用しますが、Rubin は受信 GPU がデータをロードする前にカウンター更新を採用しています。

カウント書き込みをサポートした NVLink 融合通信は、GPU 間でのデータ転送を調整するための低遅延メカニズムを提供し、計算が同期操作の完了を待たずに継続して進捗できるよう支援します。

実際、スケールアウトネットワークにおいては、2023 年頃にはすでに同様の機能を実装していました。受信側でカウンターを更新することで、往復時間 (RTT) を 1 つ節約できる仕組みです。技術的に難しいものではなく、再び NVIDIA よりも 3 年先を行く結果となりました。具体的な実装では、PTX インストラクションの fabric.try_put.async.counted::bytes や fabric.try_red.async.counted::bytes を採用しました。

PTX の fabric インストラクションは、GPU カーネルから跨 GPU での PUT 操作やリダクション処理を直接開始できます。カウント形式 (counted form) を用いると、リモート側のカウンターが「宛先でアクセスされたバイト数」に応じて自動的に累積されます。受信側はこのカウンターをポーリングして、対象データのアクセスが期待通りのバイト数に達したことを確認した後、データを消費します。これにより、独立したメモリバリア、アウェアネス（確認）、および原子フラグによる調整処理を削減できます。

もう一つの重要な事実として、カウント形式の fabric 構文はすでに PTX 9.3 で利用可能であり、最低ターゲットは sm_100+ です。つまり Blackwell アーキテクチャでもサポートされるはずです。実際にテストできるカードがいつ入手できるかは不明ですが、理論的には可能です。

2.4 Tcgen05 LUT

低ビット数のインデックスを Tensor Core の B 演算子 (B operand) に変換する機能

「tcgen05.mma.decompress::lut::b」という名前は誤解を招きやすいものです。これは任意の関数を照会できる汎用的な LUT（ルックアップテーブル）ではなく、非常に特定の「ローカルコードブック圧縮解除」を行う機能です。

B 演算子は 3 ビットのインデックスを保持します。各ローカルグループには 8 つの E4M3 形式の候補値が用意されており、MMA（マルチメディアアクセラレータ）は内部でこのインデックスを使って候補値を選択し、B の値を復元します。また、B コレクター機能により、復元された B タイルを複数の MMA 演算間で共有・再利用することが可能です。

1. その他の課題

3.1 ピーク電力の管理

AI 負荷中は電力需要が激しく変動し、瞬時のスパイクが発生します。これにより利用可能な容量が遊んでしまい、ファクトリー全体のスループットが低下する恐れがあります。Vera Rubin の電源システムは、荷電状態（SoC）に基づくインテリジェントな電力平滑化技術を採用しています。この新技術により、前世代の平滑化手法と比較して平均消費電力を約 10% 削減し、50 ミリ秒以内のピーク電力を約 20% 抑制できます。より滑らかな電力カーブを実現することで、システムは継続的な最大電力需要を下げることができ、これにより周辺供电インフラやグリッドへの負担を軽減しつつ、同じ AI ファクトリーの消費電力予算内でさらに多くの計算リソースを搭載できるようになります。

AI ファクトリーレベルでは、NVIDIA DSX MaxLPS がこの手法を GPU 間やラック間、45°C の液冷システム、そして多様な負荷にまで拡張しています。一方、DSX OS はスケジューリング、ライフサイクル管理、健康状態の自動化といった運用層を提供します。これら 2 つは協調して設計されており、遊んでいる電力を回収し、固定されたメガワット規模の範囲内で有効な計算能力を増やすことを目指しています。DSX MaxLPS を導入することで、運用者は負荷性能への影響を最小限に抑えつつ、同じ消費電力予算で最大 40% の追加 GPU を高効率な動作点で展開することが可能になります。

具体的な実装については、実は私が少年時代に父親と共に高低圧配電設備の製作に関わっていた経験があります。その手法は概ねシンプルです。

• ラック内のコンデンサがサブ秒レベルの電力スパイクを吸収します。
• Rack power steering（ラック電源ステアリング）により、CPU、GPU、NVLink スイッチの間で動的に電力予算を移動させます。
• DSX MaxLPS はワークロード、テレメトリデータ、冷却状況、そしてファクトリー全体の電力上限に基づいて、GPU の電力制限とタスクリソースを再配分します。

一言付け加えるなら、Blackwell の電力制御は実際には非常に不十分でした。先日 B 系列 GPU の逆アセンブリ解析を行った際、特に MegaMoE のようなハードウェアを極限まで駆使する演算子に対してテストを実施した結果がそれを物語っています。

8 枚の B300 を EP8 モードで接続し、V4-Pro 構成（エクスパート数 384、トップ 6、隠れ層サイズ 7168、中間層サイズ 3072）とした場合の降頻分析は以下の表の通りです。

さらに、いくつかの消費電力に関するアボレーション実験（単一変数の影響を調べる実験）を行いました。結果、単一のサブシステムだけでは DVFS（動的電圧・周波数調整）がトリガーされないことが分かりました。例えば、cuBLAS の GEMM 演算のみでピーク時の 1072W に達しても TDP の 97% を超えるわけではなく、降頻は起こりません。

しかし、MegaMoE は TC（Tensor Core）、HBM、NVLink の三者が同時に飽和状態に陥ることで閾値を超え、結果として降頻が発生します。

Rubin においても同様の降頻現象が見られるかどうかは不明ですが、Nvidia のこの部分の対応については、あまり良くない印象を受けます。

3.2 NV-HBI を介したクロスダイ間のネットワーク問題

Blackwell においては、過去に多くの記事で「Cross-Die（ダイ間）へのメモリアクセスにおける遅延の不一致」が指摘されてきました。しかし、これらの記述は実際のテスト結果とは乖離しているように思えます。

Hopper アーキテクチャにおいて Cross L2 Partition の間に大きな遅延差があることは事実ですが、Blackwell では各ダイに 2 つの L2 パーティションがあり、チップ全体で 2 つのダイを持つため、遅延テストでは最大で 4 つのピークが現れるはずです。しかし実際にはそのような現象は見られませんでした。

実は NV-HBI の遅延は、多くの人が考えているほど数百サイクルも異なるような大きな差ではなく、むしろ非常に低い値です。

一方、Rubin では HBM バンド幅が 22TB/s に達します。L2 キャッシュの容量や NV-HBI の帯域を考慮すると、何らかのアフィニティ（親和性）処理を行わない場合、全体として HBM バンド幅を最大限に活用できず、ピーク性能の 60〜70% 程度しか発揮できない可能性が高いと考えられます。

4. まとめ

総じて Rubin アーキテクチャは、Blackwell をベースにしつつもごく僅かな変更を加えた製品と言えます。I/O 機能を専用の I/O ダイへ移設したことは、優れたトレードオフの選択でした。これによりメインダイの面積をより多くの SM（ストリーミングマルチプロセッサ）に充てることが可能になりました。

さらに K 次元の拡大、周波数の向上、そして HBM バンド幅の増強が相まって、性能は大幅に向上しました。ランタイムレベルでも多数の微細な最適化が施され、SM のアイドル状態（空泡）を低減することに成功しています。特に ThreadBlock レベルでの依存関係管理の改善により、比較的低遅延な Megakernel を構築しやすくなっています。

しかし、消費電力や放熱の制約を考慮すると、チップ全体を最大限に駆動した際の電力制限（パワーウォール）が実際に存在するかどうかについては、まだ懐疑的な見方をしています。また、アテンション機構においてスパースなGEMM演算が活用できるかについても、同様に疑問が残ります。

より詳細な分析については、もしスポンサーの方からRカードを提供していただければ、エージェントを用いて約1週間以内に完全な逆エンジニアリング解析を完了させることができるでしょう。 :)

参考資料

[1] Inside NVIDIA Rubin GPU Architecture: Powering the Era of Agentic AI: https://developer.nvidia.com/blog/inside-nvidia-rubin-gpu-architecture-powering-the-era-of-agentic-ai/

[2] PTX 9.4: https://docs.nvidia.com/cuda/developer-preview/13.4/pdf/ptx_isa_9.4.pdf

微信で開くにはこちらへ