100 DAYS OF GPU Course preview
Day 55Module 6
in-technical-review

Stream-ordered allocation and memory pools

Give each iteration of day 54's pipeline a fresh scratch buffer, and its throughput falls. The kernels and sizes have not changed. The only new calls are cudaMalloc(&d_tmp, bytes) at the top of the loop and cudaFree(d_tmp) at the bottom.

Those calls stop the overlap from day 51 and the pipeline from day 54. A timer shows the cost but not its cause. This page compares three versions of the loop and uses an Nsight Systems timeline to find the call responsible.

What cudaFree waits for

cudaMalloc and cudaFree are not stream operations. They take no stream argument and act on the whole device. A block being freed may still be in use by queued kernels, so the allocator must wait until they can no longer write to it.

NVIDIA's walkthrough of this API says: "the first cudaFree call has to wait for kernelA to finish, so it synchronizes the device before freeing the memory" (https://developer.nvidia.com/blog/using-cuda-stream-ordered-memory-allocator-part-1/ , checked 2026-09-01).

Put that inside a loop and each iteration queues two kernels, then waits for the device to finish. This has the same effect as a device synchronise. The queue never holds more than one iteration, so the program cannot overlap iterations.

cudaMallocAsync and cudaFreeAsync change the ordering rules rather than the cost of the bytes. Each takes a stream, so the host queues the allocation and free like kernel launches, then continues.

The memory comes from a memory pool owned by the device, and "Each call to cudaFreeAsync returns memory to the pool, which is then available for re-use on subsequent cudaMallocAsync requests" (same post, checked 2026-09-01). The second iteration's allocation is a bookkeeping entry against a block the pool already holds, not a request to the OS. The host does not wait for the device.

Diagram: one iteration, three allocation strategies. Three horizontal bands, each a CUDA API row above a GPU kernel row on one time axis, showing three loop iterations. Band 1, malloc per iteration: each iteration is a cudaMalloc bar, two short launch bars, then a wide cudaFree bar spanning the kernels below, with a gap on the kernel row before the next iteration starts. Caption "100 allocations, 100 device drains." Band 2, hoisted: the API row is a burst of narrow launch bars that ends early; the kernel row is solid. Caption "1 allocation, 0 drains." Band 3, cudaMallocAsync: same solid kernel row as band 2, and the API row adds narrow alloc and free bars but still finishes ahead of the device. Caption "100 allocations, 0 drains." Alt text: "Freeing with cudaFree drains the device every iteration and leaves gaps on the kernel row. A hoisted buffer and a hundred stream-ordered allocations both keep the kernel row solid."

Host malloc does not track GPU work

Host malloc and free are calls into a user-space allocator that mostly touches a free list; nothing about them cares what your other threads are doing with unrelated data, and putting them in a loop is rarely what kills a C program. So cudaMalloc in a loop looks like a style problem, not a performance problem.

On the GPU, the free depends on the work queue because the caller returns before the kernels run. Only the driver knows when the block is idle, so cudaFree waits for all queued work.

cudaFreeAsync resolves it by making the free part of the queue itself: the block is reclaimed when the stream reaches the free, after the kernels that use it and before whatever is queued next. This keeps the same safety without stopping the host.

Day 51 gave you the rule that order lives in streams rather than in source code; this page extends that rule to the lifetime of memory.

Three loops, one accumulator

Full program in code/day55-mallocasync/mempool.cu. Three properties matter.

Every loop does identical work and must produce identical results. One hundred iterations, each writing 2 * in[i] into a 16 MiB scratch buffer and accumulating it. The three accumulators are checked against the host reference, so "the fast version skipped work" is a failure, not a footnote.

The events bracket the whole loop, on the stream. Device time the GPU spends idle while the host stands inside cudaMalloc or cudaFree is inside the measurement, because that idle time is the cost being tested. A warm-up before any timing covers both kernels and both allocators, so the first cudaMallocAsync does not pay for creating the pool's first block inside a measured phase.

Every phase is an NVTX range, named the way day 41 named them, so nsys stats --report nvtx_sum prints the comparison as a table without the GUI.

Here is the loop everyone writes first, and the one this page exists to replace:

// The loop everyone writes first. cudaMalloc and cudaFree are not stream
// operations: cudaFree cannot hand back memory the queued kernels still
// use, so it waits for the device to drain, once per iteration.
static void loopMallocPerIter(const float* d_in, float* d_scratch, float* d_acc,
                              cudaStream_t stream, int blocks) {
    (void)d_scratch;  // this variant allocates its own
    nvtxRangePushA("malloc-per-iter");
    for (int iter = 0; iter < kIters; ++iter) {
        float* d_tmp = nullptr;
        CUDA_CHECK(cudaMalloc(&d_tmp, kScratchBytes));
        writeScaled<<<blocks, kThreadsPerBlock, 0, stream>>>(d_in, d_tmp,
                                                             kElems);
        accumulate<<<blocks, kThreadsPerBlock, 0, stream>>>(d_tmp, d_acc,
                                                            kElems);
        CUDA_CHECK(cudaFree(d_tmp));
    }
    CUDA_CHECK(cudaGetLastError());
    nvtxRangePop();
}

The stream-ordered version differs by exactly two calls, which is the point: this is a mechanical replacement, not a rewrite. The hoisted variant between them, one buffer allocated before the loop, is in the repo and is the baseline both others are judged against.

// Same shape as the first loop, but allocation and free are operations in
// stream order against the device's default pool. The host queues them and
// moves on; the free returns the block to the pool, and the next iteration's
// cudaMallocAsync reuses it without touching the OS.
static void loopMallocAsync(const float* d_in, float* d_scratch, float* d_acc,
                            cudaStream_t stream, int blocks) {
    (void)d_scratch;  // this variant allocates its own
    nvtxRangePushA("malloc-async");
    for (int iter = 0; iter < kIters; ++iter) {
        float* d_tmp = nullptr;
        CUDA_CHECK(cudaMallocAsync(&d_tmp, kScratchBytes, stream));
        writeScaled<<<blocks, kThreadsPerBlock, 0, stream>>>(d_in, d_tmp,
                                                             kElems);
        accumulate<<<blocks, kThreadsPerBlock, 0, stream>>>(d_tmp, d_acc,
                                                            kElems);
        CUDA_CHECK(cudaFreeAsync(d_tmp, stream));
    }
    CUDA_CHECK(cudaGetLastError());
    nvtxRangePop();
}

The pool has one knob this lesson touches: the release threshold. CUDA 12.6 documents it as the "Amount of reserved memory in bytes to hold onto before trying to release memory back to the OS" with a default of zero (https://docs.nvidia.com/cuda/archive/12.6.2/cuda-runtime-api/group__CUDART__MEMORY__POOLS.html , checked 2026-09-01). Default zero means every synchronisation empties the pool, and the next allocation after a sync goes back to the OS as if the pool were not there.

// The release threshold is the amount of memory the pool may keep across a
// synchronisation. At the default of zero, any stream or event synchronise
// hands everything back to the OS; at UINT64_MAX the pool keeps what it has.
static void raiseReleaseThreshold(cudaMemPool_t pool) {
    uint64_t threshold = UINT64_MAX;
    CUDA_CHECK(cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold,
                                       &threshold));
}

The program proves the knob works by reading cudaMemPoolAttrReservedMemCurrent twice: after a synchronise at the default threshold, and after a synchronise with the threshold raised.

Note. Inside a loop with no synchronisation, the default threshold costs nothing, because releasing happens at sync points and the loop has none. The threshold matters to a program that syncs between batches, which is every real program, and that is why the demo reads the pool after a sync rather than inside the loop.

Results

Re-verified 2026-09-02 on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). All correctness and pool-reservation gates reproduced. The loop ratios moved only from 2.56x to 2.52x and from 1.01x to 1.01x.

Both transcripts and the fresh CUDA 13 report are listed in front matter. The published CUDA 12.6 report the exercise needs is at content/profiles/stream-ordered-allocation/: day55-mempool.nsys-rep with its nvtx_sum and cuda_api_sum CSV exports beside it.

Loop Total (ms) us per iteration vs hoisted
malloc-per-iter 85.588 855.9 2.52x
hoisted 33.977 339.8 1.00x
malloc-async 34.248 342.5 1.01x
Release threshold Pool reserved after sync (bytes)
0 (default) 0
UINT64_MAX 33,554,432

Predictions

  1. malloc-per-iter costs at least 3 times hoisted, and the excess sits in cudaFree, not cudaMalloc: in cuda_api_sum, cudaFree's total time exceeds cudaMalloc's, because the free carries the drain and the malloc only carries the OS bookkeeping. cuda_api_sum aggregates the whole capture, and here that grades the loop: 100 of the capture's 104 calls to each of the two functions come from the malloc-per-iter loop, the other four from setup, teardown and one warm-up. If the two totals come out even, the mechanism section above is half wrong and gets rewritten.
  2. malloc-async stays within 1.25 times hoisted. A hundred pool reuses should cost queue entries, not memory traffic. Over 1.5 and the pool is doing per-iteration work I have not accounted for.
  3. The pool holds nothing at the default threshold and at least 16,779,660 bytes after the threshold is raised, both read after a stream synchronise. That number is one scratch buffer; the pool may legally hold more.
  4. The timeline shapes differ more than the totals. Inside malloc-per-iter the kernel row shows a gap every iteration under a cudaFree API bar; inside malloc-async the API row finishes ahead of the kernel row, the day 51 signature of a host running free of its device.

What the run said

  1. Both parts failed. The loop cost is smaller than predicted: 2.52x hoisted, under the 3x bar. In the published CUDA 12.6 cuda_api_sum, the excess sits in cudaMalloc (104 calls, 81.143 ms total) rather than cudaFree (104 calls, 49.310 ms). On this driver, the allocation call after the free accounts for most of the wait, so the prediction assigned the cost to the wrong call.
  2. Held. malloc-async landed at 1.01x hoisted: a hundred pool reuses cost queue entries, not memory traffic.
  3. Held. The pool reports 0 bytes at the default threshold and 33,554,432 bytes after raising it, both read after a stream synchronise; that is two scratch buffers' worth, which the prediction's "may legally hold more" allowed.
  4. Left to the report. The two timeline shapes (a gap under an API bar per iteration against an API row running ahead of the kernel row) are the shipped capture's to show; open day55-mempool.nsys-rep and compare the malloc-per-iter and malloc-async ranges.

The ratio matters more than the measured milliseconds. The drain cost varies by GPU, but cudaFree inside a loop must wait for queued uses of the block. Stream-ordered freeing avoids that wait.

Run it yourself

Use a CUDA GPU that reports cudaDevAttrMemoryPoolsSupported as 1.

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o mempool mempool.cu
nsys profile -t cuda,nvtx -o day55-mempool ./mempool

The program checks whether cudaMallocAsync is supported and reports an unsupported device instead of failing. If you cannot run the profiler, read /setup/learn-cuda-without-a-gpu and do the exercise from the shipped report.

Exercise

Write the three loops yourself, or take your own pipeline from day 54 and give each iteration a per-iteration scratch buffer, then convert it to cudaMallocAsync. Measure all three with events, profile the malloc and async versions with -t cuda,nvtx, and name the API call that cost the difference.

Time: 30 to 45 minutes. Submit: the three per-iteration times, the malloc-per-iter over malloc-async ratio, and one sentence naming the call the time came from.

Check: a ratio, since absolute times differ card to card. malloc-per-iter over hoisted should come out well above 1, and malloc-async over hoisted close to 1. If your async loop is no faster than your malloc loop, either the pool is releasing between iterations because something inside your loop synchronises, or the allocations are not on the stream the kernels are on; the cuda_api_sum table says which, and the shipped report at content/profiles/stream-ordered-allocation/ gives the same table if you cannot run nsys.

Hint 1

The cost you are removing is not the bytes; the same bytes get allocated a hundred times in both loops. Which call in the slow loop cannot return until the device has finished everything, and what is the device doing while the host waits in it?

Hint 2

Sort cuda_api_sum by total time. Compare the total for cudaFree against the total for cudaMalloc, then against the sum of everything else in the loop. One of those three comparisons is the whole lesson.

Solution

The call is cudaFree. It cannot release a block the queued kernels might still touch, so it waits for the device to drain, once per iteration, and the drain costs more than the allocation bookkeeping ever did. cudaFreeAsync queues the free into the stream instead, after the kernels that use the block, so the pool reclaims it at exactly the right moment with nobody waiting; the next cudaMallocAsync reuses the block without going near the OS.

Queue a resource's lifetime operations in the stream that uses it. A host call that waits for the device serialises the pipeline. A stream-ordered operation keeps that wait in the queue.

Pitfalls

Your async loop is no faster because something inside it synchronises. A cudaStreamSynchronize, a blocking event query or a synchronous cudaMemcpy inside the loop drains the queue, and at the default release threshold every drain also empties the pool back to the OS. Hoist the sync out of the loop, or raise the threshold and keep the sync.

You freed on a different stream than the one using the memory. cudaFreeAsync orders the free against the stream you pass it, not against every stream. Freeing on stream B while stream A's kernels still read the block is a use-after-free that no compiler warns about. Free on the stream that used the memory, or connect the two with an event first, the day 52 pattern.

You timed the async loop with a host clock and no sync. cudaMallocAsync, the launches and cudaFreeAsync all return immediately, so a wall clock around the loop measures submission, not execution. Synchronise the stream before reading the clock. The program times with events on the stream and synchronises on the stop event, which is the same discipline day 9 taught with one kernel.

You expected the first allocation to get faster. The pool starts empty, and its first block still comes from the OS; what gets cheap is reuse. A program that allocates once at startup gains nothing from cudaMallocAsync, which is why day 5 taught cudaMalloc and this page has no reason to unteach it there.

You read the pool holding memory as a leak. After the threshold is raised, nvidia-smi shows the process holding the pool's reserved bytes even though every allocation was freed. That memory is the pool doing its job. cudaMemPoolAttrReservedMemCurrent against cudaMemPoolAttrUsedMemCurrent tells you how much is cache versus how much your code still holds.

Go deeper

Next

Day 56 captures a whole launch sequence as a CUDA graph and replays it for less than the launch overhead day 48 measured, and stream-ordered allocation is what makes memory legal inside one: a graph can own allocation nodes, but never a cudaMalloc. Day 60's frame pipeline then uses today's pool under a real frames-per-second target, where one accidental drain is a dropped frame you can see.