0Pricing
CUDA Academy · Lekcja

Uruchamianie kernelów z poziomu kernela

Rekurencja i udoskonalanie po stronie urządzenia

Uruchamianie kernelów z poziomu kernela to bezpłatna lekcja CUDA Academy na CoddyKit. To lekcja 1 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.

Kernele uruchamiające kernele

Dzięki dynamic parallelism działający kernel GPU może sam uruchomić kolejny kernel, bez konieczności powrotu do CPU. 🚀

Ta sama składnia po stronie urządzenia

Uruchomienie wygląda tak samo jak po stronie hosta: znane potrójne nawiasy ostre działają bezpośrednio w kodzie urządzenia.

__global__ void child(int* d);
__global__ void parent(int* d) {
    child<<<1, 32>>>(d);
}

Obliczanie przyszłych obliczeń

Wątek nadrzędny w czasie działania określa, ile pracy jest potrzebne, a następnie uruchamia kernel potomny o dokładnie dopasowanym rozmiarze.

child<<<blocks, threads>>>(buf);

Stworzone dla nieregularnych kształtów

Rozwiązanie sprawdza się doskonale, gdy praca zależy od danych: wątek, który wykryje więcej zadań, może od razu uruchomić więcej wątków.

Rekurencja po stronie urządzenia

Kernel może nawet uruchomić sam siebie, zapewniając prawdziwą rekurencję na GPU dla problemów typu dziel i zwyciężaj.

__global__ void solve(int lo, int hi) {
    if (hi - lo > 1) solve<<<1, 2>>>(lo, mid);
}

Strumienie nadrzędne i potomne

Każde uruchomienie potomne dołącza do strumienia. Domyślnie elementy potomne korzystają ze strumienia bloku wątków nadrzędnych, ale można wskazać własny strumień.

child<<<g, b, 0, myStream>>>(d);

Synchronizacja wewnątrz kernela

Element nadrzędny może zaczekać na elementy potomne za pomocą cudaDeviceSynchronize wywołanego w kodzie urządzenia, a następnie odczytać ich wyniki.

child<<<1, 64>>>(d);
cudaDeviceSynchronize();

Niejawne zakończenie elementów potomnych

Nawet bez ręcznego oczekiwania wszystkie uruchomione elementy potomne na pewno zakończą działanie, zanim zwróci się sam kernel nadrzędny.

Pamięć widoczna dla obu stron

Element nadrzędny i potomny współdzielą pamięć globalną, więc przekazane wskaźniki pozostają prawidłowe. Dane lokalne i współdzielone nie są jednak dostępne po drugiej stronie uruchomienia.

Kompilacja na potrzeby uruchamiania po stronie urządzenia

Należy skompilować program z użyciem relocatable device code i dołączyć środowisko uruchomieniowe urządzenia, aby zagnieżdżone uruchomienia faktycznie działały.

nvcc -rdc=true -lcudadevrt prog.cu

Twardy limit zagnieżdżania uruchomień

Zagnieżdżanie ma ograniczenie: istnieje maksymalna głębokość uruchamiania, a jej przekroczenie powoduje błąd zamiast uruchomienia kolejnych kerneli.

Szybkie sprawdzenie

Skąd pochodzi kernel uruchomiony dynamicznie?

Podsumowanie: uruchamianie po stronie urządzenia

Kernel może uruchamiać inne kernele, a nawet sam siebie, za pomocą tej samej składni potrójnych nawiasów. Udostępnij pamięć globalną i skompiluj program z opcją -rdc=true. Świetna praca! 🎉

Często zadawane pytania

Czy lekcja „Uruchamianie kernelów z poziomu kernela” jest bezpłatna?

Tak — pełny tekst „Uruchamianie kernelów z poziomu kernela” 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 „Uruchamianie kernelów z poziomu kernela”?

Rekurencja i udoskonalanie po stronie urządzenia Ć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 1 z 4.

Ile czasu zajmuje lekcja „Uruchamianie kernelów z poziomu kernela”?

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. Uruchamianie kernelów z poziomu kernela
  2. Kiedy opłaca się dynamic parallelism
  3. Przechwytywanie pracy do grafu
  4. Ponowne odtwarzanie grafów w celu ograniczenia narzutu
← Powrót do CUDA Academy