Shared Memory

Shared memory is allocated per thread block and is only visible to the threads within that block. This provides a fast, private workspace for cooperative algorithms. Threads in different blocks cannot directly access each other's shared memory, enforcing a natural data isolation that simplifies synchronization.

• Example: A block of threads processing a tile of a matrix can load the tile into shared memory, perform all calculations there, and then write the final result back to global memory.

Unlike the automatic L1/L2 hardware caches, shared memory is explicitly managed by the programmer. You control what data is placed into it and when. This requires more effort but allows for predictable, high-performance data reuse patterns that are optimal for a specific algorithm.

• Contrast: The hardware cache is transparent but subject to unpredictable eviction. Shared memory provides deterministic, algorithm-aware caching.

Shared memory resides on-chip, physically close to the streaming multiprocessor (SM) cores. This gives it latency comparable to registers and bandwidth far exceeding global (device) memory. Accessing shared memory is typically orders of magnitude faster than accessing global DRAM.

• Performance Impact: Kernels that stage data in shared memory for repeated access can see throughput improvements of 10x or more compared to kernels that read repeatedly from global memory.

To achieve high bandwidth, shared memory is divided into equally sized, independently accessible modules called memory banks. Concurrent accesses to different banks can be serviced simultaneously, but accesses to the same bank cause a bank conflict, serializing the operations and reducing throughput.

• Optimal Pattern: To avoid conflicts, ensure threads in a warp access addresses that map to different banks (e.g., stride-1 access for 32-bit words).

Shared memory acts as a rendezvous point, requiring explicit synchronization to ensure data written by some threads is visible to others. The __syncthreads() barrier is used to synchronize all threads in a block, guaranteeing memory consistency before proceeding.

• Critical Use Case: In a parallel reduction or scan algorithm, threads write partial results to shared memory, call __syncthreads(), and then read the results written by their peers for the next computation stage.

The amount of shared memory per streaming multiprocessor is a fixed hardware resource (e.g., 64KB or 128KB on modern NVIDIA GPUs). This capacity is statically allocated per thread block at kernel launch. Exceeding this limit will prevent the kernel from launching.

• Trade-off: The programmer must balance shared memory usage against the number of concurrent thread blocks that can reside on an SM (occupancy). Using more shared memory per block can reduce occupancy and vice-versa.

The classic use case for shared memory is in optimizing General Matrix Multiply (GEMM) operations, the foundation of neural network layers. The algorithm proceeds in tiles:

• Global memory loads a tile of matrix A and a tile of matrix B into shared memory.
• All threads in the block cooperatively compute partial results from these fast, on-chip tiles.
• This process repeats, accumulating results, minimizing costly accesses to slower global memory. This tiling strategy is fundamental to high-performance linear algebra libraries like cuBLAS and is directly analogous to the blocked computation in a transformer's attention or feed-forward layers.

Shared memory enables efficient parallel reduction algorithms, essential for operations like computing a sum, maximum, or softmax normalization across a dataset.

• Threads first load data from global memory into shared memory.
• A tree-based reduction is then performed within the block, where threads at each step combine partial results stored in shared memory.
• This avoids repeated global memory accesses, drastically reducing latency. For example, computing the sum of logits for a softmax function across a vocabulary dimension heavily relies on this pattern.

In CNNs, applying a filter (kernel) to an input image or feature map involves reusing input pixels for multiple output computations. Shared memory is used to stage a region of the input image:

• A block of threads loads an input tile, plus a surrounding halo region (for the filter's overlap), from global memory into shared memory.
• Each thread then repeatedly reads from this fast, on-chip cache to compute its output pixel, reusing the loaded input data for multiple convolutions.
• This eliminates redundant global memory fetches, which is critical for the performance of vision models in real-time inference.

While the full KV cache resides in global memory, shared memory acts as a critical staging area within the attention computation kernel for a single layer and head.

• During the attention score calculation (Q @ K^T), tiles of the Key cache are loaded into shared memory to be reused by all queries in a block.
• Similarly, during the output calculation (scores @ V), tiles of the Value cache are loaded.
• This tiled approach via shared memory is essential for achieving high memory bandwidth utilization when computing self-attention for long sequences, a bottleneck in transformer inference.

High-performance sorting algorithms on the GPU, like radix sort or merge sort, use shared memory as a temporary workspace for data being rearranged within a block.

• Bitonic sort and other parallel sorting networks are implemented within a thread block using shared memory for compare-and-swap operations.
• Scan (prefix sum) operations, another fundamental primitive, are implemented using efficient shared memory algorithms before writing final results back to global memory. These primitives are building blocks for data preprocessing and reordering tasks in ML pipelines.

To avoid kernel launch overhead and intermediate global memory writes, ML engineers write custom CUDA kernels that fuse multiple operations. Shared memory is the glue in these kernels.

• Example: A kernel that performs layer normalization immediately after a matrix multiply can hold the intermediate results in shared memory, compute the mean and variance across threads, and apply the normalization—all without a round-trip to global memory.
• Example: Fusing an activation function (like GELU or SiLU) with a preceding linear layer. This operator fusion, enabled by shared memory staging, is a key technique in inference engines like TensorRT and OpenAI's Triton to minimize latency.

Global memory is the large, high-bandwidth DRAM physically located on the GPU, accessible by all threads across all blocks. It serves as the primary workspace for kernels but has higher latency than on-chip memories.

• Contrast with Shared Memory: While shared memory is a small, software-managed cache for thread block cooperation, global memory is the main, large-capacity pool for all GPU data.
• Access Patterns are Critical: Performance depends heavily on achieving coalesced memory access, where consecutive threads in a warp access consecutive addresses, maximizing bandwidth.

The memory hierarchy organizes GPU memory into multiple levels with different capacities, latencies, and bandwidths, optimized for data locality.

• Levels (Fastest to Slowest): Registers → Shared Memory / L1 Cache → L2 Cache → Global Memory (HBM/GDDR) → Host Memory (DDR) → Storage (NVMe).
• Programming Implication: High-performance algorithms explicitly manage data movement up this hierarchy. A common pattern is to load data from global memory into shared memory, perform cooperative computations, then write results back.

Coalesced memory access is an optimal pattern where consecutive threads in a GPU warp access consecutive memory addresses in a single, aligned transaction.

• Why it Matters: It maximizes effective bandwidth from global memory, which is crucial for performance. Uncoalesced access can serialize transactions, reducing throughput by an order of magnitude.
• Relationship to Shared Memory: Shared memory is often used to enable coalescing. Threads can perform uncoalesced gathers from global memory into shared memory, then perform fast, bank-conflict-free accesses from shared memory for the core computation.

A bank conflict occurs in GPU shared memory when two or more threads within a warp attempt to access different data words stored in the same memory bank simultaneously, forcing serialized access.

• Shared Memory Structure: It is divided into equally sized memory banks (e.g., 32 banks). Concurrent accesses to different addresses in the same bank cause a conflict.
• Performance Impact: Severe bank conflicts can reduce shared memory bandwidth to a single access per cycle, negating its low-latency advantage. Avoiding conflicts through careful data layout (e.g., padding) is a key optimization.

Unified Virtual Memory is a memory management architecture that creates a single, contiguous virtual address space shared between a CPU and GPU, allowing both processors to access the same memory pointers.

• Contrast with Explicit Management: UVM abstracts the explicit cudaMemcpy calls required in traditional programming. The system handles page migration (demand paging) between host and device memory on access faults.
• Trade-off: Simplifies programming by enabling zero-copy access and oversubscription, but can introduce performance variability due to page fault overhead if data locality is poor.

L1 and L2 caches are hardware-managed, transparent caches on the GPU that store frequently accessed data from global memory to reduce latency.

• L1 Cache: Per-SM (Streaming Multiprocessor), often configurable with shared memory (e.g., 64KB partitioned as 48KB shared/16KB L1 or vice versa). It caches local memory accesses.
• L2 Cache: Unified cache shared across all SMs, backing global memory accesses.
• vs. Shared Memory: Caches are automatic and general-purpose. Shared memory is explicitly managed by the programmer, offering predictable, low-latency storage for specific data reuse patterns, acting as a software-controlled cache.