Instructor guide

Why we teach this lesson

On GPUs the bottleneck is rarely arithmetic — it is getting the data there. A single transfer across PCIe can easily cost more time than the kernel that consumes it. CUDA offers several strategies for this — device, managed, pinned and constant memory — and picking the wrong one costs more performance than kernel tuning can win back afterwards. The underlying trade-offs (explicit copies versus automatic migration, page-locked transfers, keeping read-only data close to the compute units) are not vendor-specific and transfer to other accelerators.

Intended learning outcomes

After completing this module, learners will be able to:

• Choose between device memory, managed memory, pinned memory and constant memory for a given workload

• Allocate and free managed memory using cudaMallocManaged in C/C++ (and the managed attribute in CUDA Fortran), and explain how data migration works

• Allocate device memory with cudaMalloc or device in CUDA Fortran and copy data between host and device

• Allocate pinned host memory with cudaMallocHost (or pinned in CUDA Fortran) and explain when pinned memory is required (e.g. for asynchronous copies)

• Use constant memory for small read-only data accessed by many threads, and stay within its 64 KB limit

• Insert cudaDeviceSynchronize() correctly before reading GPU-modified data on the host, and recognise the race conditions that occur otherwise

Timing

Episode	Lecture	Exercises	Total
Software setup	10 min	—	10 min
1. Unified memory	45 min	15 min	60 min
2. Manual memory management	30 min	30 min	60 min
3. Pinned and constant memory	10 min	0 min	10 min
4. Synchronization and summary	10 min	0 min	10 min
Total	105 min	45 min	150 min

These numbers are estimates from slide-based deliveries of the course: the material has not yet been taught in notebook form, which may speed things up, and the order of the material was different. They may be underestimates and should be re-measured at the first delivery.

Hardware requirements

An NVIDIA GPU with compute capability 7.5 or newer is recommended; compute capability 7.0 still works but its support is being phased out. Use CUDA Toolkit 12.x or newer: CUDA 13 drops support for these older GPUs, so 12.x is both the floor and, for CC 7.0 hardware, the ceiling. Profiling on compute capability 7.0 is being phased out as well: Nsight Compute 2025.3 and Nsight Systems 2025.4 dropped Volta support, so to profile on such a device use the Nsight versions bundled with HPC SDK 25.7 or with the CUDA 12.x toolchain of later HPC SDKs (profiling is not needed for this module). CUDA Fortran requires the NVIDIA HPC SDK (nvfortran). Providing GPU access for a course run (workstations, a training cluster or scheduler-allocated nodes) is the responsibility of the site delivering the course; the number of participants is limited by the number of available GPUs and supervisors.

Participants need a machine with an NVIDIA GPU — either their own workstation or a node of a cluster — to compile and run the exercises; the software-setup episode tells them so in its first paragraph.

Learner personas

Learners need sufficient knowledge of C/C++ or Fortran. Basic familiarity with a text editor or IDE, and with the Linux shell when working on a cluster, is also advisable. Some knowledge about parallel programming is a plus. This module assumes the content of the Introduction to CUDA module.

How to teach each episode

The lists below are derived from the material of each episode and from the questions that came up in earlier deliveries. Every teaching episode opens with a short introduction, a questions box and an objectives box; the questions work well as a two-minute warm-up before the lecture part.

Software setup (10 min)

• Emphasise: the exercises are compiled and run by the participants themselves on a machine with an NVIDIA GPU; nvcc for CUDA C/C++ and nvfortran for CUDA Fortran; the three flag families (optimisation, debug symbols, -lineinfo for profiling); the C/C++ templates need cuda_utils.h next to the .cu file.

• Demo live: nvcc --version / nvfortran --version, then compile and run one program on the GPU machine so that everybody has seen the full cycle once.

• Common questions: which module to load or container to start, and which queue of the batch system to use — all of this is site-specific; refer to the documentation of the system the course runs on.

Episode 1: Unified memory (45 min + 15 min)

• Emphasise: the two physically separate memory spaces and the three ways to make data accessible on the GPU; managed memory migrates data automatically, but synchronisation stays the programmer’s job; the four unified-memory paradigms and how to query which one a system supports; prefetching is a pure performance optimisation that never changes correctness.

• Run: the type-along (managed-memory add_kernel in both languages), the discussion When is managed memory the wrong choice?, and the three-step synchronisation slideshow — step through it slowly, the “lucky” run in step 1 is the point.

• Demo live: the vector-scalar exercise, starting from <<<1, 1>>> and then moving to many threads and blocks; show that the single-thread version is correct but slow.

• Common questions: what good execution-configuration defaults are — the move from <<<1, 1>>> to (size + 255) / 256 blocks of 256 threads in the exercise is the natural place to answer this; let participants time a few configurations. For very large arrays remind them of the grid-stride loop from the Introduction module, which decouples the grid size from the data size. Also: which cudaMemPrefetchAsync signature to use — the interface changed with CUDA 13, both are in the episode.

Episode 2: Manual memory management (30 min + 30 min)

• Emphasise: explicit allocation with cudaMalloc / the device attribute; the direction constants of cudaMemcpy and cudaMemcpyDefault; cudaMemcpy blocks the host and therefore acts as a built-in synchronisation point; in CUDA Fortran an array assignment between host and device arrays is a copy whose synchronisation behaviour is the same as that of cudaMemcpy.

• Run: both exercises. The first is the manual-memory counterpart of the episode 1 exercise; the second (vector addition) is about the execution configuration only since allocation and copies are already filled in.

• Demo live: compiling a template together with cuda_utils.h; the bandwidth and kernel time reported by the vector-addition template for the trivial <<<4, 4>>> launch versus a tuned configuration.

• Common questions: good defaults for the block size and other tuning points — let participants compare a few configurations and look at the reported bandwidth rather than giving one number; for vector sizes beyond what one grid can cover, point out that the template’s kernel already uses the grid-stride loop from the Introduction module; whether cudaDeviceSynchronize() is needed before the cudaMemcpy of the result (not strictly, see the solution, but it makes the dependency explicit and is needed for the timer).

Episode 3: Pinned and constant memory (10 min)

• Emphasise: the two transfer paths in the slideshow (staging buffer versus direct transfer); pinned memory is required for asynchronous copies; constant memory is for small data that all threads read and nobody writes, 64 KB per device.

• Run: the discussion Which memory for which array? — it is the only interactive element of the episode.

• Demo live: nothing new to compile; point back to the episode 2 vector-addition template, which already allocates its host arrays with cudaMallocHost, and explain now what that bought.

• Common questions: “why not pin everything?” — pinned pages can neither be swapped out nor relocated by the operating system, so pinning large amounts takes memory away from the OS and other processes and can slow down the whole system (best-practices guide); then the caveats box (PCIe access is slow, race conditions on mapped memory, Fortran restrictions).

Episode 4: Synchronization and summary (10 min)

• Emphasise: the summary table; local memory is per-thread storage that physically lives in global memory and is no faster than it; the pointer rules for host and device code; the rule of thumb for cudaDeviceSynchronize().

• Run: the discussion Pick the strategy, then find the synchronisation — a good closing exercise before the quiz, because it combines every memory type of the module with the synchronisation question.

• Common questions: which pointers are safe to dereference on the host — the answer is the pointer-rules section; use the segmentation fault as the concrete symptom.

Additional teaching recommendations

Preparing exercises

The day before the course:

• Test that all exercises and their solutions compile and run in both languages.

• Check the software environment, e.g. that cross-compilation works as intended. Profiling tools are not needed for this module.

• Make sure all materials (module pages, exercise templates and solutions) are reachable by the participants.

Exercise logistics

Learners download the exercise templates from the episode pages, edit them in an editor or IDE of their choice and compile them with nvcc (CUDA C/C++) or nvfortran (CUDA Fortran) on a machine with an NVIDIA GPU; solutions are offered for download next to each template.

The C/C++ exercises use the error-checking macro from cuda_utils.h; the header is offered for download next to every C template and solution, lives in both exercises/templates/ and exercises/solutions/, and has to sit in the same directory as the .cu file that is being compiled.

The episode 1 exercise starts from the CPU snippet shown in the episode; episode 2 offers downloadable templates and solutions for both of its exercises. Episodes 3 and 4 have no exercises, only discussions.

Delivery formats

The module can be delivered on-site or online. The lecture parts are given from the notebook-based episode pages, the exercises are done by the participants on the GPU machine, in a shell or an IDE.

Other practical aspects

• The number of participants is limited by the number of available GPUs and supervisors, as long as exercises are part of the delivery.

• See the hardware requirements above for the GPU and toolkit versions to provide.

Interesting questions you might get

The module’s own discussion prompts are the questions that come up most often:

• Episode 1, When is managed memory the wrong choice? — The runtime migrates pages for you, but it reacts to access patterns rather than anticipating them. Consider two workloads: a lookup table written once on the host and then read by every kernel for the rest of the run, and an array that host and device touch alternately inside a loop. Which of the two would you expect the automatic migration to handle badly, and why? What would you measure before deciding to manage the transfers by hand instead?

• Episode 3, Which memory for which array? — Your kernel uses two arrays: a 100-entry table of coefficients that every thread reads and nobody writes, and a large input buffer copied to the device once per timestep. One belongs in constant memory, the other is a candidate for pinned host memory. Which is which — and what goes wrong if you swap them?

• Episode 4, Pick the strategy, then find the synchronisation — Three short scenarios: (a) 100 floats, read by every thread, never written; (b) an 8 GB field that host and device touch alternately inside a loop; (c) an input buffer for the next timestep that you want to copy to the device while a kernel is still working on the current one — the kernel does not touch this buffer. For each: which memory type would you reach for, and at which point in the program must you synchronise before the data can be trusted? Expected answers: (a) constant memory, no extra synchronisation after cudaMemcpyToSymbol; (b) managed memory with prefetching, cudaDeviceSynchronize() before each host phase; (c) pinned host memory + asynchronous copy, synchronise the stream/device before the kernel that consumes the buffer.

Typical pitfalls

The instructor’s experience is that learners run into three kinds of problems:

• An incomplete or incorrect understanding of the SIMT model leads to poor performance, incorrect results or crashes.

• Incorrect synchronisation between host and device.

• Segmentation faults caused by incorrect memory handling.

The last two are the ones that bite in this module (the SIMT model is the subject of the Introduction to CUDA module). Where they show up, episode by episode:

• Incorrect synchronisation — episode 1: the host reads managed memory before the kernel has finished (the three-step slideshow shows the lucky run, the wrong run and the fixed run); episode 2: the “Synchronization needed?” line in the templates — cudaMemcpy synchronises implicitly, cudaDeviceSynchronize() is still needed for the timer; episode 4: the rule of thumb.

• Segmentation faults from incorrect memory handling — episode 1: on Linux with Pascal or newer GPUs (concurrentManagedAccess = 1) host access to managed memory while a kernel runs does not crash — it silently returns stale or half-written data, which is why the “lucky run” in the slideshow is so instructive; a segmentation fault occurs only on systems without concurrent managed access (Windows, WSL, GPUs older than Pascal); episode 2: a host pointer passed to a kernel or a device pointer dereferenced in host code; episode 3: race conditions on mapped pinned memory, pointer arithmetic on constant-memory symbols; episode 4: the pointer rules.

Restrictions stated in the episodes’ own caution, attention and important boxes:

• Unified memory (episode 1): const-qualified variables and C++ references cannot be declared as managed memory; C++ classes/structs with __managed__ members have many restrictions; managed memory can only be allocated and freed in host code.

• Pinned memory (episode 3): race conditions when accessing mapped pinned memory from both host and device; CPU locks are not visible to the GPU and GPU atomics may not work on pinned host memory; access over PCIe is slow compared to device memory; CUDA Fortran does not allow pinned memory in kernel arguments and only allocatable arrays may be pinned.

• Constant memory (episode 3): cannot be dynamically allocated; maximum 64 KB per device; must not be written from device code; in C/C++ host code, constant memory variables are symbols — do not use pointer arithmetic on them.

• Synchronisation (episode 4): always call cudaDeviceSynchronize() before the host accesses data that was modified by a kernel.