Declaring __shared__ Arrays
Fast scratchpad memory per block.
Declaring __shared__ Arrays is a free CUDA Academy lesson on CoddyKit — lesson 1 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.
A Scratchpad Per Block
Every block gets a tiny, fast pool of on-chip memory called shared memory. Think of it as a scratchpad the whole block can write to and read from together.
Why It Is So Fast
Shared memory sits right on the streaming multiprocessor, so it is roughly 100x faster than global memory. Use it to avoid hammering slow off-chip DRAM. 🚀
The __shared__ Keyword
You declare it with the __shared__ qualifier inside a kernel. This one array is then shared by every thread in the block.
__global__ void kern() {
__shared__ float tile[256];
}Fixed Size at Compile Time
When you give a size in brackets, that is static shared memory. The compiler must know the size, so it has to be a constant, not a runtime value.
__shared__ int counts[128];One Copy, Not One Per Thread
This is the key idea: a __shared__ array is created once per block, not once per thread. All threads see the exact same array.
Threads Cooperate Through It
Because every thread sees the same data, shared memory lets threads cooperate. One thread can stash a value and a neighbor can pick it up.
Block Scope, Not Beyond
Its lifetime matches the block. The array is born when the block starts and gone when it ends, so it has block scope only. 🧱
Each Thread Owns a Slot
A common pattern is one slot per thread, indexed by threadIdx.x. Each thread loads its element into shared memory in parallel.
__shared__ float s[256];
s[threadIdx.x] = input[i];A Tiny but Precious Resource
Shared memory is small, often just 48 to 100 KB per SM. Asking for too much per block reduces how many blocks can run at once.
The Classic Use: Staging Tiles
The most common job is staging a tile of global data so the block can reuse it many times without going back to slow DRAM.
Not Visible to Other Blocks
Remember the boundary: shared memory is private to its block. Two different blocks each get their own separate copy and cannot peek at each other.
Quick Check
Let us check how shared memory is scoped.
Recap
You learned that __shared__ gives each block a fast on-chip scratchpad, created once per block and ideal for staging reusable data. Next: keeping threads in step. 🎯
Frequently asked questions
Is the “Declaring __shared__ Arrays” lesson free?
Yes — the full text of “Declaring __shared__ Arrays” 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 “Declaring __shared__ Arrays”?
Fast scratchpad memory per block. 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 1 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Declaring __shared__ Arrays” 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
- Declaring __shared__ Arrays
- Synchronizing with __syncthreads
- Avoiding Bank Conflicts
- Dynamic Shared Memory