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.
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.
- 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.
- 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.
- 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
- 01Prepare and verify
Build a CPU reference, learn the required C++, and check the driver, compiler, toolkit, and device before allocating GPU memory.
- 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.
- 03Coordinate memory
Choose data layouts, reuse tiles in shared memory, and add the block, stream, or host synchronization that a dependency requires.
- 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
- 01What is CUDA? How a GPU helpsUp next
- 02Set up and verify a CUDA workstationOpen
- 03The C++ you need before your first kernelOpen
- 04Write and verify your first CUDA kernelOpen
- 05Map arrays to threads, blocks, and gridsOpen
- 06Own device memory and move data safelyOpen
- 07Memory coalescing: make neighboring threads read neighborsOpen
- 08Shared memory: load once and cooperate inside a blockOpen
- 09Synchronization: make dependencies explicitOpen
- 10Reductions and atomics: many values become oneOpen
- 11Matrix multiplication: tile a complete computationOpen
- 12Debug CUDA: find memory errors before measuring speedOpen
- 13Measure completed work and find the bottleneckOpen
- 14Streams, events, and safe asynchronous pipelinesOpen
- 15Floating-point accuracy and useful correctness testsOpen
- 16Use CUDA libraries and understand Tensor CoresOpen
- 17Replay CUDA graphs and package a reproducible programOpen
- 18Capstone: build, validate, and measure a GPU pipelineOpen
- QFinal graded quizComplete all units to unlock
Optional field project
Build and measure a GPU data pipeline
- 01Establish the environment and a reference
- 02Add a correct parallel summary
- 03Measure and investigate one bottleneck
- 04Package a reproducible decision
View optional project brief
Sign in to read the project brief.