Stream Multiprocessor (SM)

CUDA Cores are the primary scalar execution units for integer and single-precision floating-point (FP32) arithmetic. An SM contains hundreds of these cores, organized into groups that share instruction dispatch. For example, an NVIDIA GA102 SM (Ampere architecture) contains 128 FP32 CUDA Cores. These cores are grouped with specialized pipelines for:

• Tensor Cores: Dedicated hardware for mixed-precision matrix multiply-accumulate operations (e.g., FP16, BF16, INT8, INT4), accelerating deep learning workloads.
• FP64 Cores: Double-precision floating-point units for scientific computing.
• Special Function Units (SFUs): Handle transcendental operations like sine, cosine, and square root.

The register file is the fastest memory in the SM hierarchy, providing private storage for every thread. It is a massive, partitioned SRAM structure. For instance, an NVIDIA A100 SM has a 256 KB register file. Each thread in a kernel is allocated a set of registers, which hold its private variables. The size of the register file directly limits the maximum number of concurrent threads (occupancy) an SM can support, as more registers per thread means fewer threads can be resident simultaneously. Efficient register usage is a key optimization target.

Load/Store (LD/ST) Units handle memory requests from threads. Each SM has multiple such units to process addresses and move data between the register file and the memory hierarchy. They interface with:

• Shared Memory and L1 Cache (on-chip).
• L2 Cache (unified, off-chip but on the GPU die).
• Global Memory (DRAM, e.g., GDDR6/HBM). These units are responsible for coalescing memory accesses from threads within a warp into the minimum number of transactions, which is essential for achieving peak memory bandwidth. Non-coalesced access is a major performance pitfall.

The SM includes several specialized, read-only caches to optimize instruction and constant data access:

• Instruction Cache (I-Cache): Stores kernel instructions for the warp schedulers, reducing fetch latency from global memory.
• Constant Cache: A dedicated cache for the constant memory space. Constant memory is optimized for broadcast scenarios where all threads in a warp read the same address. A single read from the constant cache can service an entire warp, making it highly efficient for storing kernel parameters and lookup tables.
• Texture Cache: A separate, hardware-managed cache optimized for spatial locality in texture fetches, which is also often used for general-purpose read-only data with 2D/3D locality.

The hardware mechanism within a GPU's Stream Multiprocessor (SM) that selects which warp of threads is issued to execution units. Its primary goal is to hide instruction and memory latency by keeping the cores busy. Key strategies include:

• Greedy-Then-Oldest (GTO): Prioritizes warps with ready instructions, then the oldest ready warp.
• Round-Robin: Cycles through warps in a fixed order.
• Scoreboarding: Tracks dependencies to issue warps only when operands are ready. Effective warp scheduling is critical for achieving high occupancy and hiding the latency of global memory accesses, which can be hundreds of cycles.

The execution model that defines how a Stream Multiprocessor (SM) processes threads. A single instruction is broadcast to a warp (typically 32 threads), and each thread executes it on its own private data. The SM handles control flow divergence (e.g., if/else statements) by:

• Masking: Deactivating threads that take different paths.
• Serializing: Executing each divergent path sequentially for the active threads. This model is a key differentiator from pure SIMD architectures, as it provides the programmer with a scalar thread view while the hardware manages the underlying vectorized execution.

A programmer-defined group of threads that are guaranteed to execute concurrently on a single Stream Multiprocessor (SM). Threads within a block can:

• Synchronize using the __syncthreads() barrier.
• Communicate via fast, programmer-managed shared memory (scratchpad).
• Be arranged in 1D, 2D, or 3D grids for intuitive mapping to data structures. The SM's resources—such as registers and shared memory—are allocated per thread block. The number of blocks an SM can host simultaneously is limited by these resources, directly impacting occupancy. Blocks are independent and do not synchronize with each other.

A key GPU performance metric representing the ratio of active warps resident on a Stream Multiprocessor (SM) to the maximum number of warps it can theoretically support. It is a measure of hardware resource utilization. High occupancy helps hide latency but does not guarantee peak performance. Occupancy is limited by:

• Register usage per thread.
• Shared memory usage per thread block.
• Thread block size and grid configuration.
• Hardware limits on threads, blocks, and warps per SM. Optimizers often balance occupancy with other factors like instruction-level parallelism (ILP) and memory coalescing.

Defines the formal rules for the order in which memory operations (loads and stores) from different threads become visible to each other within a shared memory system like a GPU. The weak consistency model of modern GPUs means:

• Memory operations from a single thread are not necessarily observed in program order by other threads.
• Synchronization primitives (e.g., __syncthreads(), atomics, memory fences) are required to enforce a specific order.
• This model allows for significant hardware optimizations (e.g., write buffering, cache bypassing) but places the burden of correctness on the programmer to insert appropriate synchronization.

Indivisible read-modify-write instructions that guarantee data integrity when multiple threads concurrently access the same memory location. On a Stream Multiprocessor (SM), common atomics include:

• Atomic Add, Sub, Min, Max
• Atomic Compare-and-Swap (CAS)
• Atomic Exchange These operations are implemented in the SM's memory subsystem (e.g., in the L1 cache or shared memory) and serialize access to the target address. While essential for correctness in parallel algorithms (e.g., histogramming, reduction), overuse can lead to severe performance degradation due to serialization, known as contention.