CUDA GPU Kernel Programming: Threads, Warps & Kernels

CUDA GPU Kernel Programming

CUDA GPU programming is the transition from traditional CPU and operating-system-level programming to massively parallel computing.

A CPU usually executes a relatively small number of powerful threads, while a GPU contains thousands of lightweight execution threads designed to process large amounts of data concurrently.

NVIDIA CUDA provides a programming model where C/C++ functions called kernels are executed by many GPU threads. Threads are organized into blocks, and blocks form a grid.

The basic execution hierarchy is:

text
1CPU / Host
2    │
3    │ Launch CUDA Kernel
4    ▼
5GPU / Device
6    │
7    ├── Grid
8    │    │
9    │    ├── Block 0
10    │    │    ├── Thread 0
11    │    │    ├── Thread 1
12    │    │    └── Thread N
13    │    │
14    │    ├── Block 1
15    │    └── Block N
16    │
17    └── GPU execution

The official NVIDIA CUDA programming model is documented in the CUDA C++ Programming Guide.

What You Will Learn

By the end of this course, you will understand:

GPU architecture
CPU vs GPU execution
CUDA architecture
Host and device programming
CUDA kernels
threadIdx
blockIdx
blockDim
gridDim
Warps
SIMT execution
Thread synchronization
CUDA streams
CUDA events
Atomic operations
CUDA error handling
nvcc
CUDA Runtime API
GPU memory and execution behavior
Parallel algorithms

You will also build progressively more complex CUDA programs.

text
1Vector Addition
2      ↓
3Vector Multiplication
4      ↓
5Parallel Reduction
6      ↓
7Histogram
8      ↓
9Prefix Sum
10      ↓
11Matrix Multiplication

01. GPU Architecture

A GPU is designed for parallel workloads.

Instead of having a small number of powerful CPU cores, a GPU contains many execution units capable of running huge numbers of lightweight threads.

A simplified architecture looks like:

text
1                 GPU
2                  │
3       ┌──────────┴──────────┐
4       │                     │
5   Streaming              Streaming
6 Multiprocessor          Multiprocessor
7      (SM)                    (SM)
8       │                     │
9   ┌───┼───┐             ┌───┼───┐
10   │   │   │             │   │   │
11Threads Threads         Threads Threads

The important concept is the Streaming Multiprocessor (SM).

CUDA threads are scheduled and executed on SMs.

You should think about GPU execution at several levels:

text
1GPU
2 │
3 ├── SM
4 │    ├── Warp
5 │    │    └── Threads
6 │    │
7 │    └── Warp
8 │         └── Threads
9 │
10 └── SM
11      ├── Warp
12      └── Warp

A GPU does not simply execute one CUDA thread independently at a time. CUDA groups threads into execution groups called warps.

02. CPU vs GPU

CPU and GPU architectures are optimized for different workloads.

CPU

A CPU is optimized for:

Low-latency execution
Complex control flow
Sequential programs
Large caches
A relatively small number of powerful cores

text
1CPU
2
3Core 0
4Core 1
5Core 2
6Core 3

GPU

A GPU is optimized for:

Massive parallelism
High throughput
Data-parallel workloads
Large numbers of concurrent threads

text
1GPU
2
3Thread Thread Thread Thread
4Thread Thread Thread Thread
5Thread Thread Thread Thread
6Thread Thread Thread Thread
7...

For example, adding one million array elements is naturally parallel:

text
1C[0] = A[0] + B[0]
2C[1] = A[1] + B[1]
3C[2] = A[2] + B[2]
4...
5C[999999] = A[999999] + B[999999]

Each element can potentially be processed by a different CUDA thread.

03. CUDA Architecture

CUDA uses a host/device model.

text
1                    Application
2                        │
3                        ▼
4                  CPU / Host
5                        │
6             CUDA Runtime API
7                        │
8                        ▼
9                  GPU / Device
10                        │
11                ┌───────┴───────┐
12                │               │
13               SM              SM
14                │               │
15              Warps           Warps
16                │               │
17             Threads         Threads

The CPU is called the host.

The GPU is called the device.

The CPU controls program execution and launches GPU kernels.

The GPU executes kernels.

04. Host and Device Model

A CUDA application normally contains two types of code.

Host code

Runs on the CPU.

cpp
1int main() {
2    // CPU code
3}

Device code

Runs on the GPU.

cpp
1__global__ void kernel() {
2    // GPU code
3}

The CPU launches the GPU kernel:

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

This syntax is one of the most important parts of CUDA.

text
1kernel<<<grid, block>>>();
2       │      │
3       │      └── Threads per block
4       └───────── Number of blocks

05. Your First CUDA Kernel

Create a file:

text
1hello.cu

Example:

cpp
1#include <cstdio>
2
3__global__ void hello() {
4    printf("Hello from GPU thread!\n");
5}
6
7int main() {
8    hello<<<1, 1>>>();
9
10    cudaDeviceSynchronize();
11
12    return 0;
13}

Compile it using NVIDIA's CUDA compiler:

bash
1nvcc hello.cu -o hello

Run:

bash
1./hello

The important part is:

cpp
1hello<<<1, 1>>>();

This launches:

text
11 grid
2 └── 1 block
3      └── 1 thread

06. threadIdx

threadIdx identifies the current thread inside its block.

For one-dimensional execution:

cpp
1threadIdx.x

Example:

cpp
1#include <cstdio>
2
3__global__ void showThreadId() {
4    printf("Thread ID: %d\n", threadIdx.x);
5}
6
7int main() {
8    showThreadId<<<1, 8>>>();
9
10    cudaDeviceSynchronize();
11
12    return 0;
13}

Launch configuration:

text
11 block
2   │
3   ├── Thread 0
4   ├── Thread 1
5   ├── Thread 2
6   ├── Thread 3
7   ├── Thread 4
8   ├── Thread 5
9   ├── Thread 6
10   └── Thread 7

Each thread sees a different value of:

cpp
1threadIdx.x

07. blockIdx

blockIdx identifies the current block inside the grid.

cpp
1blockIdx.x

Example:

cpp
1#include <cstdio>
2
3__global__ void showBlockId() {
4    printf(
5        "Block: %d, Thread: %d\n",
6        blockIdx.x,
7        threadIdx.x
8    );
9}
10
11int main() {
12    showBlockId<<<4, 4>>>();
13
14    cudaDeviceSynchronize();
15
16    return 0;
17}

This creates:

text
1Grid
2│
3├── Block 0
4│   ├── Thread 0
5│   ├── Thread 1
6│   ├── Thread 2
7│   └── Thread 3
8│
9├── Block 1
10│   ├── Thread 0
11│   ├── Thread 1
12│   ├── Thread 2
13│   └── Thread 3
14│
15├── Block 2
16│
17└── Block 3

threadIdx.x tells you:

Which thread am I inside my block?

blockIdx.x tells you:

Which block am I inside the grid?

08. blockDim

blockDim contains the dimensions of the current block.

For a one-dimensional block:

cpp
1blockDim.x

Example:

cpp
1__global__ void inspectBlock() {
2    printf(
3        "Block %d has %d threads\n",
4        blockIdx.x,
5        blockDim.x
6    );
7}

If launched with:

cpp
1inspectBlock<<<4, 256>>>();

then:

cpp
1blockDim.x

is:

text
1256

The most important combination is:

cpp
1blockIdx.x
2threadIdx.x
3blockDim.x

These are used to calculate a global thread index.

cpp
1int index = blockIdx.x * blockDim.x + threadIdx.x;

This formula is fundamental to CUDA programming.

09. gridDim

gridDim tells you the dimensions of the entire grid.

Example:

cpp
1__global__ void inspectGrid() {
2    printf(
3        "Grid contains %d blocks\n",
4        gridDim.x
5    );
6}

Launch:

cpp
1inspectGrid<<<10, 256>>>();

Then:

cpp
1gridDim.x

is:

text
110

The relationship is:

text
1Grid
2 │
3 ├── blockIdx.x
4 ├── blockIdx.x
5 ├── blockIdx.x
6 └── ...

Together:

cpp
1gridDim
2blockIdx
3blockDim
4threadIdx

describe the CUDA execution hierarchy.

10. Global Thread Index

A common CUDA pattern is:

cpp
1int index =
2    blockIdx.x * blockDim.x +
3    threadIdx.x;

For example:

text
1blockIdx.x = 2
2blockDim.x = 256
3threadIdx.x = 10
4
5index = 2 × 256 + 10
6      = 522

Therefore:

text
1Block 2, Thread 10
2        ↓
3Global Thread 522

This allows thousands or millions of CUDA threads to process an array.

11. Warps

CUDA threads are grouped into warps.

On NVIDIA GPUs, a warp consists of 32 threads.

Conceptually:

text
1Block
2│
3├── Warp 0
4│   ├── Thread 0
5│   ├── Thread 1
6│   ├── ...
7│   └── Thread 31
8│
9├── Warp 1
10│   ├── Thread 32
11│   ├── ...
12│   └── Thread 63
13│
14└── ...

Warp-level execution is important for performance.

A block of:

text
1256 threads

contains:

text
1256 / 32 = 8 warps

Choosing block sizes that work well with warp organization is an important CUDA optimization technique.

12. SIMT

CUDA uses the SIMT — Single Instruction, Multiple Threads execution model.

Conceptually:

text
1          One instruction
2                │
3       ┌────────┼────────┐
4       ▼        ▼        ▼
5    Thread   Thread   Thread
6       │        │        │
7       ▼        ▼        ▼
8    Data A    Data B    Data C

A warp can execute the same instruction across multiple threads.

For example:

cpp
1C[index] = A[index] + B[index];

Many threads can execute this operation concurrently for different values of index.

Branch Divergence

Consider:

cpp
1if (index % 2 == 0) {
2    // Path A
3} else {
4    // Path B
5}

Threads inside the same warp may take different paths.

This can cause warp divergence, potentially reducing performance.

Prefer predictable and efficient control flow when optimizing CUDA kernels.

13. Synchronization

CUDA threads sometimes need to coordinate.

The basic block-level synchronization primitive is:

cpp
1__syncthreads();

Example:

cpp
1__global__ void example(float* data) {
2
3    int index = threadIdx.x;
4
5    __shared__ float sharedData[256];
6
7    sharedData[index] = data[index];
8
9    __syncthreads();
10
11    data[index] = sharedData[index] * 2.0f;
12}

The barrier means threads in the block wait until the required preceding operations have completed.

Important:

text
1Thread 0 ───────┐
2Thread 1 ───────┤
3Thread 2 ───────┼── __syncthreads()
4Thread 3 ───────┤
5Thread N ───────┘

Synchronization must be used carefully because incorrect synchronization can cause deadlocks or incorrect results.

14. CUDA Streams

A CUDA stream represents an ordered sequence of operations.

text
1Stream 0
2│
3├── Memory Copy
4├── Kernel A
5├── Kernel B
6└── Memory Copy

Multiple streams can allow independent operations to overlap when hardware and dependencies permit it.

Example:

cpp
1cudaStream_t stream;
2
3cudaStreamCreate(&stream);
4
5kernel<<<blocks, threads, 0, stream>>>();
6
7cudaStreamSynchronize(stream);
8
9cudaStreamDestroy(stream);

Streams become particularly useful when designing high-performance applications that need to overlap:

text
1CPU work
2   +
3Memory transfers
4   +
5GPU computation

15. CUDA Events

CUDA events can be used for GPU-side timing and synchronization.

Example:

cpp
1cudaEvent_t start, stop;
2
3cudaEventCreate(&start);
4cudaEventCreate(&stop);
5
6cudaEventRecord(start);
7
8kernel<<<blocks, threads>>>();
9
10cudaEventRecord(stop);
11
12cudaEventSynchronize(stop);
13
14float milliseconds = 0.0f;
15
16cudaEventElapsedTime(
17    &milliseconds,
18    start,
19    stop
20);
21
22printf(
23    "GPU time: %f ms\n",
24    milliseconds
25);
26
27cudaEventDestroy(start);
28cudaEventDestroy(stop);

This is useful when benchmarking CUDA kernels.

Do not assume that measuring only CPU wall-clock time around a kernel launch gives the actual GPU execution time, because kernel launches are generally asynchronous.

16. Atomic Operations

When multiple threads update the same memory location, a race condition can occur.

For example:

cpp
1counter++;

If thousands of threads perform this operation simultaneously, updates can be lost.

CUDA provides atomic operations.

Example:

cpp
1__global__ void increment(int* counter) {
2    atomicAdd(counter, 1);
3}

Launch:

cpp
1increment<<<100, 256>>>(counter);

Conceptually:

text
1Thread 0 ──┐
2Thread 1 ──┤
3Thread 2 ──┼── atomicAdd ──> counter
4Thread 3 ──┤
5Thread N ──┘

Atomic operations provide safe updates, but excessive atomic contention can reduce performance.

17. CUDA Error Handling

CUDA operations can fail.

A common mistake is ignoring return values.

Use a helper function to check CUDA API calls:

cpp
1#include <cstdio>
2#include <cstdlib>
3
4#define CUDA_CHECK(call)                                      \
5do {                                                          \
6    cudaError_t error = (call);                               \
7    if (error != cudaSuccess) {                               \
8        fprintf(stderr,                                       \
9                "CUDA error: %s:%d: %s\n",                   \
10                __FILE__,                                     \
11                __LINE__,                                     \
12                cudaGetErrorString(error));                   \
13        exit(EXIT_FAILURE);                                   \
14    }                                                         \
15} while (0)

Use it like:

cpp
1CUDA_CHECK(cudaMalloc(&data, size));

After launching a kernel, also check launch errors:

cpp
1kernel<<<blocks, threads>>>();
2
3CUDA_CHECK(cudaGetLastError());
4CUDA_CHECK(cudaDeviceSynchronize());

This distinguishes kernel-launch errors from errors encountered during execution.

18. nvcc

nvcc is NVIDIA's CUDA compiler driver.

A CUDA source file normally uses:

text
1.cu

Example:

bash
1nvcc vector_add.cu -o vector_add

Run:

bash
1./vector_add

You can request a specific GPU architecture when appropriate:

bash
1nvcc -arch=native vector_add.cu -o vector_add

For production projects, choose the appropriate architecture and compiler configuration for the target GPU environment rather than blindly copying compiler flags.

Check the compiler:

bash
1nvcc --version

Check the installed GPU:

bash
1nvidia-smi

19. CUDA Runtime

CUDA provides runtime APIs for interacting with the GPU.

Common functions include:

cpp
1cudaMalloc()
2cudaFree()
3cudaMemcpy()
4cudaMemcpyAsync()
5cudaDeviceSynchronize()
6cudaGetLastError()
7cudaEventCreate()
8cudaStreamCreate()

Typical memory workflow:

text
1CPU Memory
2    │
3    │ cudaMemcpy
4    ▼
5GPU Memory
6    │
7    │ Kernel
8    ▼
9GPU Result
10    │
11    │ cudaMemcpy
12    ▼
13CPU Memory

Example:

cpp
1float* deviceData;
2
3CUDA_CHECK(
4    cudaMalloc(
5        &deviceData,
6        size
7    )
8);
9
10CUDA_CHECK(
11    cudaMemcpy(
12        deviceData,
13        hostData,
14        size,
15        cudaMemcpyHostToDevice
16    )
17);

After computation:

cpp
1CUDA_CHECK(
2    cudaMemcpy(
3        hostResult,
4        deviceData,
5        size,
6        cudaMemcpyDeviceToHost
7    )
8);
9
10CUDA_CHECK(cudaFree(deviceData));

20. Project 1 — Vector Addition

Vector addition is the classic first CUDA project.

Mathematically:

text
1C[i] = A[i] + B[i]

Each element can be calculated independently.

Complete Example

cpp
1#include <cstdio>
2#include <cstdlib>
3#include <cuda_runtime.h>
4
5#define CUDA_CHECK(call)                                      \
6do {                                                          \
7    cudaError_t error = (call);                               \
8    if (error != cudaSuccess) {                               \
9        fprintf(stderr, "CUDA error: %s\n",                 \
10                cudaGetErrorString(error));                  \
11        exit(EXIT_FAILURE);                                   \
12    }                                                         \
13} while (0)
14
15__global__ void vectorAdd(
16    const float* A,
17    const float* B,
18    float* C,
19    int N
20) {
21    int i = blockIdx.x * blockDim.x + threadIdx.x;
22
23    if (i < N) {
24        C[i] = A[i] + B[i];
25    }
26}
27
28int main() {
29
30    const int N = 1 << 20;
31    const size_t size = N * sizeof(float);
32
33    float* h_A = new float[N];
34    float* h_B = new float[N];
35    float* h_C = new float[N];
36
37    for (int i = 0; i < N; i++) {
38        h_A[i] = static_cast<float>(i);
39        h_B[i] = static_cast<float>(i * 2);
40    }
41
42    float *d_A, *d_B, *d_C;
43
44    CUDA_CHECK(cudaMalloc(&d_A, size));
45    CUDA_CHECK(cudaMalloc(&d_B, size));
46    CUDA_CHECK(cudaMalloc(&d_C, size));
47
48    CUDA_CHECK(
49        cudaMemcpy(
50            d_A,
51            h_A,
52            size,
53            cudaMemcpyHostToDevice
54        )
55    );
56
57    CUDA_CHECK(
58        cudaMemcpy(
59            d_B,
60            h_B,
61            size,
62            cudaMemcpyHostToDevice
63        )
64    );
65
66    int threads = 256;
67    int blocks = (N + threads - 1) / threads;
68
69    vectorAdd<<<blocks, threads>>>(
70        d_A,
71        d_B,
72        d_C,
73        N
74    );
75
76    CUDA_CHECK(cudaGetLastError());
77    CUDA_CHECK(cudaDeviceSynchronize());
78
79    CUDA_CHECK(
80        cudaMemcpy(
81            h_C,
82            d_C,
83            size,
84            cudaMemcpyDeviceToHost
85        )
86    );
87
88    printf("C[0] = %f\n", h_C[0]);
89    printf("C[N-1] = %f\n", h_C[N - 1]);
90
91    CUDA_CHECK(cudaFree(d_A));
92    CUDA_CHECK(cudaFree(d_B));
93    CUDA_CHECK(cudaFree(d_C));
94
95    delete[] h_A;
96    delete[] h_B;
97    delete[] h_C;
98
99    return 0;
100}

Compile:

bash
1nvcc vector_add.cu -o vector_add

Run:

bash
1./vector_add

Execution Model

For:

cpp
1int threads = 256;
2int blocks = (N + threads - 1) / threads;

CUDA creates enough threads to cover all elements.

The boundary check:

cpp
1if (i < N)

prevents extra threads from accessing memory outside the array.

21. Project 2 — Vector Multiplication

The next project changes the operation:

text
1C[i] = A[i] × B[i]

Kernel:

cpp
1__global__ void vectorMultiply(
2    const float* A,
3    const float* B,
4    float* C,
5    int N
6) {
7    int i = blockIdx.x * blockDim.x + threadIdx.x;
8
9    if (i < N) {
10        C[i] = A[i] * B[i];
11    }
12}

The execution pattern remains the same.

Only the computation changes.

This demonstrates an important GPU programming concept:

Once data indexing is correct, changing the per-element operation is often straightforward.

22. Project 3 — Parallel Reduction

Reduction converts many values into fewer values.

For example:

text
1Input:
2
31 2 3 4 5 6 7 8
4
5        ↓
6
7Sum:
8
936

A CPU implementation might use:

cpp
1float sum = 0;
2
3for (int i = 0; i < N; i++) {
4    sum += data[i];
5}

A GPU reduction can divide the work among many threads.

Conceptually:

text
11 2 3 4 5 6 7 8
2│ │ │ │ │ │ │ │
3└─┴─┘ └─┴─┘ └─┴─┘ └─┴─┘
4  3     7     11    15
5
6      ↓
7
83 + 7 = 10
911 + 15 = 26
10
11      ↓
12
1336

Reduction introduces important concepts:

Shared memory
Synchronization
Tree-based computation
Warp-level optimization
Memory access patterns

Example educational reduction kernel:

cpp
1__global__ void reduceSum(
2    const float* input,
3    float* output,
4    int N
5) {
6    __shared__ float shared[256];
7
8    int tid = threadIdx.x;
9    int i = blockIdx.x * blockDim.x + tid;
10
11    shared[tid] = (i < N) ? input[i] : 0.0f;
12
13    __syncthreads();
14
15    for (int stride = blockDim.x / 2;
16         stride > 0;
17         stride /= 2) {
18
19        if (tid < stride) {
20            shared[tid] += shared[tid + stride];
21        }
22
23        __syncthreads();
24    }
25
26    if (tid == 0) {
27        output[blockIdx.x] = shared[0];
28    }
29}

Each block produces a partial sum.

A second reduction can combine those partial sums.

23. Project 4 — Histogram

A histogram counts how many values belong to each category.

Example:

text
1Input:
2
31 2 1 3 2 1
4
5Histogram:
6
70 → 0
81 → 3
92 → 2
103 → 1

A CUDA implementation can use atomic operations:

cpp
1__global__ void histogram(
2    const int* input,
3    int* bins,
4    int N,
5    int numBins
6) {
7    int i = blockIdx.x * blockDim.x + threadIdx.x;
8
9    if (i < N) {
10        int value = input[i];
11
12        if (value >= 0 && value < numBins) {
13            atomicAdd(&bins[value], 1);
14        }
15    }
16}

The atomic operation prevents multiple threads from corrupting the same histogram bin.

However, heavy contention can make this approach slower.

This project teaches the relationship between:

text
1Parallelism
2    +
3Shared data
4    +
5Atomic operations
6    +
7Contention
8    =
9Performance trade-offs

24. Project 5 — Prefix Sum

Prefix sum transforms:

text
1Input:
2
31 2 3 4

into:

text
1Output:
2
31 3 6 10

Mathematically:

text
1output[i] = input[0] + input[1] + ... + input[i]

Prefix sum is also called scan.

It is widely used in:

Parallel algorithms
Sorting
Graph processing
Stream processing
Compaction
GPU data processing

A simple sequential implementation is:

cpp
1output[0] = input[0];
2
3for (int i = 1; i < N; i++) {
4    output[i] =
5        output[i - 1] +
6        input[i];
7}

But this has sequential dependencies.

CUDA implementations use parallel scan algorithms to expose more parallelism.

This project introduces more advanced GPU algorithm design.

25. Project 6 — Matrix Multiplication

Matrix multiplication is one of the most important GPU workloads.

Given:

text
1A × B = C

each output element is:

text
1C[row][col]

calculated from a row of A and a column of B.

A simple CUDA kernel:

cpp
1__global__ void matrixMultiply(
2    const float* A,
3    const float* B,
4    float* C,
5    int N
6) {
7    int row =
8        blockIdx.y * blockDim.y +
9        threadIdx.y;
10
11    int col =
12        blockIdx.x * blockDim.x +
13        threadIdx.x;
14
15    if (row < N && col < N) {
16
17        float sum = 0.0f;
18
19        for (int k = 0; k < N; k++) {
20            sum +=
21                A[row * N + k] *
22                B[k * N + col];
23        }
24
25        C[row * N + col] = sum;
26    }
27}

Launch it with two-dimensional blocks:

cpp
1dim3 threads(16, 16);
2
3dim3 blocks(
4    (N + threads.x - 1) / threads.x,
5    (N + threads.y - 1) / threads.y
6);
7
8matrixMultiply<<<blocks, threads>>>(
9    A,
10    B,
11    C,
12    N
13);

The mapping becomes:

text
1Grid
2│
3├── Block (0,0)
4│    ├── Thread (0,0)
5│    ├── Thread (1,0)
6│    └── ...
7│
8├── Block (1,0)
9│
10├── Block (0,1)
11│
12└── ...

Each thread calculates one output matrix element.

26. One-Dimensional vs Two-Dimensional Threads

CUDA supports multidimensional execution configurations.

One-dimensional:

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

Two-dimensional:

cpp
1dim3 threads(16, 16);
2dim3 blocks(32, 32);
3
4kernel<<<blocks, threads>>>();

Three-dimensional:

cpp
1dim3 threads(8, 8, 8);
2dim3 blocks(4, 4, 4);

Corresponding thread coordinates include:

cpp
1threadIdx.x
2threadIdx.y
3threadIdx.z

and block coordinates:

cpp
1blockIdx.x
2blockIdx.y
3blockIdx.z

This is useful for:

Images
Matrices
Volumes
3D simulations

27. CUDA Memory Hierarchy

CUDA provides several memory types.

A simplified hierarchy is:

text
1Fast
2 │
3 ├── Registers
4 │
5 ├── Shared Memory
6 │
7 ├── L1 / L2 Cache
8 │
9 └── Global Memory
10 │
11Slow

Different memory spaces have different scopes.

text
1Register
2    ↓
3Individual thread
4
5Shared Memory
6    ↓
7Threads in a block
8
9Global Memory
10    ↓
11Grid / device

Example shared memory:

cpp
1__shared__ float tile[256];

Each block receives its own shared-memory region.

28. Memory Coalescing

GPU memory performance depends heavily on access patterns.

A desirable access pattern is often coalesced memory access, where neighboring threads access neighboring memory locations.

Good pattern:

text
1Thread 0 → A[0]
2Thread 1 → A[1]
3Thread 2 → A[2]
4Thread 3 → A[3]
5...

Poor access patterns can result in inefficient memory transactions.

For data-parallel CUDA code, always think about:

text
1Which thread?
2       ↓
3Which memory address?
4       ↓
5How are neighboring threads accessing memory?

29. Synchronization vs Asynchronous Execution

CUDA applications contain multiple levels of execution.

text
1CPU
2 │
3 │ launch
4 ▼
5GPU kernel
6 │
7 ├── Thread
8 ├── Thread
9 ├── Thread
10 └── Thread

A kernel launch can return control to the CPU before the GPU has completed the kernel.

Therefore:

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

does not necessarily mean:

text
1GPU finished

To wait for device completion:

cpp
1cudaDeviceSynchronize();

For streams:

cpp
1cudaStreamSynchronize(stream);

Understanding asynchronous execution is critical for both correctness and performance.

30. CUDA Programming Workflow

A typical CUDA application follows this workflow:

text
11. Initialize CPU data
2        ↓
32. Allocate GPU memory
4        ↓
53. Copy data CPU → GPU
6        ↓
74. Configure grid and blocks
8        ↓
95. Launch kernel
10        ↓
116. Check CUDA errors
12        ↓
137. Synchronize when required
14        ↓
158. Copy results GPU → CPU
16        ↓
179. Free GPU memory

In code:

cpp
1cudaMalloc()
2      ↓
3cudaMemcpy()
4      ↓
5kernel<<<...>>>()
6      ↓
7cudaGetLastError()
8      ↓
9cudaDeviceSynchronize()
10      ↓
11cudaMemcpy()
12      ↓
13cudaFree()

This workflow should become second nature.

31. Kernel Design Pattern

Most beginner CUDA kernels follow this structure:

cpp
1__global__ void kernel(
2    const float* input,
3    float* output,
4    int N
5) {
6    int index =
7        blockIdx.x * blockDim.x +
8        threadIdx.x;
9
10    if (index < N) {
11
12        // Work performed by this thread
13
14        output[index] =
15            input[index];
16    }
17}

Then:

cpp
1int threads = 256;
2
3int blocks =
4    (N + threads - 1) /
5    threads;
6
7kernel<<<blocks, threads>>>(
8    input,
9    output,
10    N
11);

Memorize the indexing pattern:

cpp
1blockIdx.x * blockDim.x + threadIdx.x

It appears throughout CUDA programming.

32. Common CUDA Mistakes

Mistake 1 — Forgetting bounds checks

Incorrect:

cpp
1output[index] = input[index];

Safer:

cpp
1if (index < N) {
2    output[index] = input[index];
3}

Mistake 2 — Ignoring errors

Bad:

cpp
1cudaMalloc(&ptr, size);

Better:

cpp
1CUDA_CHECK(cudaMalloc(&ptr, size));

Mistake 3 — Forgetting synchronization

If later code depends on GPU completion, use the appropriate synchronization mechanism.

cpp
1cudaDeviceSynchronize();

Mistake 4 — Excessive atomics

Atomics are useful but can create contention.

text
1Thousands of threads
2        │
3        ▼
4Same memory location
5        │
6        ▼
7Contention
8        │
9        ▼
10Lower performance

Mistake 5 — Ignoring memory access patterns

A mathematically correct kernel can still be slow because of poor memory access.

33. CUDA Performance Mindset

Writing CUDA code is not only about making it parallel.

You must also ask:

text
1Is the GPU busy?
2       ↓
3Are enough threads running?
4       ↓
5Are memory accesses efficient?
6       ↓
7Are warps diverging?
8       ↓
9Is synchronization excessive?
10       ↓
11Are atomics causing contention?
12       ↓
13Are CPU ↔ GPU transfers expensive?

The goal is not:

"Use the GPU."

The goal is:

"Use the GPU efficiently."

34. CPU Implementation vs CUDA Implementation

Consider vector addition.

CPU

cpp
1for (int i = 0; i < N; i++) {
2    C[i] = A[i] + B[i];
3}

CUDA

cpp
1__global__ void vectorAdd(
2    const float* A,
3    const float* B,
4    float* C,
5    int N
6) {
7    int i =
8        blockIdx.x * blockDim.x +
9        threadIdx.x;
10
11    if (i < N) {
12        C[i] = A[i] + B[i];
13    }
14}

CPU:

text
1Loop
2 │
3 ├── i = 0
4 ├── i = 1
5 ├── i = 2
6 └── ...

CUDA:

text
1Thread 0 → i = 0
2Thread 1 → i = 1
3Thread 2 → i = 2
4Thread 3 → i = 3
5...

The algorithm becomes a mapping problem:

text
1Data
2 ↓
3Threads
4 ↓
5Blocks
6 ↓
7Grid
8 ↓
9GPU

35. Recommended Project Progression

Follow the projects in this order:

text
1Level 1
2│
3└── Vector Add
4       ↓
5Level 2
6│
7└── Vector Multiply
8       ↓
9Level 3
10│
11└── Reduction
12       ↓
13Level 4
14│
15└── Histogram
16       ↓
17Level 5
18│
19└── Prefix Sum
20       ↓
21Level 6
22│
23└── Matrix Multiplication

Each project introduces a new GPU programming concept.

Project	Main Concept
Vector Add	Threads and indexing
Vector Multiply	Data parallelism
Reduction	Shared memory and synchronization
Histogram	Atomic operations
Prefix Sum	Parallel algorithms
Matrix Multiplication	2D execution and memory optimization

36. Course Capstone

The final goal is to build a small CUDA benchmark suite.

text
1CUDA Benchmark Suite
2        │
3        ├── Vector Add
4        │
5        ├── Vector Multiply
6        │
7        ├── Reduction
8        │
9        ├── Histogram
10        │
11        ├── Prefix Sum
12        │
13        └── Matrix Multiplication

For every implementation, measure:

text
1Input Size
2    ↓
3CPU Runtime
4    ↓
5GPU Runtime
6    ↓
7Speedup
8    ↓
9Memory Transfer Time
10    ↓
11Kernel Execution Time

For example:

text
1Vector Addition
2
3CPU:
4████████████████████
5
6GPU:
7████
8
9GPU Speedup:
10█████×

Do not assume the GPU will always be faster. Small workloads can be dominated by launch and memory-transfer overhead.

37. Final Mental Model

After completing this course, think about CUDA like this:

text
1                 CUDA APPLICATION
2                        │
3             ┌──────────┴──────────┐
4             │                     │
5          HOST CPU             DEVICE GPU
6             │                     │
7             │                 GRID
8             │                   │
9             │              ┌────┴────┐
10             │              │         │
11             │            BLOCK     BLOCK
12             │              │         │
13             │            WARPS     WARPS
14             │              │         │
15             │           THREADS   THREADS
16             │
17             └──── Kernel Launch ────►

The most important CUDA concepts are:

text
1Kernel
2   ↓
3Grid
4   ↓
5Block
6   ↓
7Thread
8   ↓
9Warp
10   ↓
11SIMT

And the most important indexing formula is:

cpp
1int index =
2    blockIdx.x * blockDim.x +
3    threadIdx.x;

The most important performance concepts are:

text
1Parallelism
2Memory Access
3Warps
4Occupancy
5Synchronization
6Divergence
7Atomics
8Streams
9Transfers

Once these concepts become familiar, CUDA moves from being a collection of compiler APIs to a clear execution model:

text
1CPU
2 │
3 │ Launch
4 ▼
5CUDA Kernel
6 │
7 ▼
8Grid
9 │
10 ▼
11Blocks
12 │
13 ▼
14Warps
15 │
16 ▼
17Threads
18 │
19 ▼
20Parallel Computation

What You Should Be Able to Do

After completing Course 8, you should be able to:

Explain CPU and GPU architectural differences.
Explain the CUDA host/device model.
Write and compile .cu programs.
Launch CUDA kernels.
Use threadIdx, blockIdx, blockDim, and gridDim.
Calculate global thread indices.
Explain warps and SIMT execution.
Use __syncthreads() correctly.
Allocate and transfer GPU memory.
Use CUDA streams and events.
Use atomic operations safely.
Detect CUDA errors.
Benchmark GPU kernels.
Implement vector operations.
Implement a basic reduction.
Build a histogram using atomics.
Understand prefix-sum algorithms.
Implement matrix multiplication.
Reason about GPU memory access and performance.

The next step after mastering these fundamentals is CUDA optimization: shared-memory tiling, coalesced memory access, occupancy, warp-level primitives, asynchronous copies, CUDA Graphs, profiling with Nsight, and optimized matrix/tensor computation.