0Pricing
CUDA Academy · Aula

Grupos cooperativos

Uma API moderna para escopos flexíveis de sincronização.

Grupos cooperativos é uma aula grátis de CUDA Academy no CoddyKit. Esta é a aula 4 de 4. Você pode ler a aula completa abaixo gratuitamente — depois pratica ao vivo no navegador com um editor de código integrado e um tutor de IA 24/7. Faz parte do caminho de aprendizado de CUDA Academy, e seu progresso é sincronizado entre a web e o app CoddyKit. O curso de CUDA Academy inclui 4 aulas no total.

Partes desta aula ainda não foram traduzidas e aparecem em inglês.

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. ✨

Perguntas Frequentes

A aula “Grupos cooperativos” é grátis?

Sim — o texto completo de “Grupos cooperativos” é grátis para ler aqui na web. Para praticá-la interativamente (um editor de código integrado e um tutor de IA 24/7) e desbloquear o restante do curso de CUDA Academy, atualize para CoddyKit PRO. O curso de CUDA Academy inclui 4 aulas no total.

O que vou aprender em “Grupos cooperativos”?

Uma API moderna para escopos flexíveis de sincronização. Você pratica CUDA Academy com código prático que executa diretamente no navegador, e um tutor de IA 24/7 responde suas dúvidas enquanto trabalha na aula.

Preciso ter experiência prévia para começar CUDA Academy?

Nenhuma experiência prévia é necessária. CUDA Academy no CoddyKit é estruturado para alunos iniciantes até avançados, então você pode começar aqui ou desde o início e aprender no seu ritmo. Esta é a aula 4 de 4.

Quanto tempo leva a aula “Grupos cooperativos”?

A maioria das aulas CoddyKit leva cerca de 5–10 minutos. Cada uma é compacta e interativa, então você faz progresso constante e retoma exatamente de onde parou entre web e app.

Posso escrever e executar código nesta aula de CUDA Academy?

Sim. Cada aula de CUDA Academy inclui um editor de código integrado, então você escreve e executa código real direto no navegador e recebe feedback de IA instantaneamente — nenhuma configuração local necessária.

Todas as aulas deste curso

  1. Warps, faixas e máscaras
  2. __shfl_down_sync para reduções
  3. Funções de votação e cédula
  4. Grupos cooperativos
← Voltar para CUDA Academy