cudaMemcpyAsync in a Stream
Non-blocking transfers that overlap.
cudaMemcpyAsync in a Stream is a free CUDA Academy lesson on CoddyKit — lesson 3 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.
Blocking by Default
Plain cudaMemcpy stops your CPU until the copy finishes. Your host thread just waits, doing nothing useful while bytes move.
The Async Cousin
Meet cudaMemcpyAsync. It queues the copy and returns to your CPU immediately, so the host can keep working while the transfer happens.
It Needs a Stream
The async copy takes an extra argument: a stream. The stream is the ordered queue where the transfer waits its turn to run.
cudaMemcpyAsync(d_data, h_data, bytes, cudaMemcpyHostToDevice, stream);Pinned or Bust
Here is the catch: async copies only run truly non-blocking from pinned host memory. Pass a pageable buffer and CUDA silently falls back to a blocking copy.
Returns, Does Not Finish
When the call returns, the copy has only been scheduled, not completed. The data may still be in flight, so do not read it yet.
Order Within a Stream
Work in one stream runs in order, one item after another. So a copy issued before a kernel in the same stream is guaranteed to finish first.
Wait When You Must
Before touching the results on the CPU, make sure the queue drained. Call cudaStreamSynchronize to block until that stream's work is done.
cudaStreamSynchronize(stream);The Whole Point: Overlap
Async copies let a transfer in one stream run while a kernel in another stream computes. That concurrency is how you hide transfer time.
Beware the Default Stream
If you pass stream 0, the legacy default, it can serialize against everything else. Use your own created streams to actually get overlap.
Keep the Buffer Alive
Since the copy finishes later, the host buffer must stay valid until then. Free it too early and the in-flight transfer reads garbage.
A Typical Pattern
The classic flow is async copy in, launch the kernel, then async copy out, all in one stream. Sync only at the very end.
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream);
kernel<<<grid, block, 0, stream>>>(d_in, d_out);Quick Check
What does cudaMemcpyAsync require to be genuinely non-blocking?
Recap
cudaMemcpyAsync queues a copy in a stream and returns at once, but needs pinned memory to overlap. Sync the stream before reading results. ⚡
Frequently asked questions
Is the “cudaMemcpyAsync in a Stream” lesson free?
Yes — the full text of “cudaMemcpyAsync in a Stream” 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 “cudaMemcpyAsync in a Stream”?
Non-blocking transfers that overlap. 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 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “cudaMemcpyAsync in a Stream” 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
- Why Pageable Memory Is Slow
- Pinned Memory with cudaMallocHost
- cudaMemcpyAsync in a Stream
- The Double-Buffering Pipeline