Phase 6 — CUDA Kernel Optimization for AI

This is the most important practical topic in your current CUDA roadmap. Here you combine everything you've learned:

text
1Warp/SIMT
2    ↓
3Memory hierarchy
4    ↓
5Coalescing
6    ↓
7Shared memory
8    ↓
9Occupancy
10    ↓
11Kernel optimization
12    ↓
13High-performance AI kernels

1. What Is Kernel Optimization?

Suppose you have:

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

The kernel works.

But working ≠ optimized.

Optimization means:

text
1Same result
2    ↓
3Less execution time
4    ↓
5Better GPU utilization
6    ↓
7Better throughput

For AI:

text
1Model
2 ↓
3CUDA kernels
4 ↓
5GPU hardware
6 ↓
7Latency / throughput

Better kernels can directly improve model performance.

2. The Golden Rule

The most important rule of GPU optimization is:

Measure first. Optimize second.

Don't see:

text
150% occupancy

and immediately try to make it:

text
1100%

Instead:

text
1Kernel
2 ↓
3Profile
4 ↓
5Find bottleneck
6 ↓
7Optimize bottleneck
8 ↓
9Profile again

3. The Main Kernel Bottlenecks

A CUDA kernel can be limited by:

text
1┌─────────────────────────┐
2│      Kernel Bottleneck  │
3├─────────────────────────┤
4│ Memory bandwidth        │
5│ Memory latency           │
6│ Compute throughput       │
7│ Register pressure        │
8│ Shared memory            │
9│ Synchronization          │
10│ Warp divergence          │
11│ Instruction dependency   │
12│ Low occupancy            │
13│ Launch overhead          │
14└─────────────────────────┘

Your job is to identify which one is actually limiting the kernel.

4. Optimization Strategy

Use this process:

text
11. Correctness
2      ↓
32. Benchmark
4      ↓
53. Profile
6      ↓
74. Identify bottleneck
8      ↓
95. Change one thing
10      ↓
116. Benchmark
12      ↓
137. Profile again

Don't make ten changes simultaneously.

Otherwise you won't know what actually improved performance.

5. Optimization #1 — Memory Coalescing

You already learned this.

Good:

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

Warp:

text
1T0 → A[0]
2T1 → A[1]
3T2 → A[2]
4...

This creates a contiguous access pattern.

Bad example:

cpp
1C[i] = A[i * 1024] + B[i * 1024];

Now:

text
1T0 → A[0]
2T1 → A[1024]
3T2 → A[2048]
4...

The memory requests are scattered.

Therefore:

text
1Good coalescing
2      ↓
3Efficient memory transactions
4      ↓
5Better bandwidth utilization

6. Optimization #2 — Reduce Global Memory Traffic

This is one of the biggest optimizations in AI kernels.

Suppose:

text
1Kernel A
2   ↓
3global memory
4   ↓
5Kernel B
6   ↓
7global memory
8   ↓
9Kernel C

You might instead combine operations:

text
1Kernel A + B + C
2       ↓
3registers/shared memory
4       ↓
5global memory

This is called:

Kernel Fusion

7. Example: Separate Kernels

Imagine:

cpp
1__global__ void add_bias(
2    float* x,
3    float* bias
4)
5{
6    int i =
7        blockIdx.x * blockDim.x
8        + threadIdx.x;
9
10    x[i] += bias[i];
11}

Then:

cpp
1__global__ void activation(
2    float* x
3)
4{
5    int i =
6        blockIdx.x * blockDim.x
7        + threadIdx.x;
8
9    x[i] =
10        fmaxf(x[i], 0.0f);
11}

The sequence becomes:

text
1Load X
2 ↓
3Add bias
4 ↓
5Store X
6 ↓
7Load X again
8 ↓
9Activation
10 ↓
11Store X

Extra global memory traffic occurs.

8. Fused Kernel

Instead:

cpp
1__global__ void fused(
2    float* x,
3    float* bias
4)
5{
6    int i =
7        blockIdx.x * blockDim.x
8        + threadIdx.x;
9
10    float value =
11        x[i] + bias[i];
12
13    value =
14        fmaxf(value, 0.0f);
15
16    x[i] = value;
17}

Now:

text
1Load
2 ↓
3Add
4 ↓
5Activation
6 ↓
7Store

The intermediate value remains in a register.

This can significantly reduce memory traffic.

9. Why Fusion Is Powerful

Without fusion:

text
1Global Memory
2     ↓
3     A
4     ↓
5Global Memory
6     ↓
7     B
8     ↓
9Global Memory

With fusion:

text
1Global Memory
2     ↓
3   Register
4     ↓
5     A
6     ↓
7     B
8     ↓
9Global Memory

This is especially important for:

text
1LayerNorm
2RMSNorm
3Bias
4Activation
5Softmax
6Elementwise operations
7Attention preprocessing

10. Optimization #3 — Use Registers Effectively

Registers are extremely fast.

Conceptually:

text
1Registers
2    ↓
3very fast

Example:

cpp
1float value = input[i];
2
3value = value * 2.0f;
4value = value + 1.0f;
5value = value * 3.0f;
6
7output[i] = value;

The compiler can keep value in a register.

That's good.

But:

text
1Too many registers/thread
2        ↓
3Reduced occupancy
4        ↓
5Possible register spilling
6        ↓
7Local-memory traffic

So the goal isn't:

text
1maximum registers

or:

text
1minimum registers

It is:

Use enough registers to avoid unnecessary memory traffic without causing harmful register pressure.

11. Register Spilling

Suppose the kernel requires:

text
1100+ registers/thread

but hardware/resource limits prevent efficient allocation.

Some values may spill to local memory.

Conceptually:

text
1Register
2   ↓
3not enough
4   ↓
5Local memory
6   ↓
7memory traffic

This can be very expensive.

So when optimizing a register-heavy kernel, check:

text
1Registers/thread
2Local memory
3Occupancy
4Performance

12. Optimization #4 — Shared Memory

Shared memory is useful when data is reused by threads in a block.

Suppose:

text
1Thread 0 ──┐
2Thread 1 ──┤
3Thread 2 ──┤
4            ↓
5        same data

Instead of repeatedly loading from global memory:

text
1Global Memory
2     ↓
3Thread 0
4Thread 1
5Thread 2

you can load a tile:

text
1Global Memory
2      ↓
3Shared Memory
4      ↓
5multiple threads reuse

13. Tiled Matrix Multiplication

A classic example:

text
1A × B = C

Without tiling, threads may repeatedly access global memory.

With tiling:

text
1Global A ──→ Shared A
2Global B ──→ Shared B
3                  ↓
4             Compute
5                  ↓
6                C

This increases data reuse.

14. Simplified Tiled GEMM

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

This is not a production GEMM implementation, but it teaches the fundamental optimization:

text
1Global Memory
2      ↓
3Shared Memory
4      ↓
5Reuse
6      ↓
7Registers
8      ↓
9Compute

15. Optimization #5 — Avoid Shared-Memory Bank Conflicts

You've already learned:

text
1Shared Memory
2     ↓
332 banks
4     ↓
5warp access
6     ↓
7bank mapping

Avoid patterns like:

cpp
1shared[threadIdx.x * 32]

when they cause multiple threads to map different addresses to the same bank.

Use layout changes/padding when appropriate:

cpp
1__shared__
2float tile[32][33];

instead of:

cpp
1__shared__
2float tile[32][32];

for classic transpose-style patterns.

16. Optimization #6 — Warp Divergence

Consider:

cpp
1if (threadIdx.x % 2 == 0)
2{
3    // work A
4}
5else
6{
7    // work B
8}

A warp contains:

text
132 threads

If some threads take one path and others take another, the warp may have to execute both paths under masking.

Conceptually:

text
1Warp
2 │
3 ├── Threads 0,2,4...
4 │       ↓
5 │     Path A
6 │
7 └── Threads 1,3,5...
8         ↓
9       Path B

This can reduce execution efficiency.

17. Better Branch Structure

Instead of arbitrary divergence:

cpp
1if (condition)
2{
3    ...
4}

try to organize work so that neighboring threads have similar control flow where practical.

For example, process data in separate regions:

text
1Region A → kernel A
2Region B → kernel B

rather than:

text
1Every warp
2 ↓
3half A
4half B

But don't blindly split kernels—the extra launch and memory costs may outweigh the benefit.

18. Optimization #7 — Synchronization

This instruction:

cpp
1__syncthreads();

is important when shared memory is being cooperatively used.

Example:

cpp
1shared[threadIdx.x] =
2    input[threadIdx.x];
3
4__syncthreads();
5
6float x =
7    shared[other_index];

Without synchronization, another thread may read before the data is ready.

But excessive synchronization can hurt performance.

Therefore:

text
1Correct synchronization
2        ↓
3required
4
5Unnecessary synchronization
6        ↓
7remove when safely possible

Never remove synchronization if doing so creates a race.

19. Optimization #8 — Occupancy

You've just learned this.

Remember:

text
1Occupancy =
2Active Warps
3────────────
4Maximum Warps

Occupancy can help hide:

text
1memory latency
2instruction latency

But:

text
1100% occupancy

is not necessarily optimal.

Look at:

text
1register usage
2memory stalls
3shared memory
4instruction throughput

together.

20. Optimization #9 — Block Size

Try different block sizes:

text
1128
2256
3512

For example:

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

versus:

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

The best choice depends on:

text
1register usage
2shared memory
3occupancy
4memory access
5warp behavior
6instruction workload

Benchmark them.

21. Optimization #10 — Reduce Kernel Launches

Suppose you have:

text
1Kernel 1
2Kernel 2
3Kernel 3
4Kernel 4
5Kernel 5

Each launch has overhead and may require synchronization/data movement between stages.

If operations can safely be fused:

text
1Kernel 1 + 2 + 3

you may reduce overhead and memory traffic.

This is one reason modern AI frameworks use sophisticated kernel fusion.

22. Optimization #11 — Increase Arithmetic Intensity

Remember:

text
1Arithmetic Intensity =
2FLOPs
3─────
4Bytes moved

Suppose:

text
110 FLOPs
2100 bytes

Then:

text
10.1 FLOP/byte

Very memory-heavy.

But if data can be reused:

text
1100 FLOPs
2100 bytes

then:

text
11 FLOP/byte

Higher arithmetic intensity can make better use of GPU compute resources.

23. Optimization #12 — Data Reuse

Suppose you load:

text
1X

from global memory.

If only one operation uses it:

text
1Load X
2 ↓
3Compute

not much reuse.

But:

text
1Load X
2 ↓
3Compute
4 ↓
5Compute
6 ↓
7Compute
8 ↓
9Compute

means one load supports many operations.

This is why:

text
1Registers
2Shared Memory
3Cache

are so important.

24. Optimization Hierarchy

A useful mental model:

text
1Global Memory
2      ↓
3     L2
4      ↓
5     L1
6      ↓
7Shared Memory
8      ↓
9Registers
10      ↓
11CUDA / Tensor Cores

The closer useful data is to computation, the lower the access latency generally tends to be—but each level has limited capacity and different semantics.

Your optimization goal is often:

text
1Reduce unnecessary movement

rather than simply:

text
1Make computation faster

25. Complete Optimization Example

Let's start with:

cpp
1__global__ void slow_kernel(
2    const float* x,
3    float* y,
4    int N
5)
6{
7    int i =
8        blockIdx.x * blockDim.x
9        + threadIdx.x;
10
11    if (i < N)
12    {
13        float a = x[i];
14
15        y[i] = a * 2.0f;
16
17        y[i] =
18            fmaxf(y[i], 0.0f);
19    }
20}

A better version keeps the intermediate value in a local variable:

cpp
1__global__ void optimized_kernel(
2    const float* x,
3    float* y,
4    int N
5)
6{
7    int i =
8        blockIdx.x * blockDim.x
9        + threadIdx.x;
10
11    if (i < N)
12    {
13        float value =
14            x[i] * 2.0f;
15
16        value =
17            fmaxf(value, 0.0f);
18
19        y[i] = value;
20    }
21}

The compiler may already generate similarly efficient code for the first version, so don't assume the second is faster.

This demonstrates a critical kernel-optimization lesson:

Source-code appearance does not tell you the final machine-code performance.

26. Compiler Optimization Matters

Compile with optimization enabled.

For example:

bash
1nvcc -O3 kernel.cu -o kernel

The compiler can perform:

text
1constant folding
2dead-code elimination
3instruction optimization
4register allocation
5common-subexpression elimination

Therefore, don't manually perform transformations that the compiler already handles unless profiling shows a real benefit.

27. Inspect the Generated Code

For advanced kernel work, you eventually need to understand:

text
1PTX
2SASS

The rough compilation pipeline is:

text
1CUDA C++
2   ↓
3PTX
4   ↓
5SASS
6   ↓
7GPU execution

You can inspect generated assembly with NVIDIA tools such as:

bash
1cuobjdump
2nvdisasm

This becomes useful when you're doing serious low-level optimization.

28. AI Kernel Optimization Example

Consider:

text
1LayerNorm

A naive implementation might:

text
1Load X
2 ↓
3calculate mean
4 ↓
5global memory
6 ↓
7load X again
8 ↓
9calculate variance
10 ↓
11global memory
12 ↓
13load X again
14 ↓
15normalize
16 ↓
17global memory

A highly optimized implementation can use:

text
1Warp reductions
2Shared memory
3Registers
4Vectorized loads
5Fused computation

Conceptually:

text
1Global Memory
2      ↓
3   Registers
4      ↓
5Warp Reduction
6      ↓
7Shared Memory
8      ↓
9Normalization
10      ↓
11Global Memory

This is the type of optimization you'll encounter in real AI kernels.

29. Attention Kernel Optimization

Attention is even more interesting.

Naive attention:

text
1Q
2 ↓
3QKᵀ
4 ↓
5Global Memory
6 ↓
7Softmax
8 ↓
9Global Memory
10 ↓
11× V
12 ↓
13Global Memory

Optimized attention tries to reduce intermediate memory traffic.

Conceptually:

text
1Q/K/V
2 ↓
3Tiling
4 ↓
5Shared Memory / Registers
6 ↓
7Compute
8 ↓
9Online Softmax
10 ↓
11Output

This is the fundamental motivation behind FlashAttention-style algorithms.

Later, this will connect directly to:

text
1CUDA kernels
2Shared memory
3Registers
4Warp programming
5Tensor Cores
6Memory bandwidth

30. Kernel Optimization Checklist

When you receive a slow CUDA kernel, ask:

text
1□ Is memory access coalesced?
2□ Am I transferring unnecessary data?
3□ Can data be reused?
4□ Can operations be fused?
5□ Are there shared-memory bank conflicts?
6□ Is there warp divergence?
7□ How many registers/thread?
8□ Is there register spilling?
9□ How much shared memory/block?
10□ What is occupancy?
11□ Are there unnecessary synchronizations?
12□ Is the block size appropriate?
13□ Is the kernel memory-bound?
14□ Is the kernel compute-bound?
15□ Can Tensor Cores help?
16□ What does Nsight say?

This is your practical checklist.

31. The Optimization Loop

Memorize this:

text
1              ┌──────────────┐
2              │ Write Kernel │
3              └──────┬───────┘
4                     ↓
5              ┌──────────────┐
6              │ Benchmark    │
7              └──────┬───────┘
8                     ↓
9              ┌──────────────┐
10              │ Profile      │
11              └──────┬───────┘
12                     ↓
13              ┌──────────────┐
14              │ Find Bottleneck
15              └──────┬───────┘
16                     ↓
17              ┌──────────────┐
18              │ Optimize     │
19              └──────┬───────┘
20                     ↓
21              ┌──────────────┐
22              │ Benchmark    │
23              └──────┬───────┘
24                     ↓
25                Improvement?
26                 /       \
27               YES        NO
28                ↓          ↓
29             Keep       Reconsider

This is how you should approach real GPU performance work.

32. Your AI CUDA Skill Stack Now

You've reached a strong foundation:

text
1CUDA AI Kernel Engineering
2│
3├── GPU Architecture             ✅
4├── CUDA Kernels                 ✅
5├── Thread Indexing              ✅
6├── Threads / Blocks / Grids     ✅
7├── Warp / SIMT                  ✅
8├── GPU Memory Hierarchy         ✅
9├── Memory Coalescing            ✅
10├── Memory Transactions          ✅
11├── Shared Memory                ✅
12├── Bank Conflicts               ✅
13├── Occupancy                    ✅
14└── Kernel Optimization          ✅

Now the remaining high-value topics are:

text
1├── CUDA Streams + Async Execution   ← NEXT
2├── Tensor Cores
3├── CUDA Profiling / Nsight
4└── Multi-GPU / NCCL

Next: CUDA Streams + Asynchronous Execution

This is the right next step because you've learned how to optimize one kernel. Now you'll learn how to optimize execution between kernels and memory transfers:

text
1CPU
2 │
3 ├──── Kernel A ───────────┐
4 │                         │
5 ├──── Memory Copy ────────┤
6 │                         ├── GPU
7 └──── Kernel B ───────────┘
8       ↓
9   CUDA Streams
10       ↓
11   Overlap work
12       ↓
13   Higher throughput

After that, we'll move to Tensor Cores, which is especially important for your AI/LLM direction.