Device memory and error checks: CUDA that does not hide failures
The right way to call cudaMalloc / cudaMemcpy, unified memory, and the macro you should copy once
- Know the full explicit device-memory flow and the usual traps
- Write a reusable CUDA error-check macro and understand why it is required
- Know which cases Unified Memory fits, and which it does not
- Understand why pinned memory makes copies faster
In the classic CUDA model, host and device each have their own address space. float* d_a holds a device-memory address. Dereferencing it in CPU code is a straight segfault. Passing a host pointer into a kernel crashes the same way. Keeping those two kinds of pointer apart is the basic skill of CUDA programming.
The full four-step flow
1int n = 1 << 20;2size_t bytes = n * sizeof(float);3 4std::vector<float> h_a(n, 1.0f), h_b(n, 2.0f), h_c(n, 0.0f);5 6// 1. Allocate on the device7float *d_a, *d_b, *d_c;8CUDA_CHECK(cudaMalloc(&d_a, bytes)); // note the &d_a: you are changing the pointer itself9CUDA_CHECK(cudaMalloc(&d_b, bytes));10CUDA_CHECK(cudaMalloc(&d_c, bytes));11 12// 2. Copy inputs to the device13CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));14CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));15 16// 3. Compute17int threads = 256, blocks = (n + threads - 1) / threads;18vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);19CUDA_CHECK(cudaGetLastError()); // catch launch failure (illegal config, etc.)20 21// 4. Copy results back. Synchronous cudaMemcpy waits for the kernel to finish22CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, bytes, cudaMemcpyDeviceToHost));23 24CUDA_CHECK(cudaFree(d_a));25CUDA_CHECK(cudaFree(d_b));26CUDA_CHECK(cudaFree(d_c));The macro you should copy once
Almost every CUDA Runtime API returns cudaError_t. Skipping the return value is the number one reason CUDA debugging hurts: the error stays silent and surfaces later, at a site that has nothing to do with the root cause. The macro below is standard in serious CUDA projects. Put it in your common.cuh.
1#pragma once2#include <cstdio>3#include <cstdlib>4 5#define CUDA_CHECK(call) \6 do { \7 cudaError_t err_ = (call); \8 if (err_ != cudaSuccess) { \9 fprintf(stderr, "CUDA error %s:%d: '%s' -> %s\n", \10 __FILE__, __LINE__, #call, cudaGetErrorString(err_)); \11 std::exit(EXIT_FAILURE); \12 } \13 } while (0)14 15// Kernel launch does not return an error code. Catch two kinds of error separately:16// cudaGetLastError() : synchronous errors such as an illegal config17// cudaDeviceSynchronize() : async errors during kernel execution (out of bounds, illegal instruction)18#define CUDA_CHECK_KERNEL() \19 do { \20 CUDA_CHECK(cudaGetLastError()); \21 CUDA_CHECK(cudaDeviceSynchronize()); \22 } while (0)Unified memory: less boilerplate, not free
cudaMallocManaged allocates memory that both host and device can access directly. On a page fault, the driver migrates pages in the background. The code gets much shorter, which fits prototypes and teaching. The cost is that page-migration overhead is invisible: a kernel with a bad access pattern can thrash pages in the background, and nothing in the source shows it.
1float *a, *b, *c;2CUDA_CHECK(cudaMallocManaged(&a, bytes));3CUDA_CHECK(cudaMallocManaged(&b, bytes));4CUDA_CHECK(cudaMallocManaged(&c, bytes));5 6for (int i = 0; i < n; ++i) { a[i] = 1.0f; b[i] = 2.0f; } // CPU writes directly7 8// Hint the access intent so the driver can cut page-fault thrashing9int device = 0;10CUDA_CHECK(cudaMemPrefetchAsync(a, bytes, device));11CUDA_CHECK(cudaMemPrefetchAsync(b, bytes, device));12 13vecAdd<<<blocks, threads>>>(a, b, c, n);14CUDA_CHECK(cudaDeviceSynchronize()); // CPU can safely read c only after this sync15 16printf("c[0] = %f\n", c[0]);17cudaFree(a); cudaFree(b); cudaFree(c);| Approach | Fits | Does not fit |
|---|---|---|
cudaMalloc + cudaMemcpy | production code, precise control of transfer timing, pipeline overlap | quick prototypes (lots of boilerplate) |
cudaMallocManaged | prototypes, teaching, pointer-heavy data structures, datasets larger than device memory | latency-sensitive hot paths (migration cost is hard to predict) |
cudaHostAlloc (pinned) | async copies, saturating PCIe bandwidth | huge allocations (pinning squeezes system-available memory) |
Why pinned memory is faster
Host memory from ordinary malloc is pageable. The OS can swap it out at any time. The GPU DMA engine cannot read that memory directly, so the driver copies it first into an internal pinned buffer, then DMA moves it to the device: an extra CPU copy you did not ask for. Allocate pinned memory with cudaHostAlloc and DMA can go in one hop. Bandwidth usually climbs close to the PCIe theoretical peak, and pinning is a prerequisite for cudaMemcpyAsync.
1float* h_pinned;2CUDA_CHECK(cudaHostAlloc(&h_pinned, bytes, cudaHostAllocDefault));3 4// Only pinned memory can actually copy asynchronously. Passing pageable memory to Async5// does not error, but the driver falls back to synchronous behavior and overlap is gone.6CUDA_CHECK(cudaMemcpyAsync(d_a, h_pinned, bytes, cudaMemcpyHostToDevice, stream));7 8CUDA_CHECK(cudaFreeHost(h_pinned));Self-check
Answer in your head first, then open the solution. If you cannot, reread that section.