GPU Memory & Performance Internals

High-performance CUDA programming is not only about launching thousands of GPU threads. The real performance gains come from understanding how data moves through GPU memory, how threads reuse data, how SM resources are allocated, and where a kernel becomes limited.

A CUDA kernel can be mathematically correct and still be extremely slow.

For example:

cppCopy
1__global__ void vectorAdd(
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}

This kernel is simple, but its performance depends heavily on:

textCopy
1Global-memory access
2        ↓
3Memory transactions
4        ↓
5Memory bandwidth
6        ↓
7GPU utilization
8        ↓
9Execution time

The objective of this course is to understand why kernels perform the way they do and how to optimize them systematically.

GPU Memory Hierarchy

CUDA exposes several memory spaces with different scopes and performance characteristics.

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

A useful conceptual model is:

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

The closer memory is to the executing thread, the smaller and more specialized it generally is.

Global Memory

Global memory is the large device memory used to store arrays, matrices, model weights, tensors, and other datasets.

Host code can allocate device memory:

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

Data can then be copied:

cppCopy
1cudaMemcpy(
2    d_data,
3    h_data,
4    N * sizeof(float),
5    cudaMemcpyHostToDevice
6);

A kernel can access the allocation:

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

Launch:

cppCopy
1int threads = 256;
2int blocks = (N + threads - 1) / threads;
3
4scale<<<blocks, threads>>>(d_data, N);

Why Global Memory Can Be Slow

Imagine a warp accessing:

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

This is a favorable access pattern because neighboring threads access neighboring memory locations.

Now consider:

textCopy
1Thread 0 → data[0]
2Thread 1 → data[32]
3Thread 2 → data[64]
4Thread 3 → data[96]

The access pattern is much more spread out.

This can increase the number of memory transactions required.

That leads to one of the most important CUDA concepts:

Memory access pattern matters as much as the amount of memory being accessed.

Memory Coalescing

A warp typically contains multiple threads executing the same instruction together.

For a simple array operation:

cppCopy
1int i = blockIdx.x * blockDim.x + threadIdx.x;
2
3float x = data[i];

neighboring threads access:

textCopy
1T0  → data[0]
2T1  → data[1]
3T2  → data[2]
4T3  → data[3]
5...

This is called a coalesced access pattern.

A practical example:

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

Each thread processes one adjacent element.

This is usually a good starting point for memory efficiency.

Strided Memory Access

Consider:

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

With:

textCopy
1stride = 1

the accesses are:

textCopy
10, 1, 2, 3, 4, ...

With:

textCopy
1stride = 16

they become:

textCopy
10, 16, 32, 48, 64, ...

Large strides can make global-memory accesses less efficient.

Shared Memory

Shared memory is accessible by all threads belonging to the same block.

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

The important property is:

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

This makes shared memory useful for data reuse.

Why Data Reuse Matters

Suppose 32 threads repeatedly need the same values.

Without shared memory:

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

With shared memory:

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

The data can be loaded once and reused.

This is one of the core ideas behind optimized matrix multiplication.

Synchronization

When threads cooperate through shared memory, synchronization is often necessary.

cppCopy
1__shared__ float data[256];
2
3data[threadIdx.x] = input[index];
4
5__syncthreads();
6
7float value = data[(threadIdx.x + 1) % 256];

Without synchronization, one thread could read data before another thread has finished writing it.

Think of:

textCopy
1Thread writes
2      ↓
3__syncthreads()
4      ↓
5Thread reads

The barrier establishes the required ordering for participating threads in the block.

Registers

Registers are private to individual threads.

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

Variables such as:

textCopy
1x
2y
3z

may be stored in registers when the compiler can keep them there.

Conceptually:

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

Registers are extremely valuable because they provide very fast access.

However, each SM has finite register resources.

Register Pressure

Consider a kernel with many intermediate variables:

cppCopy
1float a0, a1, a2, a3;
2float b0, b1, b2, b3;
3float c0, c1, c2, c3;
4float d0, d1, d2, d3;

If the compiler requires many registers per thread:

textCopy
1More registers/thread
2        ↓
3Fewer threads may reside simultaneously
4        ↓
5Potential occupancy reduction

If register requirements become too large, some values can spill into local memory.

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

You can inspect compiler resource usage:

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

Look for register and local-memory usage in the compiler output.

Local Memory

Despite its name, local memory should not be confused with fast on-chip memory.

It is private to a thread, but accesses can be backed by device memory.

Potential causes include:

• Register spills
• Large local arrays
• Compiler-generated temporary storage

For example:

cppCopy
1__global__ void localArray(
2    float* output
3) {
4    float values[64];
5
6    values[threadIdx.x % 64] = 1.0f;
7
8    output[threadIdx.x] =
9        values[threadIdx.x % 64];
10}

Large local arrays may increase register pressure or local-memory usage.

Therefore:

textCopy
1Local scope
2≠
3Fast memory

Constant Memory

Constant memory is designed for read-only data.

Declare:

cppCopy
1__constant__ float coefficients[64];

Copy data from the host:

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

Access it inside a kernel:

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

Constant memory can be effective when many threads read the same constant value.

L2 Cache

L2 is a GPU-wide cache shared across SMs.

textCopy
1SM 0 ─┐
2SM 1 ─┤
3SM 2 ─┼──► L2 Cache ──► Global Memory
4SM 3 ─┤
5SM N ─┘

If data is reused and remains cache-resident, subsequent accesses can benefit from caching.

However, never assume that data will always be in L2.

Profile actual behavior when performance matters.

Shared-Memory Bank Conflicts

Shared memory is organized into banks.

Conceptually:

textCopy
1Bank 0
2Bank 1
3Bank 2
4Bank 3
5...

A favorable pattern:

textCopy
1Thread 0 → Bank 0
2Thread 1 → Bank 1
3Thread 2 → Bank 2
4Thread 3 → Bank 3

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

textCopy
1Thread 0 ─┐
2Thread 1  ├──► Bank 0
3Thread 2  │
4Thread 3 ─┘

This is called a bank conflict.

For matrix tiles, padding can sometimes help.

Instead of:

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

you may use:

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

The extra column changes the address-to-bank mapping.

The correct padding depends on the access pattern.

Occupancy

Occupancy is the ratio of active warps on an SM to the maximum number of active warps supported by that SM.

For example:

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

Occupancy can help hide memory latency:

textCopy
1Warp A
2  ↓
3Waiting for memory
4
5Warp B
6  ↓
7Executes
8
9Warp C
10  ↓
11Executes
12
13Warp A
14  ↓
15Ready again

But:

100% occupancy does not guarantee maximum performance.

A kernel can perform very well at lower occupancy if it has sufficient instruction-level parallelism and good memory behavior.

What Limits Occupancy?

Several resources compete for each SM:

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

Important constraints include:

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

This is why simply increasing the block size is not always beneficial.

Latency vs Bandwidth

These concepts are different.

Latency

Latency is how long one operation takes to become available.

textCopy
1Request
2  ↓
3Wait
4  ↓
5Result

Bandwidth

Bandwidth is how much data can be transferred over time.

textCopy
1GB/s
2TB/s

A GPU hides latency by maintaining many ready warps.

textCopy
1Memory request
2      ↓
3Warp waits
4      ↓
5Another warp executes
6      ↓
7Another warp executes
8      ↓
9Original warp resumes

This is one reason GPUs are designed for massive parallelism.

Measuring Memory Bandwidth

A simple effective bandwidth estimate is:

textCopy
1Effective Bandwidth =
2Total Bytes Transferred / Execution Time

For example, if a kernel moves:

textCopy
14 GB

in:

textCopy
10.004 seconds

then:

textCopy
1Bandwidth = 4 GB / 0.004 s
2          = 1000 GB/s

This is a simplified calculation. Real benchmarking should carefully account for all memory operations.

Compute-Bound Kernels

A compute-bound kernel spends most of its time performing calculations.

Example:

cppCopy
1__global__ void computeHeavy(
2    float* output,
3    int n
4) {
5    int i = blockIdx.x * blockDim.x + threadIdx.x;
6
7    if (i < n) {
8        float x = output[i];
9
10        for (int j = 0; j < 1000; j++) {
11            x = x * 1.0001f + 0.0001f;
12        }
13
14        output[i] = x;
15    }
16}

There is substantial arithmetic relative to memory traffic.

Potential optimization areas include:

• Instruction efficiency
• Arithmetic throughput
• Instruction-level parallelism
• Specialized hardware
• Tensor Cores where applicable

Memory-Bound Kernels

Consider:

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

The computation is tiny:

textCopy
1Load
2 ↓
3Multiply
4 ↓
5Store

The kernel may therefore be limited by memory movement rather than arithmetic throughput.

Potential optimizations include:

• Coalesced accesses
• Reducing unnecessary loads
• Reducing unnecessary stores
• Increasing data reuse
• Better cache behavior

Arithmetic Intensity

Arithmetic intensity is approximately:

textCopy
1Arithmetic Intensity =
2Operations / Bytes Transferred

Low arithmetic intensity:

textCopy
1Few operations
2Many bytes
3       ↓
4Memory-bound tendency

High arithmetic intensity:

textCopy
1Many operations
2Few bytes
3       ↓
4Compute-bound tendency

This relationship is fundamental to GPU performance analysis.

Asynchronous Execution

CUDA supports asynchronous operations.

For example:

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

The CPU can continue executing while the GPU processes the kernel, subject to CUDA's synchronization semantics.

A useful mental model:

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

This allows CPU and GPU work to overlap.

CUDA Streams

A CUDA stream is an ordered sequence of operations.

Create one:

cppCopy
1cudaStream_t stream;
2
3cudaStreamCreate(&stream);

Launch a kernel into it:

cppCopy
1kernel<<<
2    blocks,
3    threads,
4    0,
5    stream
6>>>();

Synchronize:

cppCopy
1cudaStreamSynchronize(stream);

Destroy:

cppCopy
1cudaStreamDestroy(stream);

Multiple streams can be used:

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

Potentially:

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

Actual concurrency depends on GPU resources and dependencies.

Asynchronous Memory Copy

Pinned host memory can be used with asynchronous transfers.

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

Then:

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

Later:

cppCopy
1cudaStreamSynchronize(stream);

This can help create pipelines that overlap data movement with computation.

CUDA Graphs

Suppose an application repeatedly performs:

textCopy
1Copy
2 ↓
3Kernel A
4 ↓
5Kernel B
6 ↓
7Kernel C
8 ↓
9Copy Result

Launching every operation individually can introduce CPU-side launch overhead.

CUDA Graphs allow the execution pattern to be represented as a graph.

textCopy
1                 Graph
2                   │
3       ┌───────────┼───────────┐
4       ▼           ▼           ▼
5     Copy        Kernel A    Kernel B
6                               │
7                               ▼
8                            Kernel C

A simplified capture 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);

Then launch:

cppCopy
1cudaGraphLaunch(
2    graphExec,
3    stream
4);

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

CUDA Error Handling

Performance optimization is useless if the kernel is failing.

Always check CUDA errors during development.

A useful macro:

cppCopy
1#define CUDA_CHECK(call)                              \
2do {                                                  \
3    cudaError_t error = (call);                       \
4    if (error != cudaSuccess) {                       \
5        fprintf(                                      \
6            stderr,                                   \
7            "CUDA error: %s\n",                       \
8            cudaGetErrorString(error)                \
9        );                                            \
10        exit(EXIT_FAILURE);                           \
11    }                                                 \
12} while (0)

Use it:

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

After launching a kernel:

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

cudaGetLastError() is useful for detecting launch errors, while synchronization is useful when you need to observe execution errors from the kernel itself.

For performance builds, avoid unnecessary global synchronization because synchronization itself can be expensive.

CUDA Event Timing

Use CUDA events for GPU-side timing.

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

For reliable benchmarks, perform warm-up iterations and multiple timed iterations.

Profiling CUDA Kernels

A good optimization workflow is:

textCopy
1Write
2  ↓
3Compile
4  ↓
5Verify correctness
6  ↓
7Benchmark
8  ↓
9Profile
10  ↓
11Identify bottleneck
12  ↓
13Optimize
14  ↓
15Benchmark again

Never optimize based only on intuition.

Nsight Systems

Nsight Systems is useful for understanding the complete application timeline.

Conceptually:

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

It helps answer:

Where is my application spending its time?

Nsight Compute

Nsight Compute focuses on individual CUDA kernels.

It can help investigate:

textCopy
1Memory throughput
2Compute throughput
3Occupancy
4Registers
5Shared memory
6Cache behavior
7Warp behavior
8Instructions

It answers a more specific question:

Why is this particular kernel slow?

Naive GEMM

Now we can apply everything to matrix multiplication.

GEMM performs:

textCopy
1C = A × B

For each output element:

textCopy
1C[row][col] =
2    Σ A[row][k] × B[k][col]

A simple CUDA implementation:

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 + threads.x - 1) / threads.x,
5    (N + threads.y - 1) / threads.y
6);
7
8naiveGEMM<<<blocks, threads>>>(
9    A,
10    B,
11    C,
12    N
13);

This implementation is easy to understand but performs substantial redundant memory access.

Why Naive GEMM Is Slow

Imagine a block calculating a region of C.

Many threads need the same values from A and B.

textCopy
1             Global Memory
2                  │
3       ┌──────────┼──────────┐
4       ▼          ▼          ▼
5     Thread 0   Thread 1   Thread 2
6       │          │          │
7       └────── repeated data ──────┘

Instead of loading the same data repeatedly, we can load tiles into shared memory.

Tiled GEMM

Suppose:

cppCopy
1#define TILE_SIZE 16

Create shared-memory tiles:

cppCopy
1__shared__ float tileA[TILE_SIZE][TILE_SIZE];
2__shared__ float tileB[TILE_SIZE][TILE_SIZE];

Now:

textCopy
1Global Memory
2     │
3     ▼
4 ┌─────────┐
5 │ Tile A  │
6 └─────────┘
7
8 ┌─────────┐
9 │ Tile B  │
10 └─────────┘
11     │
12     ▼
13Shared Memory
14     │
15     ▼
16Repeated computation

A tiled GEMM implementation:

cppCopy
1#define TILE_SIZE 16
2
3__global__ void tiledGEMM(
4    const float* A,
5    const float* B,
6    float* C,
7    int N
8) {
9    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
10    __shared__ float tileB[TILE_SIZE][TILE_SIZE];
11
12    int row =
13        blockIdx.y * TILE_SIZE +
14        threadIdx.y;
15
16    int col =
17        blockIdx.x * TILE_SIZE +
18        threadIdx.x;
19
20    float sum = 0.0f;
21
22    int numTiles =
23        (N + TILE_SIZE - 1) / TILE_SIZE;
24
25    for (int t = 0; t < numTiles; t++) {
26
27        int aCol =
28            t * TILE_SIZE + threadIdx.x;
29
30        int bRow =
31            t * TILE_SIZE + threadIdx.y;
32
33        if (row < N && aCol < N) {
34            tileA[
35                threadIdx.y
36            ][
37                threadIdx.x
38            ] =
39                A[row * N + aCol];
40        } else {
41            tileA[
42                threadIdx.y
43            ][
44                threadIdx.x
45            ] = 0.0f;
46        }
47
48        if (bRow < N && col < N) {
49            tileB[
50                threadIdx.y
51            ][
52                threadIdx.x
53            ] =
54                B[bRow * N + col];
55        } else {
56            tileB[
57                threadIdx.y
58            ][
59                threadIdx.x
60            ] = 0.0f;
61        }
62
63        __syncthreads();
64
65        for (int k = 0; k < TILE_SIZE; k++) {
66            sum +=
67                tileA[threadIdx.y][k] *
68                tileB[k][threadIdx.x];
69        }
70
71        __syncthreads();
72    }
73
74    if (row < N && col < N) {
75        C[row * N + col] = sum;
76    }
77}

The important optimization is:

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

Register-Tiled GEMM

We can go one level further.

Instead of one thread computing:

textCopy
1C[row][col]

we can make each thread compute multiple output values.

For example:

textCopy
1Thread
2 ├── C0
3 ├── C1
4 ├── C2
5 └── C3

Those accumulators can remain in registers.

Conceptually:

textCopy
1Global Memory
2      ↓
3Shared Memory
4      ↓
5Registers
6      ↓
7Multiple output elements

This increases data reuse but also increases register pressure.

That trade-off is fundamental.

Tensor Core GEMM

Modern NVIDIA GPUs provide Tensor Cores for high-throughput matrix operations.

Conceptually:

textCopy
1CUDA Kernel
2    ↓
3Matrix Operation
4    ↓
5Tensor Cores
6    ↓
7High Throughput

Depending on the GPU architecture and operation, Tensor Cores support formats such as:

textCopy
1FP16
2BF16
3TF32
4FP8

Tensor Core programming can be accessed through CUDA libraries and APIs such as:

• cuBLAS
• cuBLASLt
• WMMA
• CUTLASS

For production GEMM, optimized libraries often outperform a hand-written naive implementation because they incorporate architecture-specific optimization.

GEMM Optimization Progression

The course project should follow:

textCopy
1Naive GEMM
2    ↓
3Tiled GEMM
4    ↓
5Shared-Memory GEMM
6    ↓
7Register-Tiled GEMM
8    ↓
9Tensor Core GEMM

Each step introduces a new performance idea.

Roofline Model

The Roofline model provides a useful way to reason about GPU performance.

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

The basic upper-bound relationship is:

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

If arithmetic intensity is low:

textCopy
1Memory bandwidth
2       ↓
3Likely memory-bound

If arithmetic intensity is high:

textCopy
1Compute throughput
2       ↓
3Likely compute-bound

This gives you a useful starting point for deciding where to investigate.

Performance Optimization Checklist

Before optimizing a CUDA kernel, ask:

Memory

textCopy
1□ Are global accesses coalesced?
2□ Are loads unnecessary?
3□ Are stores unnecessary?
4□ Can data be reused?
5□ Can caching help?
6□ Can shared memory help?

Shared Memory

textCopy
1□ Are there bank conflicts?
2□ Is synchronization necessary?
3□ Is shared memory limiting occupancy?

Registers

textCopy
1□ How many registers/thread?
2□ Is register pressure high?
3□ Are registers spilling?

Occupancy

textCopy
1□ How many active warps?
2□ What limits occupancy?
3□ Is higher occupancy actually useful?

Compute

textCopy
1□ Is the kernel compute-bound?
2□ Is arithmetic intensity high?
3□ Can specialized hardware help?

Execution

textCopy
1□ Can work overlap?
2□ Can CUDA streams help?
3□ Can CUDA Graphs reduce launch overhead?
4□ Are unnecessary synchronizations present?

Profiling

textCopy
1□ Have I measured the kernel?
2□ Have I profiled the kernel?
3□ Do I know the bottleneck?

Common CUDA Optimization Mistakes

Mistake 1: Assuming More Threads Means More Performance

Increasing:

cppCopy
1blockDim.x

does not automatically make a kernel faster.

Performance depends on:

• Register usage
• Shared memory
• Occupancy
• Memory behavior
• Instruction workload

Mistake 2: Using Shared Memory Everywhere

Shared memory is not automatically faster.

It introduces:

textCopy
1Synchronization
2+
3Instructions
4+
5Resource usage
6+
7Potential bank conflicts

Use it when it provides measurable benefit.

Mistake 3: Chasing 100% Occupancy

This assumption is incorrect:

textCopy
1100% occupancy
2      =
3Maximum performance

Instead:

textCopy
1Occupancy
2   +
3Memory behavior
4   +
5Instruction efficiency
6   +
7Latency hiding
8   +
9Compute throughput

together determine performance.

Mistake 4: Optimizing Without Profiling

Avoid:

textCopy
1"I think this is slow."
2        ↓
3Rewrite kernel

Use:

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

Complete GPU Performance Mental Model

Think about a CUDA kernel through several layers:

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

The central question is always:

What is limiting this kernel right now?

If the answer is memory:

textCopy
1Improve data movement.

If the answer is computation:

textCopy
1Improve arithmetic efficiency.

If the answer is launch overhead:

textCopy
1Consider batching or CUDA Graphs.

If the answer is latency:

textCopy
1Increase useful concurrency or data reuse.

If the answer is register pressure:

textCopy
1Reduce unnecessary state or change the algorithm.

Frontier AI Engineering Connection

These concepts are directly connected to modern AI workloads.

A simplified neural-network execution path might look like:

textCopy
1Input
2  ↓
3Embedding
4  ↓
5QKV Projection
6  ↓
7Attention
8  ↓
9MLP
10  ↓
11Normalization
12  ↓
13Output

Underneath these operations are large amounts of:

textCopy
1Matrix multiplication
2Memory movement
3Tensor operations
4Kernel launches
5Synchronization

For example:

textCopy
1Transformer Layer
2       │
3       ├── GEMM
4       ├── GEMM
5       ├── Attention
6       ├── Normalization
7       └── GEMM

Performance therefore depends on:

textCopy
1GPU Compute
2     +
3Memory Bandwidth
4     +
5Data Reuse
6     +
7Kernel Launch Overhead
8     +
9Tensor Cores
10     +
11Efficient Scheduling

This is why understanding CUDA memory and performance internals is an important foundation for high-performance AI engineering.

Final Project

Build and benchmark five versions of matrix multiplication:

textCopy
1Project 1
2Naive GEMM
3      ↓
4Project 2
5Tiled GEMM
6      ↓
7Project 3
8Shared-Memory GEMM
9      ↓
10Project 4
11Register-Tiled GEMM
12      ↓
13Project 5
14Tensor Core GEMM

For each implementation record:

textCopy
1Matrix Size
2Kernel Time
3GFLOP/s
4Effective Bandwidth
5Occupancy
6Registers / Thread
7Shared Memory / Block

Example benchmark table:

Do not focus only on the final number.

For every optimization, explain:

textCopy
1What changed?
2      ↓
3Why should it help?
4      ↓
5What resource changed?
6      ↓
7What did profiling show?
8      ↓
9Did runtime improve?

What You Should Be Able to Do

After completing this course, you should be able to:

• Explain the CUDA memory hierarchy.
• Use global memory efficiently.
• Design coalesced memory accesses.
• Use shared memory for data reuse.
• Identify shared-memory bank conflicts.
• Understand registers and register pressure.
• Identify register spills.
• Explain local memory.
• Understand constant memory.
• Understand L2 caching at a high level.
• Explain occupancy.
• Distinguish latency from bandwidth.
• Identify memory-bound kernels.
• Identify compute-bound kernels.
• Calculate basic arithmetic intensity.
• Use CUDA streams.
• Understand asynchronous memory transfers.
• Understand CUDA Graphs.
• Measure GPU kernel execution time.
• Use CUDA error checking.
• Use Nsight Systems for application-level profiling.
• Use Nsight Compute for kernel-level profiling.
• Implement naive GEMM.
• Implement tiled GEMM.
• Implement shared-memory GEMM.
• Understand register tiling.
• Understand Tensor Core GEMM.
• Optimize CUDA kernels based on measured bottlenecks.

The most important lesson is:

textCopy
1Measure
2   ↓
3Understand the bottleneck
4   ↓
5Change the memory/computation strategy
6   ↓
7Profile again
8   ↓
9Measure again

That is the foundation of serious CUDA performance engineering.