CUDA Academy · Lekcja

Warpy, lane'y i maski

Poznać jednostkę 32 wątków i maski aktywności.

Lekcja 1 z 413 kroki

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. ✨

Bezpłatny start

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

  1. Warpy, lane'y i maski
  2. __shfl_down_sync do redukcji
  3. Funkcje ballot i vote
  4. Cooperative Groups
← Powrót do CUDA Academy