🤨
RTX 4000では使えない？RTX 5000番台で追加されたThread Block Cluster解説

2025/11/21に公開
テーマ「2025年の最も大きなチャレンジ」
 Compute Capability 9.0以降の新機能 Thread Block Cluster当初HPC向けNVIDIA H100の新機能として登場したThread Block Cluster[1]ですが、ついにコンシューマー向けGPUでも使えるようになりました。個人環境でこの機能を試せるのはありがたいですね。現状RTX 5090などCompute Capability 9.0以降のGPU(5000番台)が必要です。

 対象読者本記事は相当ニッチです。対象読者は

・CUDAなどを用いたGPUプログラミングに一定の興味がある人

・GPUプログラミングを勉強するためにGPU購入を検討している人
これはなかなかにおもしろい機能でして、RTX 4000番台にはないので、GPUプログラミングを勉強するために新しいGPUを買おうとしている人は読んでおいてもいいかもしれません。

 階層が1つ増えたまず CUDA C++ プログラミングガイド[2] の説明を読むと
5.2.1.スレッドブロッククラスター

NVIDIA Compute Capability 9.0の導入により、CUDAプログラミングモデルに、スレッドブロックで構成される「スレッドブロッククラスター」と呼ばれるオプションの階層が導入されました。スレッドブロック内のスレッドがストリーミングマルチプロセッサ上で確実に同時スケジューリングされるのと同様に、クラスター内のスレッドブロックもGPU内のGPUプロセッシングクラスター (GPC) 上で確実に同時スケジューリングされます。
とあります。

つまりCC 9.0から新しい階層が1つ増えたのです。今までは「グリッド・ブロック・スレッド」の3階層だったのが「グリッド・クラスター・ブロック・スレッド」という4階層になりました。


出典: NVIDIA Corporation, CUDA C Programming Guide, Figure 6 “Memory Hierarchy”. https://docs.nvidia.com/cuda/cuda-c-programming-guide/
このBlockの集まりをThread Block Cluster(TBC)と呼びます。Blockが1つのStreaming Multiprocessor(SM)に対応していたように、このTBCは1つのGraphics Processing Cluster(GPC)に対応しています。


観点
最上位
中間 (クラスター)
ブロック単位
最小のスレッド単位


CUDA の階層
Grid
Thread Block Cluster (TBC)
Thread Block
Thread

対応するハードウェア
GPU 全体
GPC (Graphics Processing Cluster)
SM (Streaming Multiprocessor)
CUDA Core / FP ユニット

GPC/TPC/SM 構造については、PC Watch[3]に詳細な図解があります。

 Cluster内にあるThread Blockの共有メモリは互いに共有できるSM備え付けの共有メモリを別のSMからアクセスできるようになりました。これ相当すごいことだと思うんですよね。それまではSMをまたぐThread Block間通信は必ずGlobal Memoryを介さないといけなかったのが、TBC内限定ではありますがGlobal Memoryを一切介さずに通信できるようになったのです！
このメモリは分散共有メモリ（Distributed Shared Memory, DSMEM）と呼ばれ、そのアドレス空間はDistributed Shared Memory address spaceと呼ばれます。といっても新しくどこかにメモリが追加されたわけではなく分散共有メモリの実体は従来の共有メモリ(図のShared Memory)です。つまり分散共有メモリはソフトウェア的概念です。
分散共有メモリは同じCluster内のどのThreadからも読み書きやアトミック操作が可能で、Cluster内同期バリアも可能です。
Q&A
従来の3階層のCUDAコードは動く？ → 動く
TBCのサイズは自由に決められる？ → GPCの作り次第で最大16 Blockまで
RTX 4000(Ada)以前では使えない？ → 使えない。Hopper/Blackwell以降の機能
AMDや他社GPUには？ → 調べた限りそのような機構ハードは存在しない


 Thread Block Clusterの使い方「すべてのスレッドから1.0fを足しあわせて、その総和をoutに書き出す」というきわめてシンプルな処理を、TBC + 分散共有メモリ を使って実装しました。

 ここがTBC特有の部分cluster.map_shared_rank(&cluster_sum, 0)

は、「rank 0 block の共有メモリ変数 cluster_sum へのポインタ」をクラスター内から取得します。これにより、どの block からもrank 0 blockの共有メモリにアクセスできます。
cluster_sumは1要素のshared memoryです。sdata[0]にはblock内の総和があるので、それをblock 0のcluster_sumへ総計(矢印)しています。

 全Threadから1をAtomic加算するサンプル(cudaLaunchKernelExなしver)https://github.com/toropippi/ThreadBlockClusterSample/blob/main/ClusterSync/sample1_cluster_reduce_no_exlaunch.cu
コンパイル
nvcc -arch=sm_120 -std=c++17 -rdc=true sample1_cluster_reduce_no_exlaunch.cu -o sample1_cluster_reduce_no_exlaunch
結果
>sample1_cluster_reduce_no_exlaunch.exe
GPU sum = 2048  (expect 2048)

 全Threadから1をAtomic加算するサンプル(cudaLaunchKernelExありver)https://github.com/toropippi/ThreadBlockClusterSample/blob/main/ClusterSync/sample2_cluster_reduce_launch_ex.cu
コンパイル
nvcc -arch=sm_120 -std=c++17 -rdc=true sample2_cluster_reduce_launch_ex.cu -o sample2_cluster_reduce_launch_ex
結果
>sample2_cluster_reduce_launch_ex.exe
GPU sum = 2048  (expect 2048)
全体構成と解説constexpr int BLOCKS_PER_CLUSTER  = 8;    // 1クラスタあたり 8 block
constexpr int THREADS_PER_BLOCK   = 128;  // 1 block あたり 128 thread
constexpr int NUM_CLUSTERS        = 2;    // クラスタ数 2
// 合計スレッド数 = NUM_CLUSTERS * BLOCKS_PER_CLUSTER * THREADS_PER_BLOCK
したがって、総スレッド数は
2（クラスタ） × 8（block/cluster） × 128（thread/block） = 2048 thread
となります。
分散共有メモリにAtomic加算のコード// 各 block の部分和 (sdata[0]) を rank 0 block の cluster_sum に加算
if (tid == 0) {
    float* remote = cluster.map_shared_rank(&cluster_sum, 0);
    atomicAdd(remote, sdata[0]);
}

 書きかたは2パターンcuファイルをnvccでコンパイルするのは従来どおりですが、CUDA 12 現在、書きかたが2パターンあるようです。上記のように cudaLaunchKernelEx あり／なし ver です。cudaLaunchKernelEx[4]はクラスタ概念追加によりCUDAに新規実装されたものなので本来はこれを使うべきでしょう。

ただ今までの書きかたに近いのは "なしver" です。

 cudaLaunchKernelExcudaLaunchKernelExは CUDA 12 で導入された拡張カーネル起動 API であり、

従来の kernel<<<grid, block>>>() では指定できなかった 追加属性 (Launch Attributes) を起動時に動的に付与できるのが特徴です。
Thread Block Clusterサイズ(dim3)の設定もここで行います。このClusterサイズを指定するcudaLaunchAttributeClusterDimensionを設定するコードを示します。

 cudaLaunchKernelExとcudaLaunchConfig_tの使い方// 起動設定
cudaLaunchConfig_t config = {};
config.gridDim          = dim3(NUM_CLUSTERS * BLOCKS_PER_CLUSTER); // 総 block 数
config.blockDim         = dim3(THREADS_PER_BLOCK);
config.dynamicSmemBytes = THREADS_PER_BLOCK * sizeof(float);
config.stream           = 0;

// クラスタサイズを属性で指定
cudaLaunchAttribute attr[1];
attr[0].id = cudaLaunchAttributeClusterDimension;
attr[0].val.clusterDim.x = BLOCKS_PER_CLUSTER; // 1クラスタあたりの block 数
attr[0].val.clusterDim.y = 1;
attr[0].val.clusterDim.z = 1;

config.attrs   = attr;
config.numAttrs = 1;

cudaLaunchKernelEx(&config, kernel_func, args...);

 __cluster_dims__ 方式との違いこのようにcudaLaunchKernelExありとなしverで、クラスタサイズ指定に違いがあります。


方法
特徴


__cluster_dims__(X,Y,Z)

コンパイル時に固定。<<< >>> で起動可。

cudaLaunchKernelEx

起動時に動的にクラスタサイズを変更可能。ベンチ向き。

メリット起動時にクラスタサイズを自由に変えられる
これによりTBCの性能チューニングが容易
デメリットなし？(まだ慣れないだけ)

 1Clusterあたり8Blockまでのリミットと解除法デフォルトで 1 ClusterあたりのBlock数は最大8個までと決まっています。リミット解除はマナーなのでぜひ覚えておきましょう。これをcluster_reduce実行前に呼んでおきます。
    cudaFuncSetAttribute(
        cluster_reduce,
        cudaFuncAttributeNonPortableClusterSizeAllowed,
        1
    );
// 0=禁止, 1=許可
なおリミット解除せずに

BLOCKS_PER_CLUSTER = 9

にして実行するとランタイムエラーがでます。
>sample2_cluster_reduce_launch_ex.exe
cudaLaunchKernelEx error: a kernel launch error has occurred due to cluster misconfiguration
一方上記NonPortableを許可すると
>sample2_cluster_reduce_launch_ex.exe
GPU sum = 2304  (expect 2304)
ちゃんと動きます。
理由としてはNVIDIA GPUのハードウェア仕様によるもので、RTX 5090は1 GPCあたり最大16 SMを持つ構成ですがRTX 5000番台の下位クラスでは 1 GPCに16個も SM をもたせられない(歩留まりの問題などで)、そこで 1 GPCに最低 8 SM の存在は保証する、としたのだと思われます。

同時に実行可能なBlock数を8として下位GPUでも動くように保証し、RTX 5090など巨大GPCを搭載するGPUではオプションで最大Blockまで動かせる柔軟性を持たせたと考えられます。
なお実際のRTX 5090 製品では、

11 GPC / 170 SM

であり一部の GPC/SM が無効化されています。この場合、
170 SM ÷ 11 GPC ≒ 15.45 SM/GPC
となり、「全GPCが16 SMフル稼働」というわけではないことに注意してください。物理的な GPC ブロックの設計は 16 SM ですが、製品ではいくつかの SM がカットされています。
ではRTX 5000シリーズの型番と SM / GPC 構成について気になってきたのではないでしょうか。AI調査でまとめてみました。実際買う際はご自身でもちゃんと調べてください。

 Blackwell (GB20x) ダイ別構成（SM理論最大※実際とは違うので注意）この表は 「アーキ的にこのダイが最大いくつ SM を持てるか」 の一覧です。


ダイ
主な製品候補
GPC/GPU
SM/GPC
SM/GPU(理論最大)


GB202
RTX 5090 系
12
16
192

GB203
RTX 5080 / 5070 Ti 系
7
12
84

GB205
RTX 5070 系
5
10
50

GB206
RTX 5060 Ti / 5060 系
3
12
36

GB207
RTX 5050 系
2
10
20

根拠：

GB202: NVIDIA Blackwell ホワイトペーパー[5]で 12 GPC / 96 TPC / 192 SM と明記。

GB203: 同ホワイトペーパーで 7 GPC / 42 TPC / 84 SM。

GB205: 同ホワイトペーパーに 5 GPC / 25 TPC / 50 SM。

GB206 / GB207: WCCFTech[6] のBlackwell構成リーク表

GB206: 3 GPC / 18 TPC / 36 SM、GB207: 2 GPC / 10 TPC / 20 SM と記載[6:1]

 RTX 50 デスクトップ製品別 SM 有効数

製品
ダイ
SM(実際)
SM(理論最大)
備考


RTX 5090
GB202
170
192
11 GPC 構成、SM を 6 基無効

RTX 5080
GB203
84
84
フル GB203

RTX 5070 Ti
GB203
70
84
6 GPC 構成＋SM 一部無効

RTX 5070
GB205
48
50
5 GPC 構成、SM を 2 基無効

RTX 5060 Ti
GB206
36
36
フル GB206

RTX 5060
GB206
30
36
SM を 6 基無効

RTX 5050
GB207
20
20
フル GB207

根拠：

RTX 5090 (GB202)：

各種レビュー[3:1][7]で「170 SM」と明記。フル 192 から 22 SM 減だが、そのうち 6 SM は設計上の無効化、残りは歩留まりや SKU 分けとされる解説もあります。

RTX 5080 (GB203)：

公式・レビュー[7:1]で「84 SM」、GB203 フル構成と一致。

RTX 5070 Ti (GB203)：

4Gamer[7:2][8] で「6 GPC / 70 SM」と整理。

RTX 5070 (GB205)：

「5 GPC / 48 SM」＝ フル 50 SM から 2 SM 無効化[7:3][9][10]

RTX 5060 Ti (GB206)：

比較表を明示[11]

RTX 5060 (GB206)：

比較表[11:1]で「30 of 36 SM」と記載

RTX 5050 (GB207)：

比較表[12]

 Thread Block Clusterの本質はGlobal Memoryアクセス低減にある大事なことだと思うので一応書いておこうと思います。このTBCは計算ボトルネックのカーネルを魔法のように高速化する仕組みではなく、グローバルメモリ往復を削減することでメモリ帯域律速なカーネルの高速化を期待する仕組みです。
例えば巨大サイズのGEMMなんかはすでにcuBLASでほぼ演算律速(理論値80,90%超え)となっておりTBCの恩恵は薄いでしょう。一方FFTや小規模GEMM、ソート、数値流体力学といったメモリ律速の問題であればアルゴリズムの高速化が十分期待できると思われます。

 分散共有メモリ活用によるヒストグラムAtomic最適化(失敗)ここからは、Thread Block Cluster と分散共有メモリが「どのくらい効くのか」を見るために、できるだけ素朴で分かりやすいヒストグラム問題を題材にします。

題材は Wang hash の出力分布を調べるヒストグラム です。
各スレッドは自分のスレッド ID からシードを作り、Wang hash を繰り返し適用します。
hash の出力は 32bit 整数ですが、0〜262143 (2^18−1) にマスクして 262,144 bin のヒストグラムを取ります。
1 スレッドあたり 1000～10000 回 hash を生成し、同じ回数atomicAddでヒストグラム変数を更新します。
このとき、次の 2 つのカーネルを用意して性能を比較しました。

 1. ベースライン: Global メモリへの atomicAdd だけもっとも素朴な実装です。グリッド全体で共有する global_hist[262144] を Global メモリ上に 1 本だけ確保する。

各スレッドは Wang hash から 1 つ値を生成するたびに
bin = hash_value & (BINS_TOTAL - 1);  // 0〜262143
atomicAdd(&global_hist[bin], 1);
を 1 回だけ実行する。つまり 「hash 1 回につき global atomicAdd 1 回」 という、非常に単純な構造です。この実装は Thread Block Cluster も分散共有メモリも一切使っていません。

 2. TBC + 分散共有メモリ版: DSMEM 上にクラスタ共有ヒストグラムを持つ2つ目は、Thread Block Cluster と分散共有メモリをフル活用した実装です。
1 クラスタあたり BLOCKS_PER_CLUSTER = 16 block を束ねます。
各 block の共有メモリに BINS_PER_BLOCK = 16384 個の bin を持たせます。（64KB の共有メモリ）
1 クラスタ全体ではBINS_PER_BLOCK × BLOCKS_PER_CLUSTER = 16384 × 16 = 262144 となり、

これで 全 262,144 bin をクラスタ内の DSMEM 上に展開できます。
各スレッドの処理は次のような流れです。
Wang hash で bin を 1 つ生成する（global版と同じ）。
bin がどの block の担当領域かを計算する。
owner_rank   = bin / BINS_PER_BLOCK;   // どの block(rank) が担当か
owner_offset = bin % BINS_PER_BLOCK;   // その block 内のローカル index
Thread Block Cluster の API で、その block の共有メモリを DSMEM 経由で参照する。
unsigned int* remote_hist = cluster.map_shared_rank(s_hist, owner_rank);
その共有メモリ上で atomicAdd を行う。
atomicAdd(&remote_hist[owner_offset], 1);
ここまでの段階では、global メモリには一切触っていません。全ての更新はクラスタ内の共有メモリ（DSMEM）に溜め込まれています。
最後に結果を Global メモリに書き戻します。

直感的には 「global atomic の回数を大量に削減できるので、TBC 版のほうが速くなってくれそうだ」 と期待したわけです。

 測定方法と結果の概要https://github.com/toropippi/ThreadBlockClusterSample/blob/main/WangHash_Histgram/cluster_wang_hash_histogram.cu
検証はRTX 5090を使いました。結果としては、
Results (average over 3 runs, first run is warm-up):
  Global-only (one atomic per hash) : 30.777 ms
  TBC+DSMEM (local DSMEM hist)      : 195.254 ms
という形になり、TBC+DSMEM側が明確に負ける、という予想外の結果となりました。

 なぜ TBC 版が負けたのか(AIにきいた)結論だけ言うと、このワークロードが そもそも TBC/DSMEM 向きではなく、RTX 5090 の global atomic が強すぎる ためです。


global atomic のコンテンドがほぼ無い

bin 数は 2^18 (=262,144) あり、Wang hash の出力はほぼ一様です。

その結果、atomic の更新先アドレスが綺麗にバラけてしまい、L2 上の global atomic がかなり高いスループットでさばかれます。

この条件では「global atomic がボトルネック」になりにくく、削減しても旨味が出ません。


TBC 版は 1 サンプルあたりのオーバーヘッドが大きい？

各サンプルごとに

owner_rank/owner_offset の計算 → cluster.map_shared_rank → DSMEM 上の atomicAdd  という経路を通ります。

さらにクラスタ全体で cluster.sync() を 2 回叩いているため、GPC 内 block 全体を足踏みさせるコストも乗ります。

結果として、「global atomic 1 回」より「DSMEM + 追加演算 + sync」のほうが高くついてしまいます。

このAIの説明は信頼性が低いと思います。owner_offsetの計算で除算をしているから遅いと言いたいのでしょうがコンパイル時定数なので極めて軽量な処理に変換されます。cluster.sync() を2回やってるから遅いというのも長大ループの外側なので無理があります。


「Global Memory 往復削減」という TBC の強みが活きていない

TBC の本質は「クラスタ内でデータ／中間結果を共有して、GMEM 往復と global atomic を減らす」点ですが、

今回のベンチマークでは
再利用したい大きな read-only データがあるわけでもなく
block 間で共有したい中間結果も特に無く
ひたすら「独立な hash 値をカウントするだけ」

という構造になっています。

つまり 「減らせるべき GMEM トラフィック」がほとんど存在せず、TBC 側の追加コストだけが丸残りしている状態です。

このため、同じ total_samples を処理しても

global-only 版が数十 ms、TBC+DSMEM 版がその数倍 という結果になりました。

TBC はあくまで「GMEM 往復が支配的なメモリ律速カーネル向けのオプション」であって、

今回のような「低コンテンドなランダム atomic ベンチ」にそのまま投げても、むしろ遅くなる──という良い反例になったと言えます。
ということでした。今回の反省をふまえ次の記事では TBC+DSMEM をちゃんと活かした題材にしたいと思います。

 おわり、ポエム、その他
 Thread Block Clusterを扱った技術ブログは皆無これを書いている2025/11時点でThread Block Clusterを扱った日本語記事はほぼなく、世界的に見ても非常に少ないです。確かに最初はHPCユーザー向けの機能だったので個人ブログで扱うこともなかったでしょうが、Blackwell / RTX 5000番台の登場で環境がかわりつつあります。

非HPCユーザー目線でTBCを扱う技術ブログが1つはあったほうがいいと思い書きました。それもRTX 5000番台が発売されたこの2025年内になんとか書きあげたかったのです。

 今後の予想RTX 5090でさえも「全GPCが16 SMフル稼働」といかないのを見ると、やはりHPC向け機能なんだなーと思います。GPC構成についても半導体製造における歩留まりの影響で画一的に何SMと決定できない難しさがあるため、ハードウェアから歩み寄ることはできず、ソフトウェア側が頑張って抽象化なりして対応していくのでしょう。

それでもこの機能をコンシューマー向けGPUにおろしてくれたNVIDIAには感謝というか、なにか覚悟めいたものを感じます。GPUの微細化が続く限り並列化も進むわけで、キャッシュ階層も多層化していくのが自然です。その第一歩を見ている気がします。(まぁ今回キャッシュが追加されたわけではないのだけど)

一方 SM to SM 通信はまだ速くなくAtomicのヒストグラム実験は惨敗でした。NVIDIAもそこの課題はなんとなく気付いている気がするので、今後世代が進むに連れ改善していくのではないでしょうか？願望かもしれませんが。

 書ききれなかったTipsクラスタ次元とグリッド次元の整合：グリッドのブロック配置はクラスタ次元の整数倍である必要があります。
将来のGPUでもハードウェアレベルでTBCのBlock数8以上をサポートする場合がありますが、コード中ではcudaOccupancyMaxPotentialClusterSizeで実行デバイスの対応クラスタサイズを問い合わせるのが推奨されているようです。
各クラスタは独立して同期されるため、異なるクラスタ間での直接同期はできません。他クラスタとのデータ共有が必要な場合は従来通りグローバルメモリを使う必要があります。
クラスタ内同期はブロック内同期よりコストが高い可能性があるため、必要最小限に留め大きめのタスク単位で同期する方が効率的です。NVIDIAも提供資料で「クラスタ内協調はブロック内協調より粒度の大きい計算に適用すべき」と示唆しています。
ハードウェア非対応機能の差異: 現時点ではクラスタ機能はHopper以降のアーキテクチャで利用可能ですが、一部の高度機能はデータセンター向けGPU限定です。例えばグローバルメモリ→クラスタ内複数SMへの同時ブロードキャスト読込（TMAマルチキャスト）はH100ではサポートされていますが、RTX 50xxなどコンシューマGPUでは省かれているとの情報があります

 参考文献
脚注
NVIDIA Hopper Architecture In-Depth ↩︎
CUDA C++ プログラミングガイド ↩︎
PC Watch 新たに判明したGeForce RTX 5090のアーキテクチャを徹底解説 ↩︎ ↩︎
6.33. C++ API Routines ↩︎
NVIDIA RTX BLACKWELL  GPU ARCHITECTURE ↩︎
WCCFTech – Blackwell ダイ構成リーク ↩︎ ↩︎
西川善司の3DGE：GeForce RTX 50完全解説前編 Blackwell世代の構造とレイトレーシングにおける革新 ↩︎ ↩︎ ↩︎ ↩︎
Gainward　Gainward GeForce RTX 5070 Ti Phoenix ↩︎
Nvidia GeForce RTX 5070 Founders Edition review: more RTX 3090 than 4090 ↩︎
NVIDIAがGeForce RTX 5070 Ti & RTX 5070 GPUのフルスペック、GB203 & GB205 「Blackwell 」チップを発表 ↩︎
GeForce RTX 5060の性能比較＆ベンチマーク検証【2025年】 ↩︎ ↩︎
GeForce RTX 5050が発表されました ↩︎
方法	特徴
`__cluster_dims__(X,Y,Z)`	コンパイル時に固定。`<<< >>>` で起動可。
`cudaLaunchKernelEx`	起動時に動的にクラスタサイズを変更可能。ベンチ向き。
RTX 4000では使えない？RTX 5000番台で追加されたThread Block Cluster解説

Compute Capability 9.0以降の新機能 Thread Block Cluster

対象読者

階層が1つ増えた

Cluster内にあるThread Blockの共有メモリは互いに共有できる

Thread Block Clusterの使い方

ここがTBC特有の部分

全Threadから1をAtomic加算するサンプル(cudaLaunchKernelExなしver)

全Threadから1をAtomic加算するサンプル(cudaLaunchKernelExありver)

全体構成と解説

分散共有メモリにAtomic加算のコード

書きかたは2パターン

cudaLaunchKernelEx

cudaLaunchKernelExとcudaLaunchConfig_tの使い方

`__cluster_dims__` 方式との違い

メリット

デメリット

1Clusterあたり8Blockまでのリミットと解除法

Blackwell (GB20x) ダイ別構成（SM理論最大※実際とは違うので注意）

RTX 50 デスクトップ製品別 SM 有効数

Thread Block Clusterの本質はGlobal Memoryアクセス低減にある

分散共有メモリ活用によるヒストグラムAtomic最適化(失敗)

1. ベースライン: Global メモリへの atomicAdd だけ

2. TBC + 分散共有メモリ版: DSMEM 上にクラスタ共有ヒストグラムを持つ

測定方法と結果の概要

なぜ TBC 版が負けたのか(AIにきいた)

おわり、ポエム、その他

Thread Block Clusterを扱った技術ブログは皆無

今後の予想

書ききれなかったTips

参考文献

Discussion

観点	最上位	中間 (クラスター)	ブロック単位	最小のスレッド単位
CUDA の階層	Grid	Thread Block Cluster (TBC)	Thread Block	Thread
対応するハードウェア	GPU 全体	GPC (Graphics Processing Cluster)	SM (Streaming Multiprocessor)	CUDA Core / FP ユニット
ダイ	主な製品候補	GPC/GPU	SM/GPC	SM/GPU(理論最大)
GB202	RTX 5090 系	12	16	192
GB203	RTX 5080 / 5070 Ti 系	7	12	84
GB205	RTX 5070 系	5	10	50
GB206	RTX 5060 Ti / 5060 系	3	12	36
GB207	RTX 5050 系	2	10	20
製品	ダイ	SM(実際)	SM(理論最大)	備考
RTX 5090	GB202	170	192	11 GPC 構成、SM を 6 基無効
RTX 5080	GB203	84	84	フル GB203
RTX 5070 Ti	GB203	70	84	6 GPC 構成＋SM 一部無効
RTX 5070	GB205	48	50	5 GPC 構成、SM を 2 基無効
RTX 5060 Ti	GB206	36	36	フル GB206
RTX 5060	GB206	30	36	SM を 6 基無効
RTX 5050	GB207	20	20	フル GB207