Declarar arrays __shared__

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.

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

You declare it with the __shared__ qualifier inside a kernel. This one array is then shared by every thread in the block.

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.

This is the key idea: a __shared__ array is created once per block, not once per thread. All threads see the exact same array.

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.

Its lifetime matches the block. The array is born when the block starts and gone when it ends, so it has block scope only. 🧱

A common pattern is one slot per thread, indexed by threadIdx.x. Each thread loads its element into shared memory in parallel.

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 most common job is staging a tile of global data so the block can reuse it many times without going back to slow DRAM.

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.

Let us check how shared memory is scoped.

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