CUDA from zero

Start with what a GPU does, learn the C++ you need, and write your first CUDA kernel. Build correct parallel programs with threads, memory, and synchronization, then measure and improve them with coalescing, shared memory, streams, libraries, and CUDA graphs.

  • Beginner → advanced
  • 18 units · approximately 26 hours

How to work through this course

Before you begin: No CUDA, GPU, C++, or machine learning experience is assumed. You should recognize variables, loops, and functions in any language and be able to open a terminal. The C++ primer teaches the syntax, arrays, pointers, and compilation needed for the labs. Read and solve the paper exercises on any device; use a supported Linux or Windows NVIDIA GPU machine for CUDA execution. macOS and Apple GPUs cannot run these CUDA labs locally.

By the end: Write, debug, and benchmark a CUDA pipeline against a CPU reference. Explain its thread mapping, memory traffic, synchronization, and numerical tolerance, and hand over reproducible code with measured reasons for each optimization.

  1. Understand the idea. Read each section in order and study its diagrams. When a pause-and-predict question appears, commit to an answer before you reveal it. Look up unfamiliar terms in the glossary below.
  2. Try it yourself. The diagrams and prediction exercises work on any computer. Compiling and running the CUDA labs requires a supported NVIDIA GPU, locally or on a remote machine. A CPU-only C++ exercise prepares you before your first kernel.
  3. Check your understanding. Answer the knowledge check at the end of each lesson and make sure you can say why the other options are wrong. Complete all units to unlock the final quiz; the passing score is 80%.

The optional project brings the lessons together. Keep notes on your choices, results, and any failures so you can explain what you learned.

Course orientation

What is CUDA?

CUDA is NVIDIA's platform and programming model for running general computations on its graphics processing units (GPUs). In CUDA C++, a CPU program allocates data, launches a GPU function called a kernel across many threads, and collects the result. You control which work each thread performs and how threads exchange data.

The problem it solves

A loop processing millions of independent values can spend most of its time repeating the same operation. A GPU can divide that work among many threads, but transfers, launch overhead, memory access, and synchronization can erase the benefit. The course teaches you to prove correctness first and measure the whole application before choosing an optimization.

From a CPU program to a measured GPU kernel

  1. 01Prepare and verify

    Build a CPU reference, learn the required C++, and check the driver, compiler, toolkit, and device before allocating GPU memory.

  2. 02Map the work

    Launch a grid of thread blocks. Each thread computes its own index, checks the array bounds, and writes the output it owns.

  3. 03Coordinate memory

    Choose data layouts, reuse tiles in shared memory, and add the block, stream, or host synchronization that a dependency requires.

  4. 04Measure and deliver

    Validate every result, separate kernel time from application time, compare with libraries, and package a reproducible experiment.

What it covers

  • CPU and GPU fundamentals, installation checks, and a focused C++ primer.
  • Complete kernels, thread indexing, allocations, copies, error checking, and resource cleanup.
  • Coalescing, shared memory, synchronization, reductions, atomics, and tiled matrix multiplication.
  • Compute Sanitizer, event timing, Nsight tools, streams, floating-point accuracy, and library selection.
  • CUDA graph replay, deployment checks, and a capstone with correctness and performance evidence.

Outside its scope

  • A general introduction to programming. If variables and loops are new, practice them before beginning the C++ primer.
  • A GPU simulator in the browser. Diagrams explain execution; real GPU results come from running the source on supported hardware.
  • A promise of speedup. Small or transfer-heavy workloads can be faster on the CPU, and that is a valid measured result.
  • A complete treatment of multi-GPU collectives, inline assembly, or architecture-specific Tensor Core kernel design. The final lessons provide a path to further study.
Key terms in this course
CPU and GPU
The central processing unit runs the host program. The graphics processing unit executes many parallel threads on the device.
Host and device
Host refers to the CPU side of a CUDA application; device refers to the GPU side. An address must be used by a processor that can access its allocation.
Kernel
A function marked __global__ in CUDA C++ that the host launches to execute on GPU threads.
Thread, block, and grid
A thread is one execution of a kernel. Threads form blocks that can cooperate. The blocks of one kernel launch form a grid.
Warp
A group of 32 threads scheduled together on current CUDA GPUs. Divergent branches can make different lanes execute different paths.
Streaming multiprocessor (SM)
A GPU execution unit that hosts resident thread blocks and schedules their warps using shared execution and memory resources.
Global and shared memory
Global memory stores arrays accessible across blocks. Shared memory is a limited region used by cooperating threads within a block.
Coalescing
Combining a warp's nearby memory accesses into fewer memory transactions. Layout and indexing determine which addresses its lanes request.
Barrier and atomic operation
A barrier coordinates participating threads. An atomic operation makes one supported read-modify-write indivisible at its specified scope; it is not a general barrier.
Stream and event
A stream orders GPU operations. An event marks a point in that order for timing, completion checks, or dependencies between streams.
Occupancy
The fraction of the maximum supported resident warps that are active on an SM. Higher occupancy does not by itself prove lower execution time.
Bandwidth and latency
Bandwidth is data moved per second. Latency is time until an operation finishes. A workload can be limited by either.
CUDA graph
A recorded dependency graph of operations that can be instantiated and launched repeatedly. Its resources must remain valid during execution.
Tolerance
An absolute and relative error rule used to compare floating-point results, with an explicit policy for non-finite values.

Syllabus

  1. 02Set up and verify a CUDA workstation75 minProOpen
  2. 03The C++ you need before your first kernel85 minProOpen
  3. 04Write and verify your first CUDA kernel90 minProOpen
  4. 05Map arrays to threads, blocks, and grids90 minProOpen
  5. 06Own device memory and move data safely90 minProOpen
  6. 07Memory coalescing: make neighboring threads read neighbors75 minProOpen
  7. 08Shared memory: load once and cooperate inside a block85 minProOpen
  8. 09Synchronization: make dependencies explicit80 minProOpen
  9. 10Reductions and atomics: many values become one95 minProOpen
  10. 11Matrix multiplication: tile a complete computation110 minProOpen
  11. 12Debug CUDA: find memory errors before measuring speed90 minProOpen
  12. 13Measure completed work and find the bottleneck80 minProOpen
  13. 14Streams, events, and safe asynchronous pipelines90 minProOpen
  14. 15Floating-point accuracy and useful correctness tests75 minProOpen
  15. 16Use CUDA libraries and understand Tensor Cores85 minProOpen
  16. 17Replay CUDA graphs and package a reproducible program85 minProOpen
  17. 18Capstone: build, validate, and measure a GPU pipeline150 minProOpen

Optional field project

Build and measure a GPU data pipeline

  1. 01Establish the environment and a reference
  2. 02Add a correct parallel summary
  3. 03Measure and investigate one bottleneck
  4. 04Package a reproducible decision
View optional project brief

Sign in to read the project brief.