0Pricing
CUDA Academy · Урок

Функции голосования и ballot

Опрос предикатов во всём варпе.

«Функции голосования и ballot» — бесплатный урок CUDA Academy на CoddyKit. Это урок 3 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.

Части этого урока еще не переведены и отображаются на английском.

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

Часто задаваемые вопросы

Урок «Функции голосования и ballot» бесплатный?

Да — полный текст урока «Функции голосования и ballot» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.

Чему я научусь в уроке «Функции голосования и ballot»?

Опрос предикатов во всём варпе. Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.

Нужен ли мне опыт, чтобы начать CUDA Academy?

Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 3 из 4.

Сколько времени занимает урок «Функции голосования и ballot»?

Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.

Можно ли писать и запускать код в этом уроке CUDA Academy?

Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.

Все уроки этого курса

  1. Варпы, lanes и маски
  2. __shfl_down_sync для свёрток
  3. Функции голосования и ballot
  4. Кооперативные группы
← Назад к CUDA Academy