The Triple-Angle-Bracket Launch
Calling a kernel with >>.
The Triple-Angle-Bracket Launch is a free CUDA Academy lesson on CoddyKit — lesson 2 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
Launching Is Different
You do not call a kernel like a normal function. You launch it, telling the GPU how many threads to spin up at the same time. 🚀
The <<< >>> Syntax
The launch uses CUDA's special triple angle brackets. Inside them you give a launch configuration before the normal argument list.
myKernel<<<blocks, threads>>>(args);First Number: Blocks
The first value sets how many blocks launch. A block is a group of threads that run together and share fast on-chip memory.
myKernel<<<4, threads>>>(args); // 4 blocksSecond Number: Threads per Block
The second value sets threads per block. Total threads launched equals blocks times threads per block.
myKernel<<<4, 256>>>(args); // 1024 threadsArguments Come After
After the brackets, you pass the kernel's normal arguments in parentheses, just like any C++ call.
doubleIt<<<1, 256>>>(devicePtr);Launches Are Asynchronous
A launch is asynchronous: the CPU queues the work and keeps going immediately. The GPU runs the kernel in the background.
Configs Can Be Variables
The launch numbers do not have to be literals. You often compute the block count from your data size at runtime.
int blocks = (n + 255) / 256;
add<<<blocks, 256>>>(c, a, b, n);Using dim3 for 2D
For grids and blocks in 2D or 3D, you use a dim3 value. It bundles x, y, and z sizes into the launch.
dim3 threads(16, 16);
blur<<<grid, threads>>>(img);Picking Threads per Block
A safe default for threads per block is 128 or 256. Values must be multiples of 32 and stay at or below 1024.
kernel<<<blocks, 128>>>(data);A Bad Config Fails
If you ask for too many threads per block, the launch fails silently. You must check for errors right after launching.
kernel<<<1, 2048>>>(p); // too big, failsRead It Left to Right
Read a launch as: run this kernel across this many blocks of this many threads, with these arguments. Simple once it clicks.
vectorAdd<<<grid, block>>>(c, a, b, n);Quick Check
Time to check the launch configuration.
Recap: The Launch
You launch with kernel<<
Frequently asked questions
Is the “The Triple-Angle-Bracket Launch” lesson free?
Yes — the full text of “The Triple-Angle-Bracket Launch” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “The Triple-Angle-Bracket Launch”?
Calling a kernel with >>. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 2 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “The Triple-Angle-Bracket Launch” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Anatomy of a Kernel
- The Triple-Angle-Bracket Launch
- printf Inside a Kernel
- cudaDeviceSynchronize Explained