0Pricing
CUDA Academy · Lekcja

Budowanie histogramu

Używać operacji atomowych z prywatyzacją w pamięci współdzielonej.

Budowanie histogramu 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.

Co zlicza histogram

Histogram zlicza, ile danych wejściowych trafia do każdego przedziału. Wiele wątków będzie próbować zwiększyć ten sam przedział, więc jest to problem dla operacji atomowych. 📊

Podejście naiwne

Każdy wątek odczytuje jeden element, znajduje jego przedział i zwiększa jego licznik. Bez ochrony popularne przedziały tracą zliczenia z powodu wyścigów.

Wersja z globalną operacją atomową

Najprostsza poprawna poprawka to jedno atomicAdd na element, wykonywane bezpośrednio w pamięci globalnej. Działa, ale często używane przedziały wykonują wątki szeregowo.

atomicAdd(&hist[bin], 1);

Problem rywalizacji

Gdy dane skupiają się w kilku przedziałach, tysiące wątków kierują operacje pod ten sam adres. Taka rywalizacja może boleśnie spowolnić globalne operacje atomowe.

Prywatyzacja na ratunek

Prywatyzacja daje każdemu blokowi jego własny prywatny histogram w szybkiej pamięci współdzielonej. Wątki kolidują tylko wewnątrz swojego bloku, a nie w całej siatce.

Zadeklaruj histogram współdzielony

Każdy blok deklaruje współdzieloną tablicę o rozmiarze równym liczbie przedziałów. Znajduje się ona na układzie, więc operacje atomowe są tam znacznie tańsze niż w pamięci globalnej.

__shared__ int local[NBINS];

Krok 1: wyczyść przedziały

Wątki wspólnie zerują histogram współdzielony, a następnie wywołują __syncthreads, aby nikt nie zliczał przed zakończeniem czyszczenia.

local[tid] = 0;
__syncthreads();

Krok 2: zliczaj lokalnie

Teraz każdy wątek atomowo zwiększa licznik swojego przedziału we współdzielonej pamięci. To ten sam atomicAdd, ale wykonywany na szybkiej kopii na układzie zamiast w pamięci globalnej.

atomicAdd(&local[bin], 1);

Krok 3: scal z pamięcią globalną

Po synchronizacji wątki dodają każdy współdzielony przedział do globalnego histogramu, wykonując jedno atomicAdd na przedział. To znacznie mniej globalnych operacji atomowych niż wcześniej.

atomicAdd(&hist[i], local[i]);

Dlaczego to jest szybsze

Operacje atomowe w pamięci współdzielonej są szybkie, a kosztowne operacje atomowe w pamięci globalnej są teraz wykonywane raz na przedział na blok, zamiast raz na element.

Kontroluj liczbę przedziałów

Prywatny histogram musi zmieścić się w pamięci współdzielonej. Przy zbyt dużej liczbie przedziałów należy podzielić pracę na przebiegi lub wrócić do globalnych operacji atomowych.

Szybkie sprawdzenie

Jedno pytanie dotyczące strategii tworzenia histogramu.

Podsumowanie: tworzenie histogramu

Zbudowałeś histogram za pomocą globalnych operacji atomowych, a następnie przyspieszyłeś go dzięki prywatyzacji w pamięci współdzielonej: wyczyść, zliczaj lokalnie, scal. ✅

Często zadawane pytania

Czy lekcja „Budowanie histogramu” jest bezpłatna?

Tak — pełny tekst „Budowanie histogramu” 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 „Budowanie histogramu”?

Używać operacji atomowych z prywatyzacją w pamięci współdzielonej. Ć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 „Budowanie histogramu”?

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. Warunki wyścigu na GPU
  2. atomicAdd i funkcje pokrewne
  3. Budowanie histogramu
  4. Własne operacje atomowe z atomicCAS
← Powrót do CUDA Academy