CUDA Academy · Lekcja

Cooperative Groups

Poznać nowoczesne API do elastycznej synchronizacji.

Cooperative Groups to bezpłatna lekcja CUDA Academy na CoddyKit. To lekcja 4 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.

Czytelniejsze API synchronizacji

Surowe maski i intrinsics działają, ale ich użycie bywa kłopotliwe. Cooperative groups opakowują wątki w obiekty, które można nazwać, określić ich rozmiar i jawnie synchronizować.

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

Pobierz cały blok

Najpierw pobierz uchwyt do swojego bloku wątków. this_thread_block zwraca grupę, którą można synchronizować podobnie jak za pomocą syncthreads, ale w postaci obiektu.

cg::thread_block block = cg::this_thread_block();

Synchronizuj przez grupę

Wywołanie sync dla bloku jest nowoczesną barierą. Działa dokładnie tak jak syncthreads, ale jasno wskazuje metodę grupy, której Pan używa.

block.sync();

Wydziel tile warpa

Blok można podzielić na tiles o stałym rozmiarze. Tile złożony z 32 lane'ów daje grupę wielkości warpa, z czytelnymi metodami zamiast surowych masek shuffle.

auto warp = cg::tiled_partition<32>(block);

Metody zastępują maski

Grupa kafelkowa udostępnia shfl_down i podobne metody bez argumentu maski. Grupa zna już swoich członków, więc API pozostaje krótkie i bezpieczne.

val += warp.shfl_down(val, offset);

Poznaj swoją pozycję

Każda grupa udostępnia swoich członków oraz pozycję bieżącego wątku. thread_rank zwraca indeks wątku w grupie, a size zwraca liczbę należących do niej wątków.

int rank = warp.thread_rank();

Mniejsze tile również

Tile nie muszą mieć szerokości 32. tiled_partition o rozmiarze 8 lub 16 tworzy grupy mniejsze niż warp, przydatne, gdy dane naturalnie układają się w małe zestawy.

Redukcje na poziomie grupy

Biblioteka udostępnia gotowe operacje kolektywne. Grupowe reduce sumuje tile jednym wywołaniem, ukrywając pętlę przesunięć, którą wcześniej trzeba było napisać ręcznie.

int total = cg::reduce(warp, val, cg::plus<int>());

Synchronizujące się gridy

Najbardziej zaawansowaną grupą jest grid group. Przy użyciu uruchomienia kooperatywnego każdy blok może synchronizować się na jednej barierze — czego nie potrafi zwykły kernel.

Wymagane uruchomienie kooperatywne

Synchronizacja całego gridu działa tylko wtedy, gdy kernel zostanie uruchomiony za pomocą cudaLaunchCooperativeKernel i gdy GPU ją obsługuje. Zwykłe uruchomienie na to nie pozwoli.

Po co to robić

Cooperative groups sprawiają, że kod warpów jest czytelny i przenośny: nie trzeba ręcznie zarządzać maskami, zakresy są jasne, a wielokrotnego użytku operacje kolektywne odpowiadają rzeczywistym zamiarom.

Szybkie sprawdzenie

Proszę przypomnieć sobie, jak utworzyć kooperatywną grupę rozmiaru warpa z bloku.

Podsumowanie

Cooperative groups zamieniają maski w nazwane obiekty: można podzielić blok na tile, wywołać reduce, a nawet synchronizować cały grid. Teraz ma Pan warp-level CUDA pod kontrolą. ✨

Często zadawane pytania

Czy lekcja „Cooperative Groups” jest bezpłatna?

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

Poznać nowoczesne API do elastycznej synchronizacji. Ć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 4 z 4.

Ile czasu zajmuje lekcja „Cooperative Groups”?

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