100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryMemory
CC 7.5

What is unified memory in CUDA?

One pointer valid on both the host and the device, with the driver moving pages between them when the other side touches them.

Swap cudaMalloc for cudaMallocManaged and a lot of code disappears. The cudaMemcpy calls go, the direction argument that people get backwards goes with them, and the h_ and d_ pointer pair collapses into one address both sides can dereference. The bytes did not stop crossing PCIe, though. What changed is who schedules the crossing: the allocation does not decide where the memory lives, first touch does, and "managed memory is usually migrated when it is used by a processor other than the processor where it currently resides".

The unit of that migration is a page, and the trigger is a fault. A thread loads an address whose page is on the other side, the load cannot retire, the driver moves the page, and all of that happens inside the kernel launch, on a clock you thought was timing arithmetic. That is page migration, and it is why the first belief to drop is that managed memory costs what the copy cost. A cudaMemcpy is one bulk transfer you issue once. A fault storm is thousands of small moves, each started by a stalled thread, each with its own bookkeeping. Same bytes, different schedule, and the schedule is the expensive part.

The second belief is that cudaMemPrefetchAsync makes the whole thing free. It does not. A prefetch is the copy written again with a different spelling and one pointer instead of two, and what you actually bought was the single pointer plus permission to skip the prefetch wherever it does not matter. cudaMemAdvise is not a substitute either: setting a preferred location "does not cause data to migrate to that location immediately", it only steers the policy when a fault happens, and on a cold buffer that means the same faults with different bookkeeping.

What managed memory definitely does not buy is a licence to stop synchronising. Read a managed pointer on the host straight after a launch and the CPU races the kernel for the same pages. Day 5's explicit version killed the process on that line, because a device pointer is not an address the host can read. Here the read succeeds and hands you whatever the page held, so the failure mode moved from a crash to a wrong answer. Put a cudaDeviceSynchronize() in between.

Measured

Six rows, one kernel, 16,777,827 floats at 64 MiB a buffer, with the pages reset to the host between rows so that "cold" means the same thing every time. Day 19 on a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), nvcc -O3 -arch=sm_75.

Where the pages are when the kernel starts Time
explicit cudaMalloc, resident 0.776 ms
managed, resident 0.776 ms
managed, resident, one launch 0.798 ms
managed, cold, one launch 69.916 ms
managed, cold, preferred location set 75.640 ms
managed, prefetched, one launch 0.797 ms

The first two rows are means of 10 runs after three warm-ups; the other four are single launches, because a page migrates once and asking again measures the resident case. Read them against each other rather than against the means. The cold row is the entire cost of unified memory on this workload, and one prefetch call removes it: 0.797 ms, within noise of the explicit allocation at 0.776. Advice made it worse rather than better. Moving the same three buffers deliberately took 17.349 ms, which is the honest floor for getting this data onto the card once.

This device reported managedMemory 1, concurrentManagedAccess 1, pageableMemoryAccess 1 and pageableMemoryAccessUsesHostPageTables 0, which is the full software-coherence case. On WSL 2 those attributes come back differently, because full managed memory support "is not available on Windows native and therefore WSL 2 will not support it for the foreseeable future". Nothing on this page reproduces there, and that is your machine rather than your code.

Code

From code/day19-unified-memory/unified_memory.cu. Three allocations and a host loop, with nothing between them that names a transfer.

    float* u_a = nullptr;
    float* u_b = nullptr;
    float* u_out = nullptr;
    CUDA_CHECK(cudaMallocManaged(&u_a, bytes));
    CUDA_CHECK(cudaMallocManaged(&u_b, bytes));
    CUDA_CHECK(cudaMallocManaged(&u_out, bytes));

    for (size_t i = 0; i < kElems; ++i) {
        u_a[i] = h_a[i];
        u_b[i] = h_b[i];
    }

That loop is the first touch, so all three buffers now live in host DRAM, including u_out, which the kernel only writes. A store to an absent page faults exactly like a load to one, and nothing in the machine knows the kernel is about to overwrite every byte, so a write-only output buffer still gets fetched before it is written over. Leave it out of your prefetch and it is a third of the traffic you did not remove.

Diagram

An original SVG: host DRAM on the left, the PCIe link down the middle, device DRAM on the right, and the kernel drawn as a bar whose length is its time. Band one, cold, with fault arrows breaking the bar up. Band two, prefetched, with one wide arrow before an unbroken bar. Band three, resident, with nothing crossing.

Alt text: "The same kernel over sixty-four mebibytes per buffer, three ways. Cold, the page traffic hides inside the kernel bar and it runs for 69.916 milliseconds. Prefetched, one call moves the pages first and the kernel takes 0.797. Resident, nothing crosses the link."

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.