Ballot and Vote Functions
Polling predicates across a warp.
Ballot and Vote Functions is a free CUDA Academy lesson on CoddyKit — lesson 3 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
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. ✨
Frequently asked questions
Is the “Ballot and Vote Functions” lesson free?
Yes — the full text of “Ballot and Vote Functions” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Ballot and Vote Functions”?
Polling predicates across a warp. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Ballot and Vote Functions” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Warps, Lanes, and Masks
- __shfl_down_sync for Reductions
- Ballot and Vote Functions
- Cooperative Groups