0Pricing
CUDA Academy · レッスン

BallotとVote関数

ワープ全体で述語を評価します。

「BallotとVote関数」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン3/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これは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とVote関数」レッスンは無料ですか?

はい。「BallotとVote関数」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し(組み込みコードエディタと24時間対応のAIチューター)、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。

「BallotとVote関数」で何を学びますか?

ワープ全体で述語を評価します。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。

CUDA Academyを始めるのに経験は必要ですか?

事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン3/4です。

「BallotとVote関数」レッスンにはどのくらい時間がかかりますか?

ほとんどのCoddyKitレッスンは約5~10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。

このCUDA Academyレッスンでコードを書いて実行できますか?

はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

  1. ワープ、レーン、マスク
  2. リダクションに__shfl_down_syncを使う
  3. BallotとVote関数
  4. Cooperative Groups
← CUDA Academyに戻る