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
- Warps, faixas e máscaras
- __shfl_down_sync para reduções
- Funções de votação e cédula
- Grupos cooperativos