CUDA Cheatsheet

Memory Management

Use this CUDA reference while you build software engineering projects, review code for technical interview prep, or polish examples for a software engineer resume.

Allocation and Deallocation

Linear (1-D) device memory

float* d_data;
cudaMalloc(&d_data, n * sizeof(float));   // allocate on device
cudaFree(d_data);                          // free device memory

Zero-initialized

cudaMemset(d_data, 0, n * sizeof(float));            // async-safe per-stream
cudaMemsetAsync(d_data, 0, n * sizeof(float), stream);

2-D pitched allocation (avoids bank conflicts and enables 2-D copy)

float* d_mat;
size_t pitch;     // actual row stride in bytes (padded to alignment)
cudaMallocPitch(&d_mat, &pitch, width * sizeof(float), height);
// Access: d_mat[row * pitch/sizeof(float) + col]

cudaFree(d_mat);

3-D array allocation

cudaExtent extent = make_cudaExtent(W * sizeof(float), H, D);
cudaPitchedPtr pitchedPtr;
cudaMalloc3D(&pitchedPtr, extent);

// Access element (x, y, z):
// char* ptr = (char*)pitchedPtr.ptr;
// float* row = (float*)(ptr + z * pitchedPtr.ysize * pitchedPtr.pitch
//                            + y * pitchedPtr.pitch);
// float val = row[x];

cudaFree(pitchedPtr.ptr);

CUDA arrays (for textures / surfaces)

cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaArray_t cuArr;
cudaMallocArray(&cuArr, &desc, W, H, cudaArraySurfaceLoadStore);
cudaFreeArray(cuArr);

Host Memory

Function	Description
`cudaMallocHost(&ptr, bytes)`	Page-locked (pinned) host allocation
`cudaHostAlloc(&ptr, bytes, flags)`	Pinned with flags
`cudaFreeHost(ptr)`	Free pinned host memory

Flags for cudaHostAlloc:

Flag	Meaning
`cudaHostAllocDefault`	Standard pinned
`cudaHostAllocPortable`	Visible to all CUDA contexts
`cudaHostAllocMapped`	Zero-copy: accessible from device
`cudaHostAllocWriteCombined`	Write-combined (fast H→D, slow H reads)

float* h_data;
cudaMallocHost(&h_data, n * sizeof(float));   // pinned — DMA without CPU involvement
// ... use h_data ...
cudaFreeHost(h_data);

Why pinned? Pageable memory must be staged through a pinned bounce buffer before DMA. Pinning eliminates that copy and enables async cudaMemcpyAsync.

Memcpy

Directions

Direction constant	Meaning
`cudaMemcpyHostToDevice`	CPU → GPU
`cudaMemcpyDeviceToHost`	GPU → CPU
`cudaMemcpyDeviceToDevice`	GPU → GPU
`cudaMemcpyHostToHost`	CPU → CPU (uses CUDA DMA)
`cudaMemcpyDefault`	Infer from pointer attributes (unified memory)

1-D copy

cudaMemcpy(d_dst, h_src, bytes, cudaMemcpyHostToDevice);   // synchronous
cudaMemcpyAsync(d_dst, h_src, bytes, cudaMemcpyHostToDevice, stream);  // async

2-D copy (with pitch)

cudaMemcpy2D(
    d_dst, dstPitch,     // destination pointer + pitch
    h_src, srcPitch,     // source pointer + pitch
    width * sizeof(float), height,
    cudaMemcpyHostToDevice);

cudaMemcpy2DAsync(d_dst, dstPitch, h_src, srcPitch,
                  widthBytes, height, cudaMemcpyHostToDevice, stream);

3-D copy

cudaMemcpy3DParms p = { 0 };
p.srcPtr = make_cudaPitchedPtr(h_src, W*sizeof(float), W, H);
p.dstArray = cuArray;         // or p.dstPtr for 3-D pitched device ptr
p.extent = make_cudaExtent(W * sizeof(float), H, D);
p.kind   = cudaMemcpyHostToDevice;
cudaMemcpy3D(&p);

Copy to/from symbol (constant / `device` vars)

__constant__ float cData[64];
float h[64] = { /*...*/ };
cudaMemcpyToSymbol(cData, h, sizeof(h));                          // H→D
cudaMemcpyFromSymbol(h, cData, sizeof(h));                        // D→H
cudaMemcpyToSymbolAsync(cData, h, sizeof(h), 0, cudaMemcpyHostToDevice, stream);

Unified / Managed Memory

float* um;
cudaMallocManaged(&um, bytes);            // allocate unified

// Prefetch to device before kernel (avoid page-fault overhead)
int devId; cudaGetDevice(&devId);
cudaMemPrefetchAsync(um, bytes, devId, stream);

myKernel<<<grid, block, 0, stream>>>(um, n);

// Prefetch back to CPU
cudaMemPrefetchAsync(um, bytes, cudaCpuDeviceId, stream);
cudaStreamSynchronize(stream);

printf("%f\n", um[0]);   // safe after sync
cudaFree(um);

Memory advice hints

cudaMemAdvise(um, bytes, cudaMemAdviseSetReadMostly,      devId);  // allow read-duplication
cudaMemAdvise(um, bytes, cudaMemAdviseSetPreferredLocation, devId); // migrate toward device
cudaMemAdvise(um, bytes, cudaMemAdviseSetAccessedBy,      devId);  // map without migrating

Stream-Ordered Allocation (CUDA 11.2+)

cudaMallocAsync / cudaFreeAsync order allocation and free in a stream, avoid device-wide synchronization, and reuse freed memory from a driver-managed pool (cudaMemPool_t). This is the recommended modern allocation path for anything allocated/freed inside a processing loop.

float* d_buf;
cudaMallocAsync(&d_buf, bytes, stream);        // allocation ordered in the stream
myKernel<<<grid, block, 0, stream>>>(d_buf, n);
cudaFreeAsync(d_buf, stream);                  // free ordered in the stream
cudaStreamSynchronize(stream);

Memory pools

// Tune the device's default pool: keep freed memory cached instead of
// returning it to the OS at every synchronize (big perf win in loops)
cudaMemPool_t pool;
cudaDeviceGetDefaultMemPool(&pool, 0);
uint64_t threshold = UINT64_MAX;
cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold, &threshold);

// Explicit pool + allocate from it
cudaMemPoolProps props = {};
props.allocType     = cudaMemAllocationTypePinned;
props.location.type = cudaMemLocationTypeDevice;
props.location.id   = 0;
cudaMemPoolCreate(&pool, &props);

cudaMallocFromPoolAsync(&d_buf, bytes, pool, stream);
cudaFreeAsync(d_buf, stream);
cudaStreamSynchronize(stream);
cudaMemPoolDestroy(pool);

Gotchas: memory from cudaMallocAsync must be freed with cudaFreeAsync (or cudaFree); using it in another stream requires an event/stream dependency on the allocating stream. Check support with cudaDevAttrMemoryPoolsSupported.

Virtual Memory Management (CUDA VMM, CUDA 10.2+ driver API)

Allows reserving virtual address ranges and mapping physical memory dynamically. Check support via CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED.

#include <cuda.h>    // Driver API required

size_t granularity;
CUmemAllocationProp prop = {};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
prop.location.id   = 0;
cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM);

CUdeviceptr va;
cuMemAddressReserve(&va, totalSize, 0, 0, 0);  // reserve VA space

CUmemGenericAllocationHandle handle;
cuMemCreate(&handle, physicalSize, &prop, 0);   // allocate physical memory

cuMemMap(va, physicalSize, 0, handle, 0);       // map at va

CUmemAccessDesc access = {};
access.location = prop.location;
access.flags    = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
cuMemSetAccess(va, physicalSize, &access, 1);

// use (float*)va ...

cuMemUnmap(va, physicalSize);
cuMemRelease(handle);
cuMemAddressFree(va, totalSize);

Peer-to-Peer (Multi-GPU) Transfers

int canAccess;
cudaDeviceCanAccessPeer(&canAccess, 0, 1);   // can GPU 0 read GPU 1's memory?
if (canAccess) {
    cudaSetDevice(0);
    cudaDeviceEnablePeerAccess(1, 0);         // enable P2P

    // Direct D2D copy (NVLink or PCIe)
    cudaMemcpyPeer(d_dst_gpu0, 0, d_src_gpu1, 1, bytes);
    cudaMemcpyPeerAsync(d_dst, 0, d_src, 1, bytes, stream);
}

Memory Allocation API Summary

Function	Returns	Notes
`cudaMalloc`	`cudaError_t`	Linear device memory
`cudaMallocPitch`	`cudaError_t`	2-D pitched device memory
`cudaMalloc3D`	`cudaError_t`	3-D pitched device memory
`cudaMallocArray`	`cudaError_t`	CUDA array (texture/surface)
`cudaMallocManaged`	`cudaError_t`	Unified memory
`cudaMallocAsync`	`cudaError_t`	Stream-ordered device memory (CUDA 11.2+, recommended)
`cudaMallocFromPoolAsync`	`cudaError_t`	Stream-ordered, from an explicit `cudaMemPool_t`
`cudaMallocHost`	`cudaError_t`	Pinned host memory
`cudaHostAlloc`	`cudaError_t`	Pinned host with flags
`cudaHostRegister`	`cudaError_t`	Pin existing host allocation
`cudaHostUnregister`	`cudaError_t`	Unpin host allocation
`cudaFree`	`cudaError_t`	Free device/managed memory
`cudaFreeAsync`	`cudaError_t`	Stream-ordered free (pairs with `cudaMallocAsync`)
`cudaFreeHost`	`cudaError_t`	Free pinned host memory
`cudaFreeArray`	`cudaError_t`	Free CUDA array

Querying Pointer Attributes

cudaPointerAttributes attr;
cudaPointerGetAttributes(&attr, ptr);
// attr.type: cudaMemoryTypeUnregistered / Host / Device / Managed
// attr.device: which GPU
// attr.devicePointer, attr.hostPointer: for mapped memory

Common Pitfalls

Freeing host memory with cudaFree — use cudaFreeHost for pinned, free for regular.
Accessing device pointer from host — instant segfault; always copy back first.
Async copy without pinned source — cudaMemcpyAsync with pageable memory silently falls back to synchronous.
Forgetting cudaDeviceSynchronize before reading managed memory on the host on pre-Pascal hardware.
Not checking pitch — cudaMallocPitch sets pitch ≥ width; using width instead of pitch causes memory corruption.