CUDA streams, kernels and timing: when is the GPU actually done?CUDA 流、内核与计时：GPU 到底什么时候算完？

A Python return and a GPU completion are different eventsPython 返回与 GPU 完成，是两件事

Part 4 explained when the backward engine can schedule a node. On CUDA, scheduling an operation introduces another boundary: the host can submit device work and continue before that work finishes. Consider two dependent operations, y = x.square() and z = y + 1. A Python variable can already refer to z while its values are still being computed on the GPU.第四期解释了反向引擎何时可以调度节点。在 CUDA 上，调度操作还涉及另一道边界：主机提交设备任务后，可以先继续执行，设备任务稍后才完成。以两个依赖操作 y = x.square() 和 z = y + 1 为例，Python 变量已经引用 z 时，GPU 可能仍在计算它的值。

A CUDA stream is an ordered sequence of device work, not a dedicated CPU thread or a reserved GPU multiprocessor. Work in different streams may overlap if resources permit; separate streams alone do not guarantee concurrency. For a dependent pair, overlap is not the goal: the consumer must see the producer’s completed writes.CUDA 流是一串有顺序的设备任务，不是专属 CPU 线程，也不独占某个 GPU 多处理器。资源允许时，不同流中的任务可能重叠，但建立两条流不保证并发。对于有依赖的两步操作，目标并不是重叠执行，而是确保消费者能读取生产者已完成的写入。

CUDA Programming Guide: stream ordering and possible concurrencyCUDA 编程指南：流内顺序与可能的并发

Baseline: PyTorch 2.10.0, source commit 449b17684101; eager execution on one CUDA device. The authoring environment has only 2.10.0+cpu. The CUDA programs below include availability checks; their GPU branches and timing results have not been measured here. The timeline uses invented logical units, and the standalone C++ dependency model runs on the CPU.基线：PyTorch 2.10.0，源码提交 449b17684101；讨论单个 CUDA 设备上的 eager 执行。写作环境只有 2.10.0+cpu。下文 CUDA 程序包含可用性检查；其 GPU 分支及计时结果尚未在本环境实测。时间线使用人为设定的逻辑单位，独立的 C++ 依赖模型在 CPU 上运行。

Read the dependency on a shared timeline在同一条时间轴上看依赖

Let A produce y and B consume it. In this illustration the host finishes submission at 2, A finishes at 7, and a correctly ordered B finishes at 10. The striped region is a device-side wait: the host does not spend that interval blocked in wait_event(). The bars represent operations, not a claim that each operation launches exactly one kernel.设 A 生成 y，B 消费它。在这张示意图中，主机在 2 完成提交，A 在 7 完成，顺序正确的 B 在 10 完成。斜纹区域表示设备侧等待，主机不会在这段时间里一直阻塞于 wait_event()。条带表示操作，不代表每个操作恰好启动一个内核。

In the missing-wait view, B is allowed to begin before A completes. This is one possible hazardous ordering, not a prediction of a specific GPU scheduler. Even if one run happens to produce the expected values, the program has not established the dependency. The earlier finish is therefore not a valid performance improvement.在“缺少等待”的视图中，B 可能在 A 完成前开始。这是缺乏依赖时一种可能的危险顺序，不是对某个 GPU 调度器的预测。即使某次运行恰好得到预期值，程序也没有建立所需依赖，因此提前结束不能算有效的性能提升。

From an ATen call to the current CUDA stream从 ATen 调用走到当前 CUDA 流

For a dense two-dimensional torch.mm(a, b), the CUDA implementation enters mm_out_cuda, which reuses addmm_out_cuda_impl with beta 0 and alpha 1. The implementation handles layout and output preparation before the BLAS operation. One float32 path reaches cublasSgemm; alternative paths and library choices depend on dtype, configuration and input conditions.对于稠密二维张量的 torch.mm(a, b)，CUDA 实现进入 mm_out_cuda，以 beta 为 0、alpha 为 1 复用 addmm_out_cuda_impl。实现先处理布局和输出准备，再调用 BLAS。其中一条 float32 路径走到 cublasSgemm；其他路径及库的选择取决于数据类型、配置和输入条件。

Blas.cpp: mm_out_cuda and the addmm implementationBlas.cpp：mm_out_cuda 与 addmm 实现

CUDABlas.cpp: the float32 cuBLAS GEMM pathCUDABlas.cpp：float32 的 cuBLAS GEMM 路径

The bridge to stream ordering is visible in getCurrentCUDABlasHandle(): the handle is associated with PyTorch’s current stream before use. This small source excerpt is part of PyTorch, not a standalone compilable program.连接流顺序语义的关键位于 getCurrentCUDABlasHandle()：使用句柄前，将它关联到 PyTorch 当前流。下面是 PyTorch 内部的短源码片段，不能单独编译。

auto stream = c10::cuda::getCurrentCUDAStream();
TORCH_CUDABLAS_CHECK(cublasSetStream(handle, stream));

CublasHandlePool.cpp: associate a cuBLAS handle with the current streamCublasHandlePool.cpp：把 cuBLAS 句柄关联到当前流

The current-stream selection is thread-local and indexed by device. A with torch.cuda.stream(s) context changes that selection for operations submitted inside it, then restores the previous selection; it does not mean “wait until s is done.” Native pointwise CUDA launch helpers also obtain the current stream and pass it into the kernel launch configuration. A stream is therefore carried down to the launch boundary, rather than inferred later from the tensor’s Python name.当前流选择是线程局部的，并按设备保存。with torch.cuda.stream(s) 上下文改变内部操作提交时的当前流，退出后恢复原选择；它不表示“等 s 执行完再退出”。原生逐元素 CUDA 启动辅助函数也会取得当前流，并传入内核启动配置。因此，流信息会一路传到启动边界，而不是事后根据张量的 Python 名称推断。

CUDAStream.cpp: getCurrentCUDAStream and setCurrentCUDAStreamCUDAStream.cpp：取得与设置当前流

CUDALoops.cuh: the stream argument in a native kernel launchCUDALoops.cuh：原生内核启动中的流参数

Ordering and storage lifetime need separate guarantees执行顺序与存储生命周期，需要分别保证

The following complete program records an event after A, makes the consumer stream wait for it, then records completion after B. done.synchronize() blocks the host until that completion event. The expected vector follows directly from i² + 1: [1, 2, 5, 10, 17, 26, 37, 50]. This is an expected result, not a recorded GPU run.下面的完整程序在 A 之后记录事件，让消费者流先等待该事件，再执行 B，并在 B 之后记录完成事件。done.synchronize() 阻塞主机，直到这个完成事件结束。由 i² + 1 可直接得到预期向量 [1, 2, 5, 10, 17, 26, 37, 50]；这是预期结果，不是已记录的 GPU 运行输出。

import torch

def main():
    if not torch.cuda.is_available():
        print("CUDA unavailable; GPU example not run.")
        return
    device = torch.device("cuda:0")
    with torch.cuda.device(device):
        producer = torch.cuda.current_stream(device)
        consumer = torch.cuda.Stream(device=device)
        x = torch.arange(8, device=device, dtype=torch.float32)
        y = x.square()                         # A, on producer
        ready = torch.cuda.Event()
        ready.record(producer)
        with torch.cuda.stream(consumer):
            consumer.wait_event(ready)         # Execution dependency
            z = y + 1                         # B, on consumer
            y.record_stream(consumer)          # Allocator lifetime
            done = torch.cuda.Event()
            done.record(consumer)
        del y                                 # Safe reuse is deferred
        done.synchronize()
        actual = z.cpu()
        expected = torch.arange(8, dtype=torch.float32).square() + 1
        torch.testing.assert_close(actual, expected)
        print(actual.tolist())

if __name__ == "__main__":
    main()

wait_event prevents B from reading unfinished A output. record_stream informs the caching allocator that y is also used on the consumer stream, so dropping its last Python reference must not allow premature storage reuse. Neither API substitutes for the other. Keeping the tensor alive until the consumer has completed is another way to avoid early release; this example deliberately uses del y to expose the allocator issue.wait_event 防止 B 读取 A 尚未完成的输出。record_stream 告诉缓存分配器，y 还在消费者流上使用，因此即使最后一个 Python 引用被释放，也不能过早复用存储。两个 API 不能互相替代。把张量一直保留到消费者完成，也能避免提前释放；这里特意使用 del y 来展示分配器需要处理的问题。

Tensor.record_stream: cross-stream use and allocator safetyTensor.record_stream：跨流使用与分配器安全

consumer.wait_stream(producer) is convenient when waiting for the producer’s already-submitted work. Its implementation records an event on the producer and waits on it. It does not make the consumer wait for every future operation that will ever be submitted to the producer. For a specific boundary, an explicitly named event makes the dependency easier to inspect.需要等待生产者已提交的任务时，consumer.wait_stream(producer) 很方便：实现会在生产者上记录事件，再等待该事件。它不会让消费者等待生产者今后提交的所有任务。若依赖对应某个明确边界，显式命名的事件更便于检查。

streams.py: wait_event and the submission boundary of wait_streamstreams.py：wait_event 与 wait_stream 的提交边界

Choose the synchronization boundary deliberately明确选择同步边界

In CUDAEvent.h, recording reaches cudaEventRecordWithFlags, a stream wait reaches cudaStreamWaitEvent, and host synchronization reaches cudaEventSynchronize. An event marks a point in a stream; it does not imply that unrelated work in other streams has completed.在 CUDAEvent.h 中，记录事件走到 cudaEventRecordWithFlags，流等待走到 cudaStreamWaitEvent，主机同步走到 cudaEventSynchronize。事件标记一条流上的位置，不代表其他流上无关的任务也已完成。

CUDAEvent.h: record, block, elapsed_time and synchronizeCUDAEvent.h：记录、等待、计时与同步

torch.cuda.synchronize: the device-wide boundarytorch.cuda.synchronize：设备范围的同步边界

Reading a CUDA scalar with .item() requires its value on the CPU. Putting that read inside a benchmark loop can add a synchronization point to every iteration. Likewise, asynchronous errors can surface at a later synchronization call rather than the launch that caused them. CUDA_LAUNCH_BLOCKING=1 is useful for debugging the origin of an error, but changes execution behavior and should not silently remain enabled for a performance report.通过 .item() 读取 CUDA 标量，需要把它的值提供给 CPU。把这个读取放在基准测试循环内，可能让每次迭代都多出一个同步点。同样，异步错误可能在后续同步调用处暴露，而不是在引发错误的启动位置暴露。CUDA_LAUNCH_BLOCKING=1 有助于调试错误来源，但会改变执行行为，不应在性能报告中未经说明地保持启用。

PyTorch CUDA semantics: asynchronous execution and debuggingPyTorch CUDA 语义：异步执行与调试

Three timers answer three different questions三种计时，回答三个不同问题

Use host-clock readings before submission, after submission and after the completion wait. Device events provide a separate pair of timestamps, and ; never subtract a host timestamp from a device timestamp.用主机时钟记录提交前的 、提交后的 和完成等待后的 。设备事件提供另一对时间戳 与 ；不要把主机时间戳与设备时间戳直接相减。

Host interval: how long did Python, dispatch and submission take within the chosen region? It may include blocking by the runtime; it is not guaranteed to be pure launch overhead. Event interval: elapsed device time between two recorded events on the measured path. It can include idle gaps and interference from other work, so it is not necessarily the sum of kernel durations. Synchronized wall interval: how long until this submitted batch is known to be complete, including submission and the final host wait?主机区间：选定区域内，Python、调度与提交用了多久？其中可能包含运行时阻塞，不保证只是启动开销。事件区间：测量路径上两个已记录事件之间的设备经过时间。它可能包含空闲间隙和其他任务的干扰，因此不一定等于各内核时长之和。同步后的墙钟区间：从开始提交，到确认这批任务完成一共多久？它包含提交和最后的主机等待。

CUDA events need enable_timing=True, and both must have completed before querying elapsed time. The event API returns milliseconds. For work on several streams, put explicit joins before the end event; two events recorded only on the main stream do not automatically bracket independent work on a worker stream.CUDA 事件需要设置 enable_timing=True，并在两个事件都完成后查询经过时间。事件 API 返回毫秒。多流任务需要在结束事件前显式汇合；仅在主流上记录两个事件，不会自动把工作流上的独立任务包含在内。

torch.cuda.Event: timing, recording and completiontorch.cuda.Event：计时、记录与完成

For the timeline above, a valid cross-stream measurement places the start event before A, uses the A→B dependency, records a completion event after B, and makes the measuring stream wait for that completion before recording its end event. The measured interval then covers the dependent path, rather than just the act of enqueueing B.对于前面的时间线，有效的跨流测量应在 A 前放置开始事件，保留 A→B 的依赖，在 B 后记录完成事件，并让测量流等待这个完成事件后再记录结束事件。这样测到的是有依赖的执行路径，而不只是把 B 入队的动作。

A reproducible matrix-multiplication benchmark可复现的矩阵乘法基准测试

This program reuses an output buffer, disables autograd, warms up the selected workload, checks an 8×8 output tile against CPU float64, and records seven rounds. Each round contains 50 matrix multiplications. The reported medians are medians of batch means, not per-request p50 latency. Allocation, input creation, correctness checking and CPU output transfer are outside the timed region.这个程序复用输出缓冲区、关闭自动微分、预热选定工作负载，并用 CPU float64 校验一个 8×8 输出块，再记录七轮数据。每轮包含 50 次矩阵乘法。报告的中位数是各批次平均值的中位数，不是单次请求延迟的 p50。分配、输入生成、正确性检查与输出传回 CPU 均不在计时区域内。

import json
import os
import statistics
from time import perf_counter
import torch

@torch.inference_mode()
def main():
    if not torch.cuda.is_available():
        print("CUDA unavailable; benchmark not run.")
        return
    device = torch.device("cuda:0")
    with torch.cuda.device(device):
        torch.manual_seed(0)
        # PyTorch 2.10 precision API; do not mix with allow_tf32.
        torch.backends.cuda.matmul.fp32_precision = "ieee"
        n, warmup, repeats, rounds = 1024, 10, 50, 7
        a = torch.randn(n, n, device=device, dtype=torch.float32)
        b = torch.randn_like(a)
        out = torch.empty_like(a)
        def work():
            torch.mm(a, b, out=out)

        for _ in range(warmup):
            work()
        torch.cuda.synchronize(device)
        reference = a[:8].cpu().double() @ b[:, :8].cpu().double()
        torch.testing.assert_close(out[:8, :8].cpu().double(), reference,
                                   rtol=1e-4, atol=1e-3)
        stream = torch.cuda.current_stream(device)
        start = torch.cuda.Event(enable_timing=True)
        end = torch.cuda.Event(enable_timing=True)
        # Create the lazily initialized events before collecting samples.
        start.record(stream)
        end.record(stream)
        end.synchronize()
        samples = []
        for _ in range(rounds):
            torch.cuda.synchronize(device)
            t0 = perf_counter()
            start.record(stream)
            for _ in range(repeats):
                work()
            end.record(stream)
            submitted = perf_counter()
            end.synchronize()
            completed = perf_counter()
            samples.append({
                "host_region_ms_per_call": 1000 * (submitted - t0) / repeats,
                "stream_interval_ms_per_call": start.elapsed_time(end) / repeats,
                "batch_wall_ms_per_call": 1000 * (completed - t0) / repeats,
            })
        print(json.dumps({
            "torch": torch.__version__, "built_cuda": torch.version.cuda,
            "gpu": torch.cuda.get_device_name(device),
            "capability": torch.cuda.get_device_capability(device),
            "dtype": str(a.dtype), "shape": list(a.shape),
            "a_stride": list(a.stride()), "b_stride": list(b.stride()),
            "matmul_fp32_precision": torch.backends.cuda.matmul.fp32_precision,
            "CUDA_LAUNCH_BLOCKING": os.getenv("CUDA_LAUNCH_BLOCKING", "unset"),
            "warmup": warmup, "repeats": repeats, "rounds": rounds,
            "samples": samples,
            "median_of_batch_means": {
                key: statistics.median(row[key] for row in samples)
                for key in samples[0]
            },
        }, indent=2))

if __name__ == "__main__":
    main()

The host region also contains the event-record calls, so treat it as a measured region rather than an isolated launch-cost estimate. The wall measurement ends after the end-event wait, not after a full application request: a real end-to-end test must place input transfers, preprocessing and output retrieval inside its boundary when those are part of the request. Do not add these three measurements together; they describe overlapping intervals.主机区域还包含事件记录调用，因此应把它视为选定区域的耗时，而不是独立的启动成本估计。墙钟测量在结束事件等待完成后停止，并不覆盖完整应用请求：若输入传输、预处理和输出取回属于请求流程，真正的端到端测试需要把它们放进测量边界。不要把三项结果相加，它们对应重叠的时间区间。

Record the driver version from nvidia-smi alongside the JSON; torch.version.cuda identifies the CUDA version used to build PyTorch, not the driver. Also record competing GPU workloads and power/clock conditions. Keeping dtype, shape and precision fixed matters: allowing TF32 changes the numerical contract as well as possible performance. This example selects IEEE float32 through the PyTorch 2.10 precision API.发布 JSON 时，还应附上 nvidia-smi 中的驱动版本；torch.version.cuda 表示构建 PyTorch 所用的 CUDA 版本，不是驱动版本。同时记录其他 GPU 负载及功耗、频率条件。固定数据类型、形状和精度很重要：允许 TF32 会同时改变数值约定与可能的性能。此例通过 PyTorch 2.10 的精度 API 选择 IEEE float32。

PyTorch CUDA semantics: controlling float32 and TF32 precisionPyTorch CUDA 语义：float32 与 TF32 精度控制

Make the ordering constraints executable in C++把顺序约束写成可运行的 C++

This CPU-only model computes the earliest allowed start under the stated ordering constraints. A starts at 1 and takes 6 units; B is submitted at 2 and takes 3. Same-stream order and an explicit event both force B to start at 7. Without either, the model permits 2. It models a dependency bound, not SM occupancy, kernel throughput or actual CUDA scheduling. Edit the durations to see which constraint determines the critical path.这个纯 CPU 模型计算给定顺序约束下允许的最早开始时刻。A 在 1 开始、持续 6 个单位；B 在 2 提交、持续 3 个单位。同流顺序和显式事件都会把 B 的开始限制到 7；两者都没有时，模型允许在 2 开始。它建模的是依赖边界，不是 SM 占用率、内核吞吐率或真实 CUDA 调度。修改持续时间，可以观察关键路径由哪条约束决定。

#include <algorithm>
#include <iostream>
#include <stdexcept>
#include <string>

struct Stream {
    int available = 0;
    void wait(int event_time) {
        available = std::max(available, event_time);
    }
    int enqueue(int submitted, int duration) {
        const int start = std::max(submitted, available);
        available = start + duration;
        return start;
    }
};

int main(int argc, char** argv) {
    const int a_duration = argc > 1 ? std::stoi(argv[1]) : 6;
    const int b_duration = argc > 2 ? std::stoi(argv[2]) : 3;
    if (a_duration < 1 || a_duration > 100 ||
        b_duration < 1 || b_duration > 100)
        throw std::invalid_argument("durations must be in [1, 100]");
    for (const std::string mode : {"same stream", "event wait", "missing wait"}) {
        Stream producer, consumer;
        producer.enqueue(1, a_duration);
        const int ready = producer.available;
        Stream& target = mode == "same stream" ? producer : consumer;
        if (mode == "event wait") target.wait(ready);
        const int start = target.enqueue(2, b_duration);
        std::cout << mode << ": A_done=" << ready
                  << ", B=[" << start << ',' << target.available << "]"
                  << ", ordered=" << std::boolalpha << (start >= ready) << '\n';
    }
}

Output checked locally and with GCC 14.2 on Godbolt:以下输出已在本地及 Godbolt 的 GCC 14.2 中核对：

same stream: A_done=7, B=[7,10], ordered=true
event wait: A_done=7, B=[7,10], ordered=true
missing wait: A_done=7, B=[2,5], ordered=false

If A is shortened enough that it finishes by 2, the missing-wait scenario may look ordered in this particular schedule. That observation does not fix the program: correctness needs an ordering guarantee across permitted schedules, not one successful run.如果把 A 缩短到在 2 之前完成，“缺少等待”在这一组时序下也可能显得有序。但这不能修复程序：正确性需要对允许的时序建立顺序保证，而不是依赖某一次恰好成功的运行。

Use a trace to explain the intervals用运行轨迹解释时间区间

A timing table cannot show where the gaps occur. This separate profiler program exports cuda-mm-trace.json. Inspect the CPU operator and CUDA runtime lanes alongside the GPU kernel lanes in a compatible trace viewer: the CPU span shows host work, while the kernel spans show device execution. One ATen operation need not correspond to one kernel. Profiling adds overhead, so collect the trace separately from the benchmark.计时表无法显示间隙发生在哪里。下面独立的 profiler 程序导出 cuda-mm-trace.json。在兼容的轨迹查看器中，结合 CPU 算子、CUDA 运行时与 GPU 内核轨道阅读：CPU 区段显示主机工作，内核区段显示设备执行。一个 ATen 操作不一定对应一个内核。采集轨迹会增加开销，因此应与基准测试分开运行。

import torch
from torch.profiler import profile, ProfilerActivity, record_function

@torch.inference_mode()
def main():
    if not torch.cuda.is_available():
        print("CUDA unavailable; trace not recorded.")
        return
    with torch.cuda.device(0):
        a = torch.randn(1024, 1024, device="cuda:0")
        b, out = torch.randn_like(a), torch.empty_like(a)
        for _ in range(10):
            torch.mm(a, b, out=out)
        torch.cuda.synchronize()
        with profile(activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA]) as prof:
            with record_function("measured_mm_batch"):
                for _ in range(5):
                    torch.mm(a, b, out=out)
            torch.cuda.synchronize()
        prof.export_chrome_trace("cuda-mm-trace.json")
        print(prof.key_averages().table(sort_by="self_cuda_time_total", row_limit=10))

if __name__ == "__main__":
    main()

torch.profiler: CPU/CUDA activities and Chrome trace exporttorch.profiler：CPU／CUDA 活动与 Chrome 轨迹导出

Overlapping a host-to-device copy with computation needs more than non_blocking=True: pinned host memory, suitable hardware, independent work and appropriate streams must also be present. A dependent consumer must still wait for the copy. Keep a pinned source unchanged until its asynchronous copy completes, and wait before reading an asynchronously copied device-to-host result on the CPU.要让主机到设备的复制与计算重叠，仅有 non_blocking=True 不够：还需要锁页主机内存、合适的硬件、独立任务与恰当的流。有依赖的消费者仍必须等待复制。异步复制完成前，不要改写锁页源缓冲区；设备到主机的异步复制完成前，也不要在 CPU 上读取结果。

PyTorch transfer tutorial: pinned memory, overlap and buffer lifetimePyTorch 传输教程：锁页内存、重叠与缓冲区生命周期

The same reasoning applies to part 4’s backward graph. CUDA backward operations use the streams of their corresponding forward operations; the engine handles necessary internal dependencies. If another stream consumes the resulting gradients, establish the required external dependency before that use. A Python backward() return is not a general device-wide completion barrier. The dependency graph determines what can be submitted; CUDA streams determine how the submitted device work is ordered.同样的推理也适用于第四期的反向图。CUDA 反向操作使用对应前向操作的流，引擎处理必要的内部依赖。如果另一条流要消费得到的梯度，使用前仍需建立所需的外部依赖。Python 的 backward() 返回，不是通用的设备全局完成屏障。依赖图决定哪些任务可以提交，CUDA 流决定已提交的设备任务如何排序。

PyTorch CUDA semantics: streams used by backward passesPyTorch CUDA 语义：反向传播使用的流