05_shared_memory

CUDA Memory Hierarchy

To understand shared memory effectively, one must first situate it within the broader context of the GPU memory architecture.

CUDA Thread can access data from several memory spaces during its execution, each with unique properties of scope, lifetime, and speed.

Thread-Private Memory

The fastest and most local memory.

• Registers: The fastest memory on the GPU, private to a single thread and managed by the compiler. Used for local variables, with near-zero access latency.
• Local Memory: Functionally has the same scope as registers (thread-private) but is physically located in off-chip global memory. It is used for “register spilling” when registers are exhausted or to store large automatic arrays. It is critical to understand that “local” refers to its scope, not its speed; local memory is as slow as global memory.

Block-Shared Memory

• Shared Memory: A low-latency on-chip memory space shared by all threads in a block. Its lifetime is that of the block’s execution. It is the key enabler for fast data exchange between threads and for data reuse.

Grid-Wide & Application-Lifetime Memory

Memory accessible by all threads in all blocks.

• Global Memory: The largest memory space but also the slowest. It physically resides in off-chip DRAM, and its data persists between kernel launches within the same application.
• Constant Memory: A read-only memory space accessible by all threads, which is cached and broadcast efficiently to threads within a warp.
• Texture Memory: Another read-only memory space with specialized hardware for spatial locality, filtering, and various addressing modes, making it useful for graphics and image processing.

Cluster-Shared Memory

Introduced with Compute Capability 9.0, Distributed Shared Memory allows threads in a block cluster to access the shared memory of all participating thread blocks in the cluster, breaking the traditional block-level barrier.

Memory Type	Location	Scope	Lifetime	Access	Relative Speed	Typical Size	Managed By
Registers	On-chip	Thread	Thread	R/W	Fastest	Tens to hundreds of KB per thread	Compiler
Local Memory	Off-chip	Thread	Thread	R/W	Slowest	-	Compiler
Shared Memory	On-chip	Block	Block	R/W	Very Fast	Up to 99KB per SM	Programmer
L1 Cache	On-chip	Block	Block	R/W	Very Fast	Tens of KB per SM	Hardware
Constant Memory	Off-chip	Grid	Application	Read-Only	Fast (cached)	Tens of KB	Programmer
Texture Memory	Off-chip	Grid	Application	Read-Only	Fast (cached)	-	Programmer
Global Memory	Off-chip	Grid	Application	R/W	Slowest	Several GB to tens of GB	Programmer

Programmer-Managed Cache vs. Hardware-Managed Cache

• Shared Memory as a “Scratchpad”: It is crucial to emphasize that shared memory is explicitly controlled by the programmer. The programmer decides what data to put in it, when to put it there, and when it is valid. This provides immense power but also responsibility.
• L1 Cache as an Automatic System: This contrasts with the hardware-managed L1 cache, which automatically caches global memory loads based on access patterns and heuristics. While helpful, it can be suboptimal for algorithms with complex or irregular data reuse patterns that the programmer knows about but the hardware cannot predict.

The CUDA memory hierarchy reflects a core philosophy of high-performance computing: exposing hardware control to the programmer for maximum performance. The distinction between programmer-managed shared memory and hardware-managed L1 cache is the clearest example of this philosophy. It is a deliberate design choice that trades ease of use (letting the hardware cache everything) for the potential of higher, more predictable performance when the programmer explicitly manages data locality. The programmer has domain knowledge of the algorithm’s data reuse patterns (e.g., which matrix tile will be needed for the next calculation) and can therefore make more intelligent caching decisions than the hardware. Shared memory is thus not just “fast memory”; it is a tool for the programmer to enforce a specific data locality and reuse strategy, bypassing the unpredictability of a hardware cache. It is a powerful concept that elevates the programmer from a mere instruction writer to a manager of data flow on the silicon itself.

How Shared Memory Works

Declaration and Allocation: Static vs. Dynamic

• Static Allocation: Uses the __shared__ keyword for arrays whose size is known at compiler time (e.g., __shared__ float tile;). This is simple and efficient for fixed-size problems. However, statically declared shared memory may have a limit of 48KB.
• Dynamic Allocation: Uses the extern __shared__ float tile; syntax. The size is not specified in the kernel code but is provided as the third execution configuration parameter upon kernel launch: kernel_name<<<gridDim, blockDim, size_in_bytes>>>(...);. This is essential for writing flexible kernels whose memory requirements may change based on input data or the target GPU.
• Combining Static and Dynamic: If a kernel uses both, the dynamically allocated memory begins after all statically allocated variables. Only one dynamic allocation is permitted per kernel, but it can be partitioned using pointers.

Scope, Lifetime, and Synchronization

• The Per-Block Paradigm: Shared memory is private to a thread block. Its contents are visible to all threads within that block but are completely inaccessible to threads in other blocks. The data exists only for the lifetime of the block.
• The Critical Role of __syncthreads(): This is the synchronization primitive for shared memory. It acts as a barrier, ensuring no thread in the block proceeds past that point until all threads have reached it. Crucially, it includes a memory fence, guaranteeing that any shared memory writes by any thread are visible to all other threads after the barrier.
• A Common Pitfall: Synchronization in Divergent Code: Placing __syncthreads() inside a conditional branch (if/else) that is not uniform across the entire block is a grave risk. It can cause a deadlock, as some threads may wait forever at a barrier that other threads never reach. This is a classic and difficult-to-debug error.

Architecture-Specific Capacity and Configuration

• Evolution Across Architectures: Discuss how shared memory capacity has grown over several GPU generations. This is critical for performance portability.
• Configuring the L1/Shared Memory Split: On modern architectures (Volta and newer), the on-chip memory can be configured on a per-kernel basis to be used as either L1 cache or shared memory. Explain how to use cudaFuncSetAttribute() to provide a hint to the driver about the preferred shared memory size to optimize for occupancy and performance. On Volta (CC 7.0), the unified data cache is 128 KB, and shared memory capacity can be set to 0, 8, 16, 32, 64, or 96 KB. On Turing (CC 7.5), it can be set to 32 KB or 64 KB for a 96 KB cache.

Dynamic shared memory allocation is a cornerstone of writing high-performance, portable CUDA code. It decouples the algorithm’s logic from the hard-coded constraints of a specific GPU architecture. Static allocation hard-codes the memory size into the compiled binary. However, the optimal tile size may depend on the total shared memory available on the target GPU (e.g., 48 KB vs. 96 KB). By using dynamic allocation, the host code can query the device properties at runtime with cudaGetDeviceProperties and then launch the same kernel with a size_in_bytes parameter that is optimized for that specific device. This means a single compiled kernel can achieve near-optimal performance on a wide range of hardware without recompilation. It transforms a static, compile-time decision into a more powerful and flexible runtime one—a key strategy for building robust, future-proof GPU applications.

Anatomy of a Bottleneck: Understanding and Mitigating Bank Conflicts

The Banked Architecture Explained

To achieve its high bandwidth, shared memory is not a single monolithic block of memory. It is divided into 32 independent modules called banks. Successive 32-bit words are mapped to successive banks. That is, address 0 is in bank 0, address 4 bytes is in bank 1, … address 124 bytes is in bank 31, and address 128 bytes wraps back around to bank 0. The system can service one request from each of the 32 banks simultaneously in a single clock cycle. This is the source of shared memory’s immense bandwidth.

Defining and Visualizing Bank Conflicts

A bank conflict occurs when two or more threads within the same warp attempt to access different addresses that fall into the same memory bank in the same instruction cycle. When a conflict occurs, the hardware must serialize the accesses. A 2-way bank conflict takes two cycles instead of one, halving the effective bandwidth. A 32-way conflict (the worst case) takes 32 cycles, reducing bandwidth by a factor of 32.

• The Critical Exception: Broadcast. If multiple threads in a warp access the exact same 32-bit address, this is not a conflict. The value is simply broadcast to all requesting threads in one cycle.

Conflict Avoidance Strategies

Padding

The most common technique. Add dummy elements to one dimension of a 2D array to change the stride between consecutive rows, thereby offsetting the bank indices.

Data Layout Transformation

When the wasted memory from padding is unacceptable (as it can lower occupancy), more advanced techniques can be used. Memory Swizzling is a technique that uses bitwise operations (e.g., XOR) to remap logical array indices to physical memory locations in a way that breaks up conflicting access patterns. It is more complex to implement but avoids memory waste.

Algorithmic Redesign

Sometimes the access pattern itself can be changed. For example, instead of all threads processing column j, the work can be rearranged so they access data along a diagonal, which may be conflict-free.

The existence of bank conflicts reveals that on a GPU, algorithms and data structures cannot be designed in isolation from the hardware’s microarchitecture to achieve peak performance. A theoretically efficient algorithm can be slow in practice if its memory access patterns are hostile to the underlying banked memory structure. An algorithm designer working in the abstract might assume all accesses to a fast cache are equally fast. The CUDA shared memory model violates this assumption. The performance of an access is not just a property of the memory, but a property of the collective access pattern of the 32 threads in a warp. This forces the programmer to think not just in threads, but in warps. When designing a data layout in shared memory, one must ask, “if 32 threads with threadIdx.x from 0 to 31 execute this instruction, which banks will they hit?” Thie means data structures must be designed with “bank-awareness”. The decision to pad an array is a direct modification of a data structure solely to satisfy a hardware constraint. This tight coupling between algorithm, data structure, and microarchitecture is a defining characteristic of high-performance GPU programming.

Key Programming Patterns

Pattern 1: Tiled Matrix Multiplication (Data Reuse)

• The Problem with a Naive Approach: A simple kernel where each thread computes one element of the output matrix C results in massive, redundant global memory reads. Each element of A is read width_B times, and each element of B is read height_A times.
• The Tiled Solution: The key idea is to partition the large matrices into smaller sub-matrices (tiles) that can fit into shared memory.
• The “Load-Sync-Compute” Pattern:
  • Load: Each thread in a block cooperates to load one tile of A and one tile of B from global memory into two __shared__ arrays. This is done in a coalesced fashion for efficiency.
  • Sync: __syncthreads() is called to ensure the entire tile is loaded before computation begins.
  • Compute: Each thread now performs a partial dot product using data that is in fast shared memory. All threads in the block reuse the same data from the shared memory tiles.
  • Loop: The process repeats, loading the next pair of tiles and accumulating the result, until the full matrix product is computed.

Pattern 2: Efficient Parallel Reduction (Inter-thread Communication)

• The Goal: To aggregate a large array of values into a single value (e.g., finding a sum, max, or min).
• Shared Memory Implementation:
  1. Each thread loads one or more elements from global memory into a shared memory array.
  2. In a synchronized loop, threads combine values. For example, thread i adds the value from thread i + s to its own, where the stride s is halved each iteration (s = blockDim.x / 2; s > 0; s /= 2).
  3. __syncthreads() within the loop is necessary to ensure that reads in one iteration see the writes from the previous iteration.
• Modern Alternative: Warp-Level Primitives: For modern GPUs, intra-wrap reductions can be done much more efficiently using warp shuffle instructions like __shfl_down_sync(). These allow threads within a warp to exchange register data directly without using shared memory at all. This is faster and reduces shared memory pressure.
• The Hybrid Approach: The most optimal reductions often use a hybrid strategy: perform a fast reduction within each warp using shuffle instructions, have the warp’s leader” (thread 0) write its partial result to shared memory, and then perform a final, small reduction on those partial results in shared memory.

Pattern 3: Staging for Coalesced Access (Data Reordering)

• The Problem: Sometimes an algorithm needs to read from global memory in a non-coalesced pattern (e.g., accessing columns of a row-major matrix) but write in a coalesced way.
• The Solution: Use shared memory as a temporary staging area.
  1. Threads perform a non-coalesced read from global memory into a shared memory tile. This is slow but unavoidable.
  2. Call __syncthreads().
  3. Threads then read the data from shared memory in a different pattern and write it to global memory in a fully coalesced fashion. The cost of the initial non-coalesced read is amortized by the benefit of the coalesced write.

The different usage patterns of shared memory —data reuse, communication, and reordering— all contribute to a single, unifying purpose: improving the arithmetic intensity of a kernel. Arithmetic intensity is the ratio of floating-point operations to memory operations. By using shared memory to reduce the number of slow global memory accesses, all of these patterns increase this ratio, ensuring the GPU spends more of its time computing and less time waiting for data. In matrix multiplication, shared memory enables data reuse: instead of N global reads for an element, we do one. The number of floating-point operations is the same, but memory operations are reduced. Arithmetic intensity increases. In reduction, shared memory facilitates communication, turning many global reads/writes into fast shared memory operations. Floating-point operations are constant, memory operations are reduced. Arithmetic intensity increases. Shared memory is therefore not just a cache; it is a strategic tool to fundamentally restructure a kernel’s relationship with the memory system to favor computation over data movement. This is the key to unlocking GPU performance.

Advanced Performance Tuning and Architectural Considerations

The Occupancy vs. Performance Trade-off

• Defining Occupancy: The ratio of active warps on a Streaming Multiprocessor (SM) to the maximum number of warps that SM can support. High occupancy is crucial for hiding memory and instruction latency.
• Shared Memory as a Limiting Factor: The total amount of shared memory on an SM is a fixed resource. If each block requests a large amount of shared memory, it reduces the number of blocks that can reside on the SM simultaneously, thus lowering occupancy.
• Striking a Balance: This creates a critical trade-off. Using more shared memory might improve the performance of an individual block (e.g., by allowing for larger tiles), but it can hurt overall kernel performance by reducing occupancy and thus the ability to hide latency. The goal is not to maximize shared memory usage or occupancy in isolation, but to find the “sweet spot” that yields the best performance.
• Using the Occupancy Calculator: Advise readers to use NVIDIA’s Nsight tools and the occupancy calculator to model how block size, register usage, and shared memory usage effect the number of active blocks.

Register Spilling: The Hidden Performance Killer

• What is Register Spilling? When a kernel requires more variables per thread than the number of physical registers available, the compiler “spills” the excess variables into local memory.
• The Performance Impact: This is catastrophic for performance, because local memory is not a special on-chip memory but simply an abstraction for thread-private data in slow, off-chip global memory (DRAM). A kernel that spills registers is effectively performing many high-latency global memory operations where it should be performing zero-latency register operations.
• Shared Memory as a “Better Spill Location”: A key advanced optimization is to recognize when register pressure is high and proactively use shared memory as a programmer-controlled cache for variables that would otherwise be spilled. Manually loading/storing a variable from shared memory is much faster than letting the compiler automatically spill it to local memory.
• Controlling Register Count: Mention __launch_bound__ and the -maxregcount compiler flag as tools to explicitly manage the trade-off between register usage and occupancy.

Performance optimization in CUDA is a multi-dimensional problem of managing scarce resources; registers, shared memory, and threads. Optimizing for one resource in isolation can negatively impact another. For example, reducing register usage to increase occupancy might cause spills, which is worse. Increasing shared memory usage to improve data reuse might decrease occupancy. The expert programmer understands these interdependencies and uses profiling tools to identify the true bottleneck before applying an optimization. Occupancy is limited by both registers-per-thread and shared-memory-per-block. An attempt to improve performance by increasing the shared memory tile size might push the kernel over a resource “cliff”, causing a sharp drop in concurrent blocks that hurts performance more than the larger tile helps. This shows that performance tuning is not a checklist but a holistic analysis. The question is not “How can I use more shared memory?” but “What is limiting my performance right now?” If it’s global memory latency, then using more shared memory for better tiling is likely the answer. If it’s low occupancy in a latency-bound kernel, then reducing both register and shared memory usage per block might be the answer. The interconnectedness is the essence of advanced CUDA optimization.

A Summary of Core Principles

• Profile First
• Maximize Data Reuse
• Mind the Banks
• Synchronize Correctly
• Check Occupancy
• Beware of Spills: Check compiler output (-Xptxas -v) for register spilling. If it occurs, it is a top-priority problem to fix, potentially by using shared memory as a manual cache.

Mastering its use is the key to unlocking the true potential of the GPU. The ability to minimize data movement and maximize computation rests on the judicious management of shared memory, an essential skill in every domain of high-performance computing.