This article explores the computational essence of SIMD (Single Instruction Multiple Data) and SIMT (Single Instruction Multiple Thread) in the context of AI chips, detailing their relationship and how NVIDIA CUDA implements these concepts at a low level.

The Core of SIMD Computation

SIMD applies the same operation to multiple data elements simultaneously, leveraging data-level parallelism rather than concurrency. It involves multiple computations but only a single process in execution. With a single command, SIMD acts on multiple data values, often used to boost CPU capability for data parallelism. This requires wider ALU (Arithmetic Logic Unit) units and simpler control logic.

SIMD remains single-threaded, not multi-threaded. The hardware needs exactly one compute core but processes multiple data points at once. This is distinct from GPU multi-threading. The essence of SIMD is hardware-level parallel operation on multiple data with identical instructions.

Consider adding two vectors: to multiply vector A and vector B elementwise, each with four elements, to produce C:

C[0:3] = A[0:3] × B[0:3]

For SIMD to compute multiple elements in a single multiplication cycle, the hardware must increase the number of ALU units and data paths. A control unit dispatches data to multiple processing units, boosting throughput in a single clock cycle.

Practical SIMD has trade-offs:

• Disadvantage: SIMD uses one thread that processes many data elements concurrently, but due to ALU width limits, data types, formats, and sizes must be strictly aligned.
• Advantage: Performance improves by leveraging memory bandwidth; multiple data points can be read and written simultaneously. For C[0:3] = A[0:3] × B[0:3], code size and execution cycles shrink to one-quarter of the original.

Without SIMD, the computer executes separate instructions per element:

LD t1, B[i]
LD t2, C[i]
ADD t3, t1, t2
ST A[i], t3

LD t1, B[i+1]
LD t2, C[i+1]
ADD t3, t1, t2
ST A[i+1], t3

LD t1, B[i+2]
LD t2, C[i+2]
ADD t3, t1, t2
ST A[i+2], t3

LD t1, B[i+3]
LD t2, C[i+3]
ADD t3, t1, t2
ST A[i+3], t3

With SIMD, a single store instruction processes four elements:

v1 = LD B[i], 4
v2 = LD C[i], 4
v3 = ADD v1, v2, 4
ST A[i], 4, v3

Intel introduced SIMD with MMX (Multimedia Extensions) in 1996; ARM added NEON for its Cortex architecture. NEON uses 128-bit registers (16 total, also usable as 32 64-bit registers) that hold aligned same-type vectors. For instance, one register (128 bits) can store four 32-bit elements. The following code multiplies two such vectors:

// Perform multiplication on four elements simultaneously
C[0:3] = A[0:3] * B[0:3]

// Register s15 holds vector B
vldmia.32 r0!, {s15}
// Register s14 holds vector A
vldmia.32 r1!, {s14}
// s15 = s15 * s14
vmul.f32 s15, s15, s14
// Store result
vstmia.32 r2!, {s15}

MMX (MultiMedia eXtensions) was Intel's 1996 SIMD extension for identical operations on multiple data, accelerating multimedia and image processing. Later SSE and AVX expanded this. ARM NEON, introduced in 2004, includes load/store, integer, and floating-point SIMD instructions for parallel compute efficiency.

Thus, SIMD fundamentally changes the count of ALU units and data pathways, exposing more instructions to higher layers. Programmers rarely manipulate SIMD instructions directly.

The Core of SIMT Computation

SIMT (Single Instruction Multiple Threads) is a concept introduced by NVIDIA for GPUs. Like SIMD, both broadcast the same instruction to multiple execution units for data parallelism. The key difference is that SIMD mandates all vector elements execute synchronously within a single thread, whereas SIMT allows multiple threads within a warp to execute independently.

SIMT resembles CPU multithreading: multiple compute cores exist, each with its own register file (RF) and ALU, but no separate instruction cache, decoder, or program counter. Instructions are broadcast from a unified cache to all SIMT cores. While SIMT cores are independent and operate on different data with identical instructions, SIMD shares a single ALU.

For the vector multiplication example C[0:3] = A[0:3] × B[0:3], SIMT assigns one thread per element. Each thread performs a single multiplication, and when all threads finish in parallel, the entire multiplication is complete.

Hardware-wise, SIMT uses a multi-core cluster (SIMT Core Cluster). A CPU loads a kernel into this cluster. Each SIMT core has its own RF, ALU, and data cache but shares a single program counter and instruction decoder. Instructions are broadcast to all cores for execution. A GPU comprises multiple such clusters, each containing multiple SIMT cores; a SIMT core holds multiple thread blocks.

GPU SIMT can be viewed as a special SIMD structure with a pipeline split into a SIMT front-end and a SIMD back-end. The pipeline includes three scheduling loops: Fetch, Issue, and Register Access.

• Fetch loop: Fetch, I-Cache, Decode, I-Buffer.
• Issue loop: I-Buffer, Score Board, Issue, SIMT-Stack.
• Register Access loop: Operand Collector, ALU, Memory.

These three loops form the SIMT hardware core pipeline. The fetching places instructions into a stack; at runtime, the stack distributes threads to ALUs, executing in SIMD fashion. SIMT handles thread control at the front-end.

Key similarities and differences between SIMD and SIMT:

• Both are based on "single instruction, multiple data."
• SIMT appears multi-threaded but ultimately executes single-threaded at the hardware level, using multiple cores for parallelism.
• SIMT is more flexible, allowing each thread to independently address data, while SIMD requires contiguous, aligned, same-type data.

Hence, SIMT generalizes SIMD, offering more flexible and developer-friendly programming.

NVIDIA CUDA Implementation

GPU thread hierarchy partitions images: a grid represents the total task, subdivided into blocks. Each block contains many threads. Threads execute independently within blocks, processing pixel data, and can share data and synchronize.

CUDA's parallel programming model follows SPMD (Single Program Multiple Data), distinct from SIMT. In CUDA, a grid is an array of thread blocks mapped to streaming multiprocessors (SMs). Each block contains multiple warps; block size affects kernel performance. The smallest execution unit is a thread. Threads in a block share a fast, synchronized memory.

Unlike SIMD, SIMT lets you write thread-level parallel code for scalar threads or cooperative data-parallel code. To ensure correctness, developers can often ignore SIMT behavior; branching within a warp is rarely needed. Maintaining simple code yields significant hardware parallelism. All threads in a block run the same kernel; each thread has a unique index (threadIdx.x) for memory addresses and control decisions.

Multiple thread blocks form a grid. A block is the basic scheduling unit for an SM. SM is hardware; a block is an abstraction. Because blocks and threads span multiple dimensions, thread indexing uses block index (blockIdx) and intra-block thread index (threadIdx).

blockIdx.x and blockDim.x access the block's x-index and its x-dimension size in CUDA. blockIdx and blockDim are 3D vectors.

Hardware-software mapping: threads map to CUDA cores; thread blocks are assigned to SMs. SM maintains block and thread IDs, schedules execution. Each warp contains 32 threads executing in SIMD. A block is scheduled on one SM via warps and stays until kernel completion. An SM can hold multiple blocks concurrently.

AI framework development: define a neural network, write code using the framework, which builds a forward graph and a reverse graph via automatic differentiation. Matrix multiplication is a key operator. Here’s a CUDA example for C = A × B:

#include <stdio.h>

#define N 4

__global__ void matMul(int *a, int *b, int *c) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    int sum = 0;
    for (int k = 0; k < N; ++k) {
        sum += a[row * N + k] * b[k * N + col];
    }
    c[row * N + col] = sum;
}

int main() {
    int a[N][N], b[N][N], c[N][N];
    int *da, *db, *dc;

    cudaMalloc((void**)&da, N * N * sizeof(int));
    cudaMalloc((void**)&db, N * N * sizeof(int));
    cudaMalloc((void**)&dc, N * N * sizeof(int));

    for (int i = 0; i < N; ++i)
        for (int j = 0; j < N; ++j) {
            a[i][j] = i * N + j;
            b[i][j] = j * N + i;
        }

    cudaMemcpy(da, a, N * N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(db, b, N * N * sizeof(int), cudaMemcpyHostToDevice);

    dim3 block(2, 2);
    dim3 grid(N/block.x, N/block.y);

    matMul<<<grid, block>>>(da, db, dc);

    cudaMemcpy(c, dc, N * N * sizeof(int), cudaMemcpyDeviceToHost);

    for (int i = 0; i < N; ++i) {
        for (int j = 0; j < N; ++j)
            printf("%d ", c[i][j]);
        printf("\n");
    }

    cudaFree(da); cudaFree(db); cudaFree(dc);
    return 0;
}

Programming vs. Hardware Execution

Programming Model is an abstraction enabling programmers to organize and control code without considering hardware details. It includes language, data structures, algorithms, and concurrency mechanisms.

Hardware Execution Model describes how hardware executes code, including architecture, instruction set, registers, memory hierarchy, caches, and parallelism strategies. It determines actual execution (e.g., SIMD and SIMT).

• Difference: Programming model focuses on logical structure and behavior from the programmer’s perspective; hardware model focuses on low-level execution details.
• Link: The programming model defines behavior, and the compiler maps it to the hardware model. Understanding both helps optimize performance and resource utilization.

A compiler transforms the programming model into the hardware execution model, making the two conceptually distinct.