Choose C++, Python, or Rust
Write CUDA C++, Python with Numba, or Rust with cuda-oxide. Investigate occupancy, memory behavior, PTX, SASS, and timelines where the lab and GPU support them.
GPUFlight Performance Lab
Learn GPU performance on NVIDIA and AMD
Guided labs with a target and measured evidence. Write CUDA kernels for NVIDIA or HIP C++ for AMD, run on managed hardware, and see why the result changed. No local GPU setup required.
- Guided learning paths and focused labs
- CUDA: C++ · Python · Rust (beta)
- AMD: HIP C++ (beta)
- Free: at most 5 hosted runs per day, no credit card
Now with AMD support · Beta
CUDA on NVIDIA. HIP on AMD.
Explore GPU programming across two platforms. Choose a supported lab, select NVIDIA or AMD, and work with the languages and GPUs offered for that lab.
Write HIP C++ on AMD GPUs
Run supported exercises on managed AMD GPUs with HIP and ROCm. Check correctness, inspect kernel timing and traces, and compare the effect of your code changes.
Choose AMD in a supported lab, then work on the kernel beside the problem statement. This example explores shared-memory tiling.
Open the image to see the code at full size.
The result separates kernel-only time from GPU end-to-end time. Open the session details or select saved runs to compare your changes.
Actual MI325X run from the shared-memory lab.AMD support is in beta, and lab coverage varies by platform. Python and Rust currently target NVIDIA CUDA. The screenshots above show AMD HIP; the video walkthroughs below show CUDA examples.
Watch a real lab · 2:15
From a code change to a measured result
Follow a vector-add kernel through Performance Lab. Adjust the block size, check correctness, and explore occupancy and the profiling timeline.
Try the occupancy labWhat you'll learn
Build GPU performance intuition, from kernels to models
Explore how parallelism, data movement, and fusion shape performance. Write code, measure a change, and learn to explain the result. Available exercises and profiling targets vary between CUDA and HIP.
Elementwise kernels
Map a CPU loop to GPU threads, improve memory access, and fuse operations to reduce data movement. Apply those ideas to Transformer workloads.
Matrix multiplication
Understand how thread layout, shared memory, and tiling affect matrix multiplication. Compare your implementation with cuBLAS and learn where the performance gap comes from.
Reductions and softmax
Turn a serial sum into a parallel reduction, then build toward stable softmax and online normalization. Learn how threads cooperate without unnecessary synchronization.
Transformer decoder blocks
Connect individual kernel optimizations in a decoder block. Explore attention, tiling, and intermediate storage while checking results against a CPU reference.
Explore focused labs on the concepts behind fast kernels:
- Occupancy
- Memory coalescing
- Shared memory and bank conflicts
- Stencils
- Tensor Cores and TMA
- CUB
- LayerNorm and inference
- PyTorch
- Telemetry
See memory optimizations in action
Two short walkthroughs, with code changes and before-and-after profiling in the app.
App walkthrough · 2:30
Fix memory coalescing
Change how threads access a SAXPY kernel’s inputs, then compare memory efficiency and kernel time in the profiler.
Try the coalescing labApp walkthrough · 3:05
Reduce memory traffic with fusion
In this warmup, fuse two loops and remove a temporary array. Read the memory operations in the code, then measure the traffic they generate.
Try the fusion labCUDA language choices
CUDA C++, Python, or Rust: pick one per run
The same vector addition, expressed in each language. Start with a kernel you can read, then use the lab's profiling views to understand its behavior.
CUDA C++
CUDA kernels
solution.cu
__global__ void add(const float* a,
const float* b,
float* c, int n) {
int i = blockIdx.x * blockDim.x
+ threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}Python
Numba CUDA
solution.py
@cuda.jit
def add(a, b, c, n):
i = cuda.grid(1)
if i < n:
c[i] = a[i] + b[i]Rust Beta
cuda-oxide
solution.rs
#[kernel]
pub fn add(a: &[f32], b: &[f32],
mut c: DisjointSlice<f32>) {
let idx = thread::index_1d();
if let Some(out) = c.get_mut(idx) {
*out = a[idx.get()] + b[idx.get()];
}
}Kernel excerpts only; the lab supplies the surrounding setup. These three examples target NVIDIA CUDA; AMD labs use HIP C++. Available languages depend on the lab. CPU baselines and PyTorch exercises keep their own language.
Compare the compiler output · 6:43
C++ vs. Rust for CUDA: what the PTX reveals
Watch the same vector addition in CUDA C++ and Rust. Select both runs, open Compare, and follow the PTX—the intermediate instructions produced by the compiler—to see how raw pointers and checked slices affect the output.
Compare your own kernelsCompare across GPUs
Run it on two GPUs, then compare the evidence
Compare two runs from the same lab, including CUDA on NVIDIA and HIP on AMD where both are supported. Select the runs in your history to inspect correctness, runtime, and the source changes side by side.
- A plain-language verdict: which run was faster, and by how much
- Correctness and timing side by side, with graded metrics where available
- The source diff between the two submitted kernels
- Compare the same workload across GPUs; instruction listings and profiler counters remain specific to each platform
Share comparison turns the result into a link you can send. The example shown compares NVIDIA L4 and H100 runs.
The experiment loop
Change, run, inspect, compare
Read the target
The correctness rule and the metric you are trying to move.
Make one change
Edit the kernel or the launch configuration.
Run on the GPU
Compile and execute on managed hardware with the profiling passes the lab needs.
Explain the result
Check correctness and timing, then explore the timeline and the profiling views available for your platform, lab, and GPU.
Measurements come from real NVIDIA and AMD GPUs, so they carry normal noise. The goal is to explain why the number moved, not only to pass.
Common questions
Performance Lab FAQ
Do I need my own NVIDIA or AMD GPU?
No. Labs run on GPUFlight's managed hardware. Your computer only needs a supported browser; you do not need to install CUDA or ROCm locally.
Which languages do the labs support?
NVIDIA CUDA labs offer C++, Python with Numba, and Rust with cuda-oxide. AMD HIP labs use C++. Rust and AMD support are in beta: results and targets may still change. Available languages, GPUs, and profiling targets depend on the lab and your access.
Is this a GPU programming syntax course?
No. The labs are about how GPU programs execute and why their performance changes. The first path starts from a CPU loop; later paths assume basic kernel-writing familiarity.
What is free?
500 credits / month for Performance Lab and Workbench. At most 5 hosted runs per day, shared across Performance Lab and Workbench. Any credit cost is shown before you run. See the plan comparison.
Start with one measured change
Optimize a kernel and inspect the evidence
Have a program of your own? Open the CUDA Playground.