Why architecture matters here
Architecture matters here because the GPU's compute throughput is enormous and its global-memory bandwidth, though large, is nowhere near enough to feed the cores if every operand comes from HBM. The ratio of arithmetic to memory traffic — the arithmetic intensity — must be high for a kernel to be compute-bound rather than memory-bound, and the only way to raise it is to load a piece of data once and use it many times. Shared memory is the mechanism for that reuse: it is the staging area where a tile of data lives close to the cores so that dozens or hundreds of arithmetic operations can hit it without another trip to HBM. Without shared memory, a matrix multiply re-reads the same rows and columns from global memory over and over and runs at memory-bandwidth speed, a small fraction of the chip's arithmetic peak.
The cost of misusing shared memory is subtle precisely because the kernel still produces correct results — it just runs slowly, and the slowdown hides inside a single instruction that the source code makes look free. A bank conflict does not throw an error or corrupt data; it silently multiplies the latency of shared-memory accesses, so a kernel that should be compute-bound spends most of its time stalled on serialized scratchpad reads. Because the conflict is invisible in the code (a normal array index) and invisible in the output (correct numbers), it is one of the most common performance bugs that goes undiagnosed for the life of a kernel unless someone profiles it. Architecture literacy here is what lets you see the conflict that the source hides.
Shared memory also sits at the center of a fundamental resource trade-off that shapes the whole kernel: occupancy. Each SM has a fixed budget of shared memory (and registers), and that budget is divided among the thread blocks resident on the SM at once. A kernel that allocates a large shared-memory tile per block allows fewer blocks to be co-resident, which reduces occupancy — the number of warps available to hide memory latency by switching among them. So more shared memory per block buys more reuse but less latency-hiding; less shared memory buys more occupancy but less reuse. Every fast kernel is a deliberate point on this curve, and choosing it well requires understanding that shared memory is not free capacity but a rationed resource traded against parallelism.
Finally, shared memory is the only fast, general mechanism for threads within a block to cooperate. Threads in a warp can exchange registers directly via shuffle instructions, but cooperation across warps within a block — the essence of tiling, reductions, and prefix sums — goes through shared memory, coordinated by the __syncthreads() barrier. This makes shared memory not just a performance cache but the communication fabric of a thread block, and its correct use is inseparable from correct synchronization. The architecture ties data staging and inter-thread coordination into one resource, which is why understanding it is foundational rather than an optimization afterthought.
The architecture: every piece explained
Physically, shared memory is an on-SM SRAM scratchpad organized into 32 banks, each of which can service one 4-byte word per cycle. Successive 4-byte words map to successive banks and wrap around: word 0 is in bank 0, word 1 in bank 1, … word 31 in bank 31, word 32 back in bank 0, and so on. Because a warp is 32 threads and a shared-memory instruction is issued warp-wide, the hardware wants each of the 32 threads to hit a different bank so all 32 accesses complete in one cycle. The bank a given address lands in is simply (address / 4) mod 32 — a fact you must be able to compute in your head to reason about conflicts.
A bank conflict occurs when two or more threads in the same warp access different 4-byte words that map to the same bank. The hardware cannot read two different words from one bank in one cycle, so it serializes: an N-way conflict (N threads hitting distinct words in one bank) takes N cycles. The pathological case is a stride that aligns all 32 threads on one bank — for example, indexing tile[threadIdx.x * 32] into a tile 32 words wide, so every thread's address is a multiple of 32 words and therefore in bank 0 — a 32-way conflict that runs 32× slower. The broadcast exception: if all threads read the same address (same bank, same word), or if reads to a bank all target the identical word, the value is broadcast in a single cycle. Same-word is free; same-bank-different-word serializes.
The canonical fix is padding. Consider a 32×32 tile stored row-major as tile[32][32]; a column access tile[i][col] across the 32 threads (varying i) touches addresses 32 words apart, all in the same bank — a full 32-way conflict. Declaring the tile one column wider, tile[32][33], and ignoring the padding column shifts each row by one word, so successive rows land in successive banks and the column access spreads across all 32 banks conflict-free. The padding wastes a sliver of shared memory but eliminates the serialization — a trade that is almost always worth it. More sophisticated kernels use swizzling: an XOR-based permutation of the address that scatters accesses across banks for both row and column patterns without wasting space, which is what high-performance libraries and Tensor-Core kernels use.
Correct use also requires the synchronization primitives. __syncthreads() is a barrier: every thread in the block must reach it before any proceeds, which is how you guarantee that all threads have finished writing the shared tile before any thread starts reading it, and vice versa when the tile is reused for the next stage. Getting the barriers right is a correctness requirement, not a performance one: omit a barrier and threads read data that other threads have not written yet — a race that produces wrong, non-deterministic results. Modern architectures also add asynchronous copy instructions that stream data from global to shared memory in the background, overlapping the load of the next tile with computation on the current one, which is the basis of the deeply-pipelined kernels that reach peak throughput.
End-to-end flow
Trace a tiled matrix-multiply kernel computing C = A × B, the archetypal shared-memory workload. The output matrix is divided into tiles, and each thread block is responsible for one output tile. The block loops over the shared inner dimension in steps: in each step, the block cooperatively loads one tile of A and one tile of B from global memory into shared memory, with the threads arranging their loads to be coalesced (adjacent threads read adjacent global addresses) so the global reads are efficient. Each thread loads a few elements; together the block stages the two tiles into the on-SM scratchpad.
The block then calls __syncthreads() to guarantee both tiles are fully loaded before anyone reads them. Now comes the reuse that justifies everything: each thread computes its partial dot product by reading a row of the A tile and a column of the B tile from shared memory, multiplying and accumulating across the tile width. Every element loaded from global memory is read many times here — once per output element in its row or column — so the expensive global load is amortized over many cheap shared-memory reads, raising arithmetic intensity into the compute-bound regime. This is the whole game: load slow once, read fast many times.
Bank behavior determines how fast that inner loop runs. If the B tile is stored naively and threads read down a column, the accesses collide on one bank and the inner loop serializes, throwing away most of the benefit. If the tile is padded (an extra column) or swizzled, the column reads spread across all 32 banks and run at full speed. A profiler would show the naive version spending its time on 'shared memory bank conflicts' while the padded version shows near-zero conflicts and much higher throughput — the same arithmetic, the same correct result, a multiple-times difference in speed, decided entirely by the memory layout.
After the inner loop, the block calls __syncthreads() again — this time to ensure every thread has finished reading the current tiles before the next iteration overwrites them with the next pair of tiles — and repeats: load the next A and B tiles, sync, accumulate, sync. When the loop over the inner dimension completes, each thread holds its final accumulated output value and writes it back to global memory in a coalesced pattern. The kernel touched each global element the minimum number of times, did the bulk of its reads from the fast scratchpad, avoided bank conflicts through padding, and used two barriers per iteration to keep the shared tiles race-free. That combination — stage, sync, reuse, sync — is the template that every high-performance GPU kernel, up to and including the flash-attention kernels in modern LLM serving, is built on.