CUDA Academy · Урок

Запуск ядер из ядра

Рекурсия и уточнение на стороне устройства

«Запуск ядер из ядра» — бесплатный урок CUDA Academy на CoddyKit. Это урок 1 из 4. Ты можешь прочитать весь урок бесплатно ниже — а потом практиковать его прямо в браузере с встроенным редактором кода и ИИ-репетитором 24/7. Это часть пути обучения CUDA Academy, и твой прогресс синхронизируется между веб-версией и приложением CoddyKit. Курс CUDA Academy содержит 4 уроков всего.

Части этого урока еще не переведены и отображаются на английском.

Kernels Launching Kernels

With dynamic parallelism, a running GPU kernel can launch another kernel itself, no trip back to the CPU required. 🚀

Same Syntax, Device Side

The launch looks identical to a host launch: the familiar triple angle brackets work right inside device code.

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

Compute the Future Compute

A parent thread decides at runtime how much work is needed, then spawns a child kernel sized exactly for it.

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

Built for Irregular Shapes

This shines when the work is data-dependent: a thread that finds more to do can launch more threads on the spot.

Device-Side Recursion

A kernel may even launch itself, giving you true recursion on the GPU for divide-and-conquer problems.

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

Parent and Child Streams

Each child launch joins a stream. By default children use the parent thread block stream, but you can name your own stream.

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

Synchronizing Inside a Kernel

A parent can wait for its children with cudaDeviceSynchronize called from device code, then read their results.

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

Implicit Child Completion

Even without a manual wait, all launched children are guaranteed complete before the parent kernel itself returns.

Memory Both Can See

Parent and child share global memory, so pointers passed down stay valid. Local and shared data, though, do not cross the launch.

Compile for Device Launch

You must compile with relocatable device code and link the device runtime so nested launches actually work.

nvcc -rdc=true -lcudadevrt prog.cu

A Hard Launch Limit

Nesting is capped: there is a maximum launch depth, and going too deep returns an error instead of more kernels.

Quick Check

Where does a dynamically launched kernel originate?

Recap: Device Launches

A kernel can launch other kernels with the same triple-bracket syntax, even itself. Share global memory and compile with -rdc=true. Nice work! 🎉

Часто задаваемые вопросы

Урок «Запуск ядер из ядра» бесплатный?

Да — полный текст урока «Запуск ядер из ядра» бесплатно доступен здесь в веб-версии. Чтобы практиковать его интерактивно (встроенный редактор кода и ИИ-репетитор 24/7) и разблокировать остальной курс CUDA Academy, подпишись на CoddyKit PRO. Курс CUDA Academy содержит 4 уроков всего.

Чему я научусь в уроке «Запуск ядер из ядра»?

Рекурсия и уточнение на стороне устройства Ты практикуешь CUDA Academy с помощью реального кода, который запускаешь прямо в браузере, и ИИ-репетитор 24/7 отвечает на твои вопросы во время урока.

Нужен ли мне опыт, чтобы начать CUDA Academy?

Предыдущий опыт не требуется. CUDA Academy на CoddyKit структурирован для всех уровней — от новичков до продвинутых, поэтому ты можешь начать отсюда или с самого начала и учиться в своем темпе. Это урок 1 из 4.

Сколько времени занимает урок «Запуск ядер из ядра»?

Большинство уроков CoddyKit занимают около 5–10 минут. Каждый из них компактный и интерактивный, поэтому ты постоянно делаешь прогресс и продолжаешь с того же места в веб-версии и приложении.

Можно ли писать и запускать код в этом уроке CUDA Academy?

Да. Каждый урок CUDA Academy включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.

Все уроки этого курса

← Назад к CUDA Academy