0Pricing
CUDA Academy · Lesson

Cooperative Groups

A modern API for flexible sync scopes.

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();

All lessons in this course

  1. Warps, Lanes, and Masks
  2. __shfl_down_sync for Reductions
  3. Ballot and Vote Functions
  4. Cooperative Groups
← Back to CUDA Academy