CUDA programming model basics

What is CUDA?

CUDA (Compute Unified Device Architecture) is NVIDIA's parallel computing platform that allows programmers to harness GPU power for general-purpose computation, not just graphics.

WHY it exists: Graphics cards evolved to render millions of pixels in parallel. NVIDIA realized this massive parallelism could solve scientific, AI, and data problems if they gave programmers direct access.

WHAT it provides: A C/C+-like language extension + runtime libraries + compiler (nvcc) that let you write code for the GPU.

HOW it works: You write special functions (kernels) that run on the GPU. You organize work into threads, blocks, and grids. The GPU scheduler distributes these across its hardware.

The CUDA Execution Model: From First Principles

Why Paralelism Requires Structure

The fundamental problem: If you have 10,000 tasks and 10,000 workers, you need:

A way to identify which worker does task
A way to coordinate when workers need to share information
A way to map logical work organization to physical hardware

CUDA solves this with a three-level hierarchy:

Formula CUDA Hierarchy (Logical → Physical)

Logical Organization: $\text{Grid} \rightarrow \text{Blocks} \rightarrow \text{Threads}$

Physical Mapping: $\text{Grid maps to GPU} \rightarrow \text{Blocks map to SMs} \rightarrow \text{Threads map to CUDA cores}$

WHY three levels?

Thread level: Identifies individual work items
Block level: Enables cooperation among related threads (shared memory, barriers)
Grid level: Scales across the entire GPU without hardware assumptions

DERIVATION - Thread Indexing:

For a 1D problem with total N elements:

Launch configuration: <<numBlocks, threadsPerBlock>>>
Each thread computes its global index: $\text{globalIdx} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}$

WHY this formula?

blockIdx.x: Which block am I in? (0, 1, 2, ...)
blockDim.x: How many threads per block? (e.g., 256)
blockIdx.x × blockDim.x: How many threads came before my block?
threadIdx.x: My position within my block (0 to blockDim.x-1)
Sum = my unique position in the entire grid

For 2D problems (images): $\text{row} = \text{blockIdx.y} \times \text{blockDim.y} + \text{threadIdx.y}$ $\text{col} = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}$ $\text{globalIdx} = \text{row} \times \text{width} + \text{col}$

Memory Hierarchy

WHY memory matters: Moving data is expensive. A memory access can take 100× longer than a math operation. CUDA provides multiple memory types with different speed/scope tradeoffs.

Formula CUDA Memory Hierarchy (Fast → Slow)

| Memory Type | Scope | Lifetime | Speed | Size | |-------------|----------|----| | Registers | Per-thread | Thread lifetime | 1 cycle | ~64KB/SM | | Shared Memory | Per-block | Block lifetime | ~5 cycles | 48-164KB/SM | | L1/L2 Cache | Per-SM / GPU | Automatic | ~10-30 cycles | 128KB-40MB | | Global Memory | All threads | Program lifetime | ~200-400 cycles | GBs (VRAM) | | Host Memory | CPU | Persistent | ~500-1000 cycles | System RAM |

Speed reasoning from first principles:

Registers: On-chip, directly in compute units → wire distance ~mm → picoseconds
Shared memory: On-chip SRAM, shared within SM → still mm → few ns
Global memory: Off-chip GDR/HBM → wire distance ~cm + memory controller → 100+ ns
Host memory: Across PCIe bus → meters of signaling → microseconds

Access pattern formula: $\text{Effective Bandwidth} = \frac{\text{Theoretical BW} \times \text{Coalescing Efficiency}}{\text{Bank Conflicts + Unaligned Access Penalties}}$

Writing Your First CUDA Kernel

Vector Addition: The "Hello World" of CUDA

The problem: Add two arrays: C[i] = A[i] + B[i] for i = 0 to N-1

CPU version (sequential):

for (int i = 0; i < N; i++) {
    C[i] = A[i] + B[i];
}
// Time: O(N) - one element per iteration

CUDA version (parallel):

Worked example Complete CUDA Vector Addition

// KERNEL: Runs on GPU device
__global__ void vectorAdd(float *A, float *B, float *C, int N) {
   // Each thread computes its unique index
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   
   // Boundary check: only process valid indices
   if (idx < N) {
       C[idx] = A[idx] + B[idx];
   }
}
int main() {
   int N = 1000000;
   size_t bytes = N * sizeof(float);
   
   // 1. ALLOCATE HOST MEMORY
   float *h_A = (float*)malloc(bytes);
   float *h_B = (float*)malloc(bytes);
   float *h_C = (float*)malloc(bytes);
   
   // Initialize h_A, h_B with data...
   
   // 2. ALLOCATE DEVICE MEMORY
   float *d_A, *d_B, *d_C;
   cudaMalloc(&d_A, bytes);
   cudaMalloc(&d_B, bytes);
   cudaMalloc(&d_C, bytes);
   
   // 3. COPY HOST → DEVICE
   cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
   cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
   // 4. LAUNCH KERNEL
   int threadsPerBlock = 256;
   int numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
   vectorAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, N);
   
   // 5. COPY DEVICE → HOST
   cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
   
   // 6. CLEANUP
   cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
   free(h_A); free(h_B); free(h_C);
}

WHY each step? Step 1 (Host allocation): CPU memory for input/output data Step 2 (Device allocation): GPU has separate memory space—must allocate VRAM

Step 3 (Copy H→D): GPU can't access system RAM—copy data across PCIe

Step 4 (Kernel launch):

<<numBlocks, threadsPerBlock>>>: Execution configuration
numBlocks = ceil(N / 256): Need enough threads to cover all N elements
Why 256 threads/block? Hardware constraint: blocks have max threads (1024). 256 is a good balance: allows 4-8 blocks per SM, fills warps (32) evenly Formula for numBlocks: $\text{numBlocks} = \left\lceil \frac{N}{\text{threadsPerBlock}} \right\rceil = \frac{N + \text{threadsPerBlock} - 1}{\text{threadsPerBlock}}$ The ceiling division trick avoids floating-point math

Step 5 (Copy D→H): Get results back to CPU memory

Step 6 (Cleanup): Free both host and device memory

The Boundary Check: Why `if (idx < N)`?

The problem: If N = 1000and threadsPerBlock = 256:

Need 4blocks × 256 threads = 1024 threads total
But we only have 1000 elements
Last 24 threads (idx 1000-1023) are extra

Without the check: These threads would access A[1000], A[1001], ... → out of bounds! → crash or garbage data

With the check: Extra threads do nothing (idle). Tiny performance cost, prevents disaster.

Common mistake Common CUDA Mistakes

Mistake 1: Forgetting the boundary check

__global__ void bad(float *A, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   A[idx] = idx;  // CRASH if idx >= N!
}

WHY it feels right: "I launched exactly enough threads, right?" FIX: You launched blocks, and blocks come in fixed sizes. Always check if (idx < N).

Mistake 2: Wrong memory copy direction

cudaMemcpy(d_A, h_A, bytes, cudaMemcpyDeviceToHost);  // WRONG!

WHY it feels right: "d_A is destination, h_A is source..." FIX: The enum specifies the direction of data flow: DeviceToHost means GPU → CPU. For input, use HostToDevice.

Mistake 3: Not checking CUDA errors

cudaMalloc(&d_A, bytes);  // Might fail! No check!

WHY it feels right: "It's just allocation, it'll work." FIX: GPU memory can run out. Always check:

cudaError_t err = cudaMalloc(&d_A, bytes);
if (err != cudaSuccess) {
   printf("Error: %s\n", cudaGetErrorString(err));
}

Mistake 4: Using the wrong formula for2D indexing

int idx = blockIdx.x * blockDim.y + threadIdx.x;  // WRONG!

WHY it feels right: "I'm using x for threads..." FIX: Block dimensions use .x, .y, .z but must match block index dimensions. For 1D grid: use .x throughout. For 2D: blockIdx.y × blockDim.y + threadIdx.y for rows.

Synchronization and Cooperation

Thread Synchronization Within a Block

WHY needed: Threads in a block share resources (shared memory). Without synchronization, race conditions occur.

Worked example Using Shared Memory (Reduction)

Problem: Sum an array in parallel

__global__ void sumReduce(float *input, float *output, int N) {
   __shared__ float shared[256];  // Shared memory for this block
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   int tid = threadIdx.x;
   
   // Load data into shared memory
   shared[tid] = (idx < N) ? input[idx] : 0.0f;
   __syncthreads();  // Wait for all threads to finish loading
   
   // Reduction in shared memory
   for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
       if (tid < stride) {
           shared[tid] += shared[tid + stride];
       }
       __syncthreads();  // Wait before next iteration
   }
   // Thread 0 writes result
   if (tid == 0) {
       output[blockIdx.x] = shared[0];
   }
}

WHY this works:

__shared__ declaration: Allocates fast on-chip memory visible to all threads in the block
First __syncthreads(): Ensures all threads finish loading before anyone reads
Reduction loop: Halves the active threads each iteration

Stride 128: threads 0-127 add pairs
Stride 64: threads 0-63 add pairs
...until only thread 0 has the sum

Second __syncthreads(): Ensures writes from previous iteration are visible before reads in next iteration
Thread 0 writes: Only one thread writes the block's result to global memory

Derivation - Why reduction is O(log N): $\text{Iterations} = \log_2(\text{blockDim.x})$ Each iteration halves the work. For256 threads: 8 iterations instead of 255sequential additions.

Global Synchronization

Key constraint: __syncthreads() only works within a block. Threads in different blocks cannot synchronize during kernel execution.

WHY: Blocks may execute in any order, or even concurrently on different SMs. Allowing inter-block synchronization would create deadlocks.

HOW to synchronize across blocks: Launch separate kernels. The GPU guarantees all threads from kernel 1 finish before any thread from kernel 2 starts.

Performance Considerations

Memory Coalescing

Coalesced access: Threads in a warp (32 consecutive threads) access consecutive memory addresses → GPU combines into one transaction.

Uncoalesced access: Threads access scattered addresses → GPU makes many small transactions → 10-100× slower.

Formula Coalescing Rules

Perfect coalescing (one128-byte transaction for 32 threads):

float *data;
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float value = data[idx];  // Sequential access

Terrible performance (32 separate transactions):

float value = data[idx * 1000];  // Strided by 1000

WHY: GPU memory controllers work on aligned 32/64/128-byte chunks. If thread0 reads byte 0, thread 1 byte 4, thread 2 byte 8, .., thread 31 byte 124→ all in one 128-byte chunk → one transaction.

If accesses are scattered, each thread's request hits a different chunk → 32 transactions.

Occupancy

Occupancy is the ratio of active warps to maximum possible warps on an SM.

WHY it matters: Higher occupancy → more warps available → better latency hiding.

Factors limiting occupancy:

Registers per thread: More registers → fewer threads fit
Shared memory per block: More shared memory → fewer blocks fit
Block size: Too small → wastes resources; too large → limits blocks per SM

$\text{Occupancy} = \frac{\text{Active Warps per SM}}{\text{Maximum Warps per SM}}$

Typical goal: 50-75% occupancy is often enough. 100% isn't always necessary for good performance.

Recall Feynman: Explain to a 12-year-old

Imagine you have a huge math worksheet with 10,000 problems. If you do them yourself, it takes forever—one problem at a time.

Now imagine you have a classroom with 10,000 students, and you give each student exactly one problem. They all work at the same time, and suddenly the whole worksheet is done in the time it takes to solve ONE problem!

That's what a GPU does. The "students" are threads—tiny workers. CUDA is how you tell them which problem to solve.

You organize students into groups (blocks) sitting at tables. Students at the same table can share a whiteboard (shared memory) to work together. But students at different tables can't talk to each other until everyone is done.

You also have to be careful: if you have 10,000 problems but10,240 students, the last 240 don't have a problem to solve. You need to tell them "sit this one out" or they'll try to solve imaginary problems and mess things up!

Connections

GPU Architecture Overview - Hardware foundation for CUDA
Streaming Multiprocessors - Physical execution units that run blocks
Memory Hierarchy - Detailed memory performance analysis
Thread Warps and SIMT - How threads actually execute on hardware
Parallel Algorithm Design - Strategies for decomposing problems
OpenCL vs CUDA - Alternative programming models
Deep Learning Frameworks - TensorFlow/PyTorch use CUDA under the hood

Flashcards

#flashcards/hardware

What is a CUDA kernel? :: A function that runs on the GPU, executed in parallel by many threads. Declared with __global__ keyword.

What is the formula for global thread index in1D?

idx = blockIdx.x * blockDim.x + threadIdx.x

Why do we need if (idx < N) in CUDA kernels? :: Because we launch blocks of fixed size (e.g., 256 threads), we might have more threads than data elements. Extra threads must be prevented from accessing out-of-bounds memory.

What is the difference between host and device in CUDA?

Host = CPU and system RAM. Device = GPU and VRAM. They have separate memory spaces requiring explicit copies.

What does <<numBlocks, threadsPerBlock>>> specify?

The kernel launch configuration: how many blocks to launch and how many threads per block.

What is __syncthreads() used for?

Synchronizing threads within a block. Ensures all threads reach the barrier before any continue. Used when threads share data via shared memory.

Can threads in different blocks synchronize during kernel execution?

No. Blocks execute independently and may run in any order. Only way to synchronize across blocks is to launch separate kernels.

What is memory coalescing?

When threads in a warp access consecutive memory addresses, the GPU combines requests into one transaction. Non-coalesced access causes many separate transactions and is much slower.

What is __shared__ memory?

Fast on-chip memory shared by all threads in a block. Much faster than global memory (~5 cycles vs ~400 cycles) but limited in size (48-164KB per SM).

What is a warp in CUDA?

A group of 32 threads that execute in lockstep (SIMT). The fundamental execution unit on NVIDIA GPUs.

How do you calculate the number of blocks needed for N elements?

numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock (ceiling division using integer arithmetic)

What are the three levels of CUDA's execution hierarchy?

Grid → Blocks → Threads. Grid maps to GPU, blocks map to SMs, threads map to CUDA cores.

Why is 256 a common choice for threadsPerBlock?

It's a multiple of 32 (warp size), allows 4-8 blocks per SM, and stays well under the 1024 max thread limit. Good balance for occupancy.

What is GPU occupancy?

The ratio of active warps to maximum possible warps on an SM. Higher occupancy enables better latency hiding through warp switching.

What memory type is fastest in CUDA?

Registers (per-thread, ~1 cycle access), followed by shared memory (~5 cycles), then L1/L2 cache, then global memory (~400 cycles).

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho, GPU aur CPU ka core farak yahi hai ki CPU ke paas thode "smart workers" hote hain jo complex kaam fatafat karte hain, jabki GPU ke paas hazaaron "simple workers" hote hain jo ek-ek chhota kaam karte hain. CUDA basically ek management system hai jo in hazaaron workers ko batata hai ki kya karna hai, kab karna hai aur results kaise share karne hain. Matlab loop mein "yeh kaam 10,000 baar karo" likhne ke bajaye, aap kehte ho "yeh task hai, ab isse ek saath 10,000 threads mein parallel chala do." Yahi parallelism ka asli magic hai.

Ab is parallelism ko manage karne ke liye CUDA ek teen-level hierarchy use karta hai: Grid, Blocks, aur Threads. Thread sabse chhota unit hai jo ek piece of work karta hai, Block ek group hota hai jismein threads aapas mein cooperate kar sakte hain (shared memory aur synchronization ke through), aur Grid saare blocks ka collection hota hai. Physically, grid poore GPU pe map hota hai, blocks SMs pe, aur threads CUDA cores pe. Har thread apna unique global index nikaalta hai is formula se: blockIdx.x × blockDim.x + threadIdx.x — yani "mere block se pehle kitne threads aaye" plus "apne block ke andar meri position." Isse har worker ko pata chalta hai ki usse konsa data process karna hai.

Yeh sab why-matters isliye hai kyunki AI, deep learning, scientific computing aur data processing jaise kaam mein hume massive parallelism chahiye, aur CUDA hume GPU ki us power tak direct access deta hai. Ek important cheez yaad rakhna — memory movement bahut expensive hoti hai, ek memory access kabhi-kabhi math operation se 100 guna zyada time le sakta hai. Isliye CUDA alag-alag memory types deta hai (registers sabse fast, phir shared memory, phir cache, phir slow global memory) taaki aap apne data ko smartly manage karke performance maximize kar sako. Jab tum CUDA achhe se samajh loge, tab tum modern GPU-based computing ka poora foundation samajh jaoge.

What is CUDA?

The CUDA Execution Model: From First Principles

Why Paralelism Requires Structure

Memory Hierarchy

Writing Your First CUDA Kernel

Vector Addition: The "Hello World" of CUDA

The Boundary Check: Why `if (idx < N)`?

Synchronization and Cooperation

Thread Synchronization Within a Block

Global Synchronization

Performance Considerations

Memory Coalescing

Occupancy

Connections

Flashcards

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture

Connections