CUDA Academy · 课时

声明 shared 数组

每个线程块的快速暂存内存

声明 __shared__ 数组 是 CoddyKit 上的免费 CUDA Academy 课时。这是第 1 节课，共 4 节。你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。

本课时的部分内容尚未翻译，以英文显示。

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. 🎯

常见问题解答