100 DAYS OF GPU Course preview
Day 93Module 10
in-technical-review

Unified memory, second pass

Day 19 put one kernel launch on cold managed memory and measured 69.916 ms against 0.797 ms after a prefetch. The lesson people take from that number is that the fix is a single call, added once, at the top.

It is not, because a real program does not launch one kernel. It launches a kernel, reads the result on the CPU, launches another, and the pages walk back and forth across PCIe every time. The prefetch that fixed day 19 has to be repeated, and repeating a 64 MiB transfer four times is not obviously better than never having managed the memory at all.

This page runs a four-round workload five ways against an explicit cudaMemcpy baseline, and separates the call that moves bytes from the three that only change where a fault lands.

What faults, and how often

A managed allocation has no fixed home. The runtime documentation says "managed memory may not be populated when this API returns and instead may be populated on access" (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html , checked 2026-09-01). Populated on access means a fault.

The sequence, on a device where concurrentManagedAccess is 1: a warp issues a load, the address belongs to a page the GPU does not have, the load cannot retire, the driver is interrupted, the page is transferred, the warp resumes. The programming guide is direct about the price: "this can be an expensive operation, and the amount of work is proportional to the page size" (https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/unified-memory.html , checked 2026-09-01). Your CPU's pages are 4 KiB on x86; the GPU prefers 2 MiB. A 64 MiB buffer is 16,384 of the first kind.

Two things make that worse than it looks. The faults happen inside the kernel, so a profile reporting only kernel time attributes migration to your arithmetic. And page migration is a move, not a copy: the page leaves the host. Read the same array on the CPU afterwards and it walks back, once per round, in both directions, for as long as the loop runs.

Diagram: where one 64 MiB table lives across four rounds. Three bands. Host DRAM as a strip on the left, device DRAM on the right, four numbered round markers down the middle, and one arrow between the strips per crossing of the whole table. Band 1, naive: eight arrows, one into each kernel and one back into each CPU pass, the device-side ones drawn broken into many small arrows. Caption "8 crossings, 512 MiB, all of it inside a kernel or a host loop." Band 2, prefetched: the same eight, each drawn whole and placed before its kernel. Caption "8 crossings, 512 MiB, every one a call you wrote." Band 3, read-mostly: one arrow before round 1, a second copy of the table on the host strip, no arrows after that. Caption "1 crossing, 64 MiB, two read-only copies." Alt text: "One sixty-four mebibyte lookup table across four rounds. Faulting moves it eight times, prefetching moves it eight times on purpose, and read mostly duplicates it once and then moves nothing."

The hint that moves nothing

Almost everyone meets cudaMemAdvise as the other prefetch. Same argument list, same header, next to cudaMemPrefetchAsync in the same section of the manual, so it reads as the declarative version of the same operation.

It is not an operation on data at all. NVIDIA: "Setting the preferred location does not cause data to migrate to that location immediately. Instead, it guides the migration policy when a fault occurs on that memory region" (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html , checked 2026-09-01). Day 19 measured what that costs when you get it backwards: a cold buffer with SetPreferredLocation came in at 75.640 ms against 69.916 with no advice, slower, because the advice changed which policy the fault handler applied and avoided no faults.

The three pieces of advice answer three different questions, and none of them is "when do the bytes move".

SetReadMostly           // nobody writes this, so let both sides keep a copy
SetPreferredLocation    // when this faults, resolve it here
SetAccessedBy           // map it for that processor, do not migrate on access

SetReadMostly is the only one of the three that can delete traffic outright, and it does it by duplicating rather than moving: a read-only region may exist on the CPU and the GPU at once, so a host pass over it stops pulling pages off the card. The catch is in the name. Write to a read-mostly region and the duplicates collapse, which turns the saving into an extra step.

SetPreferredLocation decides where a fault resolves, so an array stops drifting back to whoever touched it last. SetAccessedBy establishes a mapping instead, so a rare reader reaches the data over the link and leaves it where it is. Neither transfers anything, which is why a program that only advises can still be slow.

Five rows over one workload

Full program in code/day93-unified-memory/managed_tuning.cu.

The workload is four rounds of out[i] += scale * table[keys[i]] with a host pass over the whole table between rounds. Three buffers, 64 MiB each, picked so each piece of advice has exactly one array it is the right answer for: u_table is read by both processors every round, u_keys by the device every round and the host once, and u_out is written by the device every round and read by the host at the end.

The rows are an explicit cudaMalloc path with two cudaMemcpy calls, then managed naive, prefetch before every kernel, advice only, and both. All five run the same kernel over the same inputs, so the answer is bit-identical across them and the program gates on that rather than hoping.

The advice, one call per access pattern:

static void applyAdvice(float* u_table, unsigned int* u_keys, float* u_out,
                        int device) {
    // Read by both processors, written by neither, so each may keep its own
    // read-only copy and the host pass stops pulling pages off the card.
    // The device argument is ignored for this advice.
    CUDA_CHECK(
        adviseFor(u_table, kTableBytes, cudaMemAdviseSetReadMostly, device));

    // Read by the device on every round, touched by the host once at fill
    // time. A preferred location does not migrate anything; it decides where
    // a fault resolves, so these pages settle on the card and stay.
    CUDA_CHECK(adviseFor(u_keys, kKeyBytes, cudaMemAdviseSetPreferredLocation,
                         device));

    // Written by the device every round and read by the host once, at the
    // end. Preferred location keeps it on the card; AccessedBy maps it for
    // the host so that last read can cross the link instead of migrating.
    CUDA_CHECK(
        adviseFor(u_out, kOutBytes, cudaMemAdviseSetPreferredLocation, device));
    CUDA_CHECK(adviseFor(u_out, kOutBytes, cudaMemAdviseSetAccessedBy,
                         cudaCpuDeviceId));
}

adviseFor is day 19's version wrapper, unchanged: CUDA 13.0 turned the int device parameter of both hint APIs into a cudaMemLocation, so one #if CUDART_VERSION >= 13000 keeps the file building on both.

The round loop, where the events go:

    for (int r = 0; r < kRounds; ++r) {
        if (prefetch) {
            CUDA_CHECK(cudaEventRecord(phaseStart));
            CUDA_CHECK(prefetchTo(u_table, kTableBytes, device));
            CUDA_CHECK(prefetchTo(u_keys, kKeyBytes, device));
            CUDA_CHECK(prefetchTo(u_out, kOutBytes, device));
            CUDA_CHECK(cudaEventRecord(phaseStop));
            row.moveMs += elapsedMs(phaseStart, phaseStop);
        }

        CUDA_CHECK(cudaEventRecord(phaseStart));
        gatherAccumulate<<<blocksFor(kElems), kThreadsPerBlock>>>(
            u_table, u_keys, u_out, kElems, scaleFor(r));
        CUDA_CHECK(cudaEventRecord(phaseStop));
        CUDA_CHECK(cudaGetLastError());

        // elapsedMs waits on phaseStop, so the launch has finished before
        // the host touches the table below. That is the required
        // synchronisation, not an accident of the timing.
        row.roundMs[r] = elapsedMs(phaseStart, phaseStop);
        row.kernelMs += row.roundMs[r];

        if (tableChecksum(u_table) != want) {
            *badRounds += 1;
        }
    }

An outer event pair wraps the whole row, from before the first kernel to after the host has read the last byte of the output, so a migration cannot escape by happening somewhere the program calls no function. Subtracting the kernel and prefetch totals leaves the column the naive row is going to lose in.

The prefetch is stream ordered, which is what makes an event bracket around it mean anything: "the migration does not begin until all prior operations in the stream have completed". Every call here is on the default stream, so nothing overlaps and nothing hides.

Results

Re-verified on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). Every row again matched the double reference and the explicit row bit for bit. The combined strategy moved from 7.7 percent slower than explicit to slightly faster, and prefetch became marginally closer to explicit than to naive. Both clean transcripts and the fresh Nsight Systems report are listed in front matter; the table is the clean CUDA 13 run.

row total ms kernel ms move ms host ms
explicit cudaMalloc 152.487 6.136 44.817 101.535
managed, naive 335.660 132.595 0.000 203.065
managed, prefetch 242.813 6.147 37.408 199.257
managed, advise 195.597 65.720 0.000 129.877
managed, advise + prefetch 150.375 6.148 18.722 125.506

The five predictions resolved as follows.

  1. Held. Naive managed memory took 335.660 ms, 2.20x the explicit row's 152.487 ms.
  2. Held narrowly under CUDA 13.0. Prefetch collapsed kernel time to 6.147 ms and made migration visible as 37.408 ms. Its 242.813 ms total was marginally closer to explicit than naive, reversing CUDA 12.6's result, though host time still dominated at 199.257 ms.
  3. Held. Advice alone took 195.597 ms, faster than naive and faster than prefetch, while its first kernel round still paid 61.102 ms.
  4. Partly held, partly refuted. Advice made rounds 2 to 4 flat at about 1.54 ms, but the naive row also fell sharply after round 1, from 65.881 to roughly 22 ms. The claim that naive rounds would not fall was wrong.
  5. Held on this card. concurrentManagedAccess printed 1; pageableMemoryAccess and pageableMemoryAccessUsesHostPageTables printed 0.

The combined row came closest to explicit at 150.375 ms and was slightly faster in this run. Milliseconds belong to this driver and card; the per-round shape is the portable diagnostic. The profiler run was deliberately kept out of the table because tracing perturbed the timings. It nevertheless records unified migration in both directions and thousands of CPU and GPU page faults; the exact totals remain in the capture transcript.

Run it yourself

A free Colab T4 is enough, and so is any discrete NVIDIA card on native Linux. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o managed_tuning managed_tuning.cu

No Compiler Explorer embed. Six 64 MiB buffers crossing the link dozens of times is more than a shared 20 second sandbox should carry, and shrinking it would shrink the migration that is the subject. On Windows or WSL 2 run it anyway: the attribute block it prints before exiting is the part your machine can still teach you.

Exercise

Give the prefetch somewhere to hide. Move it onto a second stream and issue round r + 1's copy before round r's kernel, so the next table arrives while the current kernel runs. Then measure whether the total moved.

Time: 40 to 60 minutes. Submit: total ms and move ms for the overlapped row beside the sequential prefetch row, and one sentence on which column changed.

Check: the program's four gates already cover your row. The output must sit inside the printed tolerance against the double reference and be bit-identical to the explicit row, the table checksum must match on every round, and no timed region may report zero milliseconds. A bit-for-bit failure means you moved a kernel, not a transfer. A checksum failure on round 2 or later means your prefetch and your host pass are touching the same pages at once.

Hint 1

You are not trying to make the transfer faster; it already runs at the link's rate. You are trying to make it happen while something else does. What in this loop does not need next round's pages?

Hint 2

The migration "does not begin until all prior operations in the stream have completed". Which stream are the prefetch and the kernel on now, and what does that sentence say about their order? Then ask what the host pass does to any overlap you win.

Solution

move ms should fall towards zero for rounds 2 to 4 while kernel ms stays where it was: the transfer still happens, it is just no longer alone on the clock. total ms improves by less than move ms did, and the host pass is why. The CPU loop between kernels is a serialisation point no stream can overlap, so the win is capped by the fraction of each round the GPU was busy.

Report the pair, not the total. A prefetch you overlapped and a prefetch you deleted look identical in one number, and only the first is still spending bandwidth you may need elsewhere.

Pitfalls

You advised and nothing happened. cudaMemAdvise moves no bytes. Day 19 measured a preferred location making a cold launch slower, 75.640 ms against 69.916, because it changed the fault policy and avoided no faults. cudaMemPrefetchAsync is the call that transfers.

You set SetReadMostly on an array the GPU writes. The duplicates collapse on the first write, so you paid for a duplication and threw it away. The advice is for lookup tables, weights and anything filled once and read forever. This program gates on it: the host's checksum of the table must match on every round of every row, so a stale duplicate fails a branch rather than quietly returning a wrong answer.

Your prefetch runs after the thing it was meant to help. It is stream ordered like a launch, so issued behind a long kernel on the same stream it completes after that kernel and shows up as pure added time. Day 51 is the lesson on where a stream-ordered call lands.

You profiled and timed in the same run. Nsight Systems warns that page-fault tracking "may cause significant runtime overhead" (https://docs.nvidia.com/nsight-systems/UserGuide/index.html , checked 2026-09-01). Fault counts from the profiled run, milliseconds from a clean one.

You reached for managed memory to make transfers faster. It does not. Pinned host memory raises the rate a transfer runs at, and day 53 measures that separately. Managed memory changes who schedules the crossing and when, and on a workload that keeps its data on the card neither beats not crossing at all.

Go deeper

Next

Day 94 partitions one GPU between two workloads with green contexts, MPS and MIG, the other way a program stops having the card to itself. Global memory has been a place since day 11; after this page it is a place with a policy. The LLM kernels in days 95 to 99 never set that policy, because they keep their weights on the card from load to last token.