Warpy, lane'y i maski
Poznać jednostkę 32 wątków i maski aktywności.
Warpy, lane'y i maski to bezpłatna lekcja CUDA Academy na CoddyKit. To lekcja 1 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.
Wątki wykonują się w warpach
GPU nie planuje wątków pojedynczo. Grupuje je w warp składający się z 32 wątków, które poruszają się razem i wykonują tę samą instrukcję synchronicznie.
Dlaczego 32 ma znaczenie
Na sprzęcie NVIDIA warp zawsze składa się z 32 wątków. To rzeczywista jednostka wykonawcza, dlatego dobre kernele należy projektować w grupach po 32, a nie dla pojedynczych wątków.
Każdy wątek jest lane'em
W warp każdy wątek ma pozycję od 0 do 31, nazywaną lane'em. Identyfikator lane'a pozwala prymitywom warp rozpoznać, kto komunikuje się z kim.
int lane = threadIdx.x % 32;Znajdowanie identyfikatora lane'a
Możesz odczytać lane bezpośrednio ze specjalnego rejestru, zamiast go obliczać. Rejestr laneid zawsze zawiera wartość od 0 do 31 dla bieżącego wątku.
Synchronizacja ma swoją cenę
Wątki współdzielą jeden licznik programu na warp. Gdy instrukcja if kieruje lane'y różnymi ścieżkami, warp rozgałęzia się i wykonuje każdą ścieżkę po kolei, co obniża wydajność.
Maska aktywności
Nie każdy lane jest zawsze uruchomiony. 32-bitowa maska oznacza, które lane'y są obecnie aktywne; dla każdego uczestniczącego lane'a odpowiadający bit ma wartość 1.
Dlaczego istnieją maski
Po rozgałęzieniu aktywne pozostają tylko niektóre lane'y. Prymitywy warp muszą dokładnie wiedzieć, kto jest obecny, dlatego przekazujesz im maskę aktywności, aby działały poprawnie.
Pełna maska
Gdy masz pewność, że wszystkie 32 lane'y są aktywne, maska ma wartość 0xffffffff, czyli każdy bit jest ustawiony. Ta pełna maska jest najczęściej przekazywaną wartością.
unsigned mask = 0xffffffff;Bezpieczne tworzenie maski
Wewnątrz odgałęzienia nie zgaduj maski. Wywołaj activemask, aby uchwycić dokładnie te lane'y, które właśnie dotarły do tego miejsca.
unsigned mask = __activemask();Lane'y komunikują się bez pamięci
Największą zaletą jest możliwość bezpośredniej wymiany danych między lane'ami jednego warpa za pośrednictwem rejestrów. Do wymiany lokalnej dla warpa nie są potrzebne pamięć współdzielona ani bariery.
Sync oznacza poziom warpa
Sufiks sync w tych intrinsics oznacza synchronizację na poziomie warpa, a nie barierę bloku. Koordynuje tylko lane'y wymienione w przekazanej masce.
Szybkie sprawdzenie
Przypomnij sobie, czym jest warp i jaką ma wielkość na GPU NVIDIA.
Podsumowanie
Warp to 32 lane'y wykonujące instrukcje synchronicznie, a maska śledzi, które z nich są aktywne. Ta podstawa pozwala lane'om szybko współdzielić dane. Dalej: shuffles używane do redukcji. ✨
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 „Warpy, lane'y i maski” jest bezpłatna?
Tak — pełny tekst „Warpy, lane'y i maski” 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 „Warpy, lane'y i maski”?
Poznać jednostkę 32 wątków i maski aktywności. Ć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 1 z 4.
Ile czasu zajmuje lekcja „Warpy, lane'y i maski”?
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