Host threads and the GPU

Profile a loop that fills a buffer, copies it, runs a kernel, and waits. The kernel row shows a short bar followed by a gap as long as the CPU fill. One CPU core generates the data while the GPU stays idle.

A second CPU thread can fill the next buffer while the GPU uses the current one. This lesson builds that producer-consumer pipeline with a condition variable and a callback that returns each buffer to the producer.

One context, with streams for ordering

The runtime creates one primary context per device and, as the programming guide puts it, "This context is shared among all the host threads of the application" (https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html , checked 2026-09-01). Every thread in the process sees the same device memory, modules, and streams. A pointer from cudaMalloc works from any thread.

The current device is per-thread, so each new thread should call cudaSetDevice, especially in a multi-GPU program. This selects a device's shared context rather than creating a private one.

Multiple threads can use CUDA, and runtime calls are thread-safe. Streams order their work. A stream runs operations in enqueue order, no matter which thread enqueued them.

Giving each thread its own stream keeps their ordering independent. The shared context does not provide that separation.

Two threads can enqueue into one stream without a runtime error, but the OS scheduler determines how their calls interleave.

// Thread A                              // Thread B
cudaMemcpyAsync(d_buf, h_a, n, kH2D, s); cudaMemcpyAsync(d_buf, h_b, n,
kernelA<<<g, b, 0, s>>>(d_buf);                          kH2D, s);
                                         kernelB<<<g, b, 0, s>>>(d_buf);
// One legal order: copy A, copy B, kernel A, kernel B.
// kernelA reads thread B's bytes. No error is ever reported.

Diagram: two host threads, one card, three ways. Three bands. Each band has a producer-thread row and a consumer-thread row on top, a copy row and a kernel row below, one shared time axis. Band 1, one thread: fill, copy, kernel repeat end to end on a single host row, the copy and kernel rows mostly empty. Caption "serial: the card waits out every fill." Band 2, two threads, two slots: fill bars for chunk c+1 on the producer row sit directly above the copy and kernel bars for chunk c. Caption "pipelined: the fill hides behind the copy and the kernel." Band 3, zoomed on one slot: copy bar, kernel bar, then a thin callback tick, then an arrow up to the producer row where the next fill of that slot begins. Caption "the callback returns the slot; nothing polls." Alt text: "Serially, one thread leaves the GPU idle during every fill. With two threads and two slots the next fill runs above the current copy and kernel, and a callback tick hands each slot back."

Let the callback report completion

The CPU can query GPU progress with cudaDeviceSynchronize, a stream sync, an event query in a loop. In a producer-consumer design, these checks can stop the consumer while chunk c runs. The consumer should enqueue chunk c+2 instead.

cudaLaunchHostFunc inverts the direction. You enqueue a plain host function into the stream, and the docs give the guarantee this design stands on: "The function will be called after currently enqueued work and will block work added after it" (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EXECUTION.html , checked 2026-09-01).

Enqueue the copy, kernel, and callback in that order. When the callback runs, the stream has finished using the buffer.

The same page states the two prohibitions. "The host function must not make any CUDA API calls", and it "must not perform any synchronization that may depend on outstanding CUDA work not mandated to run earlier". Both exist because later stream work waits for the callback on an internal runtime thread.

Blocking there blocks the stream. Waiting for GPU work queued after the callback causes a deadlock.

A correct callback takes a lock, changes a flag, notifies a condition variable, and returns.

Two threads, two slots, one stream

Full program in code/day59-host-threads/host_threads.cu. Three invariants run through it.

Each resource has exactly one owner at a time. Two pinned staging slots cycle through three states: the producer owns a slot while it is free, the consumer while it is ready, the stream while it is in flight. Every transition happens under one mutex. The producer thread makes no CUDA calls in its loop beyond the initial cudaSetDevice; the consumer thread owns the stream outright.

The GPU frees the slots. The consumer enqueues the async copy, the kernel, then the callback that marks the slot free. It never waits on the GPU inside the loop. The producer's condition-variable wait is the only backpressure in the program.

Threading must change nothing. The program runs the same 24 chunks serially and pipelined and requires the two outputs to be bit-identical, so "the second thread changed the answer" is a failing exit code, not a shrug.

The producer:

static void producerLoop(Handoff* handoff, float** h_staging) {
    CUDA_CHECK(cudaSetDevice(0));
    for (int chunk = 0; chunk < kChunks; ++chunk) {
        const int slot = chunk % kSlots;
        {
            std::unique_lock<std::mutex> lock(handoff->m);
            handoff->cv.wait(
                lock, [&] { return handoff->state[slot] == SlotState::kFree; });
        }
        nvtxRangePushA("fill");
        fillChunk(h_staging[slot], chunk);
        nvtxRangePop();
        {
            std::lock_guard<std::mutex> lock(handoff->m);
            handoff->state[slot] = SlotState::kReady;
        }
        handoff->cv.notify_all();
    }
}

The consumer, which is the only thread that touches the stream:

    for (int chunk = 0; chunk < kChunks; ++chunk) {
        const int slot = chunk % kSlots;
        {
            std::unique_lock<std::mutex> lock(handoff.m);
            handoff.cv.wait(
                lock, [&] { return handoff.state[slot] == SlotState::kReady; });
            handoff.state[slot] = SlotState::kInFlight;
        }
        nvtxRangePushA("enqueue");
        enqueueChunk(stream, d_in[slot], h_staging[slot], d_outPipe, chunk);
        CUDA_CHECK(cudaLaunchHostFunc(stream, slotFree, &done[slot]));
        nvtxRangePop();
    }

And the callback, which follows both prohibitions to the letter:

// Runs on CUDA's internal callback thread when the stream reaches it. It may
// not call any CUDA API, and it must return fast: work queued after it in
// this stream waits for it, and one callback thread can serve every stream
// in the process. Lock, flip the state, notify, get out. notify_all, not
// notify_one: the producer and the consumer wait on the same condition
// variable, and a single wakeup delivered to the wrong one is a deadlock.
static void CUDART_CB slotFree(void* userData) {
    SlotDone* done = static_cast<SlotDone*>(userData);
    {
        std::lock_guard<std::mutex> lock(done->handoff->m);
        done->handoff->state[done->slot] = SlotState::kFree;
    }
    done->handoff->cv.notify_all();
}

The timing uses CUDA events even for the CPU-only fill leg, by recording two events on an idle stream around the host work. Both events complete as soon as the GPU reaches them, so the elapsed time is host wall time on the GPU's clock, give or take a launch overhead measured in microseconds. It keeps all four legs on one clock; it is not how you would time a CPU function in a CPU program.

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 gates reproduced. Serial improved from 97.333 to 95.053 ms while the pipeline moved from 78.252 to 81.551 ms; the 1.17x speedup still puts the pipeline within its predicted 15 percent of the longer leg.

Both transcripts and the fresh CUDA 13 report are listed in front matter. The published CUDA 12.6 timeline is at content/profiles/host-threads-gpu/: day59-pipeline.nsys-rep with its day59-pipeline_nvtx_sum.csv and day59-pipeline_cuda_api_sum.csv exports.

All 24 chunks matched the CPU reference. The serial and pipelined outputs are bit-identical.

Leg	Total (ms)	Per chunk (ms)
fill only (CPU)	71.224	2.968
copy+kernel only	21.647	0.902
serial, one thread	95.053	3.961
pipeline, two threads	81.551	3.398

Expectations, ahead of the evidence

1. The serial leg stays within 10 percent of fill-only plus copy-plus-kernel-only. It is the same work laid end to end. If it is more than 10 percent above the sum, the per-chunk stream synchronize costs more than I think it does.
2. The pipeline leg stays within 15 percent of the longer single leg, whichever that is on this host, because two slots are enough to overlap the shorter leg with it. This is the claim being tested, and the program prints the lower bound so the run can check it directly.
3. On the timeline, the producer and consumer are two separate thread rows, NVTX fill bars for chunk c+1 sitting over the H2D copy for chunk c, and the cudaLaunchHostFunc entries in cuda_api_sum are microsecond noise, far under one percent of the run.

Claim 2 is the falsifiable one. If the pipeline comes in well above the longer leg, the overlap did not happen, and the suspects are, in order: the callback delaying the next copy (it blocks work queued after it), the condition-variable wakeup latency, or two slots being too few on this machine.

What the run said

1. Held. Serial measured 95.053 ms and remained within 10 percent of the two measured legs added together, plus the small per-chunk synchronisation cost.
2. Held, and it is the page's claim. The longer leg is the CPU fill at 71.224 ms; the pipeline landed at 81.551, 14.5 percent above it, inside the 15 percent bound. Speedup over serial: 1.17x, which is all a 2.968-to-0.902 leg ratio has to give.
3. Held in the published CUDA 12.6 capture. The capture shows the two thread rows with fill bars over the previous chunk's H2D copy, and the 24 cudaLaunchHostFunc calls total 533.8 microseconds in cuda_api_sum, 0.322 percent of the capture's API time.

The speedup depends on the ratio of fill cost to copy-plus-kernel cost, and both vary by machine. The pipeline approaches the time of its slower leg.

Run it yourself

Use a CUDA GPU and at least two CPU cores so the producer and consumer can run at the same time.

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

If you cannot run the profiler, read /setup/learn-cuda-without-a-gpu and do the exercise against the shipped report in content/profiles/host-threads-gpu/.

Exercise

Add a second consumer thread with its own stream and its own pair of staging slots and device buffers, splitting the chunks even and odd, with the single producer feeding both. Keep every gate green.

Time: 35 to 50 minutes. Submit: your host_threads.cu, its printed table, and one sentence saying which leg your second consumer shortened, if any.

Check: the program performs both checks. Gate 1 compares every chunk against the CPU reference and prints the chunk, the element index and both values for the first mismatch; gate 2 requires your pipelined output to stay bit-identical to the serial run, and prints the first differing index if threading changed the answer.

Both exit nonzero on failure. The program reports the speedup either way; the correctness gates decide whether the exercise passes.

Hint 1

Start by listing what the two consumers would actually share. The chunks partition cleanly by parity. The stream should not be shared, this page spent a section on why.

What is left, and who owns it?

Hint 2

Give each consumer its own slots, its own stream and its own SlotDone records, so the only shared object left is the producer's handoff state. Then look at your fill-only number: can one producer feed two consumers, or is the producer now the longer leg?

Solution

Duplicate the per-consumer state: slots, device input buffers, stream, callback records. The producer alternates which consumer's slots it fills; the handoff grows to four slot states under the same mutex. Each consumer loop is unchanged except for its stream and its chunk stride.

The outputs stay bit-identical because every chunk still runs the same kernel on the same bytes into its own slice of the output.

The result depends on which leg is slower. If the fill leg was slower, a second consumer adds no throughput because the producer was the limit and now feeds two queues at the same total rate. If the copy-plus-kernel leg was longer, the second stream can overlap one stream's copy with the other's kernel, and the lower bound falls toward whichever engine saturates first. The generalisation: adding a thread only helps the leg that was longest, and streams are what keep the added thread from touching anyone else's ordering.

Pitfalls

Two threads share one stream and the results scramble, but only sometimes. Every runtime call is thread-safe, so nothing crashes; the enqueue order is whatever interleaving the scheduler produced, so on some runs a kernel reads the other thread's bytes. Fix: one stream per thread, or one thread that owns all enqueueing, as this page's consumer does.

Your callback calls a CUDA function and sometimes gets cudaErrorNotPermitted. The docs say the host function "must not make any CUDA API calls"; the error is not even guaranteed, so it can appear to work. Move the CUDA call to the thread the callback wakes.

Your callback waits, and the whole program stops. The function blocks work added to the stream after it, and one runtime thread can serve every stream in the process, so a callback that waits for later GPU work deadlocks and a merely slow one stalls streams it was never enqueued on. Flip a flag, notify, return.

The producer refills a slot while the copy still reads it. An async copy from pinned memory reads the buffer when the stream reaches it, not when you enqueued it. Handing the slot back on enqueue corrupts data with no error; the callback returns it only after the stream is done.

The overlap vanished when you switched to pageable staging. A cudaMemcpyAsync from pageable memory goes through the runtime's own staging and loses most of its asynchrony. Day 53 measured it; keep the slots in cudaMallocHost memory.

Your new thread ran on the wrong GPU. The current device is per thread, and a fresh thread starts on device 0 regardless of what the spawning thread had set. Call cudaSetDevice at the top of every thread that touches CUDA. This may not change a one-GPU program, but it is required in a multi-GPU program.

Go deeper

• CUDA Runtime API, cudaLaunchHostFunc, for both prohibitions and the ordering guarantee: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EXECUTION.html (checked 2026-09-01)
• CUDA Programming Guide 2.1, "Introduction to CUDA C++", the initialization section, for the context shared among all host threads: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html (checked 2026-09-01)
• cuda-samples, Samples/0_Introduction/simpleCallback: https://github.com/NVIDIA/cuda-samples/tree/master/Samples/0_Introduction/simpleCallback (checked 2026-09-01)
• Programming Massively Parallel Processors, 4th edition, chapter 20, on overlapping computation and communication

Next

Day 60 is the module's capstone: a frame pipeline that has to hold a frames-per-second target, which takes day 51's overlap, day 54's double buffering and this page's producer-consumer handoff and makes one nsys timeline prove all three at once. After that the module hands you to the debugging tools, which is where the races this page carefully avoided get hunted for a living.