GPU Memory & Performance

GPU Memory & Performance Internals

High-performance GPU programming is not simply about launching thousands of threads.

The next level is understanding how those threads access memory, how data moves through the GPU, how much work each Streaming Multiprocessor can keep active, and where performance bottlenecks occur.

For modern AI systems, these concepts are fundamental.

A CUDA kernel can be mathematically correct and massively parallel while still performing poorly because of:

• Inefficient global-memory accesses
• Uncoalesced memory transactions
• Shared-memory bank conflicts
• Excessive register usage
• Low occupancy
• Memory latency
• Insufficient bandwidth
• Warp divergence
• Synchronization overhead
• CPU-GPU synchronization
• Poor kernel launch patterns

The central question of this course is:

Why is this GPU kernel slow, and what exactly can we change to make it faster?

GPU Memory Hierarchy

A simplified NVIDIA GPU memory hierarchy looks like this:

textCopy
1                    GPU
2                     │
3             ┌───────┴───────┐
4             │               │
5          SM 0             SM N
6             │               │
7       ┌─────┼─────┐   ┌─────┼─────┐
8       │     │     │   │     │     │
9 Registers Shared  L1  Registers Shared L1
10       │   Memory Cache     │    Memory Cache
11       │     │     │        │      │
12       └─────┴─────┴────────┴──────┘
13                     │
14                     ▼
15                  L2 Cache
16                     │
17                     ▼
18                Global Memory

Different memory spaces have different:

• Scope
• Lifetime
• Capacity
• Latency
• Bandwidth
• Access behavior

A useful conceptual hierarchy is:

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

This is a simplified performance model rather than a universal latency ranking for every access pattern.

01. GPU Memory Architecture

CUDA exposes multiple memory spaces.

The key idea is:

textCopy
1                         GPU
2                          │
3          ┌───────────────┼───────────────┐
4          │               │               │
5      Registers       Shared Memory    Global Memory
6          │               │               │
7       Thread            Block           Device

The closer data is to the executing threads, the more important it becomes to understand capacity and access constraints.

02. Global Memory

Global memory is the primary large-capacity memory space visible to CUDA kernels.

Typical allocation:

cppCopy
1float* d_data;
2
3cudaMalloc(
4    &d_data,
5    N * sizeof(float)
6);

A kernel can access it:

cppCopy
1__global__ void scale(
2    float* data,
3    int N
4) {
5    int i =
6        blockIdx.x * blockDim.x +
7        threadIdx.x;
8
9    if (i < N) {
10        data[i] *= 2.0f;
11    }
12}

Global memory is useful because it can hold large datasets.

However, it is much farther from the execution units than registers and shared memory.

Therefore, memory access patterns matter enormously.

Global Memory Access Pattern

Consider:

cppCopy
1data[i]

If neighboring threads access neighboring elements:

textCopy
1Thread 0 → data[0]
2Thread 1 → data[1]
3Thread 2 → data[2]
4Thread 3 → data[3]
5...

the GPU can generally service these accesses efficiently through coalesced memory transactions.

A poor pattern might look like:

textCopy
1Thread 0 → data[0]
2Thread 1 → data[64]
3Thread 2 → data[128]
4Thread 3 → data[192]

The exact performance impact depends on architecture, data type, cache behavior, and access pattern, but irregular access generally makes memory efficiency harder to achieve.

03. Shared Memory

Shared memory is an on-chip memory space shared by threads in the same block.

Example:

cppCopy
1__shared__ float tile[256];

Conceptually:

textCopy
1Block
2 │
3 ├── Thread 0 ──┐
4 ├── Thread 1 ──┤
5 ├── Thread 2 ──┼── Shared Memory
6 ├── Thread 3 ──┤
7 └── Thread N ──┘

This makes shared memory extremely useful for algorithms where multiple threads reuse the same data.

A typical pattern is:

cppCopy
1__global__ void example(
2    const float* input,
3    float* output
4) {
5    __shared__ float shared[256];
6
7    int tid = threadIdx.x;
8
9    shared[tid] = input[
10        blockIdx.x * blockDim.x + tid
11    ];
12
13    __syncthreads();
14
15    output[
16        blockIdx.x * blockDim.x + tid
17    ] = shared[tid];
18}

The synchronization is necessary when threads depend on data produced by other threads in the block.

Shared Memory and Data Reuse

Suppose many calculations repeatedly need the same data.

Without shared memory:

textCopy
1Global Memory
2      │
3      ├── Thread 0
4      ├── Thread 1
5      ├── Thread 2
6      └── Thread 3

With shared memory:

textCopy
1Global Memory
2      │
3      ▼
4Shared Memory
5      │
6 ┌────┼────┐
7 ▼    ▼    ▼
8T0    T1   T2

The goal is to load data once and reuse it many times.

This concept is central to optimized GEMM.

04. Registers

Registers are private to individual threads.

textCopy
1Thread 0
2 ├── Register
3 ├── Register
4 └── Register
5
6Thread 1
7 ├── Register
8 ├── Register
9 └── Register

Example:

cppCopy
1__global__ void calculation(
2    const float* input,
3    float* output
4) {
5    int i =
6        blockIdx.x * blockDim.x +
7        threadIdx.x;
8
9    float x = input[i];
10    float y = x * x;
11    float z = y + 1.0f;
12
13    output[i] = z;
14}

Variables such as:

cppCopy
1x
2y
3z

may be allocated in registers if the compiler can keep them there.

Registers provide extremely fast access, but the number available per SM is limited.

Register Pressure

Using more registers per thread can increase the amount of state each thread can maintain.

But there is a trade-off.

Suppose:

textCopy
1Kernel A
264 registers/thread
3
4Kernel B
5160 registers/thread

Kernel B may require significantly more register resources.

If register usage becomes excessive, the compiler may spill values into local memory.

Conceptually:

textCopy
1Too many variables
2        │
3        ▼
4Register pressure
5        │
6        ▼
7Register spill
8        │
9        ▼
10Local memory
11        │
12        ▼
13Potential performance loss

Therefore:

More registers can help a thread, but too many registers can reduce overall GPU efficiency.

05. Local Memory

CUDA local memory is private to a thread, but despite its name, it is not a small on-chip memory like registers.

When values cannot remain in registers, the compiler can place them in local memory.

Common causes include:

• Register spills
• Large local arrays
• Certain compiler-generated temporary storage

Conceptually:

textCopy
1Thread
2  │
3  ├── Registers
4  │
5  └── Local Memory
6          │
7          ▼
8      Device Memory

This distinction is important:

textCopy
1Local ≠ necessarily fast

"Local" primarily describes scope, not physical proximity.

You can inspect compiler resource usage with:

bashCopy
1nvcc -Xptxas=-v kernel.cu -o kernel

Compiler output can provide information about register and local-memory usage.

06. Constant Memory

Constant memory is designed for read-only data.

Example:

cppCopy
1__constant__ float coefficients[64];

Host code can copy data into it:

cppCopy
1cudaMemcpyToSymbol(
2    coefficients,
3    hostCoefficients,
4    sizeof(hostCoefficients)
5);

Kernel:

cppCopy
1__global__ void apply(
2    const float* input,
3    float* output,
4    int N
5) {
6    int i =
7        blockIdx.x * blockDim.x +
8        threadIdx.x;
9
10    if (i < N) {
11        output[i] =
12            input[i] * coefficients[0];
13    }
14}

Constant memory can be particularly effective when threads in a warp read the same constant value.

It is useful for:

• Coefficients
• Lookup constants
• Small read-only datasets
• Configuration values

07. Texture Memory

CUDA texture functionality provides specialized mechanisms for read-only data access and hardware-supported caching/filtering features.

Texture-oriented access can be useful in workloads such as:

• Image processing
• Computer vision
• Scientific visualization
• Spatial data

Modern NVIDIA architectures also provide general-purpose caching paths that make the distinction between "texture" and other read-only access patterns more nuanced than in older CUDA programming models.

For new kernels, choose the access mechanism based on the workload and current CUDA architecture rather than assuming texture memory is automatically faster.

08. L2 Cache

L2 is a GPU-wide cache shared by the SMs.

Conceptually:

textCopy
1             SM 0
2               │
3             SM 1
4               │
5             SM 2
6               │
7               ▼
8             L2 Cache
9               │
10               ▼
11          Global Memory

L2 can reduce the number of expensive accesses that reach device memory.

Consider:

textCopy
1Thread Block A
2      │
3      ▼
4    Data X
5
6Thread Block B
7      │
8      ▼
9    Data X

If the relevant data remains cache-resident, later accesses can potentially benefit from L2.

However, you should not design a kernel assuming a particular cache hit rate without measuring it.

09. Memory Coalescing

Memory coalescing is one of the most important GPU performance concepts.

Suppose a warp contains 32 threads.

Good access:

textCopy
1T0  → A[0]
2T1  → A[1]
3T2  → A[2]
4T3  → A[3]
5...
6T31 → A[31]

The threads access contiguous elements.

This is generally favorable for global-memory throughput.

A strided pattern:

textCopy
1T0  → A[0]
2T1  → A[8]
3T2  → A[16]
4T3  → A[24]
5...

can require less efficient memory transactions.

Example

Good:

cppCopy
1float value = data[index];

where:

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

Potentially poor:

cppCopy
1float value =
2    data[index * stride];

when stride is large.

The actual behavior depends on the GPU architecture, data type, cache state, and alignment.

Coalescing Mental Model

Think about a warp:

textCopy
1Warp
2 │
3 ├── T0 ──► Address A
4 ├── T1 ──► Address A + 4
5 ├── T2 ──► Address A + 8
6 └── ...

For float data, neighboring threads accessing neighboring 4-byte elements gives the GPU a favorable contiguous pattern.

This is why data layout matters so much in CUDA.

10. Shared-Memory Bank Conflicts

Shared memory is divided into banks.

Conceptually:

textCopy
1Shared Memory
2
3Bank 0
4Bank 1
5Bank 2
6Bank 3
7...

When threads access different banks efficiently, accesses can proceed in parallel.

A problematic pattern can occur when multiple threads in the same warp access different addresses that map to the same bank.

Conceptually:

textCopy
1Thread 0 ──┐
2Thread 1 ──┤
3Thread 2 ──┼──► Same bank
4Thread 3 ──┤
5Thread N ──┘

This is called a bank conflict.

Example

A simple shared-memory array:

cppCopy
1__shared__ float data[256];

can be accessed:

cppCopy
1float value = data[threadIdx.x];

For many common configurations this gives a favorable access pattern.

But multidimensional layouts can introduce conflicts.

A common optimization technique is padding.

For example:

cppCopy
1__shared__ float tile[32][33];

instead of:

cppCopy
1__shared__ float tile[32][32];

The extra column can change the mapping of addresses to banks.

The correct padding strategy depends on the access pattern and GPU architecture.

11. Occupancy

Occupancy describes how many active warps reside on an SM relative to the maximum supported active warps.

Conceptually:

textCopy
1Maximum possible active warps
2              │
3              ▼
4          Occupancy

For example:

textCopy
1Maximum = 64 warps
2Active  = 32 warps
3
4Occupancy = 50%

Higher occupancy can help hide latency.

But:

Higher occupancy does not automatically mean higher performance.

A kernel may achieve excellent performance with moderate occupancy if it has:

• Good memory access
• High instruction-level parallelism
• Efficient computation
• Low latency sensitivity

Occupancy is a resource metric, not the final performance metric.

What Limits Occupancy?

Several resources compete for SM capacity.

textCopy
1                  SM
2                  │
3        ┌─────────┼─────────┐
4        │         │         │
5     Registers Shared    Threads
6        │       Memory       │
7        └─────────┼─────────┘
8                  │
9              Active Blocks
10                  │
11                Warps

Important constraints include:

• Registers per thread
• Shared memory per block
• Threads per block
• Maximum resident blocks
• Architecture-specific limits

This is why reducing register usage or shared-memory usage can sometimes allow more blocks to become resident.

12. Register Pressure

Register pressure deserves separate attention because it is often a hidden optimization problem.

Imagine a thread computing:

cppCopy
1float a;
2float b;
3float c;
4float d;
5float e;
6float f;
7float g;
8float h;

and then maintaining many intermediate values.

The compiler may require a large number of registers.

If register demand becomes too high:

textCopy
1High register usage
2        ↓
3Fewer resident warps
4        ↓
5Lower potential occupancy

Or, if registers spill:

textCopy
1Register pressure
2        ↓
3Spilling
4        ↓
5Local-memory traffic
6        ↓
7Potential slowdown

You can inspect compiler resource usage:

bashCopy
1nvcc -Xptxas=-v kernel.cu -o kernel

Do not optimize register count blindly.

A kernel with fewer registers can sometimes run slower if reducing registers increases instruction count or memory traffic.

13. Latency

Latency is the time required for an operation to complete.

GPU memory accesses can have significant latency.

Instead of trying to eliminate every latency:

GPUs often hide latency by keeping many independent warps ready to execute.

Conceptually:

textCopy
1Warp A
2   │
3   └── Waiting for memory
4          │
5          ▼
6Warp B executes
7          │
8          ▼
9Warp C executes
10          │
11          ▼
12Warp A becomes ready

This is one reason GPU architectures rely on massive concurrency.

Latency Hiding

textCopy
1Low concurrency
2      ↓
3CPU waits
4
5High concurrency
6      ↓
7GPU switches to other ready work
8      ↓
9Latency is hidden

This is a central difference between latency-oriented and throughput-oriented architectures.

14. Bandwidth

Bandwidth describes how much data can be transferred per unit of time.

For example:

textCopy
1GB/s
2TB/s

A kernel that repeatedly loads and stores large arrays can become limited by memory bandwidth.

Suppose a kernel performs very little arithmetic:

cppCopy
1output[i] = input[i] * 2.0f;

The GPU may spend much of its time moving data rather than computing.

This can become a memory-bound workload.

Memory Bandwidth Calculation

A simple effective bandwidth estimate is:

textCopy
1Bandwidth =
2Bytes Transferred / Execution Time

For example:

textCopy
1Data moved = 1 GB
2Time       = 0.001 seconds
3
4Bandwidth = 1000 GB/s

When benchmarking real kernels, carefully count all relevant reads and writes and account for the exact execution time being measured.

15. Compute-Bound Kernels

A compute-bound kernel spends most of its time performing arithmetic or other computation rather than waiting for memory.

Conceptually:

textCopy
1Memory
2  │
3  ▼
4Data
5  │
6  ▼
7Lots of arithmetic
8  │
9  ▼
10Result

Examples may include:

• Dense mathematical transformations
• Some matrix operations
• High arithmetic-intensity kernels

A compute-bound kernel may benefit from:

• Better instruction efficiency
• Tensor Cores
• Increased arithmetic throughput
• Better instruction-level parallelism
• Reduced unnecessary operations

Arithmetic Intensity

Arithmetic intensity is approximately:

textCopy
1Arithmetic Intensity =
2Operations / Bytes Moved

High arithmetic intensity:

textCopy
1Many operations
2Few bytes moved

Low arithmetic intensity:

textCopy
1Few operations
2Many bytes moved

This is a fundamental concept in GPU performance analysis.

16. Memory-Bound Kernels

A memory-bound kernel is limited primarily by data movement.

Example:

cppCopy
1output[i] = input[i] * 2.0f;

The computation is tiny:

textCopy
11 load
21 multiplication
31 store

But the program must move the data.

Conceptually:

textCopy
1              Kernel
2                │
3        ┌───────┴───────┐
4        │               │
5     Compute          Memory
6       Low              High

If memory bandwidth is already saturated, adding more arithmetic units will not necessarily improve performance.

Optimization should instead focus on:

• Coalescing
• Reducing unnecessary loads
• Reducing unnecessary stores
• Reusing data
• Improving cache behavior
• Using shared memory where appropriate
• Reducing data movement

17. Asynchronous Execution

CUDA supports asynchronous execution between host and device and across streams.

A basic kernel launch:

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

can return control to the CPU before the GPU has finished executing.

This enables overlap.

Conceptually:

textCopy
1Time ─────────────────────────────►
2
3CPU:   Work ───── Work ───── Work
4
5GPU:       Kernel A ─────── Kernel B

The goal is to avoid unnecessary idle time.

Asynchronous Memory Transfers

Pinned host memory can be used with asynchronous copies.

Example:

cppCopy
1float* hostData;
2
3cudaMallocHost(
4    &hostData,
5    size
6);

Then:

cppCopy
1cudaMemcpyAsync(
2    deviceData,
3    hostData,
4    size,
5    cudaMemcpyHostToDevice,
6    stream
7);

Later:

cppCopy
1cudaStreamSynchronize(stream);

When used correctly, asynchronous transfers can help overlap communication and computation.

18. CUDA Streams

A stream is an ordered sequence of CUDA operations.

Example:

cppCopy
1cudaStream_t stream;
2
3cudaStreamCreate(&stream);
4
5kernel<<<
6    blocks,
7    threads,
8    0,
9    stream
10>>>();
11
12cudaStreamSynchronize(stream);
13
14cudaStreamDestroy(stream);

Multiple streams can potentially execute independent work concurrently.

textCopy
1Stream 0
2│
3├── Copy A
4├── Kernel A
5└── Copy Result A
6
7Stream 1
8│
9├── Copy B
10├── Kernel B
11└── Copy Result B

Actual overlap depends on:

• GPU capabilities
• Copy-engine availability
• Kernel resource usage
• Dependencies
• Host-memory type
• Stream semantics

Never assume that multiple streams automatically produce overlap.

Stream-Based Pipeline

A common pipeline looks like:

textCopy
1Input A ──► Copy ──► Kernel ──► Output
2Input B ──► Copy ──► Kernel ──► Output
3Input C ──► Copy ──► Kernel ──► Output

With suitable streams:

textCopy
1Time ─────────────────────────────────►
2
3Copy A   █████
4Kernel A      ███████
5Copy B      █████
6Kernel B           ███████
7Copy C            █████
8Kernel C                 ███████

The goal is to keep GPU resources and data-transfer engines busy.

19. CUDA Graphs

Launching many small kernels individually can introduce CPU-side launch overhead.

CUDA Graphs allow a sequence of operations to be captured into an executable graph.

Conceptually:

textCopy
1Without Graph
2CPU
3 │
4 ├── Launch A
5 ├── Launch B
6 ├── Launch C
7 ├── Launch D
8 └── Launch E

With a graph:

textCopy
1CPU
2 │
3 └── Launch Graph
4       │
5       ├── A
6       ├── B
7       ├── C
8       ├── D
9       └── E

A simplified example:

cppCopy
1cudaGraph_t graph;
2cudaGraphExec_t graphExec;
3
4cudaStreamBeginCapture(
5    stream,
6    cudaStreamCaptureModeGlobal
7);
8
9kernelA<<<blocks, threads, 0, stream>>>();
10kernelB<<<blocks, threads, 0, stream>>>();
11kernelC<<<blocks, threads, 0, stream>>>();
12
13cudaStreamEndCapture(
14    stream,
15    &graph
16);
17
18cudaGraphInstantiate(
19    &graphExec,
20    graph,
21    nullptr,
22    nullptr,
23    0
24);
25
26cudaGraphLaunch(
27    graphExec,
28    stream
29);
30
31cudaStreamSynchronize(stream);

CUDA Graphs are especially useful when the same execution structure repeats many times.

They can reduce CPU launch overhead and make repeated GPU workflows more efficient.

20. Profiling

Optimization without measurement is guesswork.

A good workflow is:

textCopy
1Write Kernel
2    ↓
3Verify Correctness
4    ↓
5Benchmark
6    ↓
7Profile
8    ↓
9Find Bottleneck
10    ↓
11Optimize
12    ↓
13Benchmark Again

Important NVIDIA profiling tools include:

• Nsight Systems
• Nsight Compute
• NVIDIA command-line profiling tools
• CUDA event timing

Use the official NVIDIA Nsight documentation and Nsight Compute documentation when learning profiling workflows.

CUDA Event Timing

For kernel timing:

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

Run enough iterations to obtain stable measurements.

For repeated benchmarks, warm-up iterations are often useful because the first execution can include one-time overheads.

Nsight Systems vs Nsight Compute

These tools answer different questions.

Nsight Systems

Useful for understanding the application timeline.

textCopy
1CPU
2│
3├── Data preparation
4├── Kernel launch
5├── Synchronization
6└── Other work
7
8GPU
9│
10├── Kernel A
11├── Memory Copy
12├── Kernel B
13└── Kernel C

It helps answer:

Where is time being spent across the whole system?

Nsight Compute

Provides detailed kernel-level analysis.

It helps investigate:

• Memory throughput
• Compute throughput
• Occupancy
• Warp behavior
• Cache metrics
• Instruction statistics
• Memory transactions
• Register usage

It helps answer:

Why is this specific kernel slow?

Roofline Performance Model

The Roofline model provides a useful way to reason about whether a kernel is compute-bound or memory-bound.

Conceptually:

textCopy
1Performance
2   │
3   │                 Compute Roof
4   │              ─────────────────
5   │            /
6   │          /
7   │        /
8   │      /
9   │    /
10   │  /
11   │ /
12   └──────────────────────────────►
13          Arithmetic Intensity

There are two important limits:

textCopy
1Memory Bandwidth
2       +
3Compute Throughput

A kernel's achievable performance is constrained by both.

The basic relationship is:

textCopy
1Performance ≤ min(
2    Peak Compute,
3    Arithmetic Intensity × Memory Bandwidth
4)

This is not a guarantee of achieved performance; it is a useful upper-bound model.

Naive GEMM

GEMM stands for:

General Matrix-Matrix Multiplication

The operation is commonly represented as:

textCopy
1C = A × B

For square matrices:

textCopy
1A[N][N]
2B[N][N]
3C[N][N]

A naive implementation is:

cppCopy
1__global__ void naiveGEMM(
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:

cppCopy
1dim3 threads(16, 16);
2
3dim3 blocks(
4    (N + 15) / 16,
5    (N + 15) / 16
6);
7
8naiveGEMM<<<blocks, threads>>>(
9    A,
10    B,
11    C,
12    N
13);

This is easy to understand but performs redundant global-memory accesses.

Why Naive GEMM Is Inefficient

Suppose many threads need the same data from matrix A.

Without tiling:

textCopy
1Global Memory
2     │
3     ├── Thread 0
4     ├── Thread 1
5     ├── Thread 2
6     └── Thread 3

The same data can be fetched repeatedly.

The optimization idea is:

Load reusable data into faster on-chip memory.

That leads to tiled GEMM.

Tiled GEMM

Divide matrices into tiles.

textCopy
1Matrix A
2
3┌───────┬───────┐
4│ Tile  │ Tile  │
5├───────┼───────┤
6│ Tile  │ Tile  │
7└───────┴───────┘

Each thread block loads a tile into shared memory.

Conceptually:

textCopy
1Global Memory
2      │
3      ▼
4Shared Memory Tile
5      │
6      ├── Thread 0
7      ├── Thread 1
8      ├── Thread 2
9      └── Thread N

This allows multiple threads to reuse the loaded values.

Shared-Memory GEMM

A basic tiled implementation:

cppCopy
1#define TILE 16
2
3__global__ void tiledGEMM(
4    const float* A,
5    const float* B,
6    float* C,
7    int N
8) {
9    __shared__ float As[TILE][TILE];
10    __shared__ float Bs[TILE][TILE];
11
12    int row =
13        blockIdx.y * TILE +
14        threadIdx.y;
15
16    int col =
17        blockIdx.x * TILE +
18        threadIdx.x;
19
20    float sum = 0.0f;
21
22    for (int tile = 0;
23         tile < (N + TILE - 1) / TILE;
24         tile++) {
25
26        int aCol = tile * TILE + threadIdx.x;
27        int bRow = tile * TILE + threadIdx.y;
28
29        As[threadIdx.y][threadIdx.x] =
30            (row < N && aCol < N)
31            ? A[row * N + aCol]
32            : 0.0f;
33
34        Bs[threadIdx.y][threadIdx.x] =
35            (bRow < N && col < N)
36            ? B[bRow * N + col]
37            : 0.0f;
38
39        __syncthreads();
40
41        for (int k = 0; k < TILE; k++) {
42            sum +=
43                As[threadIdx.y][k] *
44                Bs[k][threadIdx.x];
45        }
46
47        __syncthreads();
48    }
49
50    if (row < N && col < N) {
51        C[row * N + col] = sum;
52    }
53}

The major optimization is:

textCopy
1Global Memory
2      ↓
3Shared Memory
4      ↓
5Repeated reuse
6      ↓
7Arithmetic

Instead of repeatedly fetching data from global memory, threads cooperate to load tiles.

Register-Tiled GEMM

Shared-memory tiling is only one level of optimization.

The next step is register tiling.

Instead of each thread computing only one output element:

textCopy
1Thread
2  │
3  └── C[row][col]

a thread can compute a small output tile:

textCopy
1Thread
2  │
3  ├── C[row][col]
4  ├── C[row][col+1]
5  ├── C[row+1][col]
6  └── C[row+1][col+1]

The intermediate values can stay in registers.

Conceptually:

textCopy
1Global Memory
2      ↓
3Shared Memory Tile
4      ↓
5Register Tile
6      ↓
7Multiple C values

This increases data reuse and arithmetic intensity.

However, register usage increases.

Therefore:

textCopy
1Register tiling
2      ↓
3More reuse
4      +
5More registers
6      ↓
7Potential register pressure

This is a classic GPU optimization trade-off.

Tensor Core GEMM

Modern NVIDIA GPUs contain specialized Tensor Cores designed for matrix and tensor operations.

Instead of relying only on conventional CUDA arithmetic pipelines:

textCopy
1CUDA Threads
2     ↓
3Standard arithmetic

Tensor Core execution can accelerate supported matrix operations:

textCopy
1CUDA Program
2     ↓
3Tensor Core instructions
4     ↓
5Matrix computation

Tensor Core workloads commonly use lower-precision formats such as:

• FP16
• BF16
• TF32
• FP8

depending on GPU architecture and operation.

For production AI workloads, frameworks and libraries such as NVIDIA's optimized matrix-multiplication libraries can expose these capabilities without requiring every application developer to write raw Tensor Core instructions.

The important progression is:

textCopy
1Naive GEMM
2    ↓
3Tiled GEMM
4    ↓
5Shared-memory reuse
6    ↓
7Register tiling
8    ↓
9Tensor Core execution

GEMM Optimization Journey

The entire project progression can be understood as reducing unnecessary data movement and increasing useful computation.

textCopy
1Naive GEMM
2│
3│ Repeated global-memory access
4▼
5Tiled GEMM
6│
7│ Shared-memory reuse
8▼
9Shared-memory GEMM
10│
11│ Better data reuse
12▼
13Register-tiled GEMM
14│
15│ More computation per thread
16▼
17Tensor Core GEMM
18│
19│ Specialized matrix hardware
20▼
21High-performance GEMM

The goal is not simply to make the code more complicated.

The goal is:

textCopy
1More useful work
2        +
3Less unnecessary data movement
4        +
5Better hardware utilization

Memory-Bound vs Compute-Bound Example

Consider two kernels.

Kernel A

cppCopy
1output[i] = input[i] + 1.0f;

Very little arithmetic.

Likely concern:

textCopy
1Memory bandwidth

Kernel B

cppCopy
1for (int k = 0; k < 1000; k++) {
2    value =
3        value * factor +
4        bias;
5}

A lot of arithmetic.

Likely concern:

textCopy
1Compute throughput

The correct optimization strategy differs.

textCopy
1Memory-bound
2     ↓
3Optimize data movement
4
5Compute-bound
6     ↓
7Optimize computation

Performance Investigation Workflow

When a kernel is slow, do not immediately rewrite it.

Use this process:

textCopy
11. Verify correctness
2        ↓
32. Measure runtime
4        ↓
53. Identify workload size
6        ↓
74. Determine memory vs compute behavior
8        ↓
95. Profile
10        ↓
116. Inspect bottleneck
12        ↓
137. Change one thing
14        ↓
158. Benchmark again

For example:

textCopy
1Kernel = vectorAdd
2
3Observation:
4Low arithmetic intensity
5
6        ↓
7
8Profile:
9High memory throughput
10
11        ↓
12
13Diagnosis:
14Memory-bound
15
16        ↓
17
18Optimization:
19Improve access pattern
20and reduce unnecessary transfers

Performance Checklist

Before calling a CUDA kernel optimized, ask:

Memory

• Are global-memory accesses coalesced?
• Are unnecessary loads being performed?
• Are unnecessary stores being performed?
• Is data reused?
• Can shared memory improve reuse?
• Are cache effects understood?

Shared Memory

• Are there bank conflicts?
• Is synchronization necessary?
• Is shared-memory usage limiting occupancy?

Registers

• How many registers does each thread use?
• Is register pressure reducing active blocks?
• Are register spills occurring?

Occupancy

• How many warps are active?
• Is occupancy limited by registers?
• Is occupancy limited by shared memory?
• Is increasing occupancy actually necessary?

Computation

• Is the kernel compute-bound?
• Is arithmetic intensity high enough?
• Can instructions be reduced?
• Can specialized hardware be used?

Execution

• Can CPU and GPU work overlap?
• Can streams improve concurrency?
• Can repeated work use CUDA Graphs?
• Are unnecessary synchronizations present?

Profiling

• Have you measured the kernel?
• Have you profiled it?
• Do you know the actual bottleneck?

Common Optimization Mistakes

Optimizing Before Measuring

Bad workflow:

textCopy
1Guess
2 ↓
3Rewrite
4 ↓
5Assume faster

Better:

textCopy
1Measure
2 ↓
3Profile
4 ↓
5Identify bottleneck
6 ↓
7Optimize
8 ↓
9Measure again

Assuming Higher Occupancy Is Always Better

Incorrect:

textCopy
1100% occupancy = maximum performance

Not necessarily.

Occupancy helps provide latency hiding, but performance depends on the complete execution profile.

Using Shared Memory Everywhere

Shared memory is powerful, but it is not automatically beneficial.

Using it can introduce:

• Synchronization
• Extra instructions
• Bank conflicts
• Resource pressure

Use shared memory when it provides meaningful data reuse or another measurable benefit.

Ignoring Register Pressure

A kernel can become slower after aggressive register tiling if register usage becomes excessive.

Always measure.

GPU Optimization Mental Model

The most important mental model for this course is:

textCopy
1                     GPU Performance
2                           │
3          ┌────────────────┼────────────────┐
4          │                │                │
5       Memory           Compute         Concurrency
6          │                │                │
7     Bandwidth         FLOPs/s         Occupancy
8     Coalescing        Instructions     Warps
9     Cache             Tensor Cores     Latency hiding
10     Reuse
11          │                │                │
12          └────────────────┼────────────────┘
13                           │
14                           ▼
15                     Kernel Runtime

Every optimization should answer:

Which bottleneck am I reducing?

Frontier AI Engineering Connection

These concepts become especially important in modern AI workloads.

Large AI systems repeatedly execute operations such as:

textCopy
1Matrix Multiplication
2       ↓
3Attention
4       ↓
5Normalization
6       ↓
7Activation
8       ↓
9Matrix Multiplication
10       ↓
11Memory Movement

A simplified AI workload might look like:

textCopy
1Input
2  │
3  ▼
4Embedding
5  │
6  ▼
7QKV Projection
8  │
9  ▼
10Attention
11  │
12  ▼
13MLP
14  │
15  ▼
16Output

Many of these operations are dominated by combinations of:

textCopy
1Compute throughput
2+
3Memory bandwidth
4+
5Data movement
6+
7Kernel launch overhead
8+
9Specialized hardware

Understanding GPU internals allows you to reason about why optimized AI libraries and kernels are structured the way they are.

Course Project Roadmap

Your practical progression should be:

textCopy
1Project 1
2Naive GEMM
3    │
4    ▼
5Project 2
6Tiled GEMM
7    │
8    ▼
9Project 3
10Shared-Memory GEMM
11    │
12    ▼
13Project 4
14Register-Tiled GEMM
15    │
16    ▼
17Project 5
18Tensor Core GEMM
19    │
20    ▼
21Project 6
22Profile and Compare

For every implementation, record:

textCopy
1Matrix Size
2Kernel Time
3Memory Time
4GFLOP/s
5Effective Bandwidth
6Occupancy
7Register Usage
8Shared Memory Usage

Then compare:

Final Mental Model

GPU performance is fundamentally about moving data efficiently and keeping the hardware doing useful work.

Remember this hierarchy:

textCopy
1Registers
2    ↓
3Shared Memory
4    ↓
5L1 / Cache
6    ↓
7L2 Cache
8    ↓
9Global Memory

And remember these relationships:

textCopy
1Memory Access
2      ↓
3Bandwidth
4      ↓
5Data Movement
6      ↓
7Kernel Performance

textCopy
1Registers
2      ↓
3Register Pressure
4      ↓
5Occupancy / Spills
6      ↓
7Kernel Performance

textCopy
1Arithmetic Intensity
2      ↓
3Memory-bound or Compute-bound
4      ↓
5Optimization Strategy

And finally:

textCopy
1Measure
2   ↓
3Profile
4   ↓
5Find Bottleneck
6   ↓
7Optimize
8   ↓
9Measure Again

That workflow is more important than memorizing individual CUDA optimization tricks.

What You Should Be Able to Do

After completing Course 9, you should be able to:

• Explain the CUDA memory hierarchy.
• Distinguish global, shared, register, local, constant, and texture-oriented memory.
• Explain L2 cache behavior at a high level.
• Design coalesced global-memory accesses.
• Identify shared-memory bank conflicts.
• Explain occupancy and its limitations.
• Diagnose register pressure.
• Understand register spilling.
• Distinguish latency from bandwidth.
• Identify memory-bound kernels.
• Identify compute-bound kernels.
• Calculate basic arithmetic intensity.
• Understand asynchronous execution.
• Use CUDA streams.
• Understand when CUDA Graphs are useful.
• Benchmark CUDA kernels correctly.
• Use Nsight Systems for timeline analysis.
• Use Nsight Compute for kernel-level analysis.
• Implement naive GEMM.
• Implement tiled shared-memory GEMM.
• Understand register tiling.
• Understand the role of Tensor Cores.
• Reason about GPU performance from first principles.

The ultimate goal is to stop thinking of the GPU as simply "a faster CPU" and start thinking in terms of:

textCopy
1Threads
2   ↓
3Warps
4   ↓
5SM Resources
6   ↓
7Memory Hierarchy
8   ↓
9Data Reuse
10   ↓
11Arithmetic Intensity
12   ↓
13Hardware Utilization
14   ↓
15Performance

That mental model is the foundation for writing and optimizing GPU kernels used in high-performance computing and modern AI systems.