Residency Hints and the Arithmetic of Managed-Memory Thrashing

A managed allocation has no fixed home. The driver decides where each page resides, based on where it was last accessed, and moves it when the access pattern changes. That policy works well when the working set fits in device memory and badly when it does not, and the advisory interface exists so the application can supply information the driver cannot observe: which device will use a buffer, whether a buffer is read-only for the duration, and which side is allowed to touch it at all.

The advice is worth understanding precisely, because the gap between what it promises and what it guarantees is where most of the disappointment lives.

What the driver can observe on its own

The hardware records, per page, a coarse summary of accesses: whether the page was read, whether it was written, and by which processor. The granularity is a page, and the summary is sampled rather than exhaustive. From that, the driver builds a picture of recency and can approximate a least-recently-used ordering for eviction.

Two things are missing from that picture. The driver does not know the future: it cannot tell that a page evicted now will be needed by the next kernel, and it does not know the application’s structure, so it cannot distinguish a page that is genuinely cold from one that is part of a working set being cycled through faster than the sampling can track. The second case is the interesting one, because the driver’s recency estimate degrades exactly when it matters most.

The counters also lag. A page that was just migrated carries a fresh count, but the count reflects accesses since the migration, not accesses before it. A buffer that ping-pongs between host and device accumulates a misleading history on each side, which is why the pathological case is stable rather than self-correcting: each migration looks justified by the accesses that preceded it.

The advisory operations and their actual contract

Three kinds of advice are commonly available. Preferred location names the processor where a range should live when nothing forces it elsewhere. Accessed-by names a processor that is expected to access a range, which lets the driver establish a mapping in advance rather than fault on first touch. Read-mostly marks a range as not being written, which allows the driver to keep one copy on the device and one on the host rather than migrate it back and forth.

The duplication case is the one that changes the asymptotics rather than the constants. Without it, a buffer read by both the host and the device migrates on every alternation, and the cost per alternation is a full copy of the range over the bus. With read-mostly established and no writes to invalidate it, both processors read their local copy and no migration occurs. The condition is absolute: a single write from either side invalidates the arrangement, and the driver falls back to migration.

None of these are guarantees. The documentation is explicit that advice is a hint the driver may ignore when it has better information or when honouring it would conflict with correctness. A preferred location can be overridden when the device is out of memory. An accessed-by mapping can be dropped under pressure. Read-mostly is a promise the application makes, and breaking it costs more than never having made it, because the driver must then invalidate and re-establish the arrangement.

#include <cuda_runtime.h>

void prepare(int *p, size_t bytes, int dev)
{
    /* Read-mostly: allow a host copy and a device copy to coexist.
       Valid only while neither side writes to the range. */
    cudaMemAdvise(p, bytes, cudaMemAdviseSetReadMostly, dev);

    /* Preferred location biases placement when the driver has a choice.
       It does not pin the range: pressure can still evict it. */
    cudaMemAdvise(p, bytes, cudaMemAdviseSetPreferredLocation, dev);

    /* Establish device mappings ahead of the first access so the
       kernel does not take a fault per page on its first pass. */
    cudaMemAdvise(p, bytes, cudaMemAdviseSetAccessedBy, dev);

    /* Migrate now rather than waiting for a fault. */
    cudaMemPrefetchAsync(p, bytes, dev, 0);
    cudaDeviceSynchronize();
}

Distinguishing thrashing from slow migration

A program using managed memory badly has one of two problems, and they need opposite fixes. The first is too many faults: the data is in the right place eventually, but every first touch pays a fault and the fault count is proportional to the page count. The fix is prefetching and larger pages, and the symptom is a large amount of time in the driver with relatively little bus traffic after the initial migration.

The second is too much migration: the data moves repeatedly, and the bus traffic is proportional to the number of migrations times the range size. The fix is restructuring so the working set fits, or marking ranges read-mostly, or keeping one side’s copy in host-mapped memory so it never needs to move. The symptom is sustained bus traffic long after the first touch, with the traffic continuing for the duration of the kernel rather than tapering off.

The distinction is visible in the shape of the timeline, not in a single number. A program that migrates once and then runs has a burst of transfer followed by a quiet period. A program that thrashes has transfer spread evenly across the whole run, and the compute time between transfers shrinks as the working set grows relative to capacity. If the ratio of transfer time to compute time rises with problem size, capacity is the constraint; if it is flat but high, the access pattern is.

There is a third case that looks like thrashing and is not: a buffer that is genuinely shared and genuinely written by both sides. No amount of advice fixes that, because the data really does have to move. The right response is to restructure the algorithm so one side owns the buffer and the other receives results, rather than both sides writing the same range.

Sizing the working set against capacity

The arithmetic that decides whether oversubscription is workable is simple enough to do before writing code. Let the working set be the bytes touched between two reuses of the same data. Let device capacity be the usable portion of device memory after the runtime, the driver, and any other allocations have taken their share. If the working set is below capacity, migration converges: each page moves at most once per phase, and the cost is proportional to the number of phases.

If the working set exceeds capacity, migration does not converge. Pages are evicted before they are reused, and each phase re-migrates data that was migrated in the previous phase. The steady-state transfer volume per phase becomes proportional to the working set rather than to the amount of new data, and the program’s throughput is bounded by the link rather than by the device.

The useful rule is to size tiles so the working set is comfortably below capacity, by a healthy fraction rather than a hair’s breadth, because the driver’s eviction decisions are approximate and other allocations compete. A tile that exactly fills device memory leaves no room for the driver’s own structures and none for the approximation error, and it will thrash intermittently in a way that is hard to attribute.

Tile size also interacts with fault granularity. A tile whose boundaries do not align to the large page size forces the driver to mix small and large pages at the edges, and the small-page region is where the fault count concentrates. Aligning tiles to 2 MiB boundaries, and to whatever the platform’s large page size is, removes that edge effect.

What changes on a coherent platform

On an integrated device that shares physical memory with the CPU, there is nothing to migrate. A managed page is reachable by both processors at local latency, and the driver’s residency policy is close to a no-op. The same program that thrashes on a discrete GPU runs without transfer traffic on an integrated one, and the advice calls become irrelevant rather than harmful.

That portability gap is a design hazard. Code tuned for one platform encodes assumptions such as tile sizes chosen to fit a device’s memory, prefetches timed against a link’s latency, all of which are wrong on the other. A tile sized for a discrete device’s memory is far smaller than necessary on an integrated one, and the resulting loss is parallelism rather than transfer. The reverse direction is worse: a tile sized for shared memory oversubscribes the discrete device and thrashes.

The practical response is to derive tile size from a queried capacity rather than a constant, and to treat the prefetch and advice calls as no-ops when the platform reports that host and device memory are the same. The query itself is cheap and the branch costs nothing.

Where the policy ends

Advice cannot create capacity, cannot make a dependent access pattern coalesce, and cannot make a write-shared buffer single-owner. It can only remove migrations that were happening for lack of information. When a program’s transfer volume is already minimal for its algorithm, the remaining cost is the algorithm’s, and the fix is at the level of decomposition rather than memory placement, the same question that arises when splitting work across multiple devices, covered in partitioning work across devices.

The fault path that advice is trying to avoid is described in GPU page faults and demand migration, and the alternative of never migrating at all, by mapping host memory into the device permanently, is covered in host-mapped memory over PCIe. Where the kernel rather than the runtime owns the placement decision, the mechanism is the page table mirroring described in page table mirroring between CPU and device.