Under the Hood of GPU Compute: What Happens When You Run a CUDA Kernel?

Introduction: The Illusion of Instant Execution

To the modern software engineer, running a CUDA kernel feels like invoking a standard C++ function with a peculiar execution configuration syntax: my_kernel<<<grid, block>>>(args). Within microseconds, millions of parallel threads execute on billions of transistors, returning computed tensors for deep learning models or complex physical simulations.

But beneath this simple abstraction lies an incredibly complex choreography of software drivers, PCI Express transactions, hardware schedulers, and high-bandwidth memory controllers. What actually happens at the hardware level when you launch a CUDA kernel? How does a single line of host code transform into physical electrical signals across thousands of arithmetic logic units (ALUs)?

To write truly optimized high-performance code—whether you are tuning LLM inference engines, writing custom PyTorch operators, or developing high-fidelity graphics pipelines—you must understand the life cycle of a CUDA kernel from the host CPU down to the silicon of the Streaming Multiprocessor (SM).

The Host-to-Device Bridge: Launching the Kernel

The journey begins on the Host (the CPU). When your application invokes a CUDA kernel, it does not immediately execute on the Device (the GPU). Instead, it undergoes a multi-stage translation and queueing process.

1. Compilation and PTX Translation

During compilation via nvcc, your CUDA C++ code is split. The host code is compiled by a standard host compiler (like GCC or MSVC), while the device code (the kernel) is compiled into an intermediate assembly-like language called PTX (Parallel Thread Execution). At runtime, or during compilation, this PTX is further compiled into SASS (Source Assembly), the machine code specific to the target GPU architecture (e.g., Ampere, Ada Lovelace, or Hopper).

2. The CUDA Driver and Command Queues

When the host executes the <<<...>>> launch call, it interacts with the CUDA Runtime and Driver APIs. The driver translates the launch parameters (grid dimensions, block dimensions, dynamic shared memory size, and stream ID) into a command packet.

This packet is pushed into a hardware command queue (associated with a specific cudaStream_t). The GPU’s host interface—typically communicating via PCIe—reads from these queues using Direct Memory Access (DMA). Crucially, the CPU does not wait for the GPU to finish executing the kernel; the kernel launch is asynchronous. The CPU simply pushes the command onto the ring buffer and continues its execution unless an explicit synchronization barrier (like cudaDeviceSynchronize()) is met.

Inside the Silicon: Hardware Scheduling and Thread Blocks

Once the command packet crosses the PCIe bus and enters the GPU, the physical silicon takes complete control.

+-------------------------------------------------------------+
|                        GPU Chip                             |
|                                                             |
|  +-------------------------------------------------------+  |
|  |                  GigaThread Engine                    |  |
|  +-------------------------------------------------------+  |
|                             |                               |
|       +---------------------+---------------------+         |
|       |                                           |         |
|       v                                           v         |
|  +-------------------------+             +-------------------------+ |
|  | Streaming Multiprocessor|             | Streaming Multiprocessor| |
|  |          (SM)           |             |          (SM)           | |
|  |  +-------------------+  |             |  +-------------------+  | |
|  |  |  Warp Scheduler   |  |             |  |  Warp Scheduler   |  | |
|  |  +-------------------+  |             |  +-------------------+  | |
|  |  +-------------------+  |             |  +-------------------+  | |
|  |  |  Registers & L1   |  |             |  |  Registers & L1   |  | |
|  |  +-------------------+  |             |  +-------------------+  | |
|  |  +-------------------+  |             |  +-------------------+  | |
|  |  | INT32/FP32 Cores  |  |             |  | INT32/FP32 Cores  |  | |
|  |  +-------------------+  |             |  +-------------------+  | |
|  +-------------------------+             +-------------------------+ |
+-------------------------------------------------------------+

The GigaThread Engine

At the heart of the GPU's global scheduling is the GigaThread Engine. This is a dedicated hardware block responsible for distributing work across the entire chip.

The GigaThread Engine receives the kernel launch command, inspects the grid dimensions (how many thread blocks need to run), and determines which Streaming Multiprocessors (SMs) have the capacity to accept new work.

Block Allocation and Resource Constraints

A thread block is the unit of work assignment. The GigaThread Engine cannot split a single thread block across multiple SMs; an entire block must reside on a single SM because threads within a block must be able to share memory (via L1/Shared Memory) and synchronize (via __syncthreads()).

The engine evaluates each SM’s available resources:

Registers: Does the SM have enough free 32-bit registers in its physical register file to accommodate all threads in the block?
Shared Memory: Is there enough allocated shared memory space on the SM for this block?
Thread Capacity: Does adding this block exceed the maximum number of concurrent threads or blocks allowed per SM for this hardware generation?

If an SM has sufficient resources, the GigaThread Engine dispatches the block to that SM. If all SMs are full, the remaining blocks sit in a hardware queue until active blocks finish executing and release their resources.

Execution on the SM: Warps, Schedulers, and Execution Units

Once a thread block is assigned to an SM, it is partitioned into groups of 32 threads called Warps. The warp is the fundamental unit of execution in NVIDIA's SIMT (Single Instruction, Multiple Threads) architecture.

SIMT and the Warp Scheduler

Even though you write code from the perspective of an individual thread, the hardware executes instructions at the warp level. All 32 threads in a warp execute the exact same instruction at the exact same clock cycle, but on different data elements.

Within the SM, physical Warp Schedulers manage the pool of active warps. On modern architectures (like Hopper or Ada), an SM is split into multiple processing blocks (sub-cores), each containing its own warp scheduler, instruction cache, and dispatch units.

Every clock cycle, a warp scheduler selects an active warp that is "ready" to execute (meaning its next instruction's operands are available and not blocked by memory latency) and dispatches the instruction to the appropriate execution units:

FP32 Cores: For standard single-precision floating-point math.
INT32 Cores: For integer arithmetic and address calculations.
Tensor Cores: For high-throughput matrix-multiply-accumulate (MMA) operations (crucial for deep learning).
LD/ST (Load/Store) Units: For reading and writing memory.
SFUs (Special Function Units): For transcendental functions like sine, cosine, and square roots.

The Problem of Control Divergence

What happens if your code contains an "if-else" statement?

if (threadIdx.x % 2 == 0) {
    // Path A
} else {
    // Path B
}

Because a warp must execute the same instruction simultaneously, it cannot execute both Path A and Path B at the same time. This triggers warp divergence.

The hardware handles this by serializing the execution: it disables the threads that evaluate to false for Path A, runs Path A for the remaining threads, then disables the threads that took Path A, and runs Path B for the rest. This drastically reduces execution efficiency. Developers must design algorithms to minimize branch divergence within the same warp.

Memory Latency Hiding: The Secret to GPU Throughput

CPUs rely on massive, multi-megabyte caches (L1, L2, L3) and highly sophisticated branch predictors to minimize the time the processor spends waiting for data from main memory. GPUs take a completely different approach: Latency Hiding.

GPU High Bandwidth Memory (HBM) or GDDR VRAM has massive bandwidth but relatively high latency (hundreds of clock cycles). When a warp executes a global memory load instruction, it must wait. Instead of stalling the entire SM, the warp scheduler simply puts that warp to sleep and immediately context-switches to another warp that is ready to execute.

Because state switching on a GPU is instantaneous—thanks to a massive physical register file that stores the registers of all active threads simultaneously—there is zero context-switch overhead. If you have enough active warps (high occupancy), the SM can keep its execution units 100% utilized even while waiting for slow memory transactions to complete.

Tracing a Real-World Example: Vector Addition

To solidify this mental model, let's trace a simple vector addition kernel:

__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

Host Launch: The CPU invokes vectorAdd<<<256, 256>>>(d_A, d_B, d_C, N). This requests 256 thread blocks, each containing 256 threads, totaling 65,536 threads.
PCIe Transfer: The CUDA Driver packages this request and pushes it to the GPU command queue.
Global Scheduling: The GigaThread Engine reads the command. It distributes the 256 thread blocks across the available SMs.
Warp Partitioning: On each SM, the 256 threads of a block are divided into 8 warps (256 / 32 = 8).
Address Calculation: The warp scheduler selects Warp 0. The FP32/INT32 units calculate the global index i for each thread.
Memory Request: The load units initiate a read from global memory addresses A[i] and B[i]. Because the threads in the warp access contiguous memory locations, the memory controller coalesces these 32 individual reads into a single, efficient 128-byte memory transaction.
Latency Hiding: While waiting for the data to arrive from VRAM, Warp 0 is paused. The scheduler context-switches to Warp 1, then Warp 2, executing their address calculations and memory requests.
Execution: Once the data for Warp 0 arrives, it is marked as ready. The scheduler selects it, and the FP32 cores perform the addition A[i] + B[i].
Writeback: The result is written back to C[i] via a coalesced write operation.
Retirement: Once all 8 warps in a block complete, the block is retired, freeing up resources on the SM for the GigaThread Engine to dispatch more blocks.

Conclusion: Writing Hardware-Aware CUDA Code

Understanding the physical journey of a CUDA kernel shifts your programming paradigm. You realize that optimizing code isn't just about reducing instruction count; it’s about managing hardware resources. To maximize GPU throughput, keep these principles in mind:

Coalesce Memory Accesses: Ensure adjacent threads in a warp access adjacent memory locations to minimize memory transactions.
Maximize Occupancy: Keep your register usage and shared memory footprint per thread balanced so that more warps can reside on the SM simultaneously, providing better latency hiding.
Avoid Warp Divergence: Keep execution paths uniform within a single 32-thread warp.

By aligning your software design with the underlying silicon architecture, you can unlock the true, massive parallel processing potential of modern GPU hardware.