CUDA in 30 Minutes — From Zero to 100x Speedup

Overview

This guide summarizes a CUDA tutorial covering GPU programming from absolute zero to an optimized tiled matrix multiplication kernel delivering ~4,600 GFLOPS — a 100x speedup over a multi-core CPU.

The core idea of CUDA: break your problem into thousands of simple, identical parallel tasks and organize them across the GPU's thread hierarchy.

[!NOTE] Source: YouTube tutorial "CUDA in 30 minutes — make your code 100x faster". All performance numbers measured on an NVIDIA A100 GPU.

---

Performance Journey

Implementation	GFLOPS	Speed vs CPU
Multi-core CPU (baseline)	~50	1x
Naive CUDA kernel	~3,000	~60x
Tiled CUDA + Shared Memory	~4,600	~100x

The jump from naive to optimized comes from fixing the memory bottleneck — not from more computation.

---

GPU Architecture Analogy

Concept	Analogy
CPU	A single world-class master chef. Amazing at complex sequential tasks.
GPU	An army of thousands of junior chefs. Each does one simple repetitive task.
CUDA Cores	Individual junior chefs (6,912 on an A100)
Threads	Chefs executing in parallel (200,000+ on A100)
Block	A team of chefs at one cooking station. They can cooperate and share quickly.
Grid	The entire kitchen (all blocks).

Rule: For a single complex task → use CPU. For 10,000 identical tasks → use GPU.

---

Memory Hierarchy

Memory	Analogy	Size	Latency	Scope
Global	Main pantry	~80 GB	400-800 cycles	All threads in grid
Shared	Workbench	~164 KB per SM	20-30 cycles	Threads within a block
Registers	Chef's hands	per-thread	1 cycle	Single thread

[!WARNING] Global memory is the performance villain. Every access costs 400-800 cycles. The entire optimization strategy is: minimize global reads by loading data into shared memory once and reusing it.

Strategy

1. Minimize global memory access — avoid long walks to the pantry
2. Load data into shared memory — one coordinated trip per team
3. Reuse cached data on the workbench for as much computation as possible

---

CUDA Programming Model

Thread Hierarchy

Grid (entire GPU)
└── Block (one SM / cooking station)
    └── Thread (single CUDA core / chef)

Unique Thread ID

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

This single line distributes work across all threads. Each thread calculates its unique global index and processes exactly one element.

---

Five-Step CUDA Pattern

Every CUDA kernel follows this exact pattern:

Host Code (CPU — the project manager)

1. Allocate GPU memory (cudaMalloc)
2. Copy data from host to device (cudaMemcpy host→device)
3. Launch kernel (kernel<<<grid, block>>>(args))
4. Copy results from device to host (cudaMemcpy device→host)
5. Free GPU memory (cudaFree)

Device Code (GPU — the kernel)

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

Performance Trap: PCIe Transfer

Bus	Bandwidth
PCIe (host ↔ GPU)	~32 GB/s
GPU Global Memory	~2,000 GB/s
GPU Shared Memory	~19,000 GB/s

[!TIP] | Minimize host↔GPU transfers. Send data once, run as many kernels as possible, copy the final result back at the very end.

---

Naive Matrix Multiplication — The Villain

Kernel

__global__ void matmul_naive(float* A, float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; k++) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

The Problem: Redundant Global Memory Reads

Issue	Impact
Each thread reads row of A and column of B independently	Every element of A is read N times from global memory
Shared memory not used	Every access costs 400-800 cycles
1024×1024 matrices	2.1 billion redundant global memory reads

Root Cause

Threads in the same block that need the same data each make independent slow trips to global memory. The GPU spends its time waiting for data, not computing.

---

Tiling Algorithm — The Hero

Core Idea

Break the dot product into small tiles. Process each tile cooperatively: 1. One coordinated team trip to load a tile from global → shared memory 2. All threads compute using the ultra-fast shared memory workbench 3. Accumulate partial results across tiles

Memory Traffic Comparison (4×4 matrices)

Approach	Global Reads
Naive (16 threads, 8 reads each)	128
Tiled (2×2 tiles)	32

4× reduction at small scale → 16× reduction at 1024×1024 (2.1B → 134M reads).

Tiled Logic in Python (algorithm, not GPU)

def matmul_tiled(A, B, tile_size):
    N = A.shape[0]
    C = np.zeros((N, N))
    for row in range(N):
        for col in range(N):
            acc = 0.0
            for phase in range(0, N, tile_size):
                # Load a small tile of data
                for k in range(phase, min(phase + tile_size, N)):
                    acc += A[row, k] * B[k, col]
            C[row, col] = acc
    return C

---

Optimized Kernel with Shared Memory

The Three-Step Pattern

for each tile:
    1. LOAD:    global → shared memory  (coordinated team fetch)
    2. SYNC:    __syncthreads()          (barrier — wait for everyone)
    3. COMPUTE: read from shared memory  (instant access)

Kernel

__global__ void matmul_tiled(float* A, float* B, float* C, int N) {
    __shared__ float A_tile[TILE_SIZE][TILE_SIZE];
    __shared__ float B_tile[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    for (int phase = 0; phase < N / TILE_SIZE; phase++) {
        // Step 1: Load tile from global → shared (coordinated)
        A_tile[threadIdx.y][threadIdx.x] = A[row * N + (phase * TILE_SIZE + threadIdx.x)];
        B_tile[threadIdx.y][threadIdx.x] = B[(phase * TILE_SIZE + threadIdx.y) * N + col];

        // Step 2: Synchronize — ensure all data is loaded
        __syncthreads();

        // Step 3: Compute from shared memory
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += A_tile[threadIdx.y][k] * B_tile[k][threadIdx.x];
        }

        // Step 4: Sync again before next tile (avoid overwriting)
        __syncthreads();
    }

    C[row * N + col] = sum;
}

Key Lines Explained

Code	Purpose
__shared__ float A_tile[TILE][TILE]	Declares the fast workbench (shared memory)
__syncthreads()	Barrier — no thread passes until all threads in the block have arrived
A_tile[threadIdx.y][threadIdx.x] = ...	Each thread grabs exactly one element — perfect coordination

[!WARNING] __syncthreads() is critical. Without it, a fast thread could read data from shared memory that a slower thread hasn't loaded yet, producing incorrect results.

---

Compilation and Execution

# Compile
nvcc matmul_naive.cu -o matmul_naive
nvcc matmul_tiled.cu -o matmul_tiled

# Run
./matmul_naive
./matmul_tiled

Output from 1024×1024×1024 matmul:

Kernel	Performance
Naive	~3,000 GFLOPS
Tiled	~4,600 GFLOPS

---

Why This Matters for AI/ML

The tiled matrix multiplication you just implemented is not a toy problem. It is the beating heart of modern AI:

• PyTorch and TensorFlow call cuBLAS (NVIDIA's hyper-optimized BLAS library)
• cuBLAS implements the exact same tiled matmul pattern you just wrote
• Every multiplication inside GPT, Stable Diffusion, AlphaGo runs on this principle

[!NOTE] | Understanding tiled matmul means you understand how modern AI actually runs so fast. It is not magic — it is a clever algorithm + massive parallelism + deep understanding of the memory hierarchy.

---

Next Steps for Higher Performance

Technique	Description
Memory Coalescing	Optimize global memory access patterns for contiguous reads
Warp-Level Optimizations	Avoid thread divergence within warps (groups of 32 threads)
CUDA Libraries	cuBLAS, cuDNN — hand-tuned by NVIDIA engineers
Async Streams	Overlap memory transfers with computation
Tensor Cores	Specialized hardware for mixed-precision matrix multiply-accumulate

These techniques take a kernel from 4,600 GFLOPS to 20,000+ GFLOPS.

---

Glossary

Term	Definition
CUDA	NVIDIA's parallel computing platform and programming model
Kernel	A function that runs on the GPU (__global__ prefix)
Thread	A single execution unit on the GPU
Block	A group of threads that share shared memory
Grid	All blocks launched for a single kernel
SM	Streaming Multiprocessor — the compute unit inside a GPU
GFLOPS	Giga FLoating Point Operations Per Second — billions of calculations per second
Global Memory	Main GPU RAM (large, slow)
Shared Memory	On-chip SRAM per SM (small, fast)
Tiling	Algorithmic technique that breaks computation into blocks to improve data reuse
Memory Bound	Performance limited by memory bandwidth, not compute speed
Compute Bound	Performance limited by compute throughput, not memory

---

Source

• Video: CUDA in 30 Minutes — YouTube
• Category: GPU Programming / CUDA