CUDA Fundamentals

Phase 1 — CUDA Fundamentals: GPU Architecture

(with expanded practical code examples)

This is the foundation for GPU kernel programming. Before writing optimized CUDA kernels, you need to understand how your code maps onto the GPU hardware. The sections below keep the original explanations and add short, runnable-style code examples so you can practice the concepts.

1. GPU vs CPU

A CPU is designed primarily for low-latency, general-purpose sequential and moderately parallel work.
A GPU is designed primarily for massively parallel workloads.

CPU

textCopy
1CPU
2├── Few powerful cores
3├── Large caches
4├── Complex control logic
5└── Excellent for sequential/branch-heavy work

GPU

textCopy
1GPU
2├── Many SMs
3│   ├── Many warps
4│   │   └── Many threads
5│   ├── Registers
6│   └── Shared Memory
7├── Large global memory
8└── Massive parallel execution

Example task: C[i] = A[i] + B[i] for 1 000 000 elements.

CPU sequential version (for contrast):

cppCopy
1// CPU code
2void add_cpu(const float* a, const float* b, float* c, int n) {
3    for (int i = 0; i < n; ++i) {
4        c[i] = a[i] + b[i];
5    }
6}

GPU parallel version (conceptual mapping):

textCopy
1Thread 0 → C[0]
2Thread 1 → C[1]
3...
4Thread 999999 → C[999999]

GPUs excel at matrix multiplication, neural networks, image processing, simulations, scientific computing, and Transformer workloads.

2. What is CUDA?

CUDA is NVIDIA’s GPU computing platform and programming model.
You write host (CPU) code that launches kernels that run on the device (GPU).

textCopy
1CPU
2 │
3 │ launches
4 ↓
5CUDA Kernel
6 │
7 ↓
8GPU
9 ├── SM
10 │   ├── Warps
11 │   │   └── Threads
12 │   ├── Registers
13 │   └── Shared Memory
14 │
15 └── Global Memory

Minimal complete example (vector add):

cppCopy
1#include <cuda_runtime.h>
2#include <cstdio>
3
4// Device kernel
5__global__ void add_kernel(const float* a, const float* b, float* c, int n) {
6    int i = blockIdx.x * blockDim.x + threadIdx.x;
7    if (i < n) {
8        c[i] = a[i] + b[i];
9    }
10}
11
12int main() {
13    const int N = 1 << 20; // 1M elements
14    size_t bytes = N * sizeof(float);
15
16    // Host memory
17    float *h_a = new float[N];
18    float *h_b = new float[N];
19    float *h_c = new float[N];
20    for (int i = 0; i < N; ++i) {
21        h_a[i] = 1.0f;
22        h_b[i] = 2.0f;
23    }
24
25    // Device memory
26    float *d_a, *d_b, *d_c;
27    cudaMalloc(&d_a, bytes);
28    cudaMalloc(&d_b, bytes);
29    cudaMalloc(&d_c, bytes);
30
31    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
32    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
33
34    // Launch configuration
35    int threads = 256;
36    int blocks  = (N + threads - 1) / threads;
37    add_kernel<<<blocks, threads>>>(d_a, d_b, d_c, N);
38
39    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
40
41    // Quick check
42    printf("c[0] = %f (should be 3.0)\n", h_c[0]);
43
44    // Cleanup
45    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
46    delete[] h_a; delete[] h_b; delete[] h_c;
47    return 0;
48}

Compile & run (example):

bashCopy
1nvcc add.cu -o add && ./add

3. What is an SM?

SM = Streaming Multiprocessor — a parallel execution unit inside the GPU.

textCopy
1GPU
2│
3├── SM 0
4│   ├── CUDA cores
5│   ├── Tensor Cores
6│   ├── Registers
7│   ├── Shared Memory
8│   └── Warp schedulers
9├── SM 1
10│   └── ...
11└── ...

When you launch a kernel the hardware distributes thread blocks across SMs:

textCopy
1Kernel
2  │
3  ├── Block 0 ──→ SM 0
4  ├── Block 1 ──→ SM 1
5  ├── Block 2 ──→ SM 2
6  ├── Block 3 ──→ SM 0
7  └── ...

Practice: query device properties (number of SMs, etc.):

cppCopy
1#include <cuda_runtime.h>
2#include <cstdio>
3
4int main() {
5    int deviceCount = 0;
6    cudaGetDeviceCount(&deviceCount);
7
8    for (int i = 0; i < deviceCount; ++i) {
9        cudaDeviceProp prop;
10        cudaGetDeviceProperties(&prop, i);
11        printf("Device %d: %s\n", i, prop.name);
12        printf("  SMs (multiProcessorCount): %d\n", prop.multiProcessorCount);
13        printf("  Max threads per SM: %d\n", prop.maxThreadsPerMultiProcessor);
14        printf("  Warp size: %d\n", prop.warpSize);
15        printf("  Shared mem per block: %zu bytes\n", prop.sharedMemPerBlock);
16        printf("  Registers per block: %d\n", prop.regsPerBlock);
17    }
18    return 0;
19}

4. What is a CUDA Thread?

A thread is the smallest logical execution unit.

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;

Practice: print thread / block indices (small launch):

cppCopy
1__global__ void print_ids() {
2    int global_id = blockIdx.x * blockDim.x + threadIdx.x;
3    printf("Block %d, Thread %d → global id %d\n",
4           blockIdx.x, threadIdx.x, global_id);
5}
6
7int main() {
8    // 2 blocks × 4 threads = 8 threads total
9    print_ids<<<2, 4>>>();
10    cudaDeviceSynchronize();
11    return 0;
12}

Expected output (order may vary):

Block 0, Thread 0 → global id 0
Block 0, Thread 1 → global id 1
...
Block 1, Thread 3 → global id 7

5. What is a Block?

Threads are grouped into thread blocks. A kernel launch specifies the grid of blocks and the size of each block.

cppCopy
1kernel<<<blocks, threads>>>();

Example:

cppCopy
1kernel<<<100, 256>>>();   // 100 blocks × 256 threads = 25 600 logical threads

Practice: 2-D indexing (common for images / matrices):

cppCopy
1__global__ void print_2d_ids() {
2    int x = blockIdx.x * blockDim.x + threadIdx.x;
3    int y = blockIdx.y * blockDim.y + threadIdx.y;
4    printf("(%d,%d)\n", x, y);
5}
6
7int main() {
8    dim3 threads(4, 4);          // 4×4 threads per block
9    dim3 blocks(2, 2);           // 2×2 blocks
10    print_2d_ids<<<blocks, threads>>>();
11    cudaDeviceSynchronize();
12    return 0;
13}

6. What is a Warp?

NVIDIA GPUs execute threads in groups of 32 called warps (SIMT model).

textCopy
1256 threads / block → 256 / 32 = 8 warps

Practice: detect warp divergence (if/else inside a warp):

cppCopy
1__global__ void divergence_demo(int* out) {
2    int tid = threadIdx.x;
3    int warp_id = tid / 32;
4
5    if (tid % 2 == 0) {
6        out[tid] = tid * 2;   // even threads
7    } else {
8        out[tid] = tid * 3;   // odd threads  → divergence inside the warp
9    }
10}

(In real kernels avoid heavy divergence; keep warps executing the same path when possible.)

7. CUDA Cores

CUDA cores are the arithmetic units that execute + - * / etc.
Warps are scheduled onto these resources; a CUDA core is not permanently assigned to one thread.

8. Tensor Cores

Specialized units for matrix multiply-accumulate (especially important for AI).

Practice (using cuBLAS for a simple GEMM – Tensor Cores are used automatically on supported hardware):

cppCopy
1#include <cublas_v2.h>
2#include <cuda_runtime.h>
3
4// Pseudo-code sketch – real code needs proper initialization
5void gemm_example(cublasHandle_t handle,
6                  float* d_A, float* d_B, float* d_C,
7                  int M, int N, int K) {
8    const float alpha = 1.0f, beta = 0.0f;
9    // C = alpha * A * B + beta * C
10    cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
11                M, N, K, &alpha, d_A, M, d_B, K, &beta, d_C, M);
12}

On modern GPUs (Ampere+) Tensor Cores accelerate FP16 / BF16 / TF32 / FP8 matrix multiplies that dominate Transformer layers.

9–14. GPU Memory Hierarchy (with code)

textCopy
1Fastest / smallest
2        ↓
3    Registers
4        ↓
5  Shared Memory
6        ↓
7      L1 Cache
8        ↓
9      L2 Cache
10        ↓
11  Global Memory
12        ↓
13Slowest / largest

Registers

Private to each thread, extremely fast, limited.

cppCopy
1__global__ void register_example(float* out) {
2    float x = 1.0f;          // likely kept in a register
3    float y = 2.0f;
4    float result = x * y + threadIdx.x;
5    out[threadIdx.x] = result;
6}

Excessive register use → lower occupancy or spilling.

Shared Memory

Shared by all threads in a block. Perfect for data reuse / tiling.

cppCopy
1__global__ void shared_mem_example(const float* in, float* out, int n) {
2    __shared__ float tile[256];          // 256 floats shared by the block
3
4    int tid = threadIdx.x;
5    int i   = blockIdx.x * blockDim.x + tid;
6
7    if (i < n) tile[tid] = in[i];
8    __syncthreads();                     // all threads must reach here
9
10    // Now every thread can read any element of the tile
11    if (i < n) out[i] = tile[tid] * 2.0f;
12}

Global Memory

Large, accessible by all threads. Usual place for tensors.

cppCopy
1float *d_weights, *d_activations;
2cudaMalloc(&d_weights, size);
3cudaMalloc(&d_activations, size);

L1 / L2 Caches

Managed automatically by the hardware. You influence them indirectly via access patterns and cache configuration (e.g. cudaFuncSetCacheConfig).

15–16. Putting It Together & Kernel Execution Flow

cppCopy
1__global__ void add(float* a, float* b, float* c, int n) {
2    int i = blockIdx.x * blockDim.x + threadIdx.x;
3    if (i < n) c[i] = a[i] + b[i];
4}
5
6// Launch
7add<<<100, 256>>>(d_a, d_b, d_c, N);

Conceptual flow:

CPU launches kernel
  → Grid of blocks
    → Blocks scheduled onto SMs
      → Warps scheduled onto CUDA / Tensor cores
        → Threads read A/B, compute, write C

17–18. Why Memory Optimization Matters + Coalescing

Simple arithmetic is often memory-bound.

Good (coalesced) access:

cppCopy
1// Thread i reads contiguous location
2float val = a[i];

Bad (strided / scattered):

cppCopy
1float val = a[i * 100];   // large stride → poor coalescing

Practice: measure the difference (use nvprof / Nsight or simple timing):

cppCopy
1__global__ void coalesced(float* a, float* out, int n) {
2    int i = blockIdx.x * blockDim.x + threadIdx.x;
3    if (i < n) out[i] = a[i];
4}
5
6__global__ void strided(float* a, float* out, int n, int stride) {
7    int i = blockIdx.x * blockDim.x + threadIdx.x;
8    if (i < n) out[i] = a[i * stride];
9}

19. Why Shared Memory Exists (Tiling foundation)

cppCopy
1// Classic pattern for matrix multiply / attention / convolution
2__global__ void tiled_example(const float* A, float* C, int N) {
3    __shared__ float tile[TILE][TILE];
4
5    int tx = threadIdx.x, ty = threadIdx.y;
6    int row = blockIdx.y * TILE + ty;
7    int col = blockIdx.x * TILE + tx;
8
9    // Load tile from global → shared
10    if (row < N && col < N)
11        tile[ty][tx] = A[row * N + col];
12    __syncthreads();
13
14    // Reuse data from shared memory many times
15    // ... compute ...
16}

20. Connection to Transformers

textCopy
1Transformer
2  ↓
3Tensor operations (QKV projections, Attention, MLP GEMMs, RMSNorm …)
4  ↓
5CUDA kernels
6  ↓
7Threads / Warps / Shared Memory / Tensor Cores

Example mapping:

• Q = X @ Wq → GEMM kernel (Tensor Cores)
• Softmax / reductions → warp-level or block-level primitives
• RMSNorm → reduction + element-wise kernel

21. The Most Important Mental Model

textCopy
1GPU
2 │
3 ├── SM
4 │   ├── Warps (32 threads)
5 │   ├── Registers (private)
6 │   ├── Shared Memory (block-shared)
7 │   ├── CUDA Cores
8 │   └── Tensor Cores
9 │
10 └── Global Memory (+ L1/L2 caches)

Memory hierarchy (fast → slow): Registers → Shared → L1 → L2 → Global.

22. What You Should Master Before Phase 2

Be able to explain without notes:

Hardware
GPU, SM, CUDA core, Tensor Core, warp (why 32), how blocks map to SMs.

Memory
Registers, shared memory (why useful), global memory, L1/L2, coalescing.

Execution
Thread, block, grid, indexing formula, warp execution / SIMT.

AI connection
Map an LLM op → math → CUDA kernel → threads/warps → memory traffic → hardware.

Once the mental model is solid, Phase 2 can cover threads → blocks → warps → indexing → synchronization, then memory optimization.

Suggested practice order

1. Compile & run the vector-add example.
2. Print thread/block IDs.
3. Query device properties.
4. Experiment with shared-memory tile loading + __syncthreads().
5. Compare coalesced vs strided access timings.
6. Write a simple reduction (sum) using shared memory.

These short kernels give concrete experience that reinforces every concept above.