# Thread Block Cluster ## 定義 thread block cluster(スレッドブロッククラスタ、cooperative thread array [CTA] cluster)は、GPU processing cluster(GPC、近接するSM群)上に同時実行が保証された複数スレッドブロックのグループである。cooperative groups(協調グループ)がスレッド・warp・タイル・ブロック・グリッドという任意粒度を扱えるソフトウェアレベルの抽象化であるのに対し、thread block clusterはハードウェアレベルの階層であり、cooperative gridの一部SMだけを予約し残りを他カーネルに開放できる点が異なる。クラスタ内のブロックは **DSMEM(distributed shared memory、分散共有メモリ)** を通じて互いのshared memoryに読み書き・アトミック操作でき、クラスタスコープのバリア `cluster.sync()` で同期する。ポータブルな最大クラスタサイズは8ブロックだが、`cudaFuncAttributeNonPortableClusterSizeAllowed` を明示的に有効化するとBlackwellの一部で最大16ブロックまで拡張できる(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 10 Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters]])。 DSMEMは各ブロックの`__shared__`領域を論理的に1つのオンチップバッファへ結合するもので、L2キャッシュをバイパスし専用のSM-to-SMネットワークを経由する。これによりブロック間のデータ交換・部分和の集約をグローバルメモリ往復なしに行え、L2ヒット率・実効帯域を大幅に改善する。scratch memory(スクラッチメモリ)はDSMEMやクラスタ同期を支える裏方のハードウェアインフラで、バリア状態や到達フラグ等のメタデータを保持する(ユーザーコードから直接アクセスしない)。TMA(Tensor Memory Accelerator)のmulticastモードを使うと、1回のグローバルメモリ転送でクラスタ内の全ブロックのshared memoryへ同時にタイルを配信でき、DSMEM経由の「leaderが1回ロードしてmap_shared_rankで共有」方式とは別の経路として使い分けられる(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 10 Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters]])。 ## 横断的知見 - **thread block clusterは「グリッド全体を独占する `grid.sync()`」と「単一ブロック内に閉じた`__syncthreads()`」の中間に位置する、新しい同期・データ共有の粒度である**。本章は cooperative groups の `grid.sync()`(全ブロックを巻き込む重量級バリア)と thread block cluster の `cluster.sync()`(クラスタ内のみの軽量バリア)を明確に対比しており、この中間粒度の導入がGPU全体を占有せずに複数SM間でオンチップ協調を可能にする鍵だとしている(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 10 Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters]])。 - **DSMEM共有とTMA multicastは同じ「グローバルロード1回・複数ブロックで再利用」という目的を、異なるトレードオフで実現する2つの経路である**。本章は明示的な使い分け基準を示している: タイル再利用が多い、またはクラスタサイズが大きい(8〜16 CTA)場合はTMA multicast(1回書いて何度もローカル読み)が有利、共有メモリが逼迫している、またはクラスタサイズが小さく(2〜4 CTA)各フォロワーがタイルを1回しか触れない場合はDSMEM共有(フットプリントが小さい)が有利、という設計指針を提示している(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 10 Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters]])。 ## 未解決の問い - thread block clusterとDSMEMは第4章([[.raw/books/ai-systems-performance/chapters/ch-04.md|Tuning Distributed Networking Communication]]、未取り込み)で扱われるノード間・GPU間の集団通信(NCCL等)とは異なる「SM間」というより微細な粒度の通信だが、両者を貫く共通の設計原則(トポロジー認識・マルチキャストの活用等)はどの程度あるか。 - クラスタサイズを2〜4ブロックから8〜16ブロックへ拡大する際の収穫逓減点(占有SM数増加による他タスクへの悪影響とのトレードオフ)は、ワークロードの種類(GEMM、attention、疎行列演算)によってどう変わるか。 - 本章のDSMEM/thread block clusterの技法は、次章(第11章、CUDA streamによるinter-kernelパイプライン、未取り込み)のストリーム間パイプラインとどのように組み合わされ、カーネルをまたいだデータ共有の設計にどう影響するか。 ## 関連 - [[@2025__OReilly__AI Systems Performance Engineering - Chapter 10 Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters]] - [[CUDA]] - [[Warp Specialization]] - [[Persistent Kernel]] - [[wiki/entities/AI Systems Performance Engineering|AI Systems Performance Engineering]] ## 出典 - Chris Fregly, *AI Systems Performance Engineering*, O'Reilly Media, 2025, Chapter 10「Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters」.