0Pricing
CUDA Academy · Aula

Funções de votação e cédula

Consulte predicados em todas as threads de um warp.

Funções de votação e cédula é uma aula grátis de CUDA Academy no CoddyKit. Esta é a aula 3 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.

Lanes Can Vote

Beyond moving data, a warp can answer yes-or-no questions together. Vote intrinsics let every lane share a true-or-false predicate in one cheap step.

Meet ballot_sync

The most flexible vote is ballot_sync. It returns a 32-bit number where bit i is 1 exactly when lane i's predicate was true.

unsigned bits = __ballot_sync(mask, pred);

One Bit per Lane

That returned integer is a tiny bitmap of the warp. Reading bit i tells you whether lane i voted true, all packed into a single register.

Count the Trues

Want how many lanes said yes? Pop-count the ballot result. popc counts the set bits, giving the number of true votes in the warp instantly.

int yes = __popc(__ballot_sync(mask, pred));

All or Any

Two shortcuts answer common questions. all_sync returns true only if every lane agreed, while any_sync returns true if at least one lane did.

bool every = __all_sync(mask, pred);

Any Means at Least One

Use any_sync for early exits: if any lane found a match, the whole warp can react together without scanning lane by lane.

bool found = __any_sync(mask, pred);

Stream Compaction

Ballot powers compaction: each lane decides if it keeps its element, votes, and the bitmap tells lanes where to write into a packed output array.

Finding My Slot

To place its value, a lane counts the true bits below its own position. That prefix count from the ballot is its index in the compacted result.

int slot = __popc(bits & ((1u << lane) - 1));

Pick a Leader Lane

Sometimes one lane should do shared work. The lowest set bit in the ballot, found with ffs, is a natural leader the warp can all agree on.

Always Pass a Mask

Like shuffles, vote intrinsics take an active mask. Inactive lanes contribute a 0 bit, so pass the correct set to avoid counting absent lanes.

Cheap Coordination

Votes cost a single instruction and touch no memory. They let a warp coordinate decisions far cheaper than atomics or shared-memory flags would.

Quick Check

Recall what ballot_sync hands back to each lane.

Recap

You polled a warp with ballot and vote: count trues, check all or any, and compact streams cheaply. Next: cooperative groups for flexible sync. ✨

Perguntas Frequentes

A aula “Funções de votação e cédula” é grátis?

Sim — o texto completo de “Funções de votação e cédula” é 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 “Funções de votação e cédula”?

Consulte predicados em todas as threads de um warp. 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 3 de 4.

Quanto tempo leva a aula “Funções de votação e cédula”?

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