Overlapping Copy and Compute
Hiding transfers behind kernels.
Overlapping Copy and Compute is a free CUDA Academy lesson on CoddyKit — lesson 4 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
The Big Idea
Real speedups come from doing two things at once: copying one chunk of data while the GPU computes on another. 🔀
Separate Hardware Engines
A GPU has distinct copy engines and compute units. Overlap works because these can run truly in parallel, not just take turns.
Pinned Memory Is Required
Async copies need pinned host memory from cudaMallocHost. Ordinary pageable malloc forces a blocking copy that cannot overlap.
cudaMallocHost(&h_data, bytes);Async Copies
Use cudaMemcpyAsync with a stream so the transfer returns immediately and joins that stream's queue alongside kernels.
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, s);Chunk the Work
Split the array into pieces and give each piece its own stream. While stream 0 computes, stream 1 can already be copying.
Copy, Compute, Copy Back
Each chunk does the same three steps in its stream: upload input, run the kernel, download output, all without blocking the host.
cudaMemcpyAsync(d, h, n, H2D, s);
k<<<g, b, 0, s>>>(d);
cudaMemcpyAsync(h, d, n, D2H, s);How the Overlap Forms
Because chunks live in different streams, the GPU can copy chunk two while still computing chunk one. The timeline fills up. 📊
Hiding the PCIe Cost
Transfers no longer add to total time; they hide behind compute. Ideally your runtime shrinks to roughly the longer of copy or kernel.
Watch the Default Stream
One stray default-stream call mid-loop can serialize everything again. Keep every async op tied to a real, non-default stream.
Verify in a Profiler
Open Nsight Systems and look for copy and compute lanes that overlap in time. Stacked, busy lanes mean your concurrency is real.
Synchronize at the End
After issuing all chunks, call cudaDeviceSynchronize once so every stream finishes before you read the final results on the host.
cudaDeviceSynchronize();Quick Check
What does an async transfer require to overlap with compute?
Recap
By chunking work across streams with pinned memory and async copies, you hide transfers behind compute and keep the GPU truly busy. 🚀
Frequently asked questions
Is the “Overlapping Copy and Compute” lesson free?
Yes — the full text of “Overlapping Copy and Compute” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Overlapping Copy and Compute”?
Hiding transfers behind kernels. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 4 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Overlapping Copy and Compute” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- The Default Stream Trap
- Creating and Using Streams
- Events for Timing and Sync
- Overlapping Copy and Compute