CUDA Academy · レッスン

Cooperative Groups

柔軟な同期スコープに対応する最新APIを学びます。

「Cooperative Groups」はCoddyKit上の無料CUDA Academyレッスンです。これはレッスン4/4です。下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。

このレッスンの一部はまだ翻訳されておらず、英語で表示されています。

A Cleaner Sync API

Raw masks and intrinsics work, but they are fiddly. Cooperative groups wrap threads into objects you can name, size, and synchronize explicitly.

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

Grab the Whole Block

Start by getting a handle to your block of threads. this_thread_block returns a group you can sync just like syncthreads, but as an object.

cg::thread_block block = cg::this_thread_block();

Sync Through the Group

Calling sync on the block is the modern barrier. It does exactly what syncthreads does, but reads clearly as a method on the group you mean.

block.sync();

Carve Out a Warp Tile

You can split a block into fixed-size tiles. A 32-lane tile gives you a warp-sized group with clean methods instead of raw shuffle masks.

auto warp = cg::tiled_partition<32>(block);

Methods Replace Masks

A tiled group offers shfl_down and friends without a mask argument. The group already knows its members, so the API stays short and safe.

val += warp.shfl_down(val, offset);

Know Your Position

Every group exposes its members and your spot. thread_rank gives your index inside the group, and size returns how many threads it holds.

int rank = warp.thread_rank();

Smaller Tiles Too

Tiles need not be 32 wide. A tiled_partition of 8 or 16 makes sub-warp groups, handy when your data naturally clusters in small sets.

Group-Level Reductions

The library ships ready-made collectives. A group reduce sums a tile in one call, hiding the offset loop you wrote by hand earlier.

int total = cg::reduce(warp, val, cg::plus<int>());

Grids That Sync

The boldest group is the grid group. With a cooperative launch, every block can sync at one barrier, something a normal kernel cannot do.

Cooperative Launch Required

Grid-wide sync only works if you start the kernel with cudaLaunchCooperativeKernel and the GPU supports it. A normal launch will not allow it.

Why Bother

Cooperative groups make warp code readable and portable: no hand-managed masks, clear scopes, and reusable collectives that match what you mean.

Quick Check

Recall how you create a warp-sized cooperative group from a block.

Recap

Cooperative groups turn masks into named objects: tile a block, call reduce, even sync a whole grid. You now own warp-level CUDA. ✨

よくある質問

「Cooperative Groups」レッスンは無料ですか？

はい。「Cooperative Groups」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。

「Cooperative Groups」で何を学びますか？

柔軟な同期スコープに対応する最新APIを学びます。ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。

CUDA Academyを始めるのに経験は必要ですか？

事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。これはレッスン4/4です。

「Cooperative Groups」レッスンにはどのくらい時間がかかりますか？

ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。

このCUDA Academyレッスンでコードを書いて実行できますか？

はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

← CUDA Academyに戻る