0Pricing
CUDA Academy · Lesson

Ballot and Vote Functions

Polling predicates across a warp.

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);

All lessons in this course

  1. Warps, Lanes, and Masks
  2. __shfl_down_sync for Reductions
  3. Ballot and Vote Functions
  4. Cooperative Groups
← Back to CUDA Academy