Synchronization and Summary

This closing episode brings together the memory types of the module: it summarises global, constant and local memory in a single table, explains why local memory is no faster than global memory, and sets out the pointer rules that decide which pointers may be dereferenced in host code and which only in device code.

It then recaps the two synchronisation primitives used throughout the module — cudaDeviceSynchronize() and the implicit barrier of cudaMemcpy — and closes with a discussion in which you choose a memory strategy for three scenarios and find the point at which each of them must synchronise.

Questions

  • Given a workload, which memory type would you reach for first — and what would change your mind?

  • At which exact points must the host wait before it can trust what it reads?

  • Which pointers are safe to dereference on the host, and which only inside a kernel?

Objectives

  • Understand synchronisation between host and device

  • Know the characteristics and trade-offs of the different CUDA memory types

  • Understand pointer rules for host and device code in C/C++

Instructor note

  • 10 min teaching

  • 0 min exercises

Memory kinds summary

Memory

When to use

Size

Notes

Global memory

Always — default place to store data

4–384 GB

Accessible from all threads

Constant memory

Device-constant data, read by all threads

64 KB

Part of global memory with special caching

Local memory

Implicitly used for register spills

—

Part of global memory, private per thread

Despite its name, local memory is not a separate, fast memory next to the SM: it is a per-thread region whose physical location is global (device) memory, so accessing it is no faster than an ordinary global-memory access (although it is cached in L1/L2). You never allocate it explicitly — the compiler places a thread’s automatic variables there whenever they do not fit in registers, for example large per-thread arrays or structs, or when high register pressure forces register spilling.

Tip

Local-memory traffic is usually something to minimise (by keeping per-thread state small) rather than a tool you reach for deliberately.

Pointer rules (C/C++)

  • In host code, only pointers to host memory can be dereferenced without unified memory.

  • In device code with UVA, both host and device pointers can be dereferenced. Without UVA, device and host address spaces are separate.

  • However, only memory regions allocated with cudaMallocHost (or other CUDA functions) and mapped to the device address space can be accessed from the device if not unified memory.

  • Conclusion: pointers passed to kernels as arguments should (and without UVA: must) point to device or unified memory.

Warning

Segmentation faults are the typical symptom of incorrect memory handling: dereferencing a device pointer in host code crashes with a segmentation fault, and a host pointer used in a kernel (without unified memory) causes an illegal memory access that the next CUDA call reports — check every call so that it does not go unnoticed.

Synchronisation recap

Key synchronisation primitives covered in this module:

  • cudaDeviceSynchronize() — blocks the host until all preceding kernels and memory operations complete

  • cudaMemcpy — synchronous by default (acts as an implicit barrier for copies from host to device and the other way round)

Important

Rule of thumb: always call cudaDeviceSynchronize() before the host accesses data that was modified by a kernel.

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?

Keypoints

  • Global memory is the default place to store data

  • Constant memory uses a dedicated cache for small, read-only data (max 64 KB)

  • Always synchronise (cudaDeviceSynchronize) before accessing GPU-modified data on the host

  • In C/C++ host code, only host pointers (including unified memory pointer) can be dereferenced; pointers passed to kernels should point to device or unified memory

See also