CUDA Academy · レッスン

カーネルからカーネルを起動する

デバイス側の再帰と洗練

「カーネルからカーネルを起動する」はCoddyKit上の無料CUDA Academyレッスンです。これはレッスン1/4です。下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。これは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時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。

「カーネルからカーネルを起動する」で何を学びますか？

デバイス側の再帰と洗練ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。

CUDA Academyを始めるのに経験は必要ですか？

事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。これはレッスン1/4です。

「カーネルからカーネルを起動する」レッスンにはどのくらい時間がかかりますか？

ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。

このCUDA Academyレッスンでコードを書いて実行できますか？

はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

← CUDA Academyに戻る