Funkcje ballot i vote
Sprawdzać predykaty w całym warpie.
Funkcje ballot i vote to bezpłatna lekcja CUDA Academy na CoddyKit. To lekcja 3 z 4. Możesz przeczytać całą lekcję poniżej za darmo — a potem ćwiczyć ją interaktywnie w przeglądarce z wbudowanym edytorem kodu i tutorem AI dostępnym 24/7. To część ścieżki edukacyjnej CUDA Academy, a Twój postęp synchronizuje się między webem a aplikacją CoddyKit. Kurs CUDA Academy zawiera 4 lekcji w sumie.
Lane'y mogą głosować
Warp może nie tylko przenosić dane, ale także wspólnie odpowiadać na pytania typu tak lub nie. Intrinsics Vote pozwalają wszystkim lane'om przekazać predykat prawda-fałsz w jednym, tanim kroku.
Poznaj ballot_sync
Najbardziej elastycznym mechanizmem głosowania jest ballot_sync. Zwraca on liczbę 32-bitową, w której bit i ma wartość 1 dokładnie wtedy, gdy predykat lane i był prawdziwy.
unsigned bits = __ballot_sync(mask, pred);Jeden bit na lane
Zwrócona liczba całkowita to mała bitmapa warpa. Odczyt bitu i informuje, czy lane i zagłosował twierdząco — wszystkie wyniki są upakowane w jednym rejestrze.
Policz wartości prawdziwe
Chce Pan sprawdzić, ile lane'ów odpowiedziało twierdząco? Policz ustawione bity wyniku ballot. popc zlicza ustawione bity, natychmiast podając liczbę prawdziwych głosów w warpie.
int yes = __popc(__ballot_sync(mask, pred));Wszystkie czy dowolne
Dwa skróty odpowiadają na często spotykane pytania. all_sync zwraca true tylko wtedy, gdy każdy lane się zgodził, natomiast any_sync zwraca true, jeśli zgodził się przynajmniej jeden lane.
bool every = __all_sync(mask, pred);Any oznacza co najmniej jeden
Proszę używać any_sync przy wcześniejszym kończeniu działania: jeśli dowolny lane znalazł dopasowanie, cały warp może zareagować wspólnie bez sprawdzania lane'ów po kolei.
bool found = __any_sync(mask, pred);Kompaktowanie strumienia
Ballot umożliwia compaction: każdy lane decyduje, czy zachować swój element, głosuje, a bitmapa wskazuje lane'om, gdzie zapisać dane w upakowanej tablicy wynikowej.
Znajdowanie własnego miejsca
Aby umieścić swoją wartość, lane zlicza prawdziwe bity znajdujące się poniżej jego pozycji. Ta prefix count z wyniku ballot jest jego indeksem w skompaktowanym wyniku.
int slot = __popc(bits & ((1u << lane) - 1));Wybierz liderujący lane
Czasami jeden lane powinien wykonać wspólną pracę. Najniższy ustawiony bit w wyniku ballot, znaleziony za pomocą ffs, jest naturalnym liderem, co do którego cały warp może się zgodzić.
Zawsze przekazuj maskę
Podobnie jak shuffle, intrinsics głosowania przyjmują aktywną mask. Nieaktywne lane'y wnoszą bit 0, dlatego należy przekazać właściwy zbiór, aby nie zliczać nieobecnych lane'ów.
Tania koordynacja
Głosowanie kosztuje jedną instrukcję i nie korzysta z pamięci. Pozwala warpowi koordynować decyzje znacznie taniej niż atomiki czy flagi w pamięci współdzielonej.
Szybkie sprawdzenie
Proszę przypomnieć sobie, co ballot_sync zwraca każdemu lane'owi.
Podsumowanie
Przepytał Pan warp za pomocą ballot i vote: zliczał Pan wartości prawdziwe, sprawdzał, czy prawdziwe są wszystkie lub dowolne wartości, i tanio kompaktował strumienie. Dalej: cooperative groups do elastycznej synchronizacji. ✨
Ucz się C++ dzięki korepetycjom AI — za darmo
Pisz i uruchamiaj kod w przeglądarce, otrzymuj natychmiastową pomoc od korepetytora AI dostępnego 24/7 i kontynuuj naukę w sieci lub w aplikacji.
- Kursy
- 30
- Lekcje
- 120
Często zadawane pytania
Czy lekcja „Funkcje ballot i vote” jest bezpłatna?
Tak — pełny tekst „Funkcje ballot i vote” jest dostępny za darmo tutaj w sieci. Aby ćwiczyć ją interaktywnie (wbudowany edytor kodu i tutor AI dostępny 24/7) i odblokować resztę kursu CUDA Academy, przejdź na CoddyKit PRO. Kurs CUDA Academy zawiera 4 lekcji w sumie.
Co nauczysz się w „Funkcje ballot i vote”?
Sprawdzać predykaty w całym warpie. Ćwiczysz CUDA Academy z praktycznym kodem, który uruchamiasz bezpośrednio w przeglądarce, a tutor AI dostępny 24/7 odpowiada na Twoje pytania podczas pracy nad lekcją.
Czy potrzebuję doświadczenia, aby zacząć CUDA Academy?
Nie wymagamy żadnego doświadczenia. CUDA Academy w CoddyKit jest strukturyzowany dla początkujących i zaawansowanych użytkowników, więc możesz zacząć tutaj lub od początku i uczyć się w swoim tempie. To lekcja 3 z 4.
Ile czasu zajmuje lekcja „Funkcje ballot i vote”?
Większość lekcji CoddyKit trwa około 5–10 minut. Każda lekcja to mały, interaktywny krok, dzięki czemu robisz systematyczne postępy i zawsze wracasz dokładnie do tego samego miejsca — na webie i w aplikacji.
Czy mogę pisać i uruchamiać kod w tej lekcji CUDA Academy?
Tak. Każda lekcja CUDA Academy zawiera wbudowany edytor kodu, więc piszesz i uruchamiasz prawdziwy kod bezpośrednio w przeglądarce i od razu otrzymujesz sprzężenie zwrotne od AI — bez konfiguracji na komputerze.
Wszystkie lekcje w tym kursie
- Warpy, lane'y i maski
- __shfl_down_sync do redukcji
- Funkcje ballot i vote
- Cooperative Groups