CUDA Academy · Lekcja

API fragmentów WMMA

Ładowanie, mma_sync i zapisywanie fragmentów

Lekcja 3 z 413 kroki

API fragmentów WMMA 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.

API WMMA

Do bezpośredniego programowania Tensor Cores służy WMMA — API mnożenia macierzy z akumulacją dla warpu w przestrzeni nazw nvcuda::wmma. 🧩

Współpraca całego warpu

WMMA działa na poziomie całego warpu: wszystkie 32 wątki w warpie wspólnie pracują nad jednym kafelkiem. Należy myśleć o kafelkach, a nie o pojedynczych wątkach.

Poznaj fragment

Fragment to typ danych WMMA przechowujący część kafelka macierzy należącą do jednego warpu. Każdy wątek posiada niewielką część jego elementów.

Trzy role fragmentów

Fragmenty deklaruje się z oznaczeniem matrix_a, matrix_b lub accumulator. Oznaczenie informuje WMMA, jak dany kafelek uczestniczy w mnożeniu z akumulacją.

Kształty kafelków są stałe

Fragmenty używają ustalonych rozmiarów kafelków, takich jak 16 na 16 na 16. Należy wybrać obsługiwany kształt; sprzęt akceptuje tylko określone kombinacje.

Krok 1: Ładowanie

Najpierw należy wywołać load_matrix_sync, aby przenieść kafelek z pamięci do fragmentu. To ładowanie automatycznie rozdziela dane między wątki warpu.

wmma::load_matrix_sync(a_frag, ptr, ldm);

Krok 2: Wyzerowanie akumulatora

Przed dodawaniem należy wyzerować kafelek wynikowy za pomocą fill_fragment. Czysty akumulator zapewnia rozpoczęcie sumowania od znanej wartości.

wmma::fill_fragment(c_frag, 0.0f);

Krok 3: Mnożenie z akumulacją

Podstawowym wywołaniem jest mma_sync. Wykonuje ono D = A razy B plus C na załadowanych fragmentach, bezpośrednio za pomocą Tensor Cores.

wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

Krok 4: Zapis

Na koniec store_matrix_sync zapisuje fragment akumulatora z powrotem do pamięci. To zapisywanie scala części każdego wątku w jeden kafelek.

wmma::store_matrix_sync(out, c_frag, ldm, layout);

Sync oznacza synchronizację warpu

Sufiks _sync przypomina, że każde wywołanie WMMA jest synchronizowane na poziomie warpu. Wszystkie 32 lane’y muszą dotrzeć do niego razem, w przeciwnym razie zachowanie jest niezdefiniowane.

Układ i wymiar wiodący

Ładowanie i zapisywanie wymagają podania wymiaru wiodącego, czyli kroku między wierszami, a także układu row-major lub column-major, aby poprawnie odczytywać pamięć.

Szybkie sprawdzenie

Które wywołanie WMMA faktycznie wykonuje mnożenie macierzy z akumulacją na Tensor Cores?

Podsumowanie

Przeszedł Pan przez schemat działania WMMA: deklarowanie fragmentu, ładowanie kafelków, wyzerowanie akumulatora, wywołanie mma_sync, a następnie zapis. Każdy krok jest synchronizowany na poziomie warpu. 🙌

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 „API fragmentów WMMA” jest bezpłatna?

Tak — pełny tekst „API fragmentów WMMA” 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 „API fragmentów WMMA”?

Ładowanie, mma_sync i zapisywanie fragmentów Ć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 „API fragmentów WMMA”?

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. Co obliczają Tensor Cores
  2. Mieszana precyzja: FP16, BF16, TF32
  3. API fragmentów WMMA
  4. Kompromisy stabilności numerycznej
← Powrót do CUDA Academy