Stencil and Sliding Windows
Tiling neighborhoods with halo cells.
Stencil and Sliding Windows is a free CUDA Academy lesson on CoddyKit — lesson 3 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.
What a Stencil Is
A stencil computes each output from a fixed neighborhood of inputs, like averaging a pixel with the cells around it.
The Sliding Window
As you move along the array, the input window slides by one. Consecutive outputs share most of their inputs, so reuse is huge.
Perfect for Tiling
Because neighbors overlap, a tile loaded once feeds many outputs. Stencils are a textbook case where tiling shines.
Meet the Halo
To compute the edges of a tile you also need a few elements just outside it. Those extra border cells are called the halo.
Why the Halo Exists
The first thread in a tile needs its left neighbor, which belongs to the previous tile. Without halo cells that read goes wrong.
Sizing the Halo
For a radius-r stencil you load r extra cells on each side. A 3-point blur has radius 1, so one halo cell per side suffices.
A Padded Tile
Declare shared memory as blockDim plus 2 * radius so it holds the interior plus both halos in one tidy buffer.
__shared__ float tile[BLOCK + 2 * RADIUS];Loading the Interior
Each thread first loads its own element into the tile at an offset of radius, leaving room for the left halo in front.
tile[threadIdx.x + RADIUS] = in[gid];Loading the Halos
The first few threads do double duty, also fetching the left and right halo cells before the block synchronizes.
Then Sync and Stencil
After __syncthreads, each thread reads its neighbors entirely from the tile, never touching global memory again for that step.
Reuse Multiplier
Every interior value now serves 2r + 1 outputs from one shared-memory load. That is exactly the redundant traffic tiling removes.
Quick Check
Why does a tiled stencil need halo cells?
Recap
Stencils slide overlapping windows, so a halo of r extra cells per side lets a tile serve every output with one load each. ✅
Frequently asked questions
Is the “Stencil and Sliding Windows” lesson free?
Yes — the full text of “Stencil and Sliding Windows” 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 “Stencil and Sliding Windows”?
Tiling neighborhoods with halo cells. 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 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Stencil and Sliding Windows” 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
- The Data Reuse Problem
- The Load-Sync-Compute Pattern
- Stencil and Sliding Windows
- Handling Edge Tiles