LESSON READING // 02
TASK CONCURRENCY
Use multiple streams to express independence between chunks, giving copies, kernels, and returns a chance to overlap while preserving each chunk's order.
Write down answers as you go; the estimate includes code tracing and practice.
PLAY THE INTERACTIVE LESSON→- 01ORIENT + OBJECTIVES05 MIN
- 02MENTAL MODEL08 MIN
- 03DEEP DIVES15 MIN
- 04CODE TRACE12 MIN
- 05HISTORY + ECOSYSTEM08 MIN
- 06PRACTICE + REVIEW12 MIN
Treat the GPU program as a workflow and use timelines to find gaps among CPU submission, movement, and computation.
The game teaches stream dependencies; the lab turns them into a real trace that separates an asynchronous API from actual overlap.
- Lessons 00–01
- H2D / D2H basics
- In-stream ordering
- Nsight Systems
- CUDA Streams / Events
- Pinned Host Memory
ORIENT + OBJECTIVES
REMEMBER THIS FIRST
Concurrency comes from independent work and correct dependencies; an asynchronous API only communicates that possibility.
BY THE END, YOU SHOULD BE ABLE TO
- 01
Draw dependency timelines for one stream and several streams.
- 02
Identify operations that may overlap and those needing stream order or events.
- 03
Explain how pinned host memory, copy engines, and default-stream semantics affect overlap.
- 04
Move synchronization to actual result-consumption boundaries.
BUILD THE MENTAL MODEL
- 01QUEUE
Commands in one stream execute in submission order, forming a natural dependency chain.
- 02SPLIT
A stream per independent chunk lets H2D, kernel, and D2H work from different chunks interleave.
- 03JOIN
Submit all independent work first, then synchronize only before results are actually consumed.
CONCEPT DEEP DIVES
A STREAM MEANS ORDER
It describes an ordered command flow. Different streams may execute concurrently; they are not promised to do so.
DEPENDENCIES FOLLOW DATA
Keeping each chunk's copy, kernel, and return in one stream preserves local ordering without serializing other chunks.
OVERLAP HAS PREREQUISITES
Real overlap can require device capability, pinned host memory, sufficient resources, and freedom from implicit synchronization.
STREAMS DEFINE ORDER BEFORE THEY OFFER CONCURRENCY
Commands in one stream retain submission order, so H2D → kernel → D2H naturally describes one chunk. Placing three chunks in that same stream is correct but serializes independent chunks. A stream per chunk preserves each local chain while exposing concurrency across chains.
Different streams may execute concurrently; they are not guaranteed to. Resource pressure, implicit synchronization, default-stream behavior, or real dependencies can serialize them. Streams declare independence and the runtime realizes what current conditions permit.
PAUSE AND REASONWhy can one H2D stream, one kernel stream, and one D2H stream break correctness?＋
REFERENCE ANSWERA chunk's kernel may run before its H2D, or D2H before its kernel. Events would be needed to rebuild data dependencies.
TRANSFER OVERLAP DEPENDS ON BOTH HOST AND DEVICE
The host behavior and DMA path of cudaMemcpyAsync depend on memory type and direction. Pinned host memory provides stable direct DMA access and is commonly required for H2D/D2H overlap. Pageable memory can trigger staging or synchronization.
The device also needs available copy engines, and kernel resource use must leave concurrency possible. Verify with a device timeline rather than an API name. Too many tiny chunks can lose to launch and scheduling overhead.
PAUSE AND REASONDoes a quickly returning API prove the copy overlapped a kernel?＋
REFERENCE ANSWERNo. It only shows the host did not wait for the full operation. A device timeline must show actual DMA/kernel overlap.
EVENTS EXPRESS LOCAL DEPENDENCIES
cudaDeviceSynchronize waits broadly for prior device work. If stream B depends only on one operation in stream A, record an event in A and make B wait on that event. Unrelated streams and the host can keep moving.
CUDA Graphs can reduce repeated submission overhead for a stable dependency graph, but they do not invent independence. Design correct dependencies first, then choose events or graphs for precision and overhead.
PAUSE AND REASONWhy is an event better when B depends only on A's H2D?＋
REFERENCE ANSWERIt creates only the required edge, while a device-wide synchronization waits for unrelated work too.
PUT IT BACK INTO CODE
01for (int s = 0; s < 3; ++s) {02 cudaMemcpyAsync(d[s], h[s], bytes, H2D, stream[s]);03 kernel<<<grid, block, 0, stream[s]>>>(d[s]);04 cudaMemcpyAsync(out[s], d[s], bytes, D2H, stream[s]);05}06cudaDeviceSynchronize(); // before consuming resultsTRACE THE CODE
- 01LOOPONE ITERATION OWNS ONE CHUNK
The same s selects host buffer, device buffer, and stream, preventing accidental storage sharing.
- 02H2DSUBMIT INPUT ASYNCHRONOUSLY
The host buffer still needs a valid lifetime until transfer completion.
- 03KERNELSTREAM ORDER PRESERVES LOCAL DEPENDENCY
The kernel follows its chunk's H2D without an intervening host wait.
- 04JOINWAIT BEFORE RESULTS ARE CONSUMED
Submitting all chunks first preserves the widest copy/compute/return overlap window.
KNOWLEDGE ATLAS
OPTIONAL REFERENCE · REVISIT AS NEEDED · OUTSIDE THE CORE 60 MINUTES
CORE VOCABULARY
- Stream
- An ordered device command sequence; different streams expose possible independence but do not promise concurrency.
- Event
- A recorded stream completion point used for timing or a precise cross-stream dependency.
- Pinned Memory
- Host memory that is not paged out, enabling a direct DMA path and supporting asynchronous transfer.
- Copy Engine
- Hardware that performs movement; count and direction support constrain overlap with compute or another copy.
- CUDA Graph
- A reusable graph of operations and dependencies that reduces repeated CPU submission overhead.
WORKED EXAMPLE
- 01SERIAL BASELINE
3 × (2 + 5 + 2) = 27 ms when every chunk waits for the previous one.
- 02FIND CADENCE
The 5 ms kernel is the slowest stage, so steady-state cadence is one chunk per 5 ms.
- 03ADD FILL + DRAIN
The first chunk needs 9 ms; two more cadences add 10 ms: 9 + 2 × 5 = 19 ms.
DIAGNOSTIC PLAYBOOK
Multiple streams still form one line
- INSPECT FIRST
- Default stream, host-memory type, synchronization, and device capability
- EVIDENCE
- An aligned Nsight Systems view of host API, copies, and kernels
More streams make it slower
- INSPECT FIRST
- Tiny chunks, launch cost, contention, and bandwidth saturation
- EVIDENCE
- A stream-count sweep with CPU submission and GPU utilization
A consumer sometimes reads stale data
- INSPECT FIRST
- Cross-stream producer edges, buffer lifetime, and host reuse
- EVIDENCE
- An event dependency graph and explicit buffer ownership
HARDWARE → ECOSYSTEM
- 01ASYNC KERNELS + STREAMS
The host can continue after dispatch.
Programs can expose host, copy, and GPU overlap. - 02KEPLER HYPER-Q
More hardware work connections reduced false serialization.
Software-expressed concurrency reached hardware more often. - 03EVENTS + CUDA GRAPHS
Broad waits became precise, reusable dependency graphs.
The focus moved from more queues to better dependencies and lower launch cost.
HARDWARE + ECOSYSTEM COORDINATE
CUDA streams predate Kepler. Hyper-Q mattered because more hardware work connections reduced false serialization between independent streams, giving software-expressed concurrency a better route into hardware.
PRACTICE + REVIEW
Why avoid cudaDeviceSynchronize immediately after every cudaMemcpyAsync?
- A. It corrupts the copy
- B. It blocks the host and closes the overlap window too early
- C. It clears device memory
REVEAL ANSWER＋
The placement of synchronization defines the concurrency boundary. Waiting too early serializes otherwise independent commands.
Three chunks share the default stream and each runs H2D→kernel→D2H. How many ordered chains result?
HINT＋
All commands are submitted to one stream.
REFERENCE ANSWER＋
One chain of nine operations. Chunk independence is not exposed.
The host calls cudaStreamSynchronize after every chunk submission. Why is it correct but slow?
HINT＋
When can the next chunk be submitted?
REFERENCE ANSWER＋
The host waits for a full chunk before submitting the next, closing cross-chunk overlap. Submit first, then wait at consumption.
A kernel in B consumes data produced in A. Avoid a device-wide wait.
HINT＋
Record a local completion signal.
REFERENCE ANSWER＋
Record an event in A after production and call cudaStreamWaitEvent in B before submitting the consumer.
Multiple streams and async copies show no overlap. Name three checks.
HINT＋
Host memory, device capability, synchronization.
REFERENCE ANSWER＋
Check pinned host buffers, copy-engine/concurrency support, and explicit or implicit/default-stream synchronization; also inspect tiny task sizes and resource saturation.
OPTIONAL HANDS-ON LAB
OUTSIDE THE CORE 60 MINUTES · REQUIRES A SUITABLE CUDA / GPU ENVIRONMENT
LAB GOAL
BUILD A MEASURABLE COPY–COMPUTE PIPELINE
Transform nine default-stream tasks into a stream-per-chunk pipeline and verify overlap and dependencies.PROCEDURE
- 01
Capture a pageable-memory, default-stream baseline.
- 02
Use pinned memory and three non-blocking streams with H2D→kernel→D2H per chunk.
- 03
Move synchronization after submission and add one event dependency.
- 04
Annotate host API, copy engines, and kernels in Nsight Systems.
Submit before/after timelines, synchronization reasoning, and end-to-end time with hardware/task-size limits.
Capture the stable sequence as a CUDA Graph and compare CPU submission overhead.