GPU Architecture and CUDA Basics

In the late 1990s and early 2000s, the computer industry responded to the demanding graphics requirements of video games and 3D animation by designing dedicated Graphics Processing Units (GPUs).

Because GPUs were engineered to render millions of pixels simultaneously, they possessed massive computational throughput. By the early 2000s, researchers began harnessing this processing power for general scientific computing—giving rise to GPGPU (General-Purpose computing on Graphics Processing Units).

6.1 The Evolution of GPGPU: From Shaders to CUDA

Early GPGPU programming was notoriously difficult: programmers were forced to rephrase arithmetic computations as graphical operations using APIs like OpenGL or Direct3D. Algorithms had to be disguised as textures, polygons, and fragment shaders.

To overcome these hurdles, specialized computing platforms emerged:

OpenCL (Open Computing Language): An open standard designed for cross-platform portability across diverse hardware architectures (CPUs, GPUs, FPGAs, and DSPs). Portability comes with substantial configuration boilerplate.
CUDA (Compute Unified Device Architecture): Developed by Nvidia specifically for their GPU hardware. It integrates seamlessly into C and C++ with minimal setup, providing direct hardware access and mature compiler toolchains.

6.2 GPU Hardware Architecture: SIMD and SIMT

Traditional CPUs are optimized for low-latency execution of a single instruction stream (SISD in Flynn’s Taxonomy), featuring large caches and complex branch prediction.

In contrast, GPUs are massively parallel processors built around the SIMD (Single Instruction, Multiple Data) paradigm, allocating the majority of silicon area to arithmetic logic units (ALUs) rather than caches and control logic.

SIMD Branch Execution

In a pure SIMD architecture, a single control unit broadcasts an instruction to multiple datapaths. If a conditional branch (if-else) evaluates differently across datapaths, the hardware must serialize execution across multiple steps:

/* Datapath i executes: */
if (x[i] >= 0)
    x[i] += 1;
else
    x[i] -= 2;

Table 6.1: Execution of Conditional Branch on a SIMD System

Time Step	Datapaths with $x[i] \ge 0$	Datapaths with $x[i] < 0$
1	Evaluate condition ( $x[i] \ge 0 \to \text{True}$ )	Evaluate condition ( $x[i] \ge 0 \to \text{False}$ )
2	Execute $x[i] += 1$	Idle (Masked off)
3	Idle (Masked off)	Execute $x[i] -= 2$

Streaming Multiprocessors (SMs) and Streaming Processors (SPs)

An Nvidia GPU is organized into an array of Streaming Multiprocessors (SMs):

Each SM contains multiple control units, register files, on-chip shared memory, and dozens of datapaths called Streaming Processors (SPs) or CUDA Cores.
SMs operate asynchronously from one another. If threads on SM A take the if branch while threads on SM B take the else branch, both SMs execute simultaneously without penalizing each other.

Table 6.2: Branch Execution Across Multiple SMs

Time Step	Datapaths on SM A ( $x[i] \ge 0$ )	Datapaths on SM B ( $x[i] < 0$ )
1	Evaluate $x[i] \ge 0$	Evaluate $x[i] \ge 0$
2	Execute $x[i] += 1$	Execute $x[i] -= 2$ (Runs in parallel!)

Nvidia terms this architecture SIMT (Single Instruction, Multiple Thread). In SIMT, instructions are issued to groups of threads that execute in lockstep, but threads that stall on high-latency memory accesses can be temporarily suspended while the hardware scheduler instantly swaps in another ready group of threads, effectively hiding memory latency.

flowchart TD
  subgraph GPUArch["Figure 6.1: Simplified Architecture of a GPU"]
      direction TB
      subgraph SM0["Streaming Multiprocessor (SM 0)"]
          direction TB
          CTRL0["Control Units"]
          SP0["SP"] --- SP1["SP"] --- SP2["SP"] --- SP3["SP"]
          SMEM0["Fast On-Chip Shared Memory / L1"]
      end
      subgraph SM1["Streaming Multiprocessor (SM 1)"]
          direction TB
          CTRL1["Control Units"]
          SP4["SP"] --- SP5["SP"] --- SP6["SP"] --- SP7["SP"]
          SMEM1["Fast On-Chip Shared Memory / L1"]
      end
      L2["Shared L2 Cache"]
      GMEM["Global Device Memory (GDDR / HBM DRAM)"]
      
      SM0 <==> L2
      SM1 <==> L2
      L2 <==> GMEM
  end

6.3 Heterogeneous Computing: Host and Device

A CUDA application runs on a heterogeneous system comprising two physically distinct computing entities:

Host: The host CPU and its system RAM.
Device: The GPU and its on-board global memory.

flowchart LR
  subgraph Hetero["Figure 6.2: Heterogeneous CPU-GPU Architecture"]
      direction LR
      subgraph HostSide["Host (CPU)"]
          CPU["Host CPU"] <--> HMEM["Host System Memory (RAM)"]
      end
      PCIE["PCIe / NVLink Interconnect"]
      subgraph DeviceSide["Device (GPU)"]
          SMs["Streaming Multiprocessors (SMs)"] <--> DMEM["Device Global Memory (VRAM)"]
      end
      HostSide <==> PCIE <==> DeviceSide
  end

Developers write a single source file containing:

Host code: Standard C/C++ executed sequentially on the CPU to manage I/O, allocate memory, and launch parallel computation.
Device code (Kernels): Functions executed in parallel by thousands of GPU threads.

6.4 A First CUDA Program: Parallel Greetings

Below is Program 6.1 (cuda_hello.cu):

/* Program 6.1: CUDA program that prints greetings from threads */
#include <stdio.h>
#include <cuda.h>

/* Device code: runs on GPU */
__global__ void Hello(void) {
    printf("Hello from thread %d!\n", threadIdx.x);
} /* Hello */

/* Host code: runs on CPU */
int main(int argc, char* argv[]) {
    int thread_count;

    /* Get thread count from command line */
    thread_count = strtol(argv[1], NULL, 10);

    /* Launch thread_count threads on the GPU */
    Hello<<<1, thread_count>>>();

    /* Wait for GPU threads to complete */
    cudaDeviceSynchronize();

    return 0;
} /* main */

6.4.1 Compiling with `nvcc`

CUDA source files use the .cu extension and are compiled using Nvidia’s compiler driver, nvcc:

$ nvcc -o cuda_hello cuda_hello.cu
$ ./cuda_hello 4
Hello from thread 0!
Hello from thread 1!
Hello from thread 2!
Hello from thread 3!

nvcc separates the host code from the device code:

Host code is forwarded to the system C++ compiler (such as gcc or MSVC).
Device code is compiled into GPU machine instructions (PTX and SASS assembly).

6.4.2 Kernel Syntax and Execution Semantics

__global__ Qualifier: Declares a function as a CUDA kernel. A kernel is called from the host (CPU) and executed on the device (GPU). Kernels must always have a void return type.
Triple Angle Brackets <<<grid_dim, block_dim>>>: The execution configuration specifying how many thread blocks and threads per block to launch.
threadIdx.x: A built-in struct initialized by the hardware representing the thread’s local rank within its block.
Asynchronous Execution: Kernel launches are non-blocking (asynchronous). Control returns to the host CPU immediately after queuing the launch.
cudaDeviceSynchronize(): Blocks the host CPU until all previously issued GPU commands and kernels have completed. Without this call, main() might terminate before the GPU threads finish executing printf().

6.5 Threads, Blocks, and Grids

To scale seamlessly across hardware with varying numbers of SMs, CUDA organizes threads into a two-level hierarchy:

flowchart TD
  subgraph GridHierarchy["CUDA Thread Hierarchy"]
      direction TB
      subgraph Grid["Grid (gridDim)"]
          direction LR
          subgraph B0["Block (0)"]
              T0_0["Thread (0)"]
              T0_1["Thread (1)"]
              T0_N["Thread (blockDim.x - 1)"]
          end
          subgraph B1["Block (1)"]
              T1_0["Thread (0)"]
              T1_1["Thread (1)"]
              T1_N["Thread (blockDim.x - 1)"]
          end
          subgraph BK["Block (gridDim.x - 1)"]
              TK_0["Thread (0)"]
              TK_1["Thread (1)"]
              TK_N["Thread (blockDim.x - 1)"]
          end
      end
  end

Thread: The basic unit of execution running on a Streaming Processor (SP).
Thread Block (or Block): A group of threads that execute on a single SM, sharing on-chip shared memory and hardware barrier resources.
Grid: The entire collection of thread blocks spawned by a kernel launch.

Built-in CUDA Indexing Variables

Every thread has access to four built-in structs with unsigned fields .x, .y, and .z:

threadIdx: Coordinates of the thread within its block ( $0 \le \text{threadIdx.x} < \text{blockDim.x}$ ).
blockDim: Dimensions of the thread block.
blockIdx: Coordinates of the block within the grid ( $0 \le \text{blockIdx.x} < \text{gridDim.x}$ ).
gridDim: Dimensions of the grid.

Multi-Block Program: Program 6.2

/* Program 6.2: CUDA greetings from multiple blocks */
#include <stdio.h>
#include <cuda.h>

__global__ void Hello(void) {
    printf("Hello from thread %d in block %d\n", threadIdx.x, blockIdx.x);
}

int main(int argc, char* argv[]) {
    int blk_ct = strtol(argv[1], NULL, 10);
    int th_per_blk = strtol(argv[2], NULL, 10);

    /* Launch blk_ct blocks, each with th_per_blk threads */
    Hello<<<blk_ct, th_per_blk>>>();

    cudaDeviceSynchronize();
    return 0;
}

Running with 2 blocks of 3 threads (./cuda_hello_mult 2 3) yields:

Hello from thread 0 in block 0
Hello from thread 1 in block 0
Hello from thread 2 in block 0
Hello from thread 0 in block 1
Hello from thread 1 in block 1
Hello from thread 2 in block 1

Block Independence

CUDA mandates that thread blocks must be completely independent. The hardware scheduler may execute blocks in any order—concurrently, sequentially, or interleaved—depending on available SMs. This guarantees that CUDA programs scale automatically from mobile GPUs with 1 SM to data-center GPUs with 100+ SMs.

Multidimensional Grids with `dim3`

Grids and blocks can be configured in 1, 2, or 3 dimensions using CUDA’s dim3 type:

dim3 grid_dims(2, 3, 1);    /* 2 x 3 x 1 = 6 blocks  */
dim3 block_dims(4, 4, 4);   /* 4 x 4 x 4 = 64 threads per block */

Kernel<<<grid_dims, block_dims>>>(...);

6.6 Nvidia Compute Capabilities and Architectures

A GPU’s feature set and hardware limits are designated by its Compute Capability version $a.b$ (major $a$ , minor $b$ ):

Table 6.3: Nvidia GPU Microarchitectures and Compute Capabilities

Architecture	Compute Capability	Max Threads per Block	Max Threads per SM	Notable Features
Tesla	1.x	512	768 / 1024	First GPGPU unified shader architecture
Fermi	2.x	1024	1536	Hardware L1/L2 caches, 64-bit IEEE floating point
Kepler	3.x	1024	2048	Dynamic parallelism, Warp Shuffle instructions
Maxwell	5.x	1024	2048	High energy efficiency, improved shared memory
Pascal	6.x	1024	2048	NVLink, hardware Page Migration, 16-bit FP
Volta	7.0	1024	2048	Tensor Cores, Independent Thread Scheduling
Turing	7.5	1024	1024	RT Cores, INT32 + FP32 concurrent execution
Ampere	8.x	1024	2048	3rd Gen Tensor Cores, Asynchronous Copy