__shfl_down_sync do redukcji
Wykonywać redukcje warpów bez barier.
__shfl_down_sync do redukcji 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.
Shuffle przenosi wartości rejestrów
Shuffle pozwala jednemu lane'owi bezpośrednio odczytać wartość rejestru innego lane'a. Dane przemieszczają się między wątkami bez pamięci współdzielonej i bez bariery.
Poznaj shfl_down_sync
Podstawowym narzędziem do redukcji jest shfl_down_sync. Każdy lane pobiera wartość z lane'a znajdującego się o ustaloną liczbę pozycji dalej w warp.
float v = __shfl_down_sync(mask, val, offset);Odczytywanie argumentów
Wywołanie przyjmuje aktywną maskę, wartość do udostępnienia oraz przesunięcie. Lane i otrzymuje wartość przechowywaną przez lane i plus przesunięcie.
Dodaj otrzymaną wartość
Suma warpa polega na dodaniu zshufflowanej wartości z powrotem. Lane i pobiera liczbę sąsiada i dodaje ją do własnej sumy bieżącej.
val += __shfl_down_sync(mask, val, offset);Zmniejszanie przesunięcia o połowę
Podobnie jak w redukcji drzewiastej przesunięcie zaczyna się od 16 i w każdym kroku zmniejsza się o połowę: 16, 8, 4, 2, 1. Po pięciu krokach suma obejmuje cały warp.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);Dlaczego pięć kroków
Warp ma 32 lane'y, a 2 do potęgi piątej to 32. Dlatego dokładnie pięć kroków zmniejszania o połowę wystarcza, aby połączyć wszystkie 32 wartości w jedną.
Wynik trafia do lane'a 0
Po zakończeniu pętli pełna suma warpa znajduje się w lane 0. Pozostałe lane'y zawierają częściowe śmieciowe wartości, dlatego wynik powinien zapisać tylko lane 0.
Bariery nie są potrzebne
Ponieważ wymiana odbywa się przez rejestry w trybie lockstep, można całkowicie pominąć syncthreads. Sufiks sync już koordynuje lane'y objęte maską.
Szybciej niż przez pamięć współdzieloną
Redukcja z użyciem warp shuffle przewyższa wersję korzystającą z pamięci współdzielonej: wymaga mniej instrukcji, nie powoduje konfliktów banków i nie ma przestojów na barierze. To szybka ścieżka dla ostatnich 32 elementów.
Przekaż właściwą maskę
Jeśli niektóre lane'y już zakończyły działanie, pełna maska jest nieprawidłowa. Pobierz zbiór aktywnych lane'ów za pomocą activemask, aby każde przesunięcie dotyczyło tylko obecnych lane'ów.
Redukcje dwupoziomowe
Duże redukcje często najpierw redukują każdy warp za pomocą shuffle, a następnie łączą wyniki warpów w shared memory. Shuffle doskonale obsługuje ten tani, wewnętrzny poziom.
Szybkie sprawdzenie
Zastanów się, który lane przechowuje wynik po zakończeniu pętli.
Podsumowanie
Zsumował Pan warp za pomocą shfl_down_sync: pięciokrotnie zmniejszył Pan przesunięcie o połowę, a lane 0 przechowuje całą sumę — bez potrzeby używania barier. Dalej: ballot i vote. ✨
Często zadawane pytania
Czy lekcja „__shfl_down_sync do redukcji” jest bezpłatna?
Tak — pełny tekst „__shfl_down_sync do redukcji” 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 „__shfl_down_sync do redukcji”?
Wykonywać redukcje warpów bez barier. Ć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 „__shfl_down_sync do redukcji”?
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