How to Learn CUDA: A Path from C to Fused Kernels

How do you learn CUDA programming?

Learn CUDA in five steps, starting with the GPU execution model of threads, blocks and warps, then memory access, then classic parallel patterns such as reductions and tiled matrix multiplication, then profiling, and finally fused kernels for ML. You need working C or C++ and access to a real NVIDIA GPU to run code on, which can be a cloud or browser sandbox.

The CUDA Kernels study plan follows the same order as the table below, with every kernel run on a real GPU from your browser.

What do you need to know before learning CUDA?

You need three things before writing your first kernel.

• You need C or C++ basics, meaning pointers, arrays and manual memory allocation. CUDA C++ is C++ with a few extensions, so pointer arithmetic is everywhere.
• You need to know how memory works in a normal program, including what an array looks like in memory and why reading neighbouring values is fast.
• You need basic linear algebra, mainly vectors and matrix multiplication, because most ML kernels are built from them.

If you mainly write Python, you can start with Triton, which writes GPU kernels in Python at the level of blocks rather than single threads. The trade-offs are covered further down.

Step 1: How does the GPU execution model work?

A CUDA kernel is a function that thousands of threads run at the same time. Threads are grouped into blocks, and blocks into a grid. The hardware schedules threads in groups of 32 called warps, and each warp runs on a streaming multiprocessor (SM).

Your first job is to give each thread its own piece of the data. Vector addition is the standard first kernel, launched here with enough blocks of 256 threads to cover all n elements.

__global__ void vector_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];
    }
}

int threads = 256;
int blocks = (n + threads - 1) / threads;
vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

Two lines carry most of the lesson. The index formula maps each thread to one element. The bounds check on i matters because the number of threads launched is rounded up to whole blocks, so the last block usually has threads with nothing to do. Forgetting the check is a classic first CUDA bug, because the kernel writes past the end of the array and may not crash.

Step 2: Why does memory access matter more than arithmetic?

Beginner kernels are usually slow because of memory, not arithmetic. A GPU can do arithmetic far faster than it can fetch data from global memory, so the order in which threads read memory decides the speed. Learn these three ideas in order.

• Coalescing happens when the threads of a warp read neighbouring addresses, so the hardware combines the reads into a few transactions. When they read scattered addresses, each read costs a separate trip.
• Shared memory is a small, fast memory shared by the threads of one block. You load a tile once from global memory and let every thread reuse it.
• Bank conflicts happen because shared memory is split into banks, and threads that hit the same bank at once are serialised. Padding an array by one column often fixes it.

Matrix transpose is the best exercise here, because the naive version reads in one direction and writes in the other, so one side is always uncoalesced. A shared-memory version fixes both. The Matrix Transpose problem lets you write both and compare the timings.

Step 3: Which parallel patterns should you learn first?

Reductions and tiled matrix multiplication teach most of what ML kernels need.

A reduction such as summing an array forces you to coordinate threads. Each block sums part of the array in shared memory with a __syncthreads() call between steps, then the block results are combined. The same pattern sits under softmax, layer norm and loss functions.

Tiled matrix multiplication combines everything so far. Each block loads tiles of both matrices into shared memory, multiplies them and moves along. Simon Boehm's matmul worklog takes a naive kernel step by step towards cuBLAS-level speed, and it is the clearest public walkthrough of why each optimisation helps.

Step 4: How do you know if a kernel is fast?

You measure it, because a correct but slow kernel is the normal state of a first attempt, and guessing at the cause wastes time.

• Check correctness first against a reference, such as the same operation in PyTorch, within a floating-point tolerance.
• Time it properly by warming up first, running many iterations and comparing against a library baseline.
• Profile it with NVIDIA Nsight Compute, which reports memory throughput, occupancy and where the kernel stalls.
• Ask whether the kernel is limited by memory bandwidth or by arithmetic, which is the roofline question. An element-wise kernel is almost always memory-bound, so the goal is to get close to the GPU's peak bandwidth rather than its peak FLOPs.

Timings also change between GPUs. A kernel tuned on one card can behave differently on another with a different memory system, which is why it helps to benchmark on more than one.

Step 5: What are fused kernels, and why do ML engineers write them?

A fused kernel does several operations in one pass instead of writing intermediate results to global memory between them. For memory-bound ML operations, fusion is often the biggest speed-up available.

Softmax is the classic example. Done naively it reads the input three times, once for the maximum, once for the sum of exponentials and once to normalise. A fused version keeps the row in fast memory and does all three steps there. Layer norm works the same way.

Attention is where this matters most. FlashAttention computes attention in tiles without ever writing the full attention matrix to global memory. Once you can write a fused softmax, you have the main idea behind it.

Should you learn CUDA or Triton first?

If you stay with GPU programming, you will probably end up using both. The Triton Kernels study plan covers the Triton side.

Which resources should you use to learn CUDA?

These resources cover the path above, roughly in the order you will need them.

Read, then write. Every chapter of the book is worth more if you write its kernels yourself and time them.

Do you need your own NVIDIA GPU to learn CUDA?

No, you need access to one rather than ownership. Free cloud notebooks such as Google Colab sometimes offer an NVIDIA GPU, cloud providers rent them by the hour, and browser platforms run your code on real GPUs for you. TensorTonic runs CUDA and Triton kernels in the browser on several NVIDIA GPUs, from a T4 to an H200. LeetGPU is another browser option built around GPU programming, and our TensorTonic vs LeetGPU page compares the two. The guide to practising ML without a GPU covers the options in more detail.

Practice this

Is CUDA hard to learn?

The syntax is not hard, because it is C++ with a few additions. The hard part is the mental model of thousands of threads, memory access patterns and synchronisation running at once. The first kernel is usually easy, and the first fast kernel is not.

Do ML engineers need to know CUDA?

Not all of them. Most ML engineers work in PyTorch and never write a kernel. CUDA matters for ML infrastructure, inference and performance roles, where custom kernels and profiling are part of the job. Our CUDA interview questions guide shows what those roles test.

Can you learn CUDA with Python only?

You can start GPU programming in Python with Triton, but CUDA itself is C++. If you only know Python, learn enough C to read pointers and arrays first.

Related guides

• CUDA practice problems: from vector add to attention
• CUDA interview questions for ML engineers
• TensorTonic vs LeetGPU
• How to practise ML with no GPU, no dataset and no team