100 DAYS OF GPU Course preview
Day 19Module 2
in-technical-review

Unified memory and what it costs

Managed memory removes the part of day 5 that everyone gets wrong. Three cudaMalloc calls become three cudaMallocManaged calls, all three cudaMemcpy calls go, and one pointer replaces the h_ and d_ pair. The CPU and GPU can both dereference it.

The answer is the same. The direction argument day 5 spends a section on has nothing left to be wrong about.

The bytes did not stop moving. They move inside the kernel launch now, one page at a time, triggered by the first thread that touches an address the card does not have. No line of your code names that traffic, and no stopwatch you have written so far points at it.

This page ports day 5 to unified memory, times the same kernel with the pages in three different places, and tells you which line to add when the number is bad.

One pointer, and the runtime decides the rest

cudaMallocManaged hands back one address that is valid on both sides. It does not decide where the memory lives. The runtime documentation is blunt about that: on a device that supports concurrent access, "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-08-30).

Two sentences from the programming guide are the whole model. "Managed memory is usually allocated in the memory space of the processor where it is first touched." And: "Managed memory is usually migrated when it is used by a processor other than the processor where it currently resides." (https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/understanding-memory.html , checked 2026-08-30.)

Fill the array from a host loop and it is in your CPU's DRAM. Launch a kernel over it and it has to come across.

The unit of that move is a page, not a float. Your CPU's pages are 4 KiB on x86; the GPU prefers 2 MiB or larger. When a thread loads an address whose page is on the other side, the load cannot retire and the driver moves the page: "this can be an expensive operation, and the amount of work is proportional to the page size."

That is page migration, and it happens while your kernel is running, on the clock you thought was measuring arithmetic.

So the host to device traffic day 5 wrote by hand still crosses PCIe. What changed is who schedules the crossing and whether you can see it in your own source.

A cold kernel stalls around repeated migrations of one 64 MiB buffer; prefetch moves it first, while resident memory sends 0 bytes across PCIe.

The trade people think they are making

The belief you arrive with is that managed memory is a copy you did not have to write, so it costs what the copy cost. That is wrong twice.

It is the wrong granularity. A cudaMemcpy is one bulk transfer of one contiguous range, issued once, by you. A fault-driven migration is many small moves, each one started by a thread that stalled, each one carrying its own bookkeeping.

The byte count is the same, but the schedule costs more.

It is also the wrong place to look when it is slow. The number gets bad, and day 11 suggests checking the access pattern or the block size. Neither changes it, because the cost is not in the kernel at all.

People also hear that cudaMemPrefetchAsync fixes the cold case and conclude that managed memory is free after all. It is not.

A prefetch performs the copy through a different API and keeps one pointer instead of two. Managed memory gives you that single pointer and lets you skip the prefetch when migration cost does not matter.

The thing you did not buy is a licence to stop synchronising:

cudaMallocManaged(&u_out, bytes);
vectorAdd<<<blocks, 256>>>(u_a, u_b, u_out, n);
std::printf("%f\n", u_out[0]);  // no sync, so the CPU races the kernel

Day 5's version of those three lines compiled too, and then the process died on the last one, because d_out was not an address it could read. Here the read succeeds and hands you whatever the page held. Managed memory turned a crash into a wrong answer.

Measuring where the pages are

Full program in code/day19-unified-memory/unified_memory.cu. Four rules hold the measurement together.

The kernel does not change. It is day 5's vectorAdd, character for character. It takes pointers, and it cannot tell whether they came from cudaMalloc or cudaMallocManaged, so the allocator is the only variable.

The size is not day 5's: 611 elements would migrate in one page and measure nothing, so the array is the 16,777,827 floats day 9 used, which is 64 MiB a buffer and enough pages to measure repeated migrations.

Every managed row is warm in its code and cold in its memory. The explicit cudaMalloc row runs first, through day 9's timeKernel, which warms vectorAdd three times before it times anything. So the module is loaded before any managed row starts, and a difference between rows is migration rather than the module load day 9 found.

The reset is a prefetch to the CPU. cudaMemPrefetchAsync with cudaCpuDeviceId puts every page back on the host, so "cold" means the same thing on every row instead of meaning whatever the row before it left behind. The same call with a device id is the optimisation being measured, which is why one function does both.

Four rows are single launches, and the program says so. A page migrates once. Ask a second time and you are measuring the resident case.

    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));

    // First touch is on the host, through an ordinary pointer. No cudaMemcpy,
    // no direction argument, no second pointer to keep in step.
    for (size_t i = 0; i < kElems; ++i) {
        u_a[i] = h_a[i];
        u_b[i] = h_b[i];
    }

u_ is the course's third pointer prefix, after h_ and d_, and it means this allocator and no other. The migration itself is one call per buffer, bracketed by CUDA events because the prefetch is stream ordered like everything else:

static float movePages(cudaEvent_t start, cudaEvent_t stop,
                       float* const* buffers, int count, size_t bytes,
                       int device) {
    CUDA_CHECK(cudaEventRecord(start));
    for (int k = 0; k < count; ++k) {
        CUDA_CHECK(prefetchTo(buffers[k], bytes, device));
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));

    float ms = 0.0f;
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    return ms;
}

One limit explains why the program asks the device before it asks the clock. Not every machine can run the table, so the first thing main() does is read four attributes, the three the programming guide's own decision tree uses plus managedMemory, and name which unified memory model you are using:

    int managedMemory = 0;
    int concurrent = 0;
    int pageable = 0;
    int hostPageTables = 0;
    CUDA_CHECK(cudaDeviceGetAttribute(&managedMemory, cudaDevAttrManagedMemory,
                                      device));
    CUDA_CHECK(cudaDeviceGetAttribute(
        &concurrent, cudaDevAttrConcurrentManagedAccess, device));
    CUDA_CHECK(cudaDeviceGetAttribute(&pageable,
                                      cudaDevAttrPageableMemoryAccess, device));
    CUDA_CHECK(cudaDeviceGetAttribute(
        &hostPageTables, cudaDevAttrPageableMemoryAccessUsesHostPageTables,
        device));

Hardware. If day 3 put you on WSL 2, this program prints concurrentManagedAccess 0, skips the migration table, runs the correctness check and exits 0. That is correct behaviour, not a broken build. NVIDIA: "Full Managed Memory Support is not available on Windows native and therefore WSL 2 will not support it for the foreseeable future ... CUDA queries will say whether it is supported or not and applications are expected to check this." (https://docs.nvidia.com/cuda/wsl-user-guide/index.html , section 5.1, checked 2026-08-30.) For the numbers, use a free Colab T4, which is native Linux. For the lesson, note that your machine migrates managed memory in bulk at launch and back at synchronize: a different cost with the same cause.

Results

Originally measured on a Tesla T4 with driver 595.84 and CUDA 12.6. The exact shipped program was re-verified with driver 580.173.02 and CUDA 13.0 (V13.0.88) on 2026-09-02. Correctness and the resident, cold and prefetch behaviors held, but pageableMemoryAccess changed from 1 to 0 and the timing values moved.

Both transcripts remain in evidence; the original 12.6 table follows.

GPU: Tesla T4 (compute capability 7.5)
n = 16777827 floats, 64.0 MiB per buffer, 256 threads/block, 65539 blocks

What this device says about unified memory
  managedMemory                          1
  concurrentManagedAccess                1
  pageableMemoryAccess                   1
  pageableMemoryAccessUsesHostPageTables 0
  paradigm  full, for every allocation, software coherence

Where the pages are when the kernel starts
  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 migration on its own
  3 buffers, host to device                    17.349 ms
  1 buffer, device to host                      9.922 ms

The first two rows are means of 10 runs after 3 warm-ups. The
other four are single launches, because a page migrates once and
you cannot get it back, so they carry more noise. Read those
four against each other, not against the means.

all 16777827 elements match the CPU reference
managed and explicit agree bit for bit

Cold managed memory remained about two orders of magnitude slower than the resident case. It was 69.916 against 0.798 ms under CUDA 12.6, and 83.336 against 0.793 ms under CUDA 13.0. Nothing about the arithmetic changed.

The difference is pages migrating from host to device while the kernel runs.

A prefetch removes essentially all of it. cudaMemPrefetchAsync brings the pages over in bulk before the launch, and the kernel then runs in 0.797 ms, within noise of the explicit cudaMalloc case at 0.776. CUDA 13.0 reproduced that at 0.792 and 0.768 ms.

Unified memory is not inherently slow. Faulting page by page during a kernel is slow, and one call fixes it.

Setting a preferred location did not remove the cold fault cost. It was 75.640 ms against 69.916 without advice under CUDA 12.6, then 80.572 against 83.336 under CUDA 13.0. Those noisy single launches reversed order, but both remain far from the roughly 0.8 ms prefetched path.

Advice is a placement hint, not an instruction to move pages, and is not a substitute for prefetch.

The migration on its own measured 17.349 ms under CUDA 12.6 and 24.214 ms under CUDA 13.0 for three buffers, which is the measured floor for moving this data once. Compare that to the 45.8 ms round trip day 9 measured with explicit copies for the same problem: unified memory is not paying a penalty for being managed, it is slower when the transfer happens during the kernel.

Both runs report managedMemory and concurrentManagedAccess as 1 and pageableMemoryAccessUsesHostPageTables as 0. The CUDA 13.0 run reports pageableMemoryAccess as 0, so its full unified-memory path applies to cudaMallocManaged allocations, not arbitrary pageable allocations.

The older driver reported that field as 1. On WSL2 unified memory is not fully supported and these numbers will not reproduce; if you are following day 3's Windows route, expect different behaviour here rather than assuming you broke something.

Run it yourself

Colab's free T4 or any card you own, on native Linux. The build line is in the repo's README:

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

There is no Compiler Explorer embed here. The program allocates six 64 MiB device buffers and moves 192 MiB of managed memory across the link once per row, which is more than a 20 second run cap and a shared sandbox should be asked to carry.

If you are on Windows, run it anyway: the attribute block at the top is the part of this lesson your machine can still teach you.

Exercise

Do the port yourself. Take day 5's vector_add.cu, raise its kElems from 611 to something in the tens of millions so there are pages to migrate, swap the three cudaMalloc calls for cudaMallocManaged, delete all three cudaMemcpy calls and the second set of pointers, and time one launch with the pages cold against one after a cudaMemPrefetchAsync to the device. Then run it once more with u_out left out of the prefetch.

Time: 30 to 45 minutes. Submit: cold divided by prefetched, and one sentence on why leaving u_out out costs you anything at all.

Check: three gates, and they are the grading. The managed result matches the CPU reference at all 16,777,827 elements; the explicit result matches the same reference; and the two agree with no tolerance at all, because one kernel over one set of inputs has to produce the same bits whichever allocator held the memory.

A tolerance failure means you broke the port. A bit-for-bit failure between two correct-looking results means you broke the memory.

Hint 1

Two things in your program can be cold: the code and the memory. Only one of them is the subject of this lesson. Which launch is the first launch of vectorAdd in your process, and what does that do to your first number?

Hint 2

Three buffers cross the link and the kernel reads only two of them. Count the pages still on the host when the launch starts, not the buffers. What does the hardware do when a thread stores to an address the card does not have?

Solution

Report the ratio, not the milliseconds. The value changes by card, but it should stay above 1. If your ratio comes out near 1, go back to hint 1.

A cold-memory row that is also the first launch of vectorAdd includes a module load, and day 9 measures that cost.

The buffer you did not prefetch still faults. A store to a page the card does not have is a fault like a load to one, and nothing in the machine knows this kernel is about to overwrite every byte of that page, so the page is fetched before it is written over.

Write-only output buffers are the case people forget, and they are a third of the traffic here.

The rule to keep: prefetch every buffer the kernel will touch. Faults are per page, so the cost of an omitted buffer depends on its size.

Pitfalls

You read a managed pointer on the host straight after the launch and get the old values. The launch is asynchronous, so the CPU races the kernel for the same pages. Put a cudaDeviceSynchronize() between them.

Where concurrentManagedAccess is 0 it is worse than stale data, and the guide is blunt: "The CPU must not access managed memory while the GPU is active."

Your code builds on one toolkit and not the other. CUDA 13.0 changed both hint APIs: cudaMemPrefetchAsync took int dstDevice and now takes a cudaMemLocation, and cudaMemAdvise did the same. NVIDIA made the same edit to its own samples (https://github.com/NVIDIA/cuda-samples/blob/master/CHANGELOG.md , checked 2026-08-30).

Put the two calls behind #if CUDART_VERSION >= 13000 wrappers once, the way this program does, rather than at seven call sites.

You call cudaMemAdvise and expect it to move data. It does not.

"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-08-30.)

cudaMemPrefetchAsync is the call that moves bytes.

You prefetch into a busy stream and it starts late. The prefetch is stream ordered. "The migration does not begin until all prior operations in the stream have completed, and completes before any subsequent operation in the stream."

Issued behind a long kernel, it finishes after the launch it was meant to help.

Your numbers do not match this page and you are on Windows or WSL 2. Nothing is broken. Managed memory there migrates in bulk when the GPU starts and back when you synchronize, oversubscription is refused, and both hint APIs refuse the device: the guide lists the differences in section 2.6.2.3.

Read the attribute block the program prints.

You profile and time 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-08-30). Take the fault counts from the profiled run and the milliseconds from a clean one.

Go deeper

Next

Day 20 closes module 2 with the convolution capstone, and it allocates explicitly rather than using anything on this page, which shows where managed memory belongs once a pipeline keeps its data on the card. Day 53 measures the pinned host memory the explicit path here left on unused, and day 93 tunes a managed workload with the full cudaMemAdvise set.

You now know that global memory can be somewhere else, and how to ask.