1. Why CUDA Streams Exist

Imagine you have:

textCopy
1Kernel A
2Kernel B
3Kernel C

A simple program might execute:

textCopy
1Kernel A
2   ↓
3Kernel B
4   ↓
5Kernel C

Sequentially.

But modern GPUs can sometimes overlap independent operations:

textCopy
1Time →
2────────────────────────────────────
3
4Stream 0:
5Kernel A ███████████
6                 Kernel C █████████
7
8Stream 1:
9       Copy B ███████

This can improve:

textCopy
1GPU utilization
2throughput
3pipeline efficiency

2. What Is a CUDA Stream?

A CUDA stream is essentially an ordered sequence of CUDA operations.

For example:

textCopy
1Stream 0
2
3Kernel A
4   ↓
5Kernel B
6   ↓
7Kernel C

Operations submitted to the same stream maintain ordering.

Conceptually:

textCopy
1A → B → C

You can create multiple streams:

textCopy
1Stream 0:
2A → B
3
4Stream 1:
5C → D

Operations in different streams may execute concurrently when the GPU, dependencies, and resource availability allow it.

3. Default Stream

If you launch:

cppCopy
1kernel<<<blocks, threads>>>();

without specifying a stream, the kernel is submitted to the default stream.

For example:

cppCopy
1kernel<<<blocks, threads>>>();

is conceptually:

textCopy
1default stream
2      ↓
3kernel

You can explicitly specify a stream:

cppCopy
1kernel<<<
2    blocks,
3    threads,
4    0,
5    stream
6>>>();

The fourth launch configuration parameter is the CUDA stream.

4. Creating a Stream

Use:

cppCopy
1cudaStream_t stream;
2
3cudaStreamCreate(&stream);

Then launch:

cppCopy
1kernel<<<
2    blocks,
3    threads,
4    0,
5    stream
6>>>();

Finally:

cppCopy
1cudaStreamDestroy(stream);

Basic lifecycle:

textCopy
1Create
2  ↓
3Use
4  ↓
5Synchronize if needed
6  ↓
7Destroy

5. Complete Working Example

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void kernel()
5{
6    int i =
7        blockIdx.x * blockDim.x
8        + threadIdx.x;
9
10    printf(
11        "Thread %d running\n",
12        i
13    );
14}
15
16int main()
17{
18    cudaStream_t stream;
19
20    cudaStreamCreate(&stream);
21
22    int threads = 32;
23    int blocks = 1;
24
25    kernel<<<
26        blocks,
27        threads,
28        0,
29        stream
30    >>>();
31
32    cudaStreamSynchronize(stream);
33
34    cudaStreamDestroy(stream);
35
36    return 0;
37}

Compile:

bashCopy
1nvcc stream.cu -o stream

Run:

bashCopy
1./stream

6. Why Synchronization Is Necessary

CUDA operations are often asynchronous with respect to the CPU.

For example:

cppCopy
1kernel<<<blocks, threads, 0, stream>>>();
2
3printf("CPU continues\n");

The CPU may continue before the GPU kernel has finished.

If you need to wait for the stream:

cppCopy
1cudaStreamSynchronize(stream);

Conceptually:

textCopy
1CPU
2 │
3 ├── Launch kernel
4 │
5 ├── Continue doing work
6 │
7 └── Synchronize
8          ↓
9       wait for GPU

7. Synchronous vs Asynchronous

Synchronous

textCopy
1CPU
2 │
3 │ launch
4 ↓
5GPU Kernel
6 │
7 │ finish
8 ↓
9CPU continues

Asynchronous

textCopy
1CPU
2 │
3 ├── launch
4 │
5 ├── continue
6 │
7 ├── do other work
8 │
9 └── synchronize later

This is useful for pipelines.

8. cudaDeviceSynchronize()

This waits for GPU work on the device to complete.

cppCopy
1kernel<<<blocks, threads>>>();
2
3cudaDeviceSynchronize();

Conceptually:

textCopy
1All relevant GPU work
2        ↓
3     complete
4        ↓
5CPU continues

But don't use it unnecessarily.

Excessive synchronization can destroy concurrency.

9. cudaStreamSynchronize()

This is more targeted.

cppCopy
1cudaStreamSynchronize(stream);

It waits for operations in that stream to finish.

Compare:

textCopy
1cudaDeviceSynchronize()
2        ↓
3wait broadly for device work

versus:

textCopy
1cudaStreamSynchronize(stream)
2        ↓
3wait for this stream

For performance-oriented code, narrower synchronization is often preferable when it is sufficient.

10. Multiple Streams

Now create two:

cppCopy
1cudaStream_t stream0;
2cudaStream_t stream1;
3
4cudaStreamCreate(&stream0);
5cudaStreamCreate(&stream1);

Then:

cppCopy
1kernelA<<<
2    blocks,
3    threads,
4    0,
5    stream0
6>>>();
7
8kernelB<<<
9    blocks,
10    threads,
11    0,
12    stream1
13>>>();

Conceptually:

textCopy
1Stream 0:
2Kernel A ███████████
3
4Stream 1:
5Kernel B ███████████

They may overlap, depending on hardware and whether the kernels have enough resources and no dependency prevents concurrency.

11. Important: Multiple Streams Do NOT Guarantee Parallel Execution

This is extremely important.

Writing:

cppCopy
1stream0
2stream1
3stream2
4stream3

doesn't automatically mean:

textCopy
14 operations execute simultaneously

The GPU considers:

textCopy
1dependencies
2resource availability
3kernel characteristics
4memory operations
5hardware concurrency
6stream semantics

For example, if Kernel A consumes most SM resources:

textCopy
1Kernel A
2████████████████████████

there may be little capacity for Kernel B.

Therefore:

textCopy
1Multiple streams
2        ≠
3Guaranteed parallelism

12. Stream Ordering

Within one stream:

cppCopy
1kernelA<<<..., stream>>>();
2kernelB<<<..., stream>>>();

the operations are ordered:

textCopy
1A
2 ↓
3B

The GPU won't arbitrarily execute B before A within that stream.

Across independent streams:

textCopy
1Stream 0: A → B
2
3Stream 1: C → D

there can be concurrency:

textCopy
1A █████████
2C ██████
3       D ███████
4          B █████████

subject to dependencies/resources.

13. Asynchronous Memory Copy

One of the most important stream use cases is:

cppCopy
1cudaMemcpyAsync()

Example:

cppCopy
1cudaMemcpyAsync(
2    device,
3    host,
4    bytes,
5    cudaMemcpyHostToDevice,
6    stream
7);

Correct argument order is:

cppCopy
1cudaMemcpyAsync(
2    dst,
3    src,
4    count,
5    kind,
6    stream
7);

For host → device:

textCopy
1Host
2 ↓
3GPU

use:

cppCopy
1cudaMemcpyHostToDevice

For device → host:

textCopy
1GPU
2 ↓
3Host

use:

cppCopy
1cudaMemcpyDeviceToHost

14. Why Pinned Host Memory Matters

For efficient asynchronous host↔device transfers, host memory generally needs to be page-locked/pinned.

Allocate it using:

cppCopy
1float* h_data;
2
3cudaMallocHost(
4    &h_data,
5    bytes
6);

Then:

cppCopy
1cudaMemcpyAsync(
2    d_data,
3    h_data,
4    bytes,
5    cudaMemcpyHostToDevice,
6    stream
7);

This enables the CUDA runtime/driver to perform asynchronous transfer behavior more effectively.

15. Complete Async Copy Example

cppCopy
1#include <stdio.h>
2#include <cuda_runtime.h>
3
4__global__ void scale(
5    float* data,
6    int N
7)
8{
9    int i =
10        blockIdx.x * blockDim.x
11        + threadIdx.x;
12
13    if (i < N)
14    {
15        data[i] *= 2.0f;
16    }
17}
18
19int main()
20{
21    int N = 1 << 20;
22
23    size_t bytes =
24        N * sizeof(float);
25
26    float* h_data;
27    float* d_data;
28
29    cudaMallocHost(
30        &h_data,
31        bytes
32    );
33
34    cudaMalloc(
35        &d_data,
36        bytes
37    );
38
39    for (int i = 0; i < N; i++)
40    {
41        h_data[i] = 1.0f;
42    }
43
44    cudaStream_t stream;
45
46    cudaStreamCreate(&stream);
47
48    cudaMemcpyAsync(
49        d_data,
50        h_data,
51        bytes,
52        cudaMemcpyHostToDevice,
53        stream
54    );
55
56    int threads = 256;
57
58    int blocks =
59        (N + threads - 1)
60        / threads;
61
62    scale<<<
63        blocks,
64        threads,
65        0,
66        stream
67    >>>(
68        d_data,
69        N
70    );
71
72    cudaMemcpyAsync(
73        h_data,
74        d_data,
75        bytes,
76        cudaMemcpyDeviceToHost,
77        stream
78    );
79
80    cudaStreamSynchronize(stream);
81
82    printf(
83        "Result: %f\n",
84        h_data[0]
85    );
86
87    cudaStreamDestroy(stream);
88
89    cudaFree(d_data);
90
91    cudaFreeHost(h_data);
92
93    return 0;
94}

Compile:

bashCopy
1nvcc async.cu -o async

Run:

bashCopy
1./async

The sequence is:

textCopy
1Host
2 ↓
3Async H2D
4 ↓
5Kernel
6 ↓
7Async D2H
8 ↓
9Synchronize

Because these operations are in the same stream, they remain ordered.

16. The Real Power: Overlap Copy + Compute

Suppose you have a large dataset.

Instead of:

textCopy
1Copy chunk 1
2Compute chunk 1
3Copy chunk 2
4Compute chunk 2
5Copy chunk 3
6Compute chunk 3

you can potentially pipeline:

textCopy
1Time →
2
3Copy 1   █████
4Compute 1      ███████
5
6Copy 2         █████
7Compute 2             ███████
8
9Copy 3                █████
10Compute 3                    ███████

Now copying and computation can overlap.

This is a major GPU throughput optimization.

17. Double Buffering

A classic technique is double buffering.

Use:

textCopy
1Buffer A
2Buffer B

Pipeline:

textCopy
1Time →
2
3Stream/Engine:
4
5Copy A     █████
6Compute A       ███████
7
8Copy B           █████
9Compute B              ███████

While the GPU computes on one buffer, another chunk can be transferred.

Conceptually:

textCopy
1Host
2 │
3 ├── Chunk A ──→ GPU Buffer A
4 │
5 └── Chunk B ──→ GPU Buffer B
6
7GPU:
8Buffer A → Compute
9Buffer B → Compute

This is common in high-throughput workloads.

18. CUDA Events

Events are useful for:

textCopy
1timing
2synchronization
3dependencies

Create:

cppCopy
1cudaEvent_t start;
2cudaEvent_t stop;
3
4cudaEventCreate(&start);
5cudaEventCreate(&stop);

Record:

cppCopy
1cudaEventRecord(start, stream);
2
3kernel<<<
4    blocks,
5    threads,
6    0,
7    stream
8>>>();
9
10cudaEventRecord(stop, stream);

Then:

cppCopy
1cudaEventSynchronize(stop);

Measure:

cppCopy
1float ms;
2
3cudaEventElapsedTime(
4    &ms,
5    start,
6    stop
7);

This is preferable to ordinary CPU wall-clock timing for measuring GPU work.

19. Stream Dependencies With Events

Suppose:

textCopy
1Stream 0:
2Kernel A

must finish before:

textCopy
1Stream 1:
2Kernel B

Instead of synchronizing the entire device, use an event.

Record:

cppCopy
1cudaEventRecord(
2    event,
3    stream0
4);

Then:

cppCopy
1cudaStreamWaitEvent(
2    stream1,
3    event,
4    0
5);

Now:

textCopy
1Stream 0
2Kernel A
3   ↓
4Event
5   │
6   └──────────────┐
7                  ↓
8Stream 1       Kernel B

This is much more flexible than global synchronization.

20. cudaStreamWaitEvent

Example:

cppCopy
1kernelA<<<
2    blocks,
3    threads,
4    0,
5    stream0
6>>>();
7
8cudaEventRecord(
9    event,
10    stream0
11);
12
13cudaStreamWaitEvent(
14    stream1,
15    event,
16    0
17);
18
19kernelB<<<
20    blocks,
21    threads,
22    0,
23    stream1
24>>>();

Execution dependency:

textCopy
1A
2 ↓
3Event
4 ↓
5B

But unrelated work in other streams may continue.

21. Avoid cudaDeviceSynchronize() Everywhere

This pattern is often bad for performance:

cppCopy
1kernelA<<<...>>>();
2
3cudaDeviceSynchronize();
4
5kernelB<<<...>>>();
6
7cudaDeviceSynchronize();
8
9kernelC<<<...>>>();
10
11cudaDeviceSynchronize();

You're forcing the CPU/GPU workflow into:

textCopy
1A → wait
2B → wait
3C → wait

Instead, if dependencies allow:

cppCopy
1kernelA<<<...>>>(...);
2kernelB<<<...>>>(...);
3kernelC<<<...>>>(...);

or use streams/events appropriately.

22. Asynchronous Does Not Mean "No Waiting"

This is an important distinction.

Asynchronous programming means:

textCopy
1CPU does not necessarily wait immediately

It does not mean:

textCopy
1operations have no dependencies

For example:

textCopy
1Copy
2 ↓
3Kernel

The kernel must not consume data before the copy is complete.

CUDA handles ordering when operations are correctly placed in a stream or explicit event dependency.

23. Streams in AI Inference

Imagine an inference pipeline:

textCopy
1Request 1
2Request 2
3Request 3
4Request 4

You might use:

textCopy
1Stream 0 → Request 1
2Stream 1 → Request 2
3Stream 2 → Request 3
4Stream 3 → Request 4

Potentially:

textCopy
1GPU
2 ├── Stream 0
3 ├── Stream 1
4 ├── Stream 2
5 └── Stream 3

But again, actual concurrency depends on whether the kernels can coexist.

If each request launches huge GEMMs that already saturate the GPU, adding streams may provide little benefit.

24. Streams in LLM Systems

For LLM serving, you may have:

textCopy
1Request A
2Request B
3Request C

Operations can involve:

textCopy
1H2D
2 ↓
3Embedding
4 ↓
5Attention
6 ↓
7MLP
8 ↓
9Sampling
10 ↓
11D2H

A sophisticated runtime can schedule work to maximize:

textCopy
1GPU utilization
2throughput
3latency

Streams are one component of that execution model.

25. Streams + Kernel Fusion

You now have two different optimization levels:

Inside a kernel

textCopy
1coalescing
2shared memory
3registers
4warps
5occupancy
6fusion

Between kernels

textCopy
1streams
2events
3async copies
4overlap
5pipeline

Think:

textCopy
1                    GPU Optimization
2                           │
3             ┌─────────────┴─────────────┐
4             ↓                           ↓
5       Kernel-level                 Execution-level
6             │                           │
7       ┌─────┼─────┐              ┌──────┼──────┐
8       ↓     ↓     ↓              ↓      ↓      ↓
9   Memory  Warp  Register      Streams  Events  Async Copy

This distinction is very useful.

26. CUDA Graphs

Once you understand streams, the next advanced concept you'll encounter is CUDA Graphs.

Instead of repeatedly launching:

textCopy
1Kernel A
2Kernel B
3Kernel C
4Kernel D

the execution graph can be captured and replayed.

Conceptually:

textCopy
1CPU
2 ↓
3Graph launch
4 ↓
5A → B → C → D

This can reduce CPU-side launch overhead for repeated execution patterns.

For AI inference/training workloads with repeated execution graphs, CUDA Graphs can be very useful.

You don't need to master CUDA Graphs before Tensor Cores, but keep it in your advanced CUDA toolbox.

27. Streams vs CUDA Graphs

Think:

Streams

textCopy
1How operations are ordered and scheduled.

CUDA Graphs

textCopy
1How a repeated execution workflow can be captured
2and launched efficiently.

They solve related but different problems.

28. Common Mistakes

Mistake 1

Creating many streams:

textCopy
1100 streams

and assuming:

textCopy
1100× performance

False.

Mistake 2

Synchronizing after every operation:

cppCopy
1cudaDeviceSynchronize();

This destroys potential concurrency.

Mistake 3

Using pageable host memory with cudaMemcpyAsync() and assuming it always provides true asynchronous host↔device behavior.

For reliable overlap, use pinned host memory:

cppCopy
1cudaMallocHost()

or another appropriate page-locked allocation mechanism.

Mistake 4

Assuming streams can make a fully saturated kernel faster.

If:

textCopy
1Kernel A
2████████████████████

already consumes essentially all useful GPU resources, another stream cannot magically create more compute capacity.

29. How to Think About Streams

Use this mental model:

textCopy
1Same Stream
2───────────
3A → B → C

Ordered.

Different streams:

textCopy
1Stream 0: A → B
2
3Stream 1: C → D

Potentially concurrent.

Dependencies:

textCopy
1A
2 ↓
3Event
4 ↓
5D

Explicitly ordered across streams.

30. Complete AI-Oriented Execution Model

You've now reached:

textCopy
1CPU
2 │
3 ├── Prepare data
4 │
5 ↓
6Pinned Memory
7 │
8 ↓
9Async H2D
10 │
11 ↓
12CUDA Stream
13 │
14 ├── Kernel 1
15 ├── Kernel 2
16 ├── Kernel 3
17 │
18 ↓
19Async D2H
20 │
21 ↓
22CPU

With multiple streams:

textCopy
1             GPU
2              │
3      ┌───────┼────────┐
4      ↓       ↓        ↓
5   Stream 0 Stream 1 Stream 2
6      │       │        │
7      ↓       ↓        ↓
8    Work A   Work B   Work C

This is the foundation of high-throughput GPU execution.

31. Practical Exercise

Create:

textCopy
1Stream 0
2Stream 1

Run two kernels:

cppCopy
1kernelA<<<..., stream0>>>();
2kernelB<<<..., stream1>>>();

Then compare:

Version 1

textCopy
1One stream

Version 2

textCopy
1Two streams

Measure with CUDA events / Nsight.

Then make Kernel A computationally heavy and observe whether the second stream still improves performance.

This teaches an important lesson:

Concurrency depends on available GPU resources, not merely the number of streams.

32. What You Should Remember

The most important concepts:

textCopy
11. A stream is an ordered sequence of CUDA operations.
2
32. Operations in different streams may execute concurrently.
4
53. Concurrency is not guaranteed.
6
74. cudaMemcpyAsync() enables asynchronous transfer behavior.
8
95. Pinned host memory is important for efficient async H2D/D2H transfers.
10
116. cudaStreamSynchronize() waits for one stream.
12
137. cudaDeviceSynchronize() waits broadly for device work.
14
158. CUDA events provide fine-grained synchronization and timing.
16
179. Streams can overlap transfers and computation.
18
1910. Too much synchronization destroys concurrency.

Your AI-Focused CUDA Roadmap

You have now completed:

textCopy
1AI CUDA Kernel Engineering
2│
3├── Warp / SIMT execution              ✅
4├── GPU memory hierarchy               ✅
5├── Memory coalescing                  ✅
6├── Shared-memory optimization         ✅
7├── Occupancy                          ✅
8├── Kernel optimization                ✅
9├── CUDA streams + async execution     ✅
10│
11├── Tensor Cores                       ← NEXT
12├── CUDA profiling / Nsight
13└── Multi-GPU / NCCL

🚀 Next: Tensor Cores

This is a major jump for AI engineering.

You'll learn:

textCopy
1CUDA Cores
2      vs
3Tensor Cores
4
5FP32
6FP16
7BF16
8TF32
9INT8
10FP8
11      ↓
12Matrix Multiply-Accumulate
13      ↓
14WMMA
15      ↓
16MMA instructions
17      ↓
18Tensor Core kernels
19      ↓
20GEMM
21      ↓
22Transformer / LLM acceleration

For your AI/LLM goal, Tensor Cores are one of the highest-value remaining CUDA topics.