动态共享内存
在启动时确定共享内存大小
动态共享内存 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 4 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
本课时的部分内容尚未翻译，以英文显示。
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! 🎉
常见问题解答
「动态共享内存」课时是免费的吗？
是的 — 「动态共享内存」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「动态共享内存」这节课中我会学到什么？
在启动时确定共享内存大小 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 4 节课，共 4 节。
「动态共享内存」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。