2.2 CUDA Programming Model & Hierarchy

CUDA Programming Model & Execution Hierarchy

The Compute Unified Device Architecture (CUDA) is NVIDIA's parallel computing platform and programming model. It exposes the massive compute capabilities of NVIDIA GPUs to standard programming languages like C, C++, Python, and Fortran.

To effectively scale AI training, inference pipelines, and cluster orchestration, infrastructure engineers must understand how CUDA structures computational workloads and maps them onto physical hardware.

---

1. Software Hierarchy: Threads, Warps, Blocks, and Grids

CUDA organizes parallel computation into a four-level hierarchical structure that balances fine-grained data parallelism with scalable coarse-grained task distribution.

+-------------------------------------------------------------------------+
|                                 GRID                                    |
|  +---------------------------+       +---------------------------+      |
|  |      THREAD BLOCK (0,0)   |  ...  |      THREAD BLOCK (M,N)   |      |
|  |  +---------------------+  |       |  +---------------------+  |      |
|  |  | Warp 0 (32 Threads) |  |       |  | Warp 0 (32 Threads) |  |      |
|  |  | Warp 1 (32 Threads) |  |       |  | Warp 1 (32 Threads) |  |      |
|  |  | ...                 |  |       |  | ...                 |  |      |
|  |  | Warp K (32 Threads) |  |       |  | Warp K (32 Threads) |  |      |
|  |  +---------------------+  |       |  +---------------------+  |      |
|  +---------------------------+       +---------------------------+      |
+-------------------------------------------------------------------------+

1. Thread

• The smallest unit of execution in CUDA.
• Executes an instance of a CUDA kernel function.
• Has access to private registers, private local memory, and built-in multidimensional coordinate variables (threadIdx.x, threadIdx.y, threadIdx.z).

2. Warp

• A collection of 32 consecutive threads within a thread block.
• Represents the fundamental hardware scheduling and execution quantum on the GPU.
• Warp creation is handled entirely by hardware without programmer intervention.

3. Thread Block (Cooperative Thread Array / CTA)

• A group of threads (up to 1,024 threads on modern architectures) that execute on a single Streaming Multiprocessor (SM).
• Threads within a block can cooperate, share high-speed on-chip Shared Memory, and synchronize execution points using the barrier primitive __syncthreads().
• Identified by multidimensional coordinates (blockIdx.x, blockIdx.y, blockIdx.z).
• Defined by dimensional bounds (blockDim.x, blockDim.y, blockDim.z).

4. Grid

• The complete collection of thread blocks spawned by a single kernel launch.
• Defined by grid dimensions (gridDim.x, gridDim.y, gridDim.z).
• All thread blocks in a grid execute the same kernel code and share access to global device memory (HBM/GDDR).

Multidimensional Thread Indexing Math

In CUDA kernels, each thread calculates its unique global data index using its block and thread coordinates.

For a standard 1D Grid of 1D Blocks:

int global_idx = blockIdx.x * blockDim.x + threadIdx.x;

For a 2D Matrix (Width $\times$ Height):

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int global_idx = row * width + col;

The Principle of Thread Block Independence

A critical requirement of the CUDA architecture is Thread Block Independence:

• Any thread block within a grid must be capable of executing in any order relative to other thread blocks—concurrently, sequentially, or arbitrarily interleaved.
• Blocks cannot synchronize with each other directly using __syncthreads() (block-level barriers only synchronize threads within that specific block).
• Why this matters for scalability: A GPU with 10 SMs can execute 10 blocks in parallel, while a massive data-center GPU with 132 SMs (such as the H100) can execute 132 blocks in parallel, running the exact same binary without modification.

2. Hardware Mapping & Resource Allocation

Understanding how software concepts map to physical silicon is crucial for troubleshooting GPU resource exhaustion and performance bottlenecks.

Software-to-Hardware Mapping Rules

1. A Grid maps to the entire GPU Device: When a kernel is launched with kernel<<<gridDim, blockDim>>>(), the GPU's hardware Work Distribution Engine distributes the grid's blocks across all available SMs.
2. A Thread Block maps to exactly ONE SM: Once assigned to an SM, a thread block executes entirely on that SM from start to finish. A block is never split or migrated across multiple SMs.
3. An SM can host MULTIPLE concurrent Thread Blocks: Depending on hardware resource availability (registers, shared memory, thread limits), an SM typically executes multiple active blocks concurrently to maximize occupancy.
4. Threads execute on CUDA Cores / Tensor Cores: Each thread lane within an active warp executes on an individual arithmetic logic unit.

Hardware Mapping Summary Table

SM Occupancy & Hardware Resource Limits

Occupancy is the ratio of active warps on an SM to the maximum theoretical warps supported by the SM microarchitecture.

Each SM has fixed hardware limits (illustrated here for the NVIDIA Hopper H100 architecture):

• Max Threads per SM: 2,048 threads (64 warps)
• Max Thread Blocks per SM: 32 blocks
• Register File Size per SM: 64K (65,536) 32-bit registers (256 KB)
• Shared Memory / L1 SRAM: Up to 228 KB per SM

Resource Bottlenecks (Occupancy Limiters):

• Register Pressure: If a kernel requires 64 registers per thread, a 1,024-thread block requires $1024 \times 64 = 65,536$ registers—consuming 100% of the SM's register file. Consequently, the SM can host only one active block, limiting occupancy to $1024 / 2048 = 50%$.
• Shared Memory Allocation: If a kernel allocates 120 KB of shared memory per block on an SM with 228 KB total shared memory, only one block can reside on that SM at a time, regardless of how few registers each thread uses.

3. CUDA Streams & Asynchronous Concurrency

In enterprise AI training and inference serving (such as NVIDIA Triton Inference Server), maximizing GPU utilization requires overlapping host-to-device data transfers with kernel computation.

The CUDA Stream Concept

A CUDA Stream is a software-managed First-In, First-Out (FIFO) queue of GPU operations (kernel executions, memory copies, and events).

• Operations placed into the same stream are guaranteed to execute strictly in the order they were submitted (sequential execution within a stream).
• Operations placed into different streams have no execution ordering dependencies and can execute concurrently on the GPU hardware (inter-stream concurrency).

Stream 1: [ H2D Copy Chunk 1 ] ---> [ Compute Kernel 1 ] ---> [ D2H Copy Chunk 1 ]
Stream 2:        [ H2D Copy Chunk 2 ] ---> [ Compute Kernel 2 ] ---> [ D2H Copy Chunk 2 ]
Stream 3:               [ H2D Copy Chunk 3 ] ---> [ Compute Kernel 3 ] ---> [ D2H Copy Chunk 3 ]
          ========================================================================>
                                      TIME (Overlapped)

Default Stream vs. Non-Default Streams

1. Default Stream (NULL Stream / Stream 0):
   • Used when no stream is explicitly passed to kernel launches (kernel<<<grid, block>>>()) or standard memory copies (cudaMemcpy).
   • By default, the legacy NULL stream acts as a synchronization barrier: it waits for all prior operations in all other streams to complete before executing, and blocks subsequent operations in other streams until it finishes.
2. Non-Default Streams (Explicit Streams):
   • Created explicitly via cudaStreamCreate(&stream). Non-blocking streams are created using cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) to bypass default stream serialization.

Asynchronous Memory Transfers (cudaMemcpyAsync)

Standard cudaMemcpy() is a synchronous, blocking function: it halts host CPU execution until the data transfer between host RAM and GPU VRAM finishes.

In contrast, cudaMemcpyAsync() is non-blocking: it enqueues the memory transfer into a specified CUDA stream and returns control to the host CPU immediately.

// Creating streams
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// Overlapping pipeline for Stream 1 and Stream 2
cudaMemcpyAsync(d_in1, h_in1, size, cudaMemcpyHostToDevice, stream1);
kernelA<<<grid, block, 0, stream1>>>(d_in1, d_out1);
cudaMemcpyAsync(h_out1, d_out1, size, cudaMemcpyDeviceToHost, stream1);

cudaMemcpyAsync(d_in2, h_in2, size, cudaMemcpyHostToDevice, stream2);
kernelA<<<grid, block, 0, stream2>>>(d_in2, d_out2);
cudaMemcpyAsync(h_out2, d_out2, size, cudaMemcpyDeviceToHost, stream2);

Critical Requirement: For cudaMemcpyAsync() to execute asynchronously and overlap with compute, the host memory pointer MUST be page-locked (pinned) using cudaHostAlloc() or cudaMallocHost(). If standard pageable memory allocated via malloc() is passed, the CUDA driver falls back to synchronous staging, eliminating all concurrency.

---

4. CUDA Synchronization Primitives & Events

Managing dependencies in asynchronous workflows requires precise synchronization mechanisms:

Synchronization APIs

• cudaDeviceSynchronize(): Host-side barrier. The host CPU thread halts until all preceding commands in all streams on the GPU device complete. (High overhead, avoid in production hot loops).
• cudaStreamSynchronize(stream): Host-side barrier for a single stream. The CPU thread waits only until operations in the specified stream finish, allowing other streams to proceed.
• cudaStreamWaitEvent(stream, event): GPU-side synchronization. Enqueues a wait command in stream that prevents further operations in that stream from executing until event has been recorded. This coordinates inter-stream dependencies entirely on the GPU without stalling the host CPU.

CUDA Events for GPU Profiling

CUDA Events provide high-precision GPU-side timestamps that are not subject to CPU-side scheduling jitter:

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream);
myKernel<<<grid, block, 0, stream>>>(d_data);
cudaEventRecord(stop, stream);

cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);