Dynamic Shared Memory
Sizing shared memory at launch time.
Dynamic Shared Memory is a free CUDA Academy lesson on CoddyKit — lesson 4 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.
When Size Is Not Known Yet
Sometimes the shared array size depends on a runtime value, so you cannot hardcode it. CUDA solves this with dynamic shared memory.
The extern Declaration
You declare it with extern __shared__ and empty brackets. No size goes in the code, so the compiler leaves it open.
extern __shared__ float sdata[];Size Set at Launch
You pass the byte count as the third launch parameter, between the thread count and the closing brackets. That sets the size at runtime.
kern<<<blocks, threads, n * sizeof(float)>>>();It Is Always in Bytes
A classic bug: that third argument is the size in bytes, not elements. Always multiply your count by sizeof to get it right.
Default Is Zero Bytes
If you skip the third argument, CUDA reserves zero dynamic shared memory. Reading sdata then is undefined behavior, so never forget it.
Static and Dynamic Together
You can use both kinds in one kernel. Static arrays keep fixed sizes while the dynamic block flexes with the launch argument.
One Buffer, Many Arrays
Dynamic shared memory gives you one flat buffer. To hold several arrays, you carve it up yourself with pointer offsets. 🔪
Carving With Offsets
Point a second array partway into the buffer. Just be careful to align each piece so a float never starts mid-word.
extern __shared__ float buf[];
float* a = buf;
float* b = &buf[blockDim.x];Why Bother With Dynamic
It lets one compiled kernel serve many tile sizes. You tune the shared size per launch without recompiling for each case.
Watch the Hardware Limit
You still cannot exceed the per-block maximum. Ask for too many bytes at launch and the kernel simply fails to start.
Opting Into Large Allocations
To go past the default cap on newer GPUs, you must opt in with cudaFuncSetAttribute before the launch. Otherwise big requests are rejected.
Quick Check
Let us check how dynamic shared memory is sized.
Recap
You learned that extern __shared__ arrays get their byte size at launch, can be sliced into several arrays, and respect the per-block limit. Course complete! 🎉
Frequently asked questions
Is the “Dynamic Shared Memory” lesson free?
Yes — the full text of “Dynamic Shared Memory” 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 “Dynamic Shared Memory”?
Sizing shared memory at launch time. 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 4 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Dynamic Shared Memory” 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