Single Instruction, Multiple Threads (SIMT)

The warp (NVIDIA) or wavefront (AMD) is the fundamental unit of execution. A single instruction is fetched and broadcast to all 32 (or 64) threads within a warp simultaneously. Threads execute the same instruction in lockstep but operate on different data elements from their private registers. This design minimizes instruction fetch and decode overhead across thousands of concurrent threads.

When threads within a warp encounter a conditional branch (e.g., an if/else statement), warp divergence occurs. The hardware serializes execution: it first executes the threads that take the 'true' path while masking off the others, then executes the 'false' path threads. This serialization can cause significant performance penalties. Optimized kernels minimize divergence by ensuring threads in a warp follow coherent control flow paths.

Unlike pure SIMD, each thread in the SIMT model has its own:

• Program counter (logically, managed per warp for lockstep execution)
• Register file (private, fast memory)
• Predication mask (for handling divergence)
• Thread ID (unique identifier for data addressing) This independent state allows threads to execute unique memory addresses and data-dependent operations, providing the flexibility of MIMD (Multiple Instruction, Multiple Data) with the efficiency of SIMD for instruction dispatch.

SIMT efficiency is critically dependent on memory coalescing. When threads in a warp access contiguous, aligned blocks of global memory, their accesses are combined into a minimal number of wide, high-bandwidth transactions. Non-coalesced access patterns (e.g., random or strided) result in many small, inefficient transactions, starving the compute units of data. Kernel design focuses on arranging data and thread indexing to maximize coalesced access.

Occupancy is the ratio of active warps on a Streaming Multiprocessor (SM) to its maximum supported warps. It is limited by three key resources:

• Register usage per thread (spilling to slower memory reduces performance)
• Shared memory usage per thread block
• Thread block size and grid configuration High occupancy helps hide memory latency by providing the hardware with more eligible warps to execute while others are stalled on memory accesses. It is a primary tuning target for kernel optimization.

• SIMD (Single Instruction, Multiple Data): A single instruction operates on a packed vector of data elements within a single thread. The programmer or compiler explicitly manages vector registers.
• SPMD (Single Program, Multiple Data): A broader parallel programming model where multiple independent processor instances run the same program on different data. SIMT is a hardware implementation of SPMD that uses warp-based lockstep execution for efficiency.
• SIMT abstracts the vector hardware, presenting a scalar programming model to the developer while the hardware handles the vectorization across a warp.

Warp divergence is the most significant performance penalty in SIMT architectures. It occurs when threads within a warp (typically 32 threads) take different execution paths due to conditional statements (e.g., if/else). The hardware must serialize execution, executing all paths taken by any thread in the warp while disabling threads not on that path. This drastically reduces parallel efficiency.

• Mitigation: Structure algorithms to have branch coherence, where all threads in a warp follow the same path. Use techniques like branch predication or computed goto to avoid conditionals. For data-dependent operations, consider sorting data by condition before processing to maximize warp coherence.

Memory coalescing is the optimization where concurrent memory accesses from threads in a warp are combined into as few transactions as possible. Modern GPUs have wide memory buses (e.g., 32 bytes). If consecutive threads access consecutive 4-byte words, their accesses can be coalesced into a single 128-byte transaction.

• Poor Pattern: Threads accessing scattered or non-sequential addresses cause multiple small transactions, wasting bandwidth.
• Best Practice: Ensure kernel memory access patterns are stride-1 and aligned. Use data layouts like Structure of Arrays (SoA) instead of Array of Structures (AoS) for SIMT access. Properly sized and aligned memory tiles in shared memory can facilitate coalesced global memory loads and stores.

Occupancy is the ratio of active warps on a Streaming Multiprocessor (SM) to the maximum possible active warps. It is limited by three key resources: threads per SM, shared memory per block, and register usage per thread.

• High occupancy helps hide latency by having more warps ready to execute when others stall.
• Trade-offs: Excessive register usage or large shared memory allocations per thread block will limit occupancy. Use compiler directives (e.g., __launch_bounds__) or manual tuning to balance. Low occupancy can still yield high performance if the kernel is compute-bound and has sufficient instruction-level parallelism (ILP) to hide latency.

SIMT hardware handles fine-grained control flow via an active mask. When divergence occurs, the mask is updated, and instructions are issued only to active threads. This is managed automatically but incurs cost.

• Predication: A compiler optimization that converts short conditional blocks into predicated (conditional) instructions. Both sides of the branch are executed, but the results are committed only for threads where the condition is true. This avoids actual divergence but increases instruction count.
• Use Case: Effective for very short if/else blocks where the branch body has few instructions. For complex branching, restructuring the algorithm is preferable.

Within a thread block, threads are synchronized using the __syncthreads() barrier. This is a critical point where all threads in the block must reach before any can proceed. Misuse causes deadlocks or performance issues.

• Warp-Level Primitives: For operations within a warp, use faster warp shuffle instructions (__shfl_sync) or warp vote functions (__any_sync, __all_sync). These avoid the overhead of a full block barrier.
• Divergence Danger: Calling __syncthreads() in divergent code is illegal and will cause deadlock. The barrier must be encountered by all threads in the block in the same execution path.

To keep the deep pipelines of GPU cores full, kernels must expose sufficient instruction-level parallelism (ILP). This allows the hardware to schedule independent instructions from a single thread to hide latency.

• Increasing ILP: Manually unroll loops, use multiple independent arithmetic operations in sequence, and reduce data dependencies within a thread.
• Latency Hiding: High warp occupancy is the primary method for hiding long-latency operations (e.g., global memory accesses). If occupancy is low, increasing ILP per thread becomes essential to maintain performance by keeping the arithmetic logic units busy while waiting for data.

A compiler optimization technique that merges multiple, separate computational kernels into a single, larger kernel. This reduces the overhead of repeated kernel launches and eliminates costly intermediate data transfers between global memory and on-chip registers or caches. For example, fusing an activation function kernel with a preceding convolution kernel saves a full round-trip to memory.

A performance penalty in SIMT architectures that occurs when threads within a single warp (a group of 32 threads in NVIDIA GPUs) follow different execution paths due to conditional branching (e.g., if/else statements). The hardware must serialize the execution of all divergent paths, significantly reducing parallel efficiency. Minimizing warp divergence is a key optimization goal.

A measure of the number of independent instructions that can be executed simultaneously by a processor's functional units within a single thread. It is exploited through:

• Superscalar execution (issuing multiple instructions per clock cycle).
• Software pipelining (overlapping instructions from different loop iterations).
• Loop unrolling (reducing loop control overhead). ILP operates within a thread, whereas SIMT exploits parallelism across many threads.

A critical optimization for parallel architectures where concurrent memory accesses from multiple threads in a warp are combined into a single, wider transaction. This maximizes effective memory bandwidth by reducing the total number of required memory operations. It requires threads to access contiguous, aligned memory addresses. Poor coalescing leads to serialized, low-bandwidth access.

A key performance metric on GPUs/NPUs that measures the ratio of active warps resident on a Streaming Multiprocessor (SM) to the maximum number of warps it can theoretically support. High occupancy helps hide instruction and memory latency. Occupancy is limited by:

• Register usage per thread.
• Shared memory usage per thread block.
• Thread block size.

A class of compiler transformations applied to nested loops to maximize data locality, parallelism, and hardware resource utilization. Key transformations include:

• Loop tiling for cache/register locality.
• Loop interchange for stride-1 access patterns.
• Loop fusion to reduce intermediate memory traffic.
• Loop unrolling to increase ILP. These are foundational for generating efficient code for compute-intensive kernels on accelerators.