Course 12 — Advanced Kernel Engineering

This course is the frontier-level continuation of CUDA kernel programming, moving from writing individual GPU kernels to designing kernels and GPU systems that exploit modern NVIDIA architectures.

You will learn how high-performance AI frameworks and inference engines approach problems such as Tensor Core utilization, asynchronous memory movement, kernel fusion, persistent execution, multi-GPU communication, and performance modeling.

The goal is not simply to learn CUDA APIs. The goal is to understand why a kernel is fast, where it becomes bottlenecked, and how to redesign it when hardware utilization is poor.

NVIDIA CUDA C++ Programming Guide

01. Warp Specialization

A CUDA kernel normally gives every warp a similar role. Warp specialization takes a different approach: different warps can specialize in different tasks.

For example:

textCopy
1Warp 0–1 → Load data
2Warp 2–5 → Compute
3Warp 6–7 → Store results

Instead of every warp repeatedly performing:

textCopy
1load → compute → store

the kernel can organize execution into specialized producer and consumer roles.

This becomes particularly interesting for AI workloads where memory movement and computation can overlap.

Why Warp Specialization Matters

Modern GPUs have enormous compute throughput, but computation can stall when data is not available quickly enough.

A specialized design can create a pipeline:

textCopy
1Global Memory
2     ↓
3Producer Warps
4     ↓
5Shared Memory
6     ↓
7Consumer Warps
8     ↓
9Tensor Cores
10     ↓
11Output

The important concept is overlapping work rather than executing every operation sequentially.

Example Concept

cppCopy
1__global__ void specialized_kernel(float* input, float* output) {
2
3    int warp_id = threadIdx.x / 32;
4
5    if (warp_id < 2) {
6        // Producer warp
7        // Load data
8    } else {
9        // Consumer warp
10        // Process data
11    }
12}

This simple example is only conceptual. Real warp-specialized kernels require careful synchronization and resource management.

Engineering Trade-off

Warp specialization can improve utilization, but it can also reduce flexibility.

You must consider:

• register usage
• shared-memory usage
• synchronization
• occupancy
• load imbalance
• producer/consumer pipeline depth

The advanced lesson is:

Do not optimize individual instructions first. Optimize the execution pipeline.

02. Cooperative Groups

CUDA traditionally provides synchronization mechanisms such as:

cppCopy
1__syncthreads();

But complex algorithms often need synchronization between more flexible groups of threads.

CUDA Cooperative Groups provides abstractions for organizing threads into groups.

Conceptually:

textCopy
1Grid
2 ├── Block
3 │    ├── Tile
4 │    ├── Tile
5 │    └── Tile
6 └── Block
7      ├── Tile
8      └── Tile

A kernel can work with:

• thread blocks
• tiled groups
• cooperative groups
• grid-level synchronization where supported

Example:

cppCopy
1#include <cooperative_groups.h>
2
3namespace cg = cooperative_groups;
4
5__global__ void kernel(float* data) {
6
7    cg::thread_block block = cg::this_thread_block();
8
9    int tid = threadIdx.x;
10
11    data[tid] *= 2.0f;
12
13    block.sync();
14}

The value of cooperative groups is that synchronization becomes more explicit and composable.

03. Asynchronous Copies

Traditional memory operations often look conceptually like:

textCopy
1Load
2 ↓
3Wait
4 ↓
5Compute
6 ↓
7Load
8 ↓
9Wait

Advanced GPU programming attempts to transform this into:

textCopy
1Load ──────────────┐
2                   ↓
3Compute ─────── Compute
4                   ↓
5Next Load ─────────┘

The objective is to hide memory latency behind useful computation.

CUDA provides asynchronous mechanisms for moving data between memory levels.

A common conceptual pipeline is:

textCopy
1Global Memory
2      ↓
3async copy
4      ↓
5Shared Memory
6      ↓
7Compute

Instead of blocking the entire computation while waiting for data, the programmer can build a pipeline where data movement and computation overlap.

04. TMA Concepts

Tensor Memory Accelerator (TMA) is a modern NVIDIA mechanism designed to efficiently move multidimensional tensor data between global memory and shared memory.

This is particularly important for AI workloads.

Consider a matrix:

textCopy
1A[M][K]

Instead of treating the matrix as thousands of independent scalar loads, a tensor-oriented mechanism can describe the multidimensional transfer.

Conceptually:

textCopy
1Global Tensor
2     │
3     │ TMA
4     ▼
5Shared Memory Tile
6     │
7     ▼
8Tensor Core Computation

This is valuable for operations such as:

textCopy
1GEMM
2Attention
3Convolution
4LayerNorm
5MoE

Why TMA Matters

AI kernels frequently operate on tiles:

textCopy
1Large Matrix
2┌─────────────────────┐
3│ ┌────┐ ┌────┐ ┌────┐│
4│ │Tile│ │Tile│ │Tile││
5│ └────┘ └────┘ └────┘│
6│ ┌────┐ ┌────┐ ┌────┐│
7│ │Tile│ │Tile│ │Tile││
8│ └────┘ └────┘ └────┘│
9└─────────────────────┘

Efficiently transferring these tiles is fundamental to achieving high Tensor Core utilization.

05. Tensor Cores

Tensor Cores are specialized hardware units designed for high-throughput matrix operations.

Traditional CUDA arithmetic might perform:

textCopy
1A × B

using ordinary CUDA cores.

Tensor Cores instead accelerate matrix multiply-accumulate operations.

Conceptually:

textCopy
1D = A × B + C

This operation is fundamental to modern AI.

For example:

textCopy
1Transformer
2    ↓
3Linear Layer
4    ↓
5Matrix Multiplication
6    ↓
7Tensor Cores

Large portions of neural-network computation ultimately reduce to matrix operations.

Tensor Core Optimization

High performance depends on more than simply using Tensor Cores.

You need to consider:

• matrix dimensions
• data type
• memory layout
• shared-memory tiling
• register fragments
• occupancy
• instruction scheduling
• memory bandwidth

This is why a theoretically Tensor-Core-enabled kernel can still perform poorly.

06. WMMA

Warp Matrix Multiply Accumulate (WMMA) provides CUDA APIs for programming matrix operations at the warp level.

A simplified conceptual structure is:

cppCopy
1#include <mma.h>
2
3using namespace nvcuda;
4
5__global__ void matmul() {
6
7    wmma::fragment<
8        wmma::matrix_a,
9        16, 16, 16,
10        half,
11        wmma::row_major
12    > a;
13
14    // Load matrix fragment
15    // Perform matrix operation
16    // Store result
17}

The important abstraction is the fragment.

Instead of manually managing every matrix element, WMMA represents pieces of matrices that participate in Tensor Core operations.

Learning Objective

You should understand the relationship:

textCopy
1CUDA Thread
2     ↓
3Warp
4     ↓
5WMMA Fragment
6     ↓
7Tensor Core
8     ↓
9Matrix Multiply

07. MMA

MMA provides a lower-level mechanism for matrix multiply-accumulate operations.

This is where CUDA kernel programming begins moving closer to the hardware instruction level.

The conceptual progression is:

textCopy
1High Level
2    │
3    ├── PyTorch
4    │
5    ├── CUDA
6    │
7    ├── WMMA
8    │
9    └── MMA
10         ↓
11     Hardware

The lower the abstraction level, the greater the control—and the greater the complexity.

At this level, you must understand:

• operand layouts
• fragments
• register allocation
• instruction shapes
• data types
• synchronization
• architecture-specific behavior

This is the territory where serious GPU kernel engineering begins.

08. Fused Kernels

Kernel fusion combines multiple operations into one GPU kernel.

Suppose an operation performs:

textCopy
1X
2 ↓
3Bias
4 ↓
5GELU
6 ↓
7Output

A naïve implementation might launch three kernels:

textCopy
1Kernel 1 → Bias
2Kernel 2 → GELU
3Kernel 3 → Store

Each launch introduces overhead and may require intermediate data to move through memory.

A fused kernel can perform:

textCopy
1             ┌──────────────┐
2X ──────────►│ Bias + GELU  │──────────► Output
3             └──────────────┘

Why Fusion Is Powerful

Without fusion:

textCopy
1GPU Memory
2   ↑ ↓
3Kernel
4   ↑ ↓
5GPU Memory
6   ↑ ↓
7Kernel

With fusion:

textCopy
1GPU Memory
2   ↓
3Fused Kernel
4   ↓
5GPU Memory

This can reduce:

• kernel-launch overhead
• global-memory traffic
• intermediate tensors
• synchronization points

Fusion is especially valuable in:

textCopy
1Transformer blocks
2Attention
3Normalization
4Activation functions
5Elementwise operations
6Inference pipelines

09. Persistent Kernels

A conventional kernel often launches a large number of thread blocks and lets the GPU scheduler execute them.

A persistent kernel takes a different approach.

Instead of repeatedly launching work, a relatively stable set of resident blocks can remain active and process multiple units of work.

Conceptually:

textCopy
1Normal:
2
3Launch → Work → Finish
4Launch → Work → Finish
5Launch → Work → Finish
6
7
8Persistent:
9
10Launch
11  ↓
12Resident GPU Workers
13  ↓
14Work
15  ↓
16Work
17  ↓
18Work
19  ↓
20Work

Persistent kernels can be useful when:

• launch overhead matters
• workloads are iterative
• synchronization is expensive
• work can be dynamically assigned
• maintaining data on-chip is beneficial

They are advanced because resource allocation becomes much more important.

10. CUDA Graphs

Repeated GPU workloads often execute the same sequence:

textCopy
1Kernel A
2   ↓
3Kernel B
4   ↓
5Kernel C
6   ↓
7Memory operation
8   ↓
9Kernel D

Launching every operation individually introduces CPU-side overhead.

CUDA Graphs allow a sequence of operations to be captured as a graph.

textCopy
1        CUDA Graph
2┌────────────────────────┐
3│ A → B → C → D          │
4└────────────────────────┘
5             ↓
6          Replay

Instead of repeatedly constructing the same execution sequence, the graph can be replayed.

This is especially useful for:

textCopy
1AI inference
2Repeated training steps
3Serving systems
4Low-latency workloads

The important principle is:

If the execution structure is repetitive, move scheduling overhead out of the critical path.

11. Triton

CUDA gives extremely detailed control, but writing highly optimized kernels can require substantial code.

Triton provides a higher-level programming model designed for writing GPU kernels efficiently, particularly for machine-learning workloads.

A conceptual Triton kernel might look like:

pythonCopy
1import triton
2import triton.language as tl
3
4@triton.jit
5def add_kernel(
6    x_ptr,
7    y_ptr,
8    output_ptr,
9    n_elements,
10    BLOCK_SIZE: tl.constexpr
11):
12    pid = tl.program_id(0)
13
14    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
15
16    mask = offsets < n_elements
17
18    x = tl.load(x_ptr + offsets, mask=mask)
19    y = tl.load(y_ptr + offsets, mask=mask)
20
21    tl.store(output_ptr + offsets, x + y, mask=mask)

The programmer describes the computational structure while Triton's compiler handles significant portions of low-level code generation.

CUDA vs Triton

textCopy
1CUDA
2 ├── Maximum control
3 ├── Complex
4 ├── Architecture-specific optimization
5 └── Excellent for low-level systems work
6
7Triton
8 ├── Higher abstraction
9 ├── Faster kernel development
10 ├── ML-oriented
11 └── Compiler-assisted optimization

A strong kernel engineer should understand both.

12. CUDA Extensions for PyTorch

PyTorch provides highly optimized operators, but sometimes a custom operation is required.

CUDA extensions allow you to connect custom CUDA kernels with PyTorch.

Conceptually:

textCopy
1Python
2  ↓
3PyTorch
4  ↓
5C++ Extension
6  ↓
7CUDA Kernel
8  ↓
9GPU

For example:

pythonCopy
1output = my_custom_op(input)

can internally invoke:

textCopy
1C++
2  ↓
3CUDA
4  ↓
5GPU Kernel

This is useful when implementing:

• custom operators
• specialized attention
• optimized reductions
• custom quantization
• research kernels
• domain-specific operations

13. Custom Autograd Kernels

Training requires both forward and backward computation.

For an operation:

textCopy
1x → f(x) → y

training also requires:

textCopy
1dy/dx

A custom autograd operation allows you to define both.

Conceptually:

textCopy
1Forward
2x ─────► Custom CUDA Kernel ─────► y
3                                      │
4                                      ▼
5Backward
6dy ◄──── Custom CUDA Kernel ◄──── dx

This becomes important when the default PyTorch implementation is not optimal.

A custom operator should consider:

• forward correctness
• backward correctness
• numerical stability
• memory usage
• saved tensors
• gradient synchronization
• mixed precision

The objective is not merely to make the forward pass fast.

A training kernel must optimize the complete forward + backward computation.

14. Distributed Kernels

Single-GPU optimization eventually reaches a physical limit.

Large AI systems therefore distribute computation across GPUs.

textCopy
1                 Model
2                  │
3        ┌─────────┼─────────┐
4        ↓         ↓         ↓
5      GPU 0     GPU 1     GPU 2
6        │         │         │
7        └─────────┼─────────┘
8                  ↓
9             Communication

This introduces a new optimization dimension:

textCopy
1Compute
2+
3Memory
4+
5Communication

A kernel can be extremely fast locally but still produce poor application performance if communication dominates execution.

15. NCCL

NCCL is NVIDIA's communication library designed for collective communication between GPUs.

Important collective operations include:

textCopy
1AllReduce
2AllGather
3ReduceScatter
4Broadcast

For example, distributed training may require:

textCopy
1GPU 0 ──┐
2GPU 1 ──┤
3GPU 2 ──┼──► AllReduce
4GPU 3 ──┘

The result must be communicated efficiently across GPUs.

Why NCCL Matters

Modern AI systems rely heavily on communication.

A useful performance model is:

textCopy
1Total Time
2=
3Computation
4+
5Communication
6+
7Synchronization

Therefore, distributed GPU engineering is not just about making kernels faster.

It is about minimizing the end-to-end critical path.

16. Multi-GPU Memory

With multiple GPUs, memory is no longer a single uniform resource.

You may have:

textCopy
1GPU 0 Memory
2GPU 1 Memory
3GPU 2 Memory
4GPU 3 Memory

and communication paths between them.

A system may need to decide:

textCopy
1Where should weights live?
2Where should activations live?
3Where should KV cache live?
4When should tensors move?
5Which GPU should execute the operation?

For large models, these decisions can determine whether the model fits and how fast inference runs.

17. CPU-GPU Synchronization

The CPU and GPU operate asynchronously.

A simplified model is:

textCopy
1CPU
2 │
3 │ Launch
4 ▼
5GPU
6 │
7 │ Execute
8 ▼
9GPU Queue

The CPU does not necessarily wait for every operation to complete.

This is why careless synchronization can destroy performance.

For example:

cppCopy
1cudaDeviceSynchronize();

forces synchronization.

Repeated synchronization can produce:

textCopy
1CPU → GPU
2     ↓
3   Wait
4     ↓
5CPU → GPU
6     ↓
7   Wait

instead of:

textCopy
1CPU → GPU → GPU → GPU
2          parallel execution

Advanced optimization therefore asks:

Where is synchronization actually necessary?

18. Performance Modeling

Performance optimization should not rely entirely on trial and error.

Before changing a kernel, create a model.

Ask:

textCopy
1How many FLOPs?
2
3How many bytes moved?
4
5What is the arithmetic intensity?
6
7Is memory bandwidth the bottleneck?
8
9Is compute throughput the bottleneck?
10
11Is synchronization limiting performance?

This gives you a hypothesis before benchmarking.

19. Roofline Analysis

The roofline model is one of the most useful tools for understanding GPU performance.

The key metric is:

textCopy
1Arithmetic Intensity
2=
3FLOPs / Bytes Transferred

Imagine a kernel performs:

textCopy
11,000,000 FLOPs

and transfers:

textCopy
1100,000 bytes

Then:

textCopy
1Arithmetic Intensity
2=
31,000,000 / 100,000
4=
510 FLOPs/byte

The roofline model compares this intensity against hardware limits.

Conceptually:

textCopy
1Performance
2   │                 Compute Roof
3   │              ────────────────
4   │            /
5   │          /
6   │        /
7   │      /
8   │____/________________________
9       Memory
10       Bandwidth

If the kernel is memory-bound, optimizing arithmetic instructions may accomplish almost nothing.

If it is compute-bound, reducing memory traffic may have less impact than improving compute utilization.

This leads to a critical engineering rule:

First identify the bottleneck. Then optimize that bottleneck.

20. Kernel Benchmarking

The final skill is measuring whether your optimization actually worked.

Never assume:

textCopy
1fewer instructions = faster

or:

textCopy
1more occupancy = faster

or:

textCopy
1Tensor Core = maximum performance

GPU performance is affected by many interacting factors.

A useful benchmarking workflow is:

textCopy
1Baseline
2   ↓
3Measure
4   ↓
5Identify bottleneck
6   ↓
7Change one thing
8   ↓
9Measure again
10   ↓
11Compare
12   ↓
13Profile
14   ↓
15Repeat

Measure metrics such as:

textCopy
1Kernel latency
2Throughput
3Bandwidth
4FLOPs/s
5Occupancy
6SM utilization
7Memory utilization
8L2 behavior
9Register usage
10Shared-memory usage
11Launch overhead

For reliable benchmarking, avoid measuring only one execution.

Instead:

textCopy
1Warmup
2Warmup
3Warmup
4
5Run
6Run
7Run
8Run
9Run
10
11Calculate statistics

This reduces the impact of initialization and transient effects.

The Advanced Kernel Engineering Workflow

By the end of this course, your optimization process should look like this:

textCopy
1                GPU Workload
2                     │
3                     ▼
4              Understand Algorithm
5                     │
6                     ▼
7              Build Baseline
8                     │
9                     ▼
10                  Profile
11                     │
12          ┌──────────┼──────────┐
13          ▼          ▼          ▼
14       Compute     Memory    Synchronization
15       Bound       Bound         Bound
16          │          │          │
17          ▼          ▼          ▼
18      Tensor      Tiling      Async
19       Cores      Fusion      Pipeline
20          │          │          │
21          └──────────┼──────────┘
22                     ▼
23              Optimize Kernel
24                     │
25                     ▼
26              Benchmark Again
27                     │
28                     ▼
29             Validate Correctness
30                     │
31                     ▼
32             Production Kernel

This is the central philosophy of advanced GPU engineering.

Capstone: Build a Production-Grade AI Kernel

The final project should combine the concepts from the entire course.

Project

Build an optimized Transformer operation such as:

textCopy
1Fused Attention / MLP Kernel

Your implementation should progressively evolve:

textCopy
1Version 1
2Naive CUDA Kernel
3       ↓
4Version 2
5Tiled Memory Access
6       ↓
7Version 3
8Shared Memory
9       ↓
10Version 4
11Asynchronous Copies
12       ↓
13Version 5
14Tensor Core / MMA
15       ↓
16Version 6
17Kernel Fusion
18       ↓
19Version 7
20CUDA Graph Integration
21       ↓
22Version 8
23PyTorch Extension
24       ↓
25Version 9
26Benchmark + Roofline Analysis

Final Benchmark

Compare:

textCopy
1PyTorch baseline
2       vs
3Naive CUDA
4       vs
5Optimized CUDA
6       vs
7Triton

Report:

textCopy
1Latency
2Throughput
3Memory bandwidth
4FLOPs/s
5GPU utilization
6Peak memory
7Speedup

For example:

textCopy
1                    Latency
2PyTorch             100 μs
3Naive CUDA           82 μs
4Triton                61 μs
5Optimized CUDA        38 μs

The exact numbers will depend on the operation, GPU architecture, tensor shapes, data types, and implementation.

The important result is not a particular number.

The important result is that you can explain why one implementation is faster than another.

What You Should Know After Course 12

After completing Advanced Kernel Engineering, you should be able to reason across the complete GPU software stack:

textCopy
1AI Model
2   ↓
3PyTorch
4   ↓
5Custom Operator
6   ↓
7Triton / CUDA
8   ↓
9Kernel
10   ↓
11Warp
12   ↓
13Tensor Core / CUDA Core
14   ↓
15Memory Hierarchy
16   ↓
17GPU Architecture
18   ↓
19Multi-GPU System

You should understand how to:

• design warp-specialized kernels
• use cooperative groups
• build asynchronous pipelines
• understand TMA-based data movement
• program Tensor Cores
• use WMMA and MMA concepts
• fuse GPU operations
• design persistent kernels
• use CUDA Graphs
• write Triton kernels
• build CUDA extensions for PyTorch
• implement custom autograd operations
• reason about distributed GPU computation
• understand NCCL collectives
• optimize multi-GPU memory movement
• reduce CPU-GPU synchronization
• model kernel performance
• perform roofline analysis
• benchmark kernels scientifically

The progression is:

textCopy
1Writing CUDA
2     ↓
3Understanding CUDA
4     ↓
5Optimizing CUDA
6     ↓
7Designing GPU kernels
8     ↓
9Designing AI kernels
10     ↓
11Designing GPU systems

That is what makes Advanced Kernel Engineering the frontier level of the CUDA/AI systems track.