In modern parallel computing architectures, GPUs (Graphics Processing Units) are no longer just “graphics chips” used for rendering, but rather powerful “multi-core processors” with exceptional general-purpose parallel computing capabilities. Among them, NVIDIA’s CUDA programming framework employs the SIMT (Single Instruction Multiple Threads) model allows us to easily write parallel programs with a “thread” mindset. However, at the hardware execution level, these massive threads are actually organized into “Warps” for SIMD (Single Instruction Multiple Data) lock-step execution. Therefore, it can be said that SIMT is essentially a SIMD based on Warp. This article will reveal the principles and significance behind this.
1. Review: What is SIMD?
SIMD stands for Single Instruction Multiple Data. Traditional vector instruction sets in CPUs (such as Intel SSE, AVX, ARM NEON, etc.) are typical representatives of SIMD. It refers to a single instruction processing multiple data at the same time. For instance, in a 128-bit vector register, we can perform addition/multiplication on 4 32-bit floating-point numbers simultaneously, significantly improving data processing throughput.
The core idea of SIMD is:A stream of instructions + Multiple data operation units = Parallel processing of multiple data within the same clock cycle..
2. The Warp Mechanism in NVIDIA GPUs
To understand the SIMT employed by NVIDIA GPUs, it is essential to first clarify how NVIDIA GPUs schedule and execute those “threads”.
1. Size of a Warp
In NVIDIA GPUs, a Warp typically consists of 32 threads. These 32 threads are treated as a single execution unit by the hardware, meaning they share instruction fetch values (PC, Program Counter) and decoding stages, executing the same instruction in the same cycle.
2. Warp Scheduling
In an NVIDIA GPU, a streaming multiprocessor (SM) may have multiple Warps existing simultaneously, and there are multiple Warp schedulers. The hardware will schedule ready Warps in turn to fully utilize execution units and hide memory access latency.
Different Warps can execute different code segments within the same kernel or execute completely different instructions at the same time. This independence allows NVIDIA GPUs to handle different tasks or divisions simultaneously.
3. Thread Divergence
Within the same Warp, if there is a conditional branch (if-else), and different threads need to execute different branch paths, NVIDIA GPUs will use a “masking + time-slicing” approach to complete each branch path separately, which may lead to performance degradation (the so-called thread divergence).
This is similar to how CPU SIMD handles data dependencies or branches, except that NVIDIA GPUs present a thread-based programming model externally, providing developers with more flexible usage.
3. SIMT: A Combination of Thread-Level Abstraction and Hardware SIMD
1. What is SIMT (Single Instruction Multiple Threads)?
From a programming model perspective, SIMT allows developers to think about parallelism in terms of “threads”: we can launch thousands, or even millions, of threads in CUDA, each with its own independent thread ID, register context, and stack space, which seems very flexible.
However, at the hardware level, these threads are allocated into several Warps, and Warps are the smallest units that truly execute “the same instruction, processing multiple data”.
2. SIMT is essentially a “Warp-based SIMD”
Warp is the SIMD container: In a Warp, 32 threads execute the same instruction (such as addition, multiplication, read/write, etc.), but operate on different registers or address indices, which is the same “lock-step” concept as traditional CPU SIMD.
Flexibility brought by thread abstraction: Unlike the “vector register” approach of CPU SIMD, NVIDIA GPUs package each “element” as a thread managed by the programming model. The hardware only needs to package these threads into Warps for SIMD execution during the execution phase, balancing flexibility and high throughput.
3. Example
When we write `<<>> kernel` in CUDA to launch multiple threads, each thread executes the same code but uses different `threadIdx`, `blockIdx`, etc., to distinguish their respective work. At the hardware level, suppose we have many blocks, each with 256 threads, these threads are divided into several Warps (each Warp has 32 threads). The “lock-step execution” occurs at the Warp level.
4. Why did NVIDIA design GPUs as SIMT instead of direct vectorization?
1. Programming Convenience
Compared to explicit vectorization on CPUs (which requires writing SSE, AVX instructions or relying on the compiler for automatic vectorization), the thread abstraction of SIMT is more intuitive: for a large number of parallel tasks, developers only need to focus on “how many threads I have and what they need to do,” without worrying about specific vector registers and data packing methods.
2. Scalability and General-Purpose Computing
GPUs typically handle massive and highly parallel computing tasks (such as rendering pixels, scientific computing, deep learning). The SIMT model is easily scalable, allowing hardware to flexibly schedule more Warps for parallel execution, significantly enhancing throughput.
3. Branching and Irregular Parallel Processing
Although Warps still execute in SIMD, through thread masking and Warp scheduling, NVIDIA GPUs can handle some irregular parallel scenarios. This means that while efficiency may be affected, NVIDIA GPUs still maintain good overall performance for large datasets and complex computations.
5. Conclusion and Outlook
Essentially, the SIMT model adopted by NVIDIA GPUs gives developers the illusion of “thread-level” control at the programming level, while the hardware still uses Warp-level SIMD (lock-step execution) to achieve high parallelism and high throughput..
This design balances usability with execution efficiency at the hardware level. This is precisely why NVIDIA’s CUDA has achieved great success in the field of general-purpose computing (GPGPU).
As GPUs continue to evolve and other heterogeneous computing architectures emerge (such as dedicated AI accelerators like NPUs), how to balance “efficient hardware execution modes” with “ease of software programming models” remains a hot research direction in the industry. Understanding the intrinsic connection between SIMT and SIMD also provides a solid foundation for us to explore parallel computing architectures and write higher-performance parallel programs.
【Final Note】
In daily GPU parallel program development, when we write `threadIdx.x` in CUDA code, we are using the SIMT abstraction provided by NVIDIA GPUs for parallel computing; however, internally, it is actually batches of 32 threads (a Warp) executing “SIMD (the same instruction, multiple data)”. Understanding this allows us to better comprehend and optimize GPU program performance, organize data and threads more rationally, avoid unnecessary thread divergence, and unleash the powerful capabilities of GPU parallel computing.