CUDA Academy · Lekcja

Eliminowanie rozbieżności warpu

Zmieniać indeksowanie, aby utrzymać warpy w pracy.

Lekcja 2 z 413 kroki

Eliminowanie rozbieżności warpu to bezpłatna lekcja CUDA Academy na CoddyKit. To lekcja 2 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.

Warpy działają synchronicznie

Warp to 32 wątki wykonujące wspólnie tę samą instrukcję. Gdy ich ścieżki są zgodne, sprzęt działa z pełną szybkością.

Koszt rozbieżności

Jeśli wątki w warpie wybierają różne gałęzie, mamy do czynienia z rozbieżnością. Sprzęt wykonuje każdą ścieżkę sekwencyjnie, pozostawiając część lane'ów bezczynnych i marnując cykle.

Naiwna redukcja powoduje rozbieżność

Prosta wersja używa tid % (2*s) do wyboru aktywnych wątków. Aktywne i bezczynne wątki przeplatają się w każdym warpie, dlatego każdy warp silnie się rozbiega.

if (tid % (2 * s) == 0)
  data[tid] += data[tid + s];

Bezczynne lane'y nadal kosztują

Mimo że połowa wątków nic nie robi, nadal zajmują one warp. Warp nie może zakończyć pracy, dopóki nie zostaną obsłużone zarówno ścieżki aktywne, jak i bezczynne.

Ponowne indeksowanie według identyfikatora wątku

Rozwiązaniem jest przypisanie aktywnej pracy do najniższych identyfikatorów wątków zamiast do rozproszonych identyfikatorów. Oblicz indeks na podstawie tid i kroku.

int index = 2 * s * tid;
if (index < blockDim.x)
  data[index] += data[index + s];

Dlaczego to pomaga

Teraz zajęte wątki są ciągłe: tid 0,1,2,... pracują, a wszystkie pozostałe pozostają bezczynne. Całe warpy są albo w pełni aktywne, albo w pełni bezczynne.

Całkowicie bezczynne warpy nic nie kosztują

Warp, w którym każdy lane jest bezczynny, kończy działanie bez wykonywania pracy. Nie ma serializacji na poziomie lane'ów, więc koszt rozbieżności w dużej mierze znika.

Pułapka operatora modulo

Ukrytym problemem był warunek z użyciem modulo. Rozpraszał on aktywne wątki w każdym warpie, co dokładnie prowadzi do rozbieżności.

Ta sama praca, lepsze przypisanie

Nie zmienili Państwo obliczeń ani liczby dodawań. Zmienili Państwo tylko to, który wątek wykonuje dane dodawanie, a warpy odwdzięczają się za to lepszą wydajnością.

Efekt narasta wraz ze skalą

W tysiącach bloków i wielu krokach usunięcie rozbieżności daje rzeczywiste przyspieszenie, często pozwalając działać kilka razy szybciej niż naiwny kernel.

Pozostał jeszcze jeden problem

Ta wersja odczytuje sąsiadów przeplatanych w pamięci współdzielonej, co może powodować konflikty banków. Następna lekcja rozwiąże również ten problem.

Szybkie sprawdzenie

Zastanów się, co powoduje rozbieżność warpa w naiwnej redukcji.

Odpowiedz w interaktywnej lekcji

Podsumowanie

Wykluczono rozbieżność, przydzielając pracę czątkom o najniższych identyfikatorach, dzięki czemu warpy są w całości aktywne albo bezczynne. Ta sama matematyka, szybsza redukcja. Następne: konflikty banków. 🚀

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

Zacznij naukę bezpłatnie Wyświetl pełny kurs

Często zadawane pytania

Czy lekcja „Eliminowanie rozbieżności warpu” jest bezpłatna?

Tak — pełny tekst „Eliminowanie rozbieżności warpu” 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 „Eliminowanie rozbieżności warpu”?

Zmieniać indeksowanie, aby utrzymać warpy w pracy. Ć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 2 z 4.

Ile czasu zajmuje lekcja „Eliminowanie rozbieżności warpu”?

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

← Powrót do CUDA Academy