From c2fc7c1645e7c414cff10718da997f2289c569d4 Mon Sep 17 00:00:00 2001 From: tlopex <820958424@qq.com> Date: Sun, 16 Aug 2026 01:14:46 -0400 Subject: [PATCH 1/4] Add GPU benchmarking and profiling appendix --- README.md | 4 +- appendix/benchmarking_gpu_kernels.md | 870 ++++++++++++++++++++++++ appendix/index.md | 3 +- appendix/nsys_example.py | 88 +++ chapter_gemm_advanced/index.md | 2 +- chapter_gemm_basics/index.md | 5 + chapter_performance/index.md | 4 + img/nsys_b200_timeline.svg | 64 ++ img/nsys_b200_timeline_zh.svg | 64 ++ img/scripts/gen_nsys_b200_timeline.py | 120 ++++ index.md | 7 +- zh/appendix/benchmarking_gpu_kernels.md | 616 +++++++++++++++++ zh/appendix/index.md | 3 +- zh/chapter_gemm_advanced/index.md | 2 +- zh/chapter_gemm_basics/index.md | 4 + zh/chapter_performance/index.md | 3 + zh/index.md | 5 +- 17 files changed, 1854 insertions(+), 10 deletions(-) create mode 100644 appendix/benchmarking_gpu_kernels.md create mode 100644 appendix/nsys_example.py create mode 100644 img/nsys_b200_timeline.svg create mode 100644 img/nsys_b200_timeline_zh.svg create mode 100644 img/scripts/gen_nsys_b200_timeline.py create mode 100644 zh/appendix/benchmarking_gpu_kernels.md diff --git a/README.md b/README.md index 7fd5993f..aaedf18e 100644 --- a/README.md +++ b/README.md @@ -25,7 +25,9 @@ vehicle is **TIRx** (Tensor IR next), a Python DSL for writing GPU kernels at th persistent scheduling, warp specialization, and 2-CTA clusters. - **Part IV — Flash Attention 4.** A complete attention kernel built from the Part III techniques: two MMAs with softmax between them, online-softmax rescaling, causal masking, and GQA. -- **Reference.** TIRx language reference and compiler internals. +- **Appendices.** TIRx language reference, [reproducible GPU benchmarking and + profiling](appendix/benchmarking_gpu_kernels.md), compiler internals, and asynchronous-kernel + debugging. ## Build the book locally diff --git a/appendix/benchmarking_gpu_kernels.md b/appendix/benchmarking_gpu_kernels.md new file mode 100644 index 00000000..942833b1 --- /dev/null +++ b/appendix/benchmarking_gpu_kernels.md @@ -0,0 +1,870 @@ +(chap_benchmarking)= +# Measuring and Analyzing GPU Kernel Performance + +GPU kernel optimization involves two separate questions: how fast the operation is, and where its +time goes. A benchmark answers the first question; a profile helps answer the second. Because a +profiler changes the execution environment, final claims about complete-operator or application +latency should be confirmed with an unprofiled measurement. + +One Python call does not necessarily correspond to one GPU kernel. It may launch several kernels, +enqueue memory copies, or wait for GPU work to finish. Before timing, define which of those steps +belong to the operation and use the same boundary for every implementation. + +In practice, first verify correctness and establish an unprofiled baseline, then use a profiler to +investigate where the time goes. After changing the implementation, repeat the same measurement. An +optimization is successful only when it improves the unprofiled baseline. + +The performance chapter ({ref}`chap_performance`) uses roofline analysis to reason about compute and +memory limits. Here the focus shifts to experiments: defining the timing boundary, choosing warm-up +and repeat budgets, and reading profiler reports. + +## Separate Measurement from Diagnosis + +The tools used in this workflow have different jobs: + +| Tool | Primary question | +|---|---| +| CUDA Events | How much GPU-stream time elapsed across the measured region? A stream is an ordered queue of GPU work. | +| Proton (provided by Triton) | Which GPU kernels were launched, how often, and which kernels account for most of the time? | +| Nsight Systems | How do host work, streams, copies, kernels, and communication overlap on a timeline? | +| Nsight Compute (`ncu`) | Why does one selected GPU kernel spend its cycles the way it does? | +| IKET (optional) | Which named phases or warp roles consume time inside one selected kernel? | + +### Three Common Profile Views + +A profile is not a single number or a single report format. The tools in this chapter produce three +complementary views: + +| View | Tools | How to read it | +|---|---|---| +| Aggregation tree | Proton | Compare call count, average duration, and total duration to locate expensive kernels. | +| Timeline | Nsight Systems; IKET inside one kernel | Read time from left to right across tracks; inspect gaps, overlap, and dependencies. | +| Per-kernel metric report | Nsight Compute | Read launch configuration, utilization, scheduler state, memory traffic, and source/SASS evidence for one launch. | + +Profiles explain where time is spent; they do not replace the performance measurement. After changing +an implementation, disable profiling and measure it again with the same timing boundary used for the +baseline. + +## Verify Correctness Before Timing + +Verify correctness separately before collecting performance: + +1. Construct representative inputs, including relevant boundary cases. +2. Run the implementation and synchronize so that GPU work has completed. +3. Compare the output with a reference under a stated tolerance. +4. If the kernel accumulates into an existing output or modifies an input in place, restore the same + initial state before each correctness check. + +Reference computation and result comparison stay outside performance timing. Design the benchmark +only after correctness passes; whether state reset is timed depends on the operation boundary defined +in the next section. + +## Define the Timing Boundary + +Before timing, define the work that constitutes one measured operation. It may contain only one +kernel, or it may include every kernel, memory copy, and state reset required to produce the complete +result. State explicitly whether compilation, input construction, allocation, or data conversion is +part of that operation. Implementations are directly comparable only when they perform the same work +inside the measured boundary. + +After fixing the scope, choose the timer: + +- **CUDA Events** record timestamps when a GPU stream reaches two points. They can measure the + device-timeline interval around one kernel or a complete operator. Kernels, memory copies, and idle + stream gaps inside that interval all count. For a multi-stream operation, measured work on every + participating stream must begin after the start event and join before the end event is recorded. +- A **synchronized wall-clock timer** starts before the host call and stops after all GPU work required + by that call has completed. It also includes Python dispatch, CUDA launch, and the wait for GPU + completion, so it is appropriate for end-to-end call latency. + +For example, if the GPU executes the start event before the host submits the next launch, that idle +stream time remains inside the CUDA Event interval. An Event interval is therefore not necessarily +the same as a kernel's start-to-finish execution interval in a profiler. Profilers are useful for +examining execution and overlap, but diagnostic profiles do not directly replace unprofiled timing at +the same boundary. + +## Measure GPU Time with CUDA Events + +CUDA launches are normally asynchronous. Python can continue after submitting work to a CUDA stream, +before the GPU has finished. A CPU timer placed immediately around that Python call can therefore +stop too early and mostly measure host submission time. Use CUDA Events to measure elapsed time on a +GPU stream. Use the synchronized wall-clock timer shown later when the boundary runs from the Python +call through GPU completion. The +[PyTorch CUDA semantics documentation](https://docs.pytorch.org/docs/stable/notes/cuda.html#asynchronous-execution) +describes this behavior in more detail. + +Start with a complete CUDA Event benchmark. The following runnable example allocates its matrices, +runs a warm-up, and then measures five rounds of one FP16 GEMM before reporting their median: + +```python +from statistics import median + +import torch + + +a = torch.randn((2048, 2048), device="cuda", dtype=torch.float16) +b = torch.randn((2048, 2048), device="cuda", dtype=torch.float16) +c = torch.empty((2048, 2048), device="cuda", dtype=torch.float16) + + +def gemm(): + torch.mm(a, b, out=c) + + +def measure_batch_ms(fn, calls): + """Return mean CUDA Event time per call over one consecutive batch, in ms.""" + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + + start.record() + for _ in range(calls): + fn() + end.record() + end.synchronize() + return start.elapsed_time(end) / calls + + +warmup_calls = 500 +repeat = 100 +rounds = 5 + +for _ in range(warmup_calls): + gemm() +torch.cuda.synchronize() + +samples_ms = [measure_batch_ms(gemm, repeat) for _ in range(rounds)] + +print(f"device: {torch.cuda.get_device_name()}") +print(f"calls per round: {repeat}") +print("round samples (ms):", [round(x, 4) for x in samples_ms]) +print(f"median CUDA Event time: {median(samples_ms):.4f} ms") +``` + +`measure_batch_ms` records start and end events in the current CUDA stream and divides their elapsed +time by the number of calls. The result is the mean GPU-stream time per GEMM during consecutive +execution. `end.synchronize()` only makes the CPU wait for that round of GPU work so that the Event +result can be read. + +Here `warmup_calls=500`, `repeat=100`, and `rounds=5` are invocation or measurement counts. They are +example values selected from measurements of this GEMM on a B200, not universal defaults. In a +ten-round calibration, 50 warm-up calls still produced a first-to-last decline from 0.01425 ms to +0.01296 ms. At 500 calls, the change narrowed to 0.01332 ms to 0.01301 ms. Results with `repeat=100` +were also more stable than with `repeat=10`. + +For another workload, increase `warmup_calls` until the first rounds no longer become consistently +faster or slower. Then increase `repeat` or `rounds` until the variation is acceptable for the +experiment. Larger counts are not automatically better: if longer runs systematically shift the +timing level, inspect temperature, power, and clock behavior and decide whether the experiment should +represent short bursts or sustained execution. Long-running kernels generally need smaller counts. + +This code reuses the same matrices, so later calls may find some data in cache. It therefore represents +a warm-cache workload. A published result should also record the GPU model, software versions, and +clock settings. + +When benchmarking the TIRx kernels in this book, there is no need to rewrite the warm-up, repeated +timing, and statistics loop for every kernel. TVM's +[`tvm.tirx.bench.bench`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/bench.py) +already provides those steps. Pass it a function that launches the prepared implementation; inputs, +outputs, and workspace remain allocated outside the measured interval. + +The helper uses a different cache policy from the manual example. The example repeatedly reuses the +same matrices, whereas `bench` evicts L2 before each measured invocation and records an independent +CUDA Event interval. Invoke it as follows: + +```python +from tvm.tirx.bench import bench + + +# run is a no-argument function that launches the operation on preallocated tensors. +result = bench( + {"tirx": run}, + timer="event", + warmup=25, + repeat=100, + rounds=5, + cooldown_s=1.0, +) + +print(result["impls"]["tirx"]) # five-round mean, in us +print(result["round_samples"]["tirx"]) # result from each round +``` + +Here `warmup=25` and `repeat=100` are time budgets in milliseconds, not invocation counts. The Event +timer first performs a short estimate, then converts the 25 ms warm-up budget and 100 ms measurement +budget into iteration counts. That estimate includes both L2 eviction and the measured call, so the +resulting counts are approximate. In the reported samples, the L2 eviction occurs before the start +event and only the invocation is timed. Short kernels therefore run more times than long kernels. +`rounds=5` repeats the complete measurement five times, while `cooldown_s=1.0` waits one second before +measuring an implementation in each round. `impls` contains the five-round mean and `round_samples` +retains the individual results. The 25/100 ms values are the Event timer defaults. Five rounds are the +default used by the TIRx-kernels CLI; `bench` itself defaults to one round. + +These values are starting points rather than a standard for every workload. Increase the warm-up +budget if results continue to drift between rounds. Increase the measurement budget or number of +rounds if the results remain noisy. Use the same timer, budgets, and rounds for every implementation, +and retain all round results instead of reporting only the fastest one. + +TIRx-kernels uses this helper in its `run_bench` entry points; see +[`tirx_kernels/attention/flash_attention4.py`](https://github.com/mlc-ai/tirx-kernels/blob/5be39749e7dfd2c4bdae9b4d396f8ec35af07126/tirx_kernels/attention/flash_attention4.py). +The local benchmark defaults to Proton when `timer` is omitted. Specify `timer="event"` as above when +the intended result is a CUDA Event interval. The two timers report different quantities, so a result +must identify which one was used. + +The function passed to `bench` is invoked repeatedly. If a kernel accumulates into its output or +modifies an input in place, restore equivalent state before every call or ensure that every measured +invocation receives fresh preallocated state. A reset inside the measured function belongs to the +operation boundary defined earlier. Otherwise later calls no longer represent the same workload. + +### Measure End-to-End Time for One Call + +Use a synchronized wall-clock timer when the target is the complete interval from one Python call +until its GPU work finishes. The following code continues with the `gemm()` defined and warmed up +above: + +```python +from statistics import median +import time + +import torch + + +def measure_single_call_ms(fn, samples=20): + values = [] + for _ in range(samples): + torch.cuda.synchronize() # exclude unfinished work from earlier calls + t0 = time.perf_counter() + fn() + torch.cuda.synchronize() # wait for this call's GPU work + values.append((time.perf_counter() - t0) * 1e3) + return values + + +host_samples_ms = measure_single_call_ms(gemm) +print("single-call samples (ms):", [round(x, 4) for x in host_samples_ms]) +print(f"median end-to-end time: {median(host_samples_ms):.4f} ms") +``` + +The first synchronization keeps unfinished earlier work outside the measurement. The second ensures +that this GEMM finishes before the timer stops. Each sample contains exactly one call, so the result +includes the Python call, CUDA launch, GPU execution, and the wait for completion. The CUDA Event +benchmark above instead reports mean GPU-stream time per GEMM during consecutive execution. + +Every implementation in a comparison must use the same timer and boundary. If both results are +reported, name them separately as *CUDA Event GPU time* and *single-call end-to-end time* rather than +placing values from different timers under one generic latency label. + +### Timing Overlapping GPU Work + +The GEMM examples above run entirely in the current CUDA stream, so their start and end events cover +all of the work. An operator that uses several streams needs additional synchronization. Events +recorded only in the current stream do not automatically include work elsewhere, which may begin +before the start event or remain unfinished after the end event. + +To time the complete operator, use the start event as a common starting signal. Every work stream +waits for start before beginning the measured work and signals completion when it finishes. The stream +that records the end event waits for all of those completion signals first. The resulting interval +then spans the operation from its earliest start through its final completion. + +Programmatic Dependent Launch (PDL) is another source of possible overlap. On GPUs with compute +capability 9.0 or newer, it can start a later kernel early in the same stream. That kernel may perform +preparation that does not depend on earlier results, then wait before consuming those results. PDL +must be enabled explicitly and follow its trigger-and-wait contract; see the +[CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html) +for the API details. + +The timing rule is the same whether overlap comes from multiple streams or PDL: place CUDA Events +around the complete operation. Overlapping kernels can cover the same time interval, so adding their +profiler durations does not produce operator latency. A Nsight Systems timeline shows the actual +ordering and overlap. Because PDL overlap is opportunistic, program correctness cannot require it to +occur. + +## Keep Experimental Conditions Consistent + +The preceding sections established the timing boundary and timer. The remaining experimental +conditions must also be held constant. The examples above start after allocation and warm-up, so they +measure subsequent repeated calls. A first call may also include CUDA initialization, JIT compilation, +autotuning, or other one-time work. If first-call latency or a complete application path is the target, +include those steps in the boundary and report the result separately from repeated-call performance. + +One measurement is not enough to establish stability. The manual CUDA Event example retains five +rounds and reports their median. `bench` reports the mean across rounds and also stores the raw values +in `round_samples`. Whichever summary is used, keep the per-round results, inspect them for trends and +outliers, and state whether the reported value is a median or mean rather than selecting only the +fastest round. When comparing implementations, repeat the experiment in a different order so that one +implementation is not always measured on a colder or hotter device. + +Cache conditions also affect the result. The manual example repeatedly uses the same matrices and +therefore measures a warm-cache workload. +The TVM 0.26 Event and Proton timers instead write a 256 MiB buffer before every measured invocation +to evict existing L2 data; that write remains outside the timed interval. Either policy can be valid. +Choose the one that represents the target application and apply it consistently to every +implementation. `torch.cuda.empty_cache()` releases unused blocks from PyTorch's caching allocator. +It does not clear GPU L2 and cannot implement a cold-L2 measurement. + +Finally, record the GPU model, driver, CUDA runtime, framework and compiler versions, and the workload +dtype and shape. Also record clock and power settings, keep unrelated processes off the device, and +watch for thermal throttling. If clocks are locked, provide the actual values and command; "fixed +clocks" alone is not enough to reproduce the experiment. + +Matching tensor shapes alone does not make two implementations comparable. Align at least three +classes of conditions: + +- **Numerical semantics:** input and output dtype, layout, transpose convention, alignment, + accumulation precision, scale, mask, epilogue, output definition, and accuracy tolerance; +- **Measured scope:** inclusion or exclusion of allocation, conversion, state reset, auxiliary + kernels, communication, and synchronization; +- **Tuning conditions:** workspace limits, whether per-shape autotuning is allowed, and the search + budget available to each implementation. + +Every implementation should also use the same cache, clock, warm-up, sampling, and timing policies. +For a library baseline, record its version, selected algorithm, and workspace. Autotuning may run +outside the timed interval, but its search budget and final configuration remain part of the +experimental record. + +## Convert Latency to Throughput + +Throughput is not measured directly by the timer. It is computed by dividing a defined amount of work +by the measured latency. A table that reports TFLOP/s, GB/s, or tokens/s should therefore retain the +original latency and explain how the work was counted. This book counts GEMM as $2MNK$ FLOPs. For +attention and fused kernels, specify whether the count represents the full dense problem, the elements +actually selected, or the work executed by the kernel. See {ref}`chap_performance` for the corresponding +formulas and roofline analysis. + +## Use Proton to Find Expensive Kernels + +The preceding benchmark tells us how long the complete operation takes. If it launches several kernels, +we still need to determine where that time is spent. Proton reports each kernel's call count, average +time, and cumulative time. + +Proton is a GPU profiler provided by the Triton project. It records CUDA kernel activity and can +therefore see TIRx kernels compiled by TVM; those kernels are not compiled by Triton. The `bench` +helper introduced above also supports `timer="proton"`: it aggregates kernel execution times for each +invocation and returns a statistic across the measured invocations. To see which kernels ran and how +often, collect a kernel tree in a separate Proton session. + +The following example reuses the matrices allocated above and combines GEMM with ReLU into a +two-kernel operation. It finishes warm-up, collects the next 100 calls, and writes `operator.hatchet` +in the current directory: + +```python +import torch +import triton.profiler as proton + + +def operation(): + torch.mm(a, b, out=c) + torch.clamp_min(c, 0, out=c) + + +def collect_proton(run, *, warmup_calls, profile_calls): + for _ in range(warmup_calls): + run() + torch.cuda.synchronize() + + session = proton.start("operator", context="shadow", data="tree") + if session is None: + raise RuntimeError("Proton session could not be created") + try: + with proton.scope("target_operation"): + for _ in range(profile_calls): + run() + torch.cuda.synchronize() + finally: + proton.finalize(session) + + +collect_proton(operation, warmup_calls=500, profile_calls=100) +``` + +Here `warmup_calls` and `profile_calls` are invocation counts, not the millisecond budgets used by +`bench`. Install a Triton version compatible with TVM and record that version with the experiment. + +First list the stored metrics. The next two commands print call counts, total time, and average time: + +```bash +proton-viewer --list operator.hatchet +proton-viewer --metrics time/ms,count --print-sorted operator.hatchet +proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet +``` + +If the viewer reports that `pandas` or `hatchet` is missing, install the optional packages with +`python -m pip install pandas llnl-hatchet`. The following is one real B200 result from this code; +kernel names are shortened for readability: + +```text +target_operation calls avg/us total/ms +├── GEMM kernel 100 14.83 1.483 +└── ReLU kernel 100 4.23 0.423 +``` + +First confirm that the expected kernels and call counts appear, then compare the leaf averages and +totals. GEMM has the largest total in this example, making it the better candidate for further +analysis with Nsight Compute. A short kernel can still accumulate substantial total time when it is +launched often. An average shown for a parent scope, however, is not the latency of one complete +operation. + +This tree is useful for finding expensive kernels, but it does not replace the benchmark above. +Overlapping kernels cover some of the same elapsed interval, while copies, synchronization, and stream +gaps may be absent from the tree. This capture also reuses the same matrices instead of reproducing the +L2-eviction policy of the earlier TVM timer, so the two sets of numbers are not directly comparable. +Obtain complete-operation time from the CUDA Event or wall-clock benchmark above. + +## Read an Application Timeline with Nsight Systems + +Proton can aggregate kernel time, but it cannot show execution order, gaps between kernels, copies, +or host waits. Use the Nsight Systems timeline to examine those relationships. + +### Capture a Reproducible Report + +The repository's `appendix/nsys_example.py` defines a small three-stage operation. It copies a +$4096\times4096$ BF16 matrix from pinned host memory to the GPU, then runs GEMM and ReLU. Inputs and +outputs are allocated before collection. The core of the script is: + +```python +import torch + + +def run(): + with torch.cuda.nvtx.range("H2D input"): + a.copy_(host_a, non_blocking=True) + with torch.cuda.nvtx.range("BF16 GEMM"): + torch.mm(a, b, out=c) + with torch.cuda.nvtx.range("ReLU"): + torch.clamp_min(c, 0, out=output) + + +def run_once_for_profiler(run, *, warmup_calls): + for _ in range(warmup_calls): + run() + torch.cuda.synchronize() + + cudart = torch.cuda.cudart() + cudart.cudaProfilerStart() + with torch.cuda.nvtx.range("target operation"): + run() + torch.cuda.synchronize() + cudart.cudaProfilerStop() +``` + +Warm-up remains outside the capture. The `target operation` NVTX range gives the invocation a clear +name, and the synchronization inside that range ensures that all three GPU operations finish before +the range closes. `cudaProfilerStart()` and `cudaProfilerStop()` control collection; they are not a +performance timer. + +The following command writes `reports/target-timeline.nsys-rep`: + +```bash +mkdir -p reports +nsys profile \ + --trace=cuda,nvtx \ + --sample=none \ + --cpuctxsw=none \ + --capture-range=cudaProfilerApi \ + --capture-range-end=stop \ + --output=reports/target-timeline \ + --force-overwrite=true \ + python appendix/nsys_example.py --profile-once +``` + +`--trace=cuda,nvtx` records CUDA APIs, GPU activity, and NVTX ranges. Disabling CPU sampling and +context-switch tracing keeps the first report focused on the CUDA timeline. If the report exposes a +long period with no GPU work, collect a separate report with the relevant host-scheduling or OS +runtime tracing enabled. + +Open the report in the GUI, or print the five most useful tables for this capture: + +```bash +nsys-ui reports/target-timeline.nsys-rep + +nsys stats \ + --format=column \ + --timeunit=us \ + --report nvtx_gpu_proj_sum \ + --report nvtx_pushpop_trace \ + --report cuda_gpu_sum \ + --report cuda_kern_exec_trace \ + --report cuda_api_sum \ + reports/target-timeline.nsys-rep +``` + +### Interpret a Real Report + +The following values come from one real capture on an NVIDIA B200 with driver 595.58.03, CUDA 13.0, +PyTorch 2.12.0+cu130, and Nsight Systems 2025.6.3. They demonstrate how to read a report and should +not be treated as reference performance for this workload. + +![A measured Nsight Systems timeline: CPU NVTX ranges and CUDA APIs appear above the H2D copy, GEMM, and ReLU on one GPU stream](../img/nsys_b200_timeline.svg) + +*This figure was redrawn from the timestamps in `cuda_api_trace`, `nvtx_pushpop_trace`, and +`cuda_gpu_trace`. Bar lengths are proportional to the measured durations.* + +`cuda_gpu_sum` reports the three GPU activities: + +| GPU activity | Count | GPU duration | Share of listed GPU time | +|---|---:|---:|---:| +| 32 MiB H2D copy | 1 | 607.230 μs | 85.4% | +| BF16 GEMM | 1 | 93.152 μs | 13.1% | +| ReLU | 1 | 11.072 μs | 1.6% | + +`cuda_kern_exec_trace` correlates each kernel with its launch API and reports API time, positive queue +time, and GPU execution separately. Positive queue time is the interval from API return to a later +kernel start; it has no positive value when the kernel starts earlier. + +| Kernel | API time | Positive queue time | GPU execution | +|---|---:|---:|---:| +| BF16 GEMM | 34.270 μs | 403.214 μs | 93.152 μs | +| ReLU | 11.064 μs | 442.166 μs | 11.072 μs | + +This report supports four concrete observations: + +1. **The H2D copy dominates this region.** It accounts for 85.4% of the three GPU-duration sum. + Looking only at `cuda_gpu_kern_sum` would omit the copy entirely, so this example uses + `cuda_gpu_sum`, which includes both kernels and memory operations. +2. **Queue time is not launch overhead.** GEMM and ReLU wait behind earlier work on the same stream. + Their queue intervals are long because they follow the H2D copy and GEMM, not because their launch + APIs took hundreds of microseconds. +3. **A long synchronization interval usually means that the host is waiting for the GPU.** + `cudaDeviceSynchronize` occupied 384.821 μs on the CPU because GPU work remained unfinished when + the host called it. That number is neither one kernel's duration nor the complete operation + latency. +4. **Intervals at different scopes cannot be added.** The three GPU durations sum to 711.454 μs. + `nvtx_gpu_proj_sum` measures from the first enclosed GPU operation's start to the last one's end, + producing 715.582 μs; the roughly 4.1 μs difference is gaps between activities. The original + `target operation` range on the CPU is 870.561 μs because it also includes dispatch and the final + synchronization wait. + +In a separate unprofiled run, 20 CUDA Event samples of the same operation had a median of 722.816 μs +and a range of 718.400–726.336 μs. The same measurement method can be rerun with: + +```bash +python appendix/nsys_example.py --event-samples 20 +``` + +That unprofiled result is the appropriate performance number to report. The single Nsight Systems +timeline explains where time went; the two numbers need not match exactly. + +These numbers also show why the timing boundary matters. If the production operation truly includes +the H2D copy, reducing or overlapping the transfer is the first place to look. If the application +already holds its input on the GPU, the copy does not belong inside the measured region. Do not +assume that GEMM is the first target merely because it is the main compute kernel. + +Apply the same reading order to other reports: first confirm that the NVTX range excludes warm-up and +initialization; inspect kernels, copies, gaps, and overlap on the GPU streams; then follow correlation +back to launch or synchronization APIs on the host. Durations on different streams cannot simply be +added, and visible overlap proves only that it occurred in this capture. Verify any claimed latency +benefit with the same unprofiled boundary. + +Report scripts vary across Nsight Systems releases. Run `nsys stats --help-reports` to list those +available in the installed version, and record `nsys --version` with the experiment. The +[Nsight Systems User Guide](https://docs.nvidia.com/nsight-systems/UserGuide/index.html) covers the +CLI and GUI; the [Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html) +explains API, queue, and kernel-execution intervals in more detail. + +## Analyze One Kernel with Nsight Compute + +Nsight Systems shows when kernels run; Nsight Compute explains why one selected kernel behaves as it +does. It collects launch configuration, occupancy, compute and memory throughput, scheduler state, +and other hardware metrics. NCU can replay the kernel several times while collecting those metrics, +so it is a diagnostic tool: do not substitute the report's `Duration` for latency measured during a +normal run. + +The timeline above showed that the H2D copy dominates the example operation. The NCU walkthrough +still selects the 93.152 μs BF16 GEMM—not because it is the operation's primary bottleneck, but to +show how the metrics from one launch determine what to inspect next. + +### Collect One Target Kernel + +Reuse the `--profile-once` path from the Nsight Systems example. Warm-up stays outside the capture +range, and only one target operation runs inside it. That operation launches a GEMM followed by +ReLU. The kernel-name filter below selects the GEMM, and `--launch-count 1` collects only the first +matching launch. Kernel replay is appropriate only when the selected kernel can be replayed in +isolation. Inspect a dependent or concurrent multi-kernel region in Nsight Systems before choosing a +different replay mode. + +Start with the `basic` section set: + +```bash +mkdir -p reports +ncu \ + --config-file off \ + --target-processes application-only \ + --profile-from-start off \ + --kernel-name-base function \ + --kernel-name 'regex:.*nvjet_sm100.*' \ + --launch-count 1 \ + --set basic \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --export reports/bf16-gemm-basic \ + --force-overwrite \ + python appendix/nsys_example.py --profile-once +``` + +`--profile-from-start off` makes NCU wait for the profiler API range in the script. The regular +expression then selects the GEMM whose name contains `nvjet_sm100`. Generated kernel names can change +with PyTorch and CUDA versions, so copy the actual name from Nsight Systems before writing a narrower +filter for another program. + +`--set basic` collects launch, occupancy, workload-distribution, and high-level throughput sections. +Cache and clock controls are explicit because they change the profiling conditions. In particular, +`--cache-control all` flushes the GPU caches that NCU can control before every replay iteration. That +helps stabilize counter collection but does not reproduce the hot-cache policy of the headline +benchmark. Section sets and defaults can change between releases, so record `ncu --version` and +inspect `ncu --config-file off --list-sets` on the collection machine. + +If the script cannot use profiler start/stop, rerun the command with `--profile-from-start on` (or +remove `--profile-from-start off`) and use `--launch-skip N --launch-count 1` to select an invocation +after warm-up. `--launch-skip` counts matching kernel launches, so a changed filter or launch order +can select a different invocation. The +[Nsight Compute CLI documentation](https://docs.nvidia.com/nsight-compute/NsightComputeCli/) +describes the kernel and launch filters in detail. + +Open the report in the GUI: + +```bash +ncu-ui reports/bf16-gemm-basic.ncu-rep +``` + +or inspect it in the terminal: + +```bash +ncu --import reports/bf16-gemm-basic.ncu-rep \ + --page details \ + --print-details header \ + --print-metric-name label-name +``` + +`header` is a compact first view. Expand the Work ID/CLC and throughput tables in the GUI, or replace +`header` with `all` to print the complete details used below. + +## Interpret a Real NCU Report + +The following values come from a real capture on the same B200 with Nsight Compute 2026.1. NCU used +nine replay passes to build the `basic` report. The percentages are NCU throughput metrics relative +to the sustained peak of the corresponding hardware subsystem; they are not application FLOPs +divided by the chip's advertised peak. + +| Metric | Measured value | +|---|---:| +| Kernel duration | 95.87 μs | +| Grid / block size | 512 blocks / 256 threads | +| Cluster size | 4 blocks | +| Registers | 255 / thread | +| Dynamic shared memory | 213.28 KB / block | +| Waves per SM | 3.46 | +| Theoretical / achieved occupancy | 12.50% / 8.98% | +| SM compute-throughput metric | 77.34% | +| Memory-throughput metric | 38.51% | +| DRAM / L2 / L1-TEX throughput metrics | 20.42% / 34.60% / 46.93% | + +NCU's 95.87 μs differs slightly from the 93.152 μs captured by Nsight Systems above. The values come +from separate profiling runs, and NCU also changes cache, clock, and replay conditions. This is why +the report's `Duration` cannot replace the headline benchmark. + +Read the table in three passes: verify the selected launch, see how its work covers the GPU, and only +then choose which throughput breakdown to expand. + +### 1. Launch Statistics and Workload Distribution + +The captured name is `nvjet_sm100_tst_128x256_64x6_2x2_2cta_h_bz_NNT`, which matches the GEMM in the +timeline above. It launches 512 blocks of 256 threads and groups four blocks into each cluster. +`Waves Per SM = 3.46` means that the grid requires three full waves and one partial wave. It describes +how the grid covers the GPU over time; it is not occupancy. + +This report also carries a Work ID/Cluster Launch Control warning. Although the nominal launch has +512 CTAs, only 380 were granted. When this warning appears, treat metrics derived from block, warp, +or thread counts cautiously rather than assuming that the nominal launch count is the executed +count. + +### 2. Occupancy + +This kernel uses 255 registers per thread and 213.28 KB of dynamic shared memory per block. Both the +register and shared-memory limits allow only one block to reside on an SM. The resulting theoretical +occupancy is 12.50%, and the measured achieved occupancy is 8.98%. + +Those values show that few warps are resident; they do not establish occupancy as the bottleneck. +This GEMM deliberately uses four-CTA clusters and an asynchronous pipeline. Reducing registers or +shared memory merely to raise occupancy can introduce spills or sacrifice tile reuse and make the +kernel slower. + +NCU also emits rule-based `Est. Speedup` suggestions. They are local upper bounds under simplified +assumptions and are useful as investigation prompts, not as expected speedups from changing the +kernel. + +### 3. Speed of Light + +The basic report shows an SM compute-throughput metric of 77.34% and a memory-throughput metric of +38.51%, with DRAM at only 20.42%. The evidence therefore does not support calling the kernel +DRAM-bound; expanding the compute pipelines is the more useful next step. + +For another kernel, compare the high-level compute and memory metrics against their respective +sustained peaks: + +- high compute and lower memory throughput suggests a compute-pipeline limit; +- high memory and lower compute throughput suggests investigating the memory hierarchy; +- both low suggests underfill, dependency latency, synchronization, imbalance, or too few eligible + warps before it suggests a peak-throughput limit. + +"Memory throughput" is not synonymous with DRAM throughput. Its limiting contributor can be L1, +L2, shared memory, or a memory-instruction pipeline. Expand the breakdown before calling a kernel +DRAM-bound. + +### Expand the report before continuing + +The `basic` report covers steps 1–3. Use its evidence to collect only the sections needed for the +next question: + +| Evidence from the basic report | Add next | +|---|---| +| Register, shared-memory, or resident-block limit | `LaunchStats`, `Occupancy` are already in `basic`; inspect their limit tables before collecting more | +| Compute path appears dominant | `ComputeWorkloadAnalysis` | +| Memory hierarchy appears dominant | `MemoryWorkloadAnalysis`; add `_Chart` for the visual breakdown or `_Tables` for detailed requests and sectors | +| Too few eligible warps or unexplained issue gaps | `SchedulerStats`, then `WarpStateStats` | +| A source or instruction location is required | `SourceCounters` | + +The compute metric is higher in this report, so add `ComputeWorkloadAnalysis` for the same isolated +launch: + +```bash +mkdir -p reports +ncu \ + --config-file off \ + --target-processes application-only \ + --profile-from-start off \ + --kernel-name-base function \ + --kernel-name 'regex:.*nvjet_sm100.*' \ + --launch-count 1 \ + --section ComputeWorkloadAnalysis \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --export reports/bf16-gemm-compute \ + --force-overwrite \ + python appendix/nsys_example.py --profile-once +``` + +The follow-up report contains these values: + +| Pipeline | Throughput metric | +|---|---:| +| TMEM | 77.23% | +| Tensor Core | 77.04% | +| Tensor FP | 76.90% | +| ALU / TMA / FMA | all below 2% | + +That evidence resolves the basic report's aggregate 77.34% compute metric to the Tensor Core and +Tensor Memory path. For another hypothesis, keep the command shape and replace the `--section` lines +rather than accumulating every section in one report. This keeps the report smaller and reduces +replay overhead. + +### 4. Compute and Memory Workload Analysis + +Compute Workload Analysis identifies which execution pipelines are active. Check the Tensor Core, +FMA, ALU, special-function, and relevant asynchronous pipelines rather than inferring Tensor Core +utilization from one aggregate compute percentage. + +Memory Workload Analysis separates DRAM, L2, L1/TEX, shared memory, and local-memory effects. Read +traffic volume together with bandwidth, cache hit rate, and local-memory spill. Detailed sector and +request tables require `MemoryWorkloadAnalysis_Tables`; source-level coalescing and shared-memory +conflict evidence can require `SourceCounters`. A high cache hit rate alone says little when the +traffic volume is small. + +### 5. Scheduler and Warp States + +Scheduler Statistics shows active, eligible, and issued warps. First determine whether schedulers +often have no eligible instruction to issue. Only then use Warp State Statistics to investigate why. +The NCU guide explicitly warns that stalls are not all avoidable and do not automatically limit +performance. + +Common states should be interpreted as clues: + +| State | Useful interpretation | Do not conclude from it alone | +|---|---|---| +| Long Scoreboard | Waiting on a dependency associated with the L1TEX path | Every wait reached DRAM | +| Short Scoreboard | Waiting on an MIO-path dependency, often involving shared memory | A bank conflict definitely exists | +| Barrier | Waiting for a synchronization dependency | The barrier is unnecessary | +| Not Selected | The warp was eligible but another warp issued | The scheduler is starved | +| Math/MIO Throttle | A pipeline or queue is under pressure | Removing arbitrary instructions will improve runtime | + +For warp-specialized kernels, aggregate stall percentages also combine roles with intentionally +different behavior. Relate the result to the producer, MMA, softmax, or writeback role before changing +synchronization. + +### 6. Source and SASS Correlation + +The SASS view and instruction attribution do not require CUDA line information. Correlating generated +CUDA source back to SASS does: the binary needs line information, and NCU must be able to find the +source file. For a TIRx module compiled through NVCC, dump the generated source before compilation: + +```bash +export TVM_CUDA_COMPILE_MODE=nvcc +export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels" +mkdir -p "$TVM_KERNEL_DUMP" +``` + +When `TVM_KERNEL_DUMP` is set, TVM retains the generated files and passes `-lineinfo` to NVCC. Add +`--import-source yes --source-folders "$TVM_KERNEL_DUMP"` to the NCU collection command. Saving +`inspect_source("cuda")` is still useful for manual comparison, but by itself it cannot add line +information to a compiled binary. A Python line may lower to many CUDA or SASS instructions, and an +asynchronous tile primitive may be understandable only in those lower-level views. + +### Advanced: NCU Changes the Experiment + +NCU collection can change the execution conditions: + +- it can replay a kernel to collect counter groups; +- its default cache control can flush GPU caches between replay iterations; +- it can control GPU clocks; +- replay can serialize or otherwise alter concurrent work; +- application replay reruns the entire program and requires deterministic execution and launch + matching; it is not a remedy for a nondeterministic launch order; +- a dependent multi-kernel region may require range replay rather than replaying one kernel in + isolation. + +Record the NCU version, replay mode, cache control, clock control, selected sections, and kernel +filter. Do not compare NCU's `Duration` column directly with an unprofiled hot-cache CUDA Event result. +Do not run Proton and NCU in the same profiling process. + +If NCU reports `ERR_NVGPUCTRPERM`, hardware-counter access is restricted. Follow NVIDIA's +[counter-permission guidance](https://developer.nvidia.com/nvidia-development-tools-solutions-err-nvgpuctrperm-nsightcompute) +or ask the system administrator to enable the required access; do not make running every experiment as +root the default solution. + +## Optional Tool: IKET + +Nsight Systems represents one kernel as a single GPU activity, while NCU aggregates hardware metrics +across the kernel. Neither produces a named timeline for load, compute, and wait phases inside a +warp-specialized kernel. A TIRx kernel with phase annotations can use IKET (In-Kernel Event Tracing) +to show when different warp roles work, wait, or overlap. + +TVM 0.26 integrates IKET for SM90-or-newer CUDA targets and validates a strict set of CUTLASS DSL, +NVRTC, and related tool versions. The instrumentation changes the generated kernel, so its timing is +useful for understanding phase relationships, not for reporting latency. See +[`python/tvm/backend/cuda/iket.py`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/backend/cuda/iket.py) +and the +[NVIDIA IKET guide](https://github.com/NVIDIA/cutlass/blob/v4.6.0/media/docs/pythonDSL/cute_dsl_general/iket_profiling.rst) +for the required versions, annotations, and Perfetto trace workflow. + +## Benchmark Checklist + +Before publishing a benchmark table or pull request, use this checklist to ensure that another person +can reconstruct the measurement: + +| Category | Record | +|---|---| +| Hardware | Exact GPU, number of devices, topology when relevant, clock and power policy | +| Software | Driver, CUDA, framework, compiler, library versions, and source commit | +| Workload | Shapes, dtype, layouts, mask, scale, epilogue, input distribution, batch/sequence details, state and reset policy | +| Correctness | Reference, tolerance, accumulation and output dtype, exceptional-input policy | +| Timing | Timer type, kernel/operator/end-to-end boundary, stream policy, CUDA Graph use, values and units for `warmup`/`repeat`, `rounds`, and raw per-round results | +| Cache | Reused inputs, rotating inputs, explicit flush policy, and whether the policy models the application | +| Statistics | Raw latency unit, median or mean, spread, independent runs, implementation order | +| Baseline | Library and algorithm, workspace, tuning budget, selected configuration | +| Profiling | Proton/IKET/Nsight Systems/NCU versions, kernel filters, IKET ranges and trace format, Nsight Systems capture options and trace, NCU sections and replay/cache/clock controls | + +The final workflow is deliberately circular. A benchmark establishes that a change matters; a profile +suggests why; the next unprofiled benchmark determines whether the explanation led to a real +improvement. diff --git a/appendix/index.md b/appendix/index.md index c0b27807..1bf77244 100644 --- a/appendix/index.md +++ b/appendix/index.md @@ -1,11 +1,12 @@ (chap_appendix)= # Overview -The main text runs through Parts I–IV. The Reference section collects material you may want to consult while reading: +The main text runs through Parts I–IV. The appendices collect material you may want to consult while reading: | Need | Where | |------|-----| | Look up a TIRx language feature | **{ref}`chap_language_reference`** | +| Measure, compare, and profile GPU kernels reproducibly | **{ref}`chap_benchmarking`** | | Compiler internals (the lowering pipeline) | **{ref}`chap_arch`** | | Debug asynchronous GEMM/FA hangs, crashes, wrong results, or slowdowns | **{ref}`chap_warp_spec_debug`** | diff --git a/appendix/nsys_example.py b/appendix/nsys_example.py new file mode 100644 index 00000000..9fe981fb --- /dev/null +++ b/appendix/nsys_example.py @@ -0,0 +1,88 @@ +"""Small multi-stage CUDA workload for the Nsight Systems appendix example.""" + +import argparse +from statistics import median + +import torch + + +def make_workload(size: int): + host_a = torch.randn((size, size), dtype=torch.bfloat16, pin_memory=True) + a = torch.empty_like(host_a, device="cuda") + b = torch.randn((size, size), dtype=torch.bfloat16, device="cuda") + c = torch.empty((size, size), dtype=torch.bfloat16, device="cuda") + output = torch.empty_like(c) + + def run(): + with torch.cuda.nvtx.range("H2D input"): + a.copy_(host_a, non_blocking=True) + with torch.cuda.nvtx.range("BF16 GEMM"): + torch.mm(a, b, out=c) + with torch.cuda.nvtx.range("ReLU"): + torch.clamp_min(c, 0, out=output) + + return run, output + + +def run_once_for_profiler(run, *, warmup_calls: int): + for _ in range(warmup_calls): + run() + torch.cuda.synchronize() + + cudart = torch.cuda.cudart() + cudart.cudaProfilerStart() + with torch.cuda.nvtx.range("target operation"): + run() + torch.cuda.synchronize() + cudart.cudaProfilerStop() + + +def measure_event_us(run, *, warmup_calls: int, samples: int): + for _ in range(warmup_calls): + run() + torch.cuda.synchronize() + + values = [] + for _ in range(samples): + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + start.record() + run() + end.record() + end.synchronize() + values.append(start.elapsed_time(end) * 1e3) + return values + + +def main(): + parser = argparse.ArgumentParser() + mode = parser.add_mutually_exclusive_group() + mode.add_argument("--profile-once", action="store_true") + mode.add_argument("--event-samples", type=int, default=0) + parser.add_argument("--size", type=int, default=4096) + parser.add_argument("--warmup-calls", type=int, default=5) + args = parser.parse_args() + + run, output = make_workload(args.size) + if args.profile_once: + run_once_for_profiler(run, warmup_calls=args.warmup_calls) + elif args.event_samples: + values = measure_event_us( + run, + warmup_calls=args.warmup_calls, + samples=args.event_samples, + ) + print( + f"median={median(values):.3f} us, " + f"min={min(values):.3f} us, max={max(values):.3f} us" + ) + else: + run() + torch.cuda.synchronize() + + if not torch.isfinite(output).all().item(): + raise RuntimeError("workload produced a non-finite output") + + +if __name__ == "__main__": + main() diff --git a/chapter_gemm_advanced/index.md b/chapter_gemm_advanced/index.md index 20982b37..9fc46b52 100644 --- a/chapter_gemm_advanced/index.md +++ b/chapter_gemm_advanced/index.md @@ -861,7 +861,7 @@ def hgemm_v9(M, N, K): ## End-to-End Results -The table below follows the progression from the naive baseline to the warp-specialized cluster kernel, with cuBLAS included as a reference. The measurements use an NVIDIA B200 with `M=N=K=4096`, fp16 inputs, locked clocks, and 1,000 timed iterations per measured version: +The table below follows the progression from the naive baseline to the warp-specialized cluster kernel, with cuBLAS included as a reference. The measurements use an NVIDIA B200 with `M=N=K=4096`, fp16 inputs, locked clocks, and 1,000 timed iterations per measured version. New measurements and reproduction attempts should follow the full protocol in {ref}`chap_benchmarking`. | Step | Technique | Time | Speedup | |------|-----------|------|---------| diff --git a/chapter_gemm_basics/index.md b/chapter_gemm_basics/index.md index 43bfdf92..b423164c 100644 --- a/chapter_gemm_basics/index.md +++ b/chapter_gemm_basics/index.md @@ -320,6 +320,11 @@ tflops = 2 * M * N * K / ms / 1e9 print(f"Performance: {ms:.3f} ms, {tflops:.1f} TFLOPS") ``` +This short timing loop is sufficient for a smoke measurement, but it is not a complete experimental +protocol. For reported results, follow {ref}`chap_benchmarking`: define the timing boundary, collect +multiple samples, state the cache and clock policy, and separate unprofiled latency measurements from +Proton or Nsight Compute analysis. + ### Limits of the Single-Tile Kernel The kernel is correct, but it still has a narrow operating range: diff --git a/chapter_performance/index.md b/chapter_performance/index.md index 906343e2..823703e7 100644 --- a/chapter_performance/index.md +++ b/chapter_performance/index.md @@ -338,3 +338,7 @@ For a memory-bound kernel, focus on reducing data movement and making transfers bandwidth ceiling. For a compute-bound kernel, focus on reducing idle time in the compute units. The roofline model does not produce the final implementation, but it prevents effort from being spent on resources that are not the bottleneck. + +Roofline interpretation starts from a trustworthy measurement. The practical workflow for timing a +kernel, locating its expensive launches with Proton, and testing a hardware hypothesis with Nsight +Compute is collected in {ref}`chap_benchmarking`. diff --git a/img/nsys_b200_timeline.svg b/img/nsys_b200_timeline.svg new file mode 100644 index 00000000..91310f96 --- /dev/null +++ b/img/nsys_b200_timeline.svg @@ -0,0 +1,64 @@ + + + + +One Real Nsight Systems Capture on B200 +4096×4096 BF16: H2D copy → GEMM → ReLU + + +0 μs + +100 μs + +200 μs + +300 μs + +400 μs + +500 μs + +600 μs + +700 μs + +800 μs + +900 μs +Outer NVTX + +Child NVTX ranges + +CUDA APIs + +GPU stream 7 + + +target operation · host NVTX 870.6 μs + +H2D input + +BF16 GEMM + +ReLU + + + + +cudaDeviceSynchronize · 384.8 μs + +memcpy API + +GEMM launch + +ReLU launch + +H2D copy · 607.2 μs + +GEMM +93.2 μs + + +ReLU · 11.1 μs +Time is relative to the outer NVTX-range start; horizontal lengths come from the measured capture. + \ No newline at end of file diff --git a/img/nsys_b200_timeline_zh.svg b/img/nsys_b200_timeline_zh.svg new file mode 100644 index 00000000..3e032218 --- /dev/null +++ b/img/nsys_b200_timeline_zh.svg @@ -0,0 +1,64 @@ + + + + +B200 上的一次真实 Nsight Systems 采集 +4096×4096 BF16：H2D copy → GEMM → ReLU + + +0 μs + +100 μs + +200 μs + +300 μs + +400 μs + +500 μs + +600 μs + +700 μs + +800 μs + +900 μs +外层 NVTX + +子 NVTX ranges + +CUDA APIs + +GPU stream 7 + + +target operation · host NVTX 870.6 μs + +H2D input + +BF16 GEMM + +ReLU + + + + +cudaDeviceSynchronize · 384.8 μs + +memcpy API + +GEMM launch + +ReLU launch + +H2D copy · 607.2 μs + +GEMM +93.2 μs + + +ReLU · 11.1 μs +时间以外层 NVTX range 的起点为 0；横向长度来自真实采集，并非示意比例。 + \ No newline at end of file diff --git a/img/scripts/gen_nsys_b200_timeline.py b/img/scripts/gen_nsys_b200_timeline.py new file mode 100644 index 00000000..8af6f988 --- /dev/null +++ b/img/scripts/gen_nsys_b200_timeline.py @@ -0,0 +1,120 @@ +"""Render the measured Nsight Systems timeline used by the benchmarking appendix.""" + +from html import escape +from pathlib import Path + + +WIDTH = 1500 +HEIGHT = 500 +LEFT = 235 +RIGHT = 1440 +T_MAX = 900.0 +BAR_HEIGHT = 34 + + +def x_pos(time_us: float) -> float: + return LEFT + (RIGHT - LEFT) * time_us / T_MAX + + +def render(*, chinese: bool, output: Path) -> None: + title = ( + "B200 上的一次真实 Nsight Systems 采集" + if chinese + else "One Real Nsight Systems Capture on B200" + ) + subtitle = ( + "4096×4096 BF16：H2D copy → GEMM → ReLU" + if chinese + else "4096×4096 BF16: H2D copy → GEMM → ReLU" + ) + rows = ( + ["外层 NVTX", "子 NVTX ranges", "CUDA APIs", "GPU stream 7"] + if chinese + else ["Outer NVTX", "Child NVTX ranges", "CUDA APIs", "GPU stream 7"] + ) + note = ( + "时间以外层 NVTX range 的起点为 0；横向长度来自真实采集，并非示意比例。" + if chinese + else "Time is relative to the outer NVTX-range start; horizontal lengths come from the measured capture." + ) + + parts = [ + '', + f'', + '', + '', + f'{escape(title)}', + f'{escape(subtitle)}', + ] + + axis_y = 88 + parts.append(f'') + for tick in range(0, 901, 100): + x = x_pos(float(tick)) + parts.append(f'') + parts.append(f'{tick} μs') + + row_y = [125, 210, 295, 380] + for label, y in zip(rows, row_y): + parts.append(f'{escape(label)}') + parts.append(f'') + + def rect(start: float, end: float, y: float, color: str, label: str = "", text_color: str = "#ffffff") -> None: + x = x_pos(start) + width = x_pos(end) - x + parts.append( + f'' + ) + if label: + parts.append( + f'{escape(label)}' + ) + + # Outer and child NVTX ranges, relative to target-operation start. + rect(0.0, 870.561, row_y[0], "#365f9d", "target operation · host NVTX 870.6 μs") + rect(12.553, 174.087, row_y[1], "#72a7d8", "H2D input") + rect(190.849, 359.104, row_y[1], "#598bc2", "BF16 GEMM") + rect(364.257, 413.265, row_y[1], "#86b6df", "ReLU", "#243247") + + # Host CUDA API intervals. + rect(116.267, 150.436, row_y[2], "#e39c45") + rect(319.529, 353.799, row_y[2], "#d7832f") + rect(397.159, 408.223, row_y[2], "#c76d25") + rect(482.407, 867.228, row_y[2], "#b85f46", "cudaDeviceSynchronize · 384.8 μs") + for time_us, label, anchor in [ + (133.351, "memcpy API", "middle"), + (336.664, "GEMM launch", "middle"), + (402.691, "ReLU launch", "start"), + ]: + x = x_pos(time_us) + dx = 12 if anchor == "start" else 0 + parts.append(f'') + parts.append( + f'{escape(label)}' + ) + + # GPU activity on the default stream. + rect(145.879, 753.109, row_y[3], "#3d9b72", "H2D copy · 607.2 μs") + rect(757.013, 850.165, row_y[3], "#6f55b5", "GEMM") + gemm_x = (x_pos(757.013) + x_pos(850.165)) / 2 + parts.append(f'93.2 μs') + rect(850.389, 861.461, row_y[3], "#a574d1") + relu_x = x_pos(855.925) + parts.append(f'') + parts.append(f'ReLU · 11.1 μs') + + parts.append(f'{escape(note)}') + parts.append('') + output.write_text("\n".join(parts), encoding="utf-8") + + +def main() -> None: + image_dir = Path(__file__).resolve().parents[1] + render(chinese=False, output=image_dir / "nsys_b200_timeline.svg") + render(chinese=True, output=image_dir / "nsys_b200_timeline_zh.svg") + + +if __name__ == "__main__": + main() diff --git a/index.md b/index.md index ff398443..6975f80b 100644 --- a/index.md +++ b/index.md @@ -39,8 +39,8 @@ This book is open source. Contributions, corrections, and examples are welcome t TMA pipelining, persistent scheduling, warp specialization, and 2-CTA clusters. - **Part IV, Flash Attention 4.** A complete attention kernel built from the Part III techniques: two MMAs with softmax between them, online-softmax rescaling, causal masking, and GQA. -- **Reference.** TIRx language reference, compiler internals, and a guide to debugging asynchronous - kernels. +- **Appendices.** TIRx language reference, a reproducible GPU benchmarking and profiling workflow, + compiler internals, and a guide to debugging asynchronous kernels. ```{toctree} :caption: Part I, Understanding the GPU @@ -82,11 +82,12 @@ chapter_flash_attention/index ``` ```{toctree} -:caption: Reference +:caption: Appendices :maxdepth: 1 appendix/index tirx_guide/language_reference/index +appendix/benchmarking_gpu_kernels tirx_guide/arch/index appendix/debugging_warp_specialized ``` diff --git a/zh/appendix/benchmarking_gpu_kernels.md b/zh/appendix/benchmarking_gpu_kernels.md new file mode 100644 index 00000000..c5fd1066 --- /dev/null +++ b/zh/appendix/benchmarking_gpu_kernels.md @@ -0,0 +1,616 @@ +(chap_benchmarking)= +# GPU Kernel 性能测量与分析 + +优化 GPU kernel 时，需要分别回答两个问题：运行一次要多久，时间主要花在哪里。Benchmark 负责测前者，profile 用来分析后者。Profiler 会改变程序的执行条件，因此涉及整个 operator 或应用路径的最终性能数字应回到关闭 profiler 后的测量中确认。 + +一次 Python 调用不一定只对应一个 GPU kernel。它可能启动多个 kernels、提交内存拷贝，或者等待 GPU 完成工作。计时前要先确定被测 operation 包含哪些步骤，并让所有实现采用相同的边界。 + +实际操作时，先验证结果并测出无 profiler 的基线，再用 profiler 查找耗时的原因。修改实现后，使用相同的计时方法重新测量；只有基线时间确实缩短，才能说明优化有效。 + +{ref}`chap_performance` 介绍了如何用 roofline 判断性能受计算吞吐还是内存带宽限制。接下来讨论实验方法：如何确定计时范围、选择 warm-up 和 repeat，以及解读 profiler 报告。 + +## 区分性能测量与性能诊断 + +这套流程中的工具各自回答不同问题： + +| 工具 | 主要回答的问题 | +|---|---| +| CUDA Events | 被测区间在 GPU stream 上经过了多长时间？Stream 是按提交顺序执行 GPU 工作的队列。 | +| Proton（Triton 提供的 profiler） | 启动了哪些 GPU kernels、各被调用多少次，哪些 kernel 占用了主要时间？ | +| Nsight Systems | Host、streams、拷贝、kernels 和通信如何在时间线上重叠？ | +| Nsight Compute（`ncu`） | 一个选定的 GPU kernel 主要受哪类硬件资源或等待限制？ | +| IKET（可选） | 一个选定的 kernel 内部，哪些命名阶段或 warp roles 占用了时间？ | + +### Profile 结果的三种常见形式 + +Profile 不是一个数字，也不只有一种报告格式。本章使用的工具会生成三种互补的视图： + +| 视图 | 工具 | 阅读重点 | +|---|---|---| +| 聚合树 | Proton | 比较调用次数、平均时间和总时间，找出主要耗时的 kernels。 | +| 时间线 | Nsight Systems；单个 kernel 内部使用 IKET | 沿横轴阅读不同 tracks 上的事件，检查空隙、重叠和依赖关系。 | +| 单 kernel 指标报告 | Nsight Compute | 查看一次 launch 的配置、利用率、scheduler 状态、内存流量以及 source/SASS 证据。 | + +这些 profile 用来解释时间花在哪里，不能替代正式的性能测量。修改实现后，应关闭 profiler，并使用与基线相同的计时边界重新测量。 + +## 计时前先验证正确性 + +性能计时前，先单独验证正确性： + +1. 构造有代表性的输入，并覆盖相关的边界情况。 +2. 运行被测实现并同步，确保 GPU 已经完成计算。 +3. 使用明确的 tolerance，将结果与 reference 比较。 +4. 如果 kernel 会在已有 output 上累加或原地修改输入，每次验证前都恢复相同的初始状态。 + +Reference 计算和结果比较不属于性能计时。正确性通过后，再开始设计 benchmark；是否把状态重置计入时间，由下一节定义的 operation 边界决定。 + +## 明确计时边界 + +计时前，先写清楚一次被测 operation 包含哪些工作。它可以只包含一个 kernel，也可以包含得到完整结果所需的全部 kernels、内存拷贝和状态重置。编译、输入构造、内存分配或数据格式转换是否属于这次 operation，也要明确说明。不同实现只有在测量相同工作时才能直接比较。 + +范围确定后，再选择计时方法： + +- **CUDA Events** 记录 GPU stream 执行到两个位置时的时间戳，适合测量一个 kernel 或整个 operator 在 device 时间线上的区间。区间内的 kernels、内存拷贝和 stream 空隙都会被计入。对于多 stream operation，所有参与计算的 streams 都必须在 start event 之后开始被测工作，并在记录 end event 前完成汇合。 +- **同步的 wall-clock timer** 从 host 发起调用前开始计时，在调用所需的 GPU 工作全部完成后停止。它还会计入 Python dispatch、CUDA launch 和等待 GPU 完成的时间，适合测量一次调用的端到端 latency。 + +例如，GPU 已经执行 start event，但 host 还没有提交下一个 launch 时，这段 stream 空闲时间仍会落在 CUDA Event 区间内。因此，CUDA Event 区间不一定等于 profiler 中某个 kernel 从开始到结束的执行区间。诊断型 profiler 适合观察 kernel 执行和重叠关系，但其结果不能直接替代相同边界下的无 profiler 计时。 + +## 使用 CUDA Events 测量 GPU 时间 + +CUDA launch 通常是异步的：Python 把工作提交到 CUDA stream 后就可以继续运行，此时 GPU 不一定已经完成。如果只用 CPU 时钟记录这次 Python 调用前后的时间，计时可能在 GPU 完成前就已经停止，得到的主要是 host 提交耗时。测量 GPU stream 上经过的时间时，应使用 CUDA Events；测量从 Python 发起调用到 GPU 完成的完整时间时，则使用后面介绍的同步 wall-clock timer。[PyTorch CUDA semantics 文档](https://docs.pytorch.org/docs/stable/notes/cuda.html#asynchronous-execution)也说明了这种异步行为。 + +先看一份可以直接运行的 CUDA Event benchmark。它在计时前分配矩阵，先执行 warm-up，再测量五轮 FP16 GEMM 并报告中位数： + +```python +from statistics import median + +import torch + + +a = torch.randn((2048, 2048), device="cuda", dtype=torch.float16) +b = torch.randn((2048, 2048), device="cuda", dtype=torch.float16) +c = torch.empty((2048, 2048), device="cuda", dtype=torch.float16) + + +def gemm(): + torch.mm(a, b, out=c) + + +def measure_batch_ms(fn, calls): + """返回连续 calls 次调用的平均 CUDA Event 时间，单位为 ms。""" + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + + start.record() + for _ in range(calls): + fn() + end.record() + end.synchronize() + return start.elapsed_time(end) / calls + + +warmup_calls = 500 +repeat = 100 +rounds = 5 + +for _ in range(warmup_calls): + gemm() +torch.cuda.synchronize() + +samples_ms = [measure_batch_ms(gemm, repeat) for _ in range(rounds)] + +print(f"device: {torch.cuda.get_device_name()}") +print(f"calls per round: {repeat}") +print("round samples (ms):", [round(x, 4) for x in samples_ms]) +print(f"median CUDA Event time: {median(samples_ms):.4f} ms") +``` + +`measure_batch_ms` 在当前 CUDA stream 中记录 start 和 end events，并用两者之间的时间除以调用次数。这样得到的是连续执行时每次 GEMM 的平均 GPU stream 时间。`end.synchronize()` 只是让 CPU 等到这轮 GPU 工作完成，以便读取 Event 结果。 + +这里的 `warmup_calls=500`、`repeat=100` 和 `rounds=5` 都表示调用或测量次数。它们不是通用标准，而是根据这个 GEMM 在 B200 上的实测结果选出的示例值。在一次 10 轮校准中，warm-up 50 次时，第一轮到最后一轮仍从 0.01425 ms 降到 0.01296 ms；增加到 500 次后，变化缩小为 0.01332 ms 到 0.01301 ms。`repeat=100` 的结果也比 `repeat=10` 更稳定。 + +选择其他 workload 的参数时，可以逐步增加 `warmup_calls`，直到前几轮不再持续变快或变慢；再增加 `repeat` 或 `rounds`，直到波动已经满足实验需要。次数也不是越多越好：如果延长实验后整体时间系统性变化，应检查温度、功耗和时钟频率，并先确定实验要表示短时运行还是持续运行。耗时较长的 kernel 通常可以使用更小的次数。 + +这段代码始终复用同一组矩阵，因此后续调用可能从 cache 中读取部分数据，测得的是 warm-cache 场景。发布结果时，还应记录 GPU 型号、软件版本和时钟设置。 + +实际测量本书中的 TIRx kernels 时，不需要为每个 kernel 重新编写 warm-up、重复计时和统计逻辑。TVM 的 [`tvm.tirx.bench.bench`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/bench.py) 已经封装了这些步骤。只需传入一个负责启动被测实现的函数，输入、输出和 workspace 仍在计时前分配。 + +这个 helper 与上例采用不同的 cache 策略：上例连续复用同一组矩阵，而 `bench` 会在每次正式调用前驱逐 L2 cache，再用一对独立的 CUDA Events 测量被测实现。调用方式如下： + +```python +from tvm.tirx.bench import bench + + +# run 是无参数函数；它只使用已经分配好的 tensors 启动被测实现。 +result = bench( + {"tirx": run}, + timer="event", + warmup=25, + repeat=100, + rounds=5, + cooldown_s=1.0, +) + +print(result["impls"]["tirx"]) # 五轮平均值，单位为 us +print(result["round_samples"]["tirx"]) # 每轮结果 +``` + +这里的 `warmup=25` 和 `repeat=100` 表示毫秒预算，不是固定的调用次数。Event timer 会先做一轮短测，再把 25 ms 的 warm-up 预算和 100 ms 的正式测量预算换算成实际次数。短测包含 L2 驱逐和被测调用，因此换算出的次数只是近似值；正式报告的 Event 时间只覆盖被测调用，L2 驱逐发生在 start event 之前。短 kernel 会自动执行更多次，长 kernel 则执行较少次。`rounds=5` 表示完整测量五轮，`cooldown_s=1.0` 表示每轮测量一个实现前暂停一秒。最终结果是五轮的平均值，每轮结果仍保存在 `round_samples` 中。25/100 ms 是 Event timer 的默认预算；五轮测量则是 TIRx-kernels 命令行工具采用的默认设置，`bench` 函数本身默认只运行一轮。 + +这些数值只是默认起点。若各轮结果仍持续漂移，应增加 warm-up 预算；若各轮波动很大，应增加正式测量预算或轮数。所有实现必须使用相同的 timer、预算和轮数，并保留每轮结果，而不是只报告最快的一次。 + +TIRx-kernels 中的 `run_bench` 也调用这个 helper，例如 [`tirx_kernels/attention/flash_attention4.py`](https://github.com/mlc-ai/tirx-kernels/blob/5be39749e7dfd2c4bdae9b4d396f8ec35af07126/tirx_kernels/attention/flash_attention4.py)。省略 `timer` 时，本地 benchmark 默认使用 Proton；需要 CUDA Event 区间时，应像上面一样显式指定 `timer="event"`。两种 timer 的结果含义不同，报告时必须注明。 + +传给 `bench` 的函数会被反复调用。如果 kernel 会累加 output 或原地修改输入，就要在每次调用前恢复相同状态，或者保证每次正式测量都使用一份尚未修改的预分配输入。若恢复操作放在被测函数中，它的时间也属于前面定义的 operation 边界。否则后一次调用面对的已经不是同一个 workload。 + +### 测量一次调用的端到端时间 + +如果关心的是从 Python 发起一次调用到 GPU 完成这次工作所经过的完整时间，可以使用同步的 wall-clock timer。下面继续测量前面已经定义并 warm-up 的 `gemm()`： + +```python +from statistics import median +import time + +import torch + + +def measure_single_call_ms(fn, samples=20): + values = [] + for _ in range(samples): + torch.cuda.synchronize() # 排除此前尚未完成的 GPU 工作 + t0 = time.perf_counter() + fn() + torch.cuda.synchronize() # 等待本次调用的 GPU 工作完成 + values.append((time.perf_counter() - t0) * 1e3) + return values + + +host_samples_ms = measure_single_call_ms(gemm) +print("single-call samples (ms):", [round(x, 4) for x in host_samples_ms]) +print(f"median end-to-end time: {median(host_samples_ms):.4f} ms") +``` + +第一个同步确保计时开始前没有更早的 GPU 工作残留；第二个同步保证计时停止前，这次 GEMM 已经完成。每个 sample 只包含一次调用，因此结果包括 Python 调用、CUDA launch、GPU 执行以及等待完成的开销。前面的 CUDA Event benchmark 测量的则是连续调用时每次 GEMM 的平均 GPU stream 时间。 + +比较多个实现时，所有实现必须使用相同的计时方法和边界。如果同时报告这两种结果，应分别命名为“CUDA Event GPU 时间”和“单次端到端时间”，而不是把使用不同 timer 得到的数字都写成同一种 latency。 + +### 重叠执行时如何计时 + +前面的 GEMM 只在当前 CUDA stream 上运行，因此 start 和 end events 可以直接包住全部工作。如果一个 operator 同时使用多个 streams，仅在当前 stream 记录 events 就不够了：其他 stream 上的工作可能在 start 之前已经开始，也可能在 end 之后仍未完成。 + +要测量整个 operator，可以把 start event 作为所有工作 streams 的共同起点：每个 stream 先等待 start，再开始被测工作，并在完成后各自记录一个 event。最后，当前 stream 等待这些完成 events，再记录 end。这样得到的区间才覆盖从最早开始到全部完成的整个 operation。 + +PDL（Programmatic Dependent Launch）是另一种可能产生重叠的情况。它允许 compute capability 9.0 或更新的 GPU 在同一 stream 中提前启动后一个 kernel：后一个 kernel 可以先完成不依赖前序结果的准备工作，在真正读取这些结果前再等待。这个过程需要显式启用，并遵守相应的 trigger 和 wait 约定；具体 API 见 [CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html)。 + +无论重叠来自多个 streams 还是 PDL，计时原则都相同：用 CUDA Events 包住完整 operation。重叠的 kernels 可能覆盖同一段时间，因此不能把 profiler 中各 kernel 的 duration 直接相加作为 operator latency。实际的执行顺序和重叠情况可以在 Nsight Systems 时间线中查看。PDL 是否产生重叠由运行时决定，程序正确性不能依赖它一定发生。 + +## 固定实验条件 + +前面确定了计时边界和计时器，接下来还要固定会影响结果的实验条件。前面的示例都在输入分配和 warm-up 完成后开始计时，测量的是后续重复调用的性能。首次调用则可能包含 CUDA 初始化、JIT、autotuning 或其他只发生一次的工作。如果研究目标是首次调用或完整应用路径，就应把相应步骤纳入计时边界并单独报告，不能与重复调用的结果混在一起。 + +一次测量不足以说明结果是否稳定。手写 CUDA Event 示例保留五轮结果并报告中位数；`bench` 则报告各轮的平均值，同时把原始结果保存在 `round_samples` 中。无论采用哪种汇总方式，都应保留每轮结果、检查是否存在趋势或异常波动，并明确报告使用的是中位数还是平均值，而不是只挑最快的一轮。比较多个实现时，还可以更换测量顺序后再运行一次，避免某个实现总是在设备较冷或较热时被测量。 + +缓存状态也会改变结果。前面的手写示例反复使用同一组矩阵，属于 warm-cache 测量。TVM 0.26 的 Event 和 Proton timers 则会在每次正式调用前写入一个 256 MiB buffer，以驱逐 L2 中已有的数据；这次写入发生在计时区间之外。两种策略都可以使用，关键是选择符合目标应用的一种，并让所有实现保持一致。`torch.cuda.empty_cache()` 只会释放 PyTorch caching allocator 中未使用的 blocks，不会清空 GPU L2 cache，因此不能用它实现 cold-L2 测量。 + +最后，记录 GPU 型号、driver、CUDA runtime、framework 和 compiler 版本，以及被测 workload 的 dtype 与 shape。还要记录时钟和功耗设置，避免其他进程占用设备，并留意热降频。若锁定时钟，应给出具体数值与命令；只写“fixed clocks”不足以复现实验。 + +相同的 tensor shape 并不代表两个实现可以直接比较。至少要对齐三类条件： + +- **数值语义：** 输入与输出的数据类型、布局、转置方式、对齐要求、累加精度、缩放、mask、epilogue、输出定义和误差阈值； +- **被测范围：** 是否包含 allocation、数据转换、状态重置、辅助 kernels、通信和同步； +- **调优条件：** workspace 上限、是否允许针对每个 shape 自动调优，以及各实现可使用的搜索预算。 + +所有实现还应采用相同的 cache、时钟、warm-up、采样和计时策略。使用库实现作为 baseline 时，需要记录版本、所选算法和 workspace；自动调优可以放在计时区间之外，但搜索预算与最终配置仍应写入实验记录。 + +## 由延迟换算吞吐率 + +吞吐率不是计时器直接测出来的，而是用约定的工作量除以延迟得到的。因此，性能表在给出 TFLOP/s、GB/s 或 tokens/s 时，也应保留原始延迟，并说明工作量如何计算。本书将 GEMM 的工作量记为 $2MNK$ FLOPs；对于 attention 和 fused kernels，则需注明统计的是完整的稠密问题、实际选中的元素，还是 kernel 真正执行的工作。相关公式和 roofline 分析见 {ref}`chap_performance`。 + +## 使用 Proton 找出耗时的 kernel + +前面的 benchmark 只告诉我们整个 operation 用了多长时间。如果它会启动多个 kernels，还需要找出时间具体花在哪些 kernels 上。Proton 可以列出每个 kernel 的调用次数、平均时间和累计时间。 + +Proton 是 Triton 项目提供的 GPU profiler。它记录 CUDA kernel 活动，因此也能看到由 TVM 编译的 TIRx kernels；这些 kernels 并不是由 Triton 编译的。前面介绍的 `bench` 也支持 `timer="proton"`：它会汇总每次调用中的 kernel 执行时间，并返回多次测量的统计结果。如果想知道其中有哪些 kernels、各调用了多少次，就需要单独采集一棵 kernel 树。 + +下面继续使用前面分配好的矩阵，并把 GEMM 和 ReLU 组成一个两-kernel operation。代码先完成 warm-up，只采集后面的 100 次调用，最后在当前目录生成 `operator.hatchet`： + +```python +import torch +import triton.profiler as proton + + +def operation(): + torch.mm(a, b, out=c) + torch.clamp_min(c, 0, out=c) + + +def collect_proton(run, *, warmup_calls, profile_calls): + for _ in range(warmup_calls): + run() + torch.cuda.synchronize() + + session = proton.start("operator", context="shadow", data="tree") + if session is None: + raise RuntimeError("Proton session could not be created") + try: + with proton.scope("target_operation"): + for _ in range(profile_calls): + run() + torch.cuda.synchronize() + finally: + proton.finalize(session) + + +collect_proton(operation, warmup_calls=500, profile_calls=100) +``` + +这里的 `warmup_calls` 和 `profile_calls` 都是调用次数，不是 `bench` 使用的毫秒预算。运行代码需要安装与 TVM 兼容的 Triton，并在实验记录中注明 Triton 版本。 + +先查看文件中有哪些 metrics，再用后两条命令分别打印调用次数、总时间和平均时间： + +```bash +proton-viewer --list operator.hatchet +proton-viewer --metrics time/ms,count --print-sorted operator.hatchet +proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet +``` + +如果 viewer 提示缺少 `pandas` 或 `hatchet`，可在 profiling 环境中运行 `python -m pip install pandas llnl-hatchet`。下面是这段代码在 B200 上的一次实际结果；为了便于阅读，缩短了 kernel 名称： + +```text +target_operation calls avg/us total/ms +├── GEMM kernel 100 14.83 1.483 +└── ReLU kernel 100 4.23 0.423 +``` + +先确认预期的 kernels 和调用次数是否正确，再比较叶节点的平均时间与总时间。这个例子中 GEMM 的总时间最大，因此它是更值得继续使用 Nsight Compute 分析的对象。一个 kernel 即使单次很短，也可能因为调用次数很多而占用大量总时间；父 scope 的平均值则不能当作一次完整 operation 的 latency。 + +这棵树适合寻找耗时的 kernel，不能替代前面的计时结果。存在重叠时，各 kernel 的时间会重复覆盖同一段区间；内存拷贝、同步和 stream 空隙也不一定出现在树中。此外，这段采集会复用同一组矩阵，没有采用前面 TVM timer 的 L2 驱逐策略，因此两处数字不能直接比较。完整 operation 的时间仍由前面的 CUDA Event 或 wall-clock benchmark 给出。 + +## 使用 Nsight Systems 阅读应用时间线 + +Proton 可以汇总各 kernel 的时间，却看不到它们以什么顺序执行，也看不到 kernel 之间的空隙、数据拷贝和 host 等待。分析这些问题时，需要使用 Nsight Systems 的时间线。 + +### 采集一份可复现的报告 + +仓库中的 `appendix/nsys_example.py` 构造了一个简单的三阶段 operation：先把一个 $4096\times4096$ 的 BF16 matrix 从 pinned host memory 复制到 GPU，再执行 GEMM 和 ReLU。输入和输出都在采集前分配。脚本中的核心代码是： + +```python +import torch + + +def run(): + with torch.cuda.nvtx.range("H2D input"): + a.copy_(host_a, non_blocking=True) + with torch.cuda.nvtx.range("BF16 GEMM"): + torch.mm(a, b, out=c) + with torch.cuda.nvtx.range("ReLU"): + torch.clamp_min(c, 0, out=output) + + +def run_once_for_profiler(run, *, warmup_calls): + for _ in range(warmup_calls): + run() + torch.cuda.synchronize() + + cudart = torch.cuda.cudart() + cudart.cudaProfilerStart() + with torch.cuda.nvtx.range("target operation"): + run() + torch.cuda.synchronize() + cudart.cudaProfilerStop() +``` + +Warm-up 位于采集范围之外，`target operation` 则为这次调用提供一个容易识别的 NVTX 名称。同步放在这个 range 内，保证三个 GPU operations 都在 range 结束前完成。`cudaProfilerStart()` 和 `cudaProfilerStop()` 只控制采集范围，不负责性能计时。 + +运行下面的命令会生成 `reports/target-timeline.nsys-rep`： + +```bash +mkdir -p reports +nsys profile \ + --trace=cuda,nvtx \ + --sample=none \ + --cpuctxsw=none \ + --capture-range=cudaProfilerApi \ + --capture-range-end=stop \ + --output=reports/target-timeline \ + --force-overwrite=true \ + python appendix/nsys_example.py --profile-once +``` + +`--trace=cuda,nvtx` 记录 CUDA API、GPU activity 和 NVTX ranges。这里关闭 CPU sampling 与 context-switch tracing，让第一份报告只聚焦 CUDA 时间线。如果报告显示 GPU 长时间没有工作，再单独采集 host scheduling 或 OS runtime 信息。 + +可以在 GUI 中打开报告，也可以从命令行打印这次采集最有用的五张表： + +```bash +nsys-ui reports/target-timeline.nsys-rep + +nsys stats \ + --format=column \ + --timeunit=us \ + --report nvtx_gpu_proj_sum \ + --report nvtx_pushpop_trace \ + --report cuda_gpu_sum \ + --report cuda_kern_exec_trace \ + --report cuda_api_sum \ + reports/target-timeline.nsys-rep +``` + +### 读懂一份真实报告 + +下面的数据来自一次真实采集：NVIDIA B200、driver 595.58.03、CUDA 13.0、PyTorch 2.12.0+cu130 和 Nsight Systems 2025.6.3。数值只用于演示读法，不能作为这个 workload 的性能基准。 + +![一次真实的 Nsight Systems 时间线：CPU 上的 NVTX ranges 和 CUDA APIs 位于上方，实际 H2D、GEMM 与 ReLU 位于同一条 GPU stream 上](../../img/nsys_b200_timeline_zh.svg) + +*时间线根据这份报告中的 `cuda_api_trace`、`nvtx_pushpop_trace` 和 `cuda_gpu_trace` 时间戳重绘。横条长度与实测时长成比例。* + +`cuda_gpu_sum` 给出三项 GPU activity 的时间： + +| GPU activity | 次数 | GPU duration | 占所列 GPU 时长总和的比例 | +|---|---:|---:|---:| +| 32 MiB H2D copy | 1 | 607.230 μs | 85.4% | +| BF16 GEMM | 1 | 93.152 μs | 13.1% | +| ReLU | 1 | 11.072 μs | 1.6% | + +`cuda_kern_exec_trace` 把每个 kernel 与对应的 launch API 关联起来，并分别给出 API time、positive queue time 和 GPU execution。这里的 positive queue time 是 launch API 返回后到 kernel 开始前的等待时间；如果 kernel 更早开始，这一项就没有正值。 + +| Kernel | API time | Positive queue time | GPU execution | +|---|---:|---:|---:| +| BF16 GEMM | 34.270 μs | 403.214 μs | 93.152 μs | +| ReLU | 11.064 μs | 442.166 μs | 11.072 μs | + +这份报告可以读出以下结论： + +1. **主要时间花在 H2D copy。** 它占三项 GPU duration 之和的 85.4%。若只看 `cuda_gpu_kern_sum`，这次 copy 会被完全漏掉；因此这里使用同时包含 kernels 和 memory operations 的 `cuda_gpu_sum`。 +2. **Queue time 不是 launch overhead。** GEMM 与 ReLU 都在同一条 stream 上等待更早提交的工作。它们分别排在 H2D copy 和 GEMM 后面，所以 queue time 远大于各自的 API time。 +3. **较长的同步 API 区间通常表示 host 正在等待 GPU。** `cudaDeviceSynchronize` 在 CPU 上持续了 384.821 μs，说明调用它时仍有 GPU 工作没有完成；这个数字不是某个 kernel 的时间，也不是整个 operation 的 latency。 +4. **不同范围的时间不能相加。** 三项 GPU duration 相加是 711.454 μs。`nvtx_gpu_proj_sum` 以该 NVTX range 中第一项 GPU work 的开始为起点、最后一项的结束为终点，得到 715.582 μs；约 4.1 μs 的差值是 activities 之间的空隙。CPU 上的原始 `target operation` NVTX range 则是 870.561 μs，其中还包含 dispatch 和最后的同步等待。 + +在另一轮关闭 profiler 的测量中，同一 operation 的 20 个 CUDA Event samples 得到 722.816 μs 的 median，范围为 718.400–726.336 μs。可以用下面的命令复现相同的测量方法： + +```bash +python appendix/nsys_example.py --event-samples 20 +``` + +这个无 profiler 结果才适合报告性能；Nsight Systems 的一次时间线用于解释时间花在哪里，两者不要求数值完全相同。 + +这组数据也说明了为什么必须先确定计时边界。如果实际应用中的 operation 确实包含 H2D copy，优化重点应首先考虑减少或重叠传输；如果输入早已位于 GPU，这次 copy 就不应放进被测范围。不能因为 GEMM 是主要的计算 kernel，就默认它是首要优化对象。 + +分析其他报告时也采用同样的顺序：先确认 NVTX range 中没有 warm-up 和初始化，再查看 GPU streams 上的 kernels、copies、空隙与重叠，随后沿 correlation 回到 host 上的 launch 或同步 API。不同 streams 上的 duration 不能直接相加；看见重叠也只说明它在这次采集中发生，是否降低了 latency 仍要通过相同边界的无 profiler 测量验证。 + +Report scripts 会随 Nsight Systems 版本变化。运行 `nsys stats --help-reports` 可以查看当前版本支持的名称，并应在实验记录中保留 `nsys --version`。[Nsight Systems User Guide](https://docs.nvidia.com/nsight-systems/UserGuide/index.html) 介绍了 CLI 和 GUI，[Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html) 则进一步解释 API、queue 和 kernel execution time。 + +## 使用 Nsight Compute 分析单个 kernel + +Nsight Systems 告诉我们各个 kernel 在什么时候运行；选定一个 kernel 后，Nsight Compute 可以继续查看它的启动配置、occupancy、计算与访存吞吐，以及 scheduler 状态。采集这些指标时，NCU 可能多次重放同一个 kernel，因此它适合诊断原因，不应使用报告中的 `Duration` 代替正常运行时测得的 latency。 + +上一节的时间线表明，示例 operation 的主要时间花在 H2D copy。下面选择其中耗时 93.152 μs 的 BF16 GEMM 演示 NCU 的读法；这并不表示 GEMM 是整个 operation 最该优化的部分。 + +### 只采集一次目标 kernel + +继续使用上一节脚本的 `--profile-once` 模式，可以把 warm-up 留在采集范围之外，并且只执行一次目标 operation。这个 operation 会依次启动 GEMM 和 ReLU；下面通过 kernel-name filter 选中 GEMM，并用 `--launch-count 1` 只采集第一个匹配的 launch。命令采用 kernel replay，因此只适合能够独立重放的 kernel。若一段工作包含跨 kernel 依赖或并发，应先查看 Nsight Systems 时间线，再决定是否需要后文介绍的其他 replay mode。 + +先收集 `basic` section set： + +```bash +mkdir -p reports +ncu \ + --config-file off \ + --target-processes application-only \ + --profile-from-start off \ + --kernel-name-base function \ + --kernel-name 'regex:.*nvjet_sm100.*' \ + --launch-count 1 \ + --set basic \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --export reports/bf16-gemm-basic \ + --force-overwrite \ + python appendix/nsys_example.py --profile-once +``` + +`--profile-from-start off` 让 NCU 等待脚本中的 profiler API；正则表达式再从该范围中选出名称包含 `nvjet_sm100` 的 GEMM。生成的完整 kernel 名称会随 PyTorch 和 CUDA 版本变化，因此分析其他程序时，应先从 Nsight Systems 抄下实际名称，再编写更精确的 filter。 + +`--set basic` 采集 launch、occupancy、workload distribution 和高层 throughput sections。Cache 与 clock controls 会改变 profiling 条件，因此命令中将它们显式写出。`--cache-control all` 会在每次 kernel replay 前清理 NCU 能够控制的 GPU caches；这有助于稳定 counter 采集，却不等同于正式 benchmark 的 hot-cache 条件。不同 NCU release 的 section sets 和 defaults 可能变化；采集时应记录 `ncu --version`，并在目标机器上运行 `ncu --config-file off --list-sets` 查看实际配置。 + +如果脚本不能调用 profiler start/stop，应把命令改为 `--profile-from-start on`（或删除 `--profile-from-start off`），再用 `--launch-skip N --launch-count 1` 选择 warm-up 后的某次调用。`--launch-skip` 只统计匹配 kernel 的 launches；filter 或 launch 顺序变化时，它可能选中另一个实例。[Nsight Compute CLI 文档](https://docs.nvidia.com/nsight-compute/NsightComputeCli/) +详细说明了 kernel 与 launch filters。 + +在 GUI 中打开报告： + +```bash +ncu-ui reports/bf16-gemm-basic.ncu-rep +``` + +也可以直接在 terminal 中查看： + +```bash +ncu --import reports/bf16-gemm-basic.ncu-rep \ + --page details \ + --print-details header \ + --print-metric-name label-name +``` + +`header` 适合先查看主要指标；后文使用的 Work ID/CLC 明细和完整 throughput breakdown 可在 GUI 中展开，或把命令中的 `header` 改为 `all` 后打印。 + +## 读懂一份真实的 NCU 报告 + +下面的数据来自同一台 B200 上的一次真实采集，使用 Nsight Compute 2026.1。NCU 用 9 个 replay passes 完成了 `basic` 报告。表中的百分比表示相应 throughput 指标占硬件子系统可持续峰值的比例，不是直接用应用 FLOPs 除以芯片标称峰值得到的利用率。 + +| 指标 | 实测值 | +|---|---:| +| Kernel duration | 95.87 μs | +| Grid / block size | 512 blocks / 256 threads | +| Cluster size | 4 blocks | +| Registers | 255 / thread | +| Dynamic shared memory | 213.28 KB / block | +| Waves per SM | 3.46 | +| Theoretical / achieved occupancy | 12.50% / 8.98% | +| SM compute throughput 指标 | 77.34% | +| Memory throughput 指标 | 38.51% | +| DRAM / L2 / L1-TEX throughput 指标 | 20.42% / 34.60% / 46.93% | + +NCU 中的 95.87 μs 与上一节 Nsight Systems 采集到的 93.152 μs 不完全相同。两者来自不同的 profiling run，而且 NCU 还改变了 cache、clock 和 replay 条件；这种差异正说明 profile 中的 `Duration` 不能替代正式 benchmark。 + +阅读这张表时，先确认捕获对象，再看工作怎样铺满 GPU，最后才判断应展开哪类指标。 + +### 1. Launch Statistics 与 Workload Distribution + +报告中的名称是 `nvjet_sm100_tst_128x256_64x6_2x2_2cta_h_bz_NNT`，与上一节时间线中的 GEMM 一致。它以 512 个 blocks、每 block 256 个 threads 启动，并把 4 个 blocks 组成一个 cluster。`Waves Per SM = 3.46` 表示整张 grid 需要三个完整 waves，再加一个不完整 wave 才能执行完；它描述的是 grid 在时间上覆盖 GPU 的方式，不是 occupancy。 + +这次报告还给出了 Work ID/Cluster Launch Control 警告：名义上启动 512 个 CTAs，报告只记录到 380 个获准执行的 CTAs。只要出现这种警告，依赖 block、warp 或 thread 数量的指标都要谨慎解释，不能把名义 launch 数量直接当作实际执行数量。 + +### 2. Occupancy + +这个 kernel 每个 thread 使用 255 个 registers，每个 block 使用 213.28 KB dynamic shared memory。报告中的 register limit 和 shared-memory limit 都只允许每个 SM 驻留一个 block，因此 theoretical occupancy 为 12.50%，实际采集到 8.98%。 + +这只能说明同时驻留的 warps 较少，不能单独证明 occupancy 是性能瓶颈。这个 GEMM 本来就采用 4-CTA clusters 和异步流水线；如果只为提高 occupancy 而减少 registers 或 shared memory，可能引入 spills、减少 tile reuse，反而变慢。 + +NCU 还会显示 `Est. Speedup` 等规则生成的提示。它们是在若干简化假设下估算的局部上限，用来提示值得调查的方向，不是修改 kernel 后可以期待的实际加速比。 + +### 3. Speed of Light + +这次 basic 报告中的 SM compute throughput 指标为 77.34%，memory throughput 指标为 38.51%，其中 DRAM 只有 20.42%。因此，现有证据不支持把它称为 DRAM-bound；下一步更合理的是展开 compute pipelines。 + +对其他 kernel，仍可先比较 compute 和 memory throughput 相对于各自 sustained peak 的比例： + +- Compute 高而 memory 较低，说明应继续检查 compute pipelines； +- Memory 高而 compute 较低，说明应继续检查 memory hierarchy； +- 两者都低，通常应先检查 underfill、dependency latency、synchronization、imbalance 或缺少 + eligible warps，而不是直接归因于 peak throughput。 + +“Memory throughput 高”不等于“DRAM-bound”。限制项也可能来自 L1、L2、shared memory 或 +memory-instruction pipeline。展开 breakdown 后才能判断。 + +### 继续阅读前，先按需扩展报告 + +`basic` 报告只覆盖前 3 步。根据其中的线索，只采集回答下一个问题所需的 sections： + +| `basic` 报告中的线索 | 下一步加入 | +|---|---| +| Registers、shared memory 或 resident blocks 构成限制 | `LaunchStats` 和 `Occupancy` 已包含在 `basic` 中；先阅读其中的 limit tables | +| Compute path 更可能构成限制 | `ComputeWorkloadAnalysis` | +| Memory hierarchy 更可能构成限制 | `MemoryWorkloadAnalysis`；用 `_Chart` 查看图形 breakdown，用 `_Tables` 查看详细 requests 与 sectors | +| Eligible warps 太少，或指令发射存在无法解释的空隙 | `SchedulerStats`，再根据结果加入 `WarpStateStats` | +| 需要定位到 source 或 instruction | `SourceCounters` | + +这次报告中的 compute 指标更高，因此对同一次隔离 launch 加入 `ComputeWorkloadAnalysis`： + +```bash +mkdir -p reports +ncu \ + --config-file off \ + --target-processes application-only \ + --profile-from-start off \ + --kernel-name-base function \ + --kernel-name 'regex:.*nvjet_sm100.*' \ + --launch-count 1 \ + --section ComputeWorkloadAnalysis \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --export reports/bf16-gemm-compute \ + --force-overwrite \ + python appendix/nsys_example.py --profile-once +``` + +这份进一步采集的报告给出： + +| Pipeline | Throughput 指标 | +|---|---:| +| TMEM | 77.23% | +| Tensor Core | 77.04% | +| Tensor FP | 76.90% | +| ALU / TMA / FMA | 均低于 2% | + +这些数据才把 basic 报告中笼统的 77.34% compute 指标落实到 Tensor Core 和 Tensor Memory 路径上。分析其他假设时沿用同一命令结构，并替换其中的 `--section` 行，不要把所有 sections 逐次累加到一份报告中。这样可以缩小报告并减少 replay overhead。 + +### 4. Compute 与 Memory Workload Analysis + +Compute Workload Analysis 会显示哪些 execution pipelines 正在工作。应分别检查 Tensor Core、 +FMA、ALU、special-function，以及相关的 asynchronous pipelines，不要从一个 aggregate compute +百分比推断 Tensor Core utilization。 + +Memory Workload Analysis 会区分 DRAM、L2、L1/TEX、shared memory 和 local-memory effects。将实际 traffic volume 与 bandwidth、cache hit rate 和 local-memory spill 放在一起看。更详细的 +sector 与 request tables 需要 `MemoryWorkloadAnalysis_Tables`；source-level coalescing 和 +shared-memory conflict 证据可能需要 `SourceCounters`。若 traffic 很小，单独一个高 cache hit +rate 并不能说明性能良好。 + +### 5. Scheduler 与 Warp States + +Scheduler Statistics 展示 active、eligible 和 issued warps。首先判断 scheduler 是否经常没有 +eligible instruction 可以发出；只有这时，才需要使用 Warp State Statistics 深挖原因。NCU +文档明确提醒，并非所有 stalls 都可避免，它们也不会自动成为性能瓶颈。 + +常见状态只应作为线索： + +| 状态 | 可以帮助判断 | 不能单独推出 | +|---|---|---| +| Long Scoreboard | 正在等待与 L1TEX 路径相关的数据依赖 | 每次等待都访问了 DRAM | +| Short Scoreboard | 正在等待 MIO 路径依赖，常见于 shared memory | 一定存在 bank conflict | +| Barrier | Warp 正在等待 synchronization dependency | 这个 barrier 没有必要 | +| Not Selected | Warp 已 eligible，但 scheduler 发出了另一个 warp | Scheduler 缺少可执行工作 | +| Math/MIO Throttle | 某个 pipeline 或 queue 承受较高压力 | 随意删掉几条指令就会变快 | + +对于 warp-specialized kernels，aggregate stall percentage 还混合了行为刻意不同的 roles。修改 +synchronization 前，应先将结果对应到 producer、MMA、softmax 或 writeback role。 + +### 6. Source 与 SASS 对应关系 + +SASS 视图和 instruction attribution 不依赖 CUDA line information；要把生成的 CUDA source +对应到 SASS，则 binary 必须包含 line information，而且 NCU 必须能找到 source 文件。通过 NVCC +编译 TIRx module 时，可以在编译前设置： + +```bash +export TVM_CUDA_COMPILE_MODE=nvcc +export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels" +mkdir -p "$TVM_KERNEL_DUMP" +``` + +设置 `TVM_KERNEL_DUMP` 后，TVM 会保留生成文件，并在 NVCC 编译时加入 `-lineinfo`。NCU 采集命令还要加入 `--import-source yes --source-folders "$TVM_KERNEL_DUMP"`。保存 +`inspect_source("cuda")` 仍便于手工对照，但它本身不能给已经编译的 binary 补上 line information。一个 Python line 可能 lower 成多条 CUDA 或 SASS instructions，异步 tile primitive 也可能只能在这些 lower-level views 中看清。 + +### 高级情况：NCU 会改变实验条件 + +NCU 采集会改变执行条件： + +- 它可能重放 kernel，以收集不同 counter groups； +- 默认 cache control 可能在 replay iterations 之间清理 GPU caches； +- 它可以控制 GPU clocks； +- Replay 可能串行化或改变原本并发的工作； +- Application replay 会重新运行整个程序，并要求各次运行的执行过程和 launch matching 具有确定性；它不能解决不确定的 launch 顺序； +- 具有跨 kernel 依赖的一段工作可能需要 range replay，而不适合只隔离重放一个 kernel。 + +报告中应记录 NCU version、replay mode、cache control、clock control、所选 sections 与 kernel +filter。不要将 NCU 的 `Duration` 直接与无 profiler 的 hot-cache CUDA Event 结果比较，也不要在同一个 profiling 进程中同时启动 Proton 和 NCU。 + +如果 NCU 报告 `ERR_NVGPUCTRPERM`，说明 hardware-counter access 受到限制。应按照 NVIDIA 的 +[counter permission 指南](https://developer.nvidia.com/nvidia-development-tools-solutions-err-nvgpuctrperm-nsightcompute) +配置权限，或请系统管理员开放所需访问；不应把所有实验长期使用 root 运行作为默认方案。 + +## 可选工具：IKET + +在 Nsight Systems 时间线中，一个 kernel 只显示为完整的 GPU 执行区间；NCU 给出的指标也覆盖整个 kernel。对于已经加入阶段标记的 warp-specialized TIRx kernel，可以使用 IKET（In-Kernel Event Tracing）查看不同 warp 分工在何时工作、等待或发生重叠。 + +TVM 0.26 已为 TIRx 接入 IKET，目前要求 SM90 或更新的 CUDA target，并对 CUTLASS DSL、NVRTC 等工具版本有严格要求。IKET 会在 kernel 中加入记录代码，因此得到的时间只适合分析阶段关系，不能作为正式 latency。具体的版本要求、标记方法和 Perfetto trace 生成流程见 [`python/tvm/backend/cuda/iket.py`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/backend/cuda/iket.py) 和 [NVIDIA IKET 文档](https://github.com/NVIDIA/cutlass/blob/v4.6.0/media/docs/pythonDSL/cute_dsl_general/iket_profiling.rst)。 + +## 性能实验检查清单 + +发布 benchmark table 或 pull request 前，可以用下面的清单确认其他人能够重建这次测量： + +| 类别 | 需要记录的内容 | +|---|---| +| Hardware | 准确的 GPU、设备数量、相关 topology、clock 与 power policy | +| Software | Driver、CUDA、framework、compiler、library versions 与 source commit | +| Workload | Shapes、dtype、layouts、mask、scale、epilogue、输入分布、batch/sequence 信息、状态与重置策略 | +| Correctness | Reference、tolerance、accumulation/output dtype、异常输入策略 | +| Timing | Timer 类型、kernel/operator/end-to-end 边界、stream policy、CUDA Graph、`warmup`/`repeat` 的数值与单位、`rounds` 以及每轮原始结果 | +| Cache | 是否复用输入、是否轮换输入、显式 flush policy，以及该策略是否符合应用 | +| Statistics | 原始 latency 单位、median/mean、spread、独立 runs、实现顺序 | +| Baseline | Library 与 algorithm、workspace、tuning budget、最终 configuration | +| Profiling | Proton/IKET/Nsight Systems/NCU versions、kernel filters、IKET ranges 与 trace 格式、Nsight Systems 采集选项与 trace、NCU sections 与 replay/cache/clock controls | + +整套流程是一个循环：benchmark 证明变化确实影响性能；profile 提供原因线索；下一次无 profiler +benchmark 再判断这个解释是否带来了真实改进。 diff --git a/zh/appendix/index.md b/zh/appendix/index.md index 0060cbce..36e1c7b0 100644 --- a/zh/appendix/index.md +++ b/zh/appendix/index.md @@ -1,11 +1,12 @@ (chap_appendix)= # 概览 -本书的主线内容位于第一至第四部分。阅读过程中需要查询具体细节时，可以使用下面的参考资料： +本书的主线内容位于第一至第四部分。附录收录了阅读过程中可能需要查询的补充内容： | 需要查询的内容 | 对应页面 | |---|---| | TIRx 语言特性的准确写法和语义 | **{ref}`chap_language_reference`** | +| 可复现地测量、比较和分析 GPU kernel 性能 | **{ref}`chap_benchmarking`** | | 编译器内部机制与 lowering 流程 | **{ref}`chap_arch`** | | 排查异步 GEMM 或 Flash Attention kernel 的卡死、崩溃、错误结果和性能下降 | **{ref}`chap_warp_spec_debug`** | diff --git a/zh/chapter_gemm_advanced/index.md b/zh/chapter_gemm_advanced/index.md index 5ecdeed6..23632f00 100644 --- a/zh/chapter_gemm_advanced/index.md +++ b/zh/chapter_gemm_advanced/index.md @@ -861,7 +861,7 @@ def hgemm_v9(M, N, K): ## 完整优化结果 -下表列出从朴素 baseline 到 warp-specialized cluster kernel 的各个阶段，并给出 cuBLAS 作为参考。测试使用 NVIDIA B200、`M=N=K=4096`、fp16 和固定 clocks，每个版本计时 1000 次： +下表列出从朴素 baseline 到 warp-specialized cluster kernel 的各个阶段，并给出 cuBLAS 作为参考。测试使用 NVIDIA B200、`M=N=K=4096`、fp16 和固定 clocks，每个版本计时 1000 次。新增测量或尝试复现这组结果时，应遵循 {ref}`chap_benchmarking` 中的完整协议。 | 步骤 | 优化方法 | 时间 | 相对第 1 步的累计加速比 | |------|----------|------|--------| diff --git a/zh/chapter_gemm_basics/index.md b/zh/chapter_gemm_basics/index.md index 6ef04e41..e294f285 100644 --- a/zh/chapter_gemm_basics/index.md +++ b/zh/chapter_gemm_basics/index.md @@ -318,6 +318,10 @@ tflops = 2 * M * N * K / ms / 1e9 print(f"Performance: {ms:.3f} ms, {tflops:.1f} TFLOPS") ``` +这段计时循环适合快速确认数量级，但还不是完整的实验协议。需要报告性能结果时，请遵循 +{ref}`chap_benchmarking`：明确计时边界，采集多组样本，说明 cache 与 clock 策略，并将 +无 profiler 的 latency 测量和 Proton、Nsight Compute 分析分开运行。 + ### 单 Tile Kernel 的限制 这个 kernel 已经能够算对，但适用范围很窄。当前仍有以下限制： diff --git a/zh/chapter_performance/index.md b/zh/chapter_performance/index.md index be9b5d6f..087e370b 100644 --- a/zh/chapter_performance/index.md +++ b/zh/chapter_performance/index.md @@ -258,3 +258,6 @@ SM 占用率受 registers、shared memory、warp slots 和 CTA slots 的限制 3. 检查实际实现离对应上限还有多远，并优化真正处于瓶颈的资源。 对于 memory-bound kernel，重点是减少数据搬运，并让传输速度尽可能接近带宽上限；对于 compute-bound kernel，重点是减少计算单元的等待时间。Roofline 模型不能直接给出最终实现，但可以避免在不构成瓶颈的部分反复调参。 + +Roofline 分析需要从可靠的测量结果出发。怎样测量 kernel 时间、怎样用 Proton 找到主要耗时的 +launch，以及怎样用 Nsight Compute 检验硬件层面的判断，统一整理在 {ref}`chap_benchmarking` 中。 diff --git a/zh/index.md b/zh/index.md index b2567331..63337fd0 100644 --- a/zh/index.md +++ b/zh/index.md @@ -17,7 +17,7 @@ - **第二部分：TIRx 概览。** 这一部分介绍 TIRx 的核心组成部分，为理解后续章节中的代码示例做准备。 - **第三部分：GEMM：从 Tiled 到 SOTA。** 这一部分完整讲解如何优化一个 tiled GEMM，并逐步加入 TMA pipelining、persistent scheduling、warp specialization 和 2-CTA cluster。 - **第四部分：Flash Attention 4。** 这一部分基于第三部分的技术构建完整的 attention kernel：两个 MMA，中间插入 softmax，并包含 online-softmax rescaling、causal mask 和 GQA。 -- **参考资料。** TIRx 语言参考、编译器内部机制，以及异步 kernel 调试指南。 +- **附录。** TIRx 语言参考、可复现的 GPU 性能测量与分析流程、编译器内部机制，以及异步 kernel 调试指南。 ```{toctree} :caption: 第一部分：理解 GPU @@ -59,11 +59,12 @@ chapter_flash_attention/index ``` ```{toctree} -:caption: 参考资料 +:caption: 附录 :maxdepth: 1 appendix/index tirx_guide/language_reference/index +appendix/benchmarking_gpu_kernels tirx_guide/arch/index appendix/debugging_warp_specialized ``` From f88aaf372ec34e3076cfe9d51e32c6514cb0adf9 Mon Sep 17 00:00:00 2001 From: tlopex <820958424@qq.com> Date: Mon, 17 Aug 2026 17:17:20 -0400 Subject: [PATCH 2/4] Refine GPU profiling workflow and examples --- appendix/benchmarking_gpu_kernels.md | 131 +++++++++++------- ...vg => nsys_b200_timeline_zh_en_tracks.svg} | 4 +- img/scripts/gen_nsys_b200_timeline.py | 8 +- zh/appendix/benchmarking_gpu_kernels.md | 76 ++++++---- 4 files changed, 132 insertions(+), 87 deletions(-) rename img/{nsys_b200_timeline_zh.svg => nsys_b200_timeline_zh_en_tracks.svg} (98%) diff --git a/appendix/benchmarking_gpu_kernels.md b/appendix/benchmarking_gpu_kernels.md index 942833b1..a1016514 100644 --- a/appendix/benchmarking_gpu_kernels.md +++ b/appendix/benchmarking_gpu_kernels.md @@ -30,7 +30,7 @@ The tools used in this workflow have different jobs: | Nsight Compute (`ncu`) | Why does one selected GPU kernel spend its cycles the way it does? | | IKET (optional) | Which named phases or warp roles consume time inside one selected kernel? | -### Three Common Profile Views +### Three Profile Views A profile is not a single number or a single report format. The tools in this chapter produce three complementary views: @@ -336,15 +336,14 @@ The preceding benchmark tells us how long the complete operation takes. If it la we still need to determine where that time is spent. Proton reports each kernel's call count, average time, and cumulative time. -Proton is a GPU profiler provided by the Triton project. It records CUDA kernel activity and can -therefore see TIRx kernels compiled by TVM; those kernels are not compiled by Triton. The `bench` -helper introduced above also supports `timer="proton"`: it aggregates kernel execution times for each -invocation and returns a statistic across the measured invocations. To see which kernels ran and how -often, collect a kernel tree in a separate Proton session. +Proton is a GPU profiler provided by the Triton project. It observes CUDA kernel activity and can +therefore analyze TIRx kernels compiled by TVM. The `bench(timer="proton")` helper introduced above +returns an aggregate kernel-time result. Here we create a separate Proton session to inspect each +kernel's call count and execution time. -The following example reuses the matrices allocated above and combines GEMM with ReLU into a -two-kernel operation. It finishes warm-up, collects the next 100 calls, and writes `operator.hatchet` -in the current directory: +The following example reuses the matrices allocated above and combines GEMM with ReLU in one +operation. It finishes warm-up, collects the next 100 calls, and writes `operator.hatchet` in the +current directory: ```python import torch @@ -387,9 +386,13 @@ proton-viewer --metrics time/ms,count --print-sorted operator.hatchet proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet ``` -If the viewer reports that `pandas` or `hatchet` is missing, install the optional packages with -`python -m pip install pandas llnl-hatchet`. The following is one real B200 result from this code; -kernel names are shortened for readability: +If `proton-viewer` reports missing optional dependencies, install them with: + +```bash +python -m pip install pandas llnl-hatchet +``` + +The following is one real B200 result from this code; kernel names are shortened for readability: ```text target_operation calls avg/us total/ms @@ -397,28 +400,32 @@ target_operation calls avg/us total/ms └── ReLU kernel 100 4.23 0.423 ``` -First confirm that the expected kernels and call counts appear, then compare the leaf averages and -totals. GEMM has the largest total in this example, making it the better candidate for further -analysis with Nsight Compute. A short kernel can still accumulate substantial total time when it is -launched often. An average shown for a parent scope, however, is not the latency of one complete +First confirm that all expected kernels appear with the correct call counts, then compare their +average and cumulative times. GEMM has the largest cumulative time in this example, so it should be +the first kernel examined with Nsight Compute. Also watch for short kernels that are launched often: +their individual calls may be inexpensive while their cumulative cost is substantial. + +Proton includes only the captured kernel times, not memory copies, synchronization, or stream gaps. +When kernels overlap, summing their durations also counts the overlapping interval more than once. +These data therefore identify kernels for further analysis; they are not the latency of the complete operation. -This tree is useful for finding expensive kernels, but it does not replace the benchmark above. -Overlapping kernels cover some of the same elapsed interval, while copies, synchronization, and stream -gaps may be absent from the tree. This capture also reuses the same matrices instead of reproducing the -L2-eviction policy of the earlier TVM timer, so the two sets of numbers are not directly comparable. -Obtain complete-operation time from the CUDA Event or wall-clock benchmark above. +This capture repeatedly uses the same matrices, whereas the earlier TVM timer evicts L2 before each +measurement, so the two experiments also have different cache conditions. Measure complete-operation +latency with CUDA Events or a synchronized wall-clock timer. -## Read an Application Timeline with Nsight Systems +## Analyze an Application Timeline with Nsight Systems Proton can aggregate kernel time, but it cannot show execution order, gaps between kernels, copies, or host waits. Use the Nsight Systems timeline to examine those relationships. -### Capture a Reproducible Report +### Capture the Target Operation Timeline -The repository's `appendix/nsys_example.py` defines a small three-stage operation. It copies a -$4096\times4096$ BF16 matrix from pinned host memory to the GPU, then runs GEMM and ReLU. Inputs and -outputs are allocated before collection. The core of the script is: +The following example shows how to restrict Nsight Systems collection to one target operation. In +`appendix/nsys_example.py`, that operation performs three steps in sequence: it copies a +$4096\times4096$ BF16 matrix from pinned host memory to the GPU in a host-to-device (H2D) copy, runs +GEMM, and applies ReLU. All required tensors are allocated before collection, so the report focuses +on these three steps. The core of the script is: ```python import torch @@ -446,12 +453,13 @@ def run_once_for_profiler(run, *, warmup_calls): cudart.cudaProfilerStop() ``` -Warm-up remains outside the capture. The `target operation` NVTX range gives the invocation a clear -name, and the synchronization inside that range ensures that all three GPU operations finish before -the range closes. `cudaProfilerStart()` and `cudaProfilerStop()` control collection; they are not a -performance timer. +`run_once_for_profiler` completes warm-up before the profiler starts and waits for the warm-up work on +the GPU to finish. It then calls `cudaProfilerStart()` and labels the measured invocation with the +`target operation` NVTX range, making it easy to locate in the timeline. The synchronization inside +the range ensures that all three GPU operations finish before `cudaProfilerStop()`. These profiler +APIs define the collection range; they do not measure performance. -The following command writes `reports/target-timeline.nsys-rep`: +The following command runs the script and writes `reports/target-timeline.nsys-rep`: ```bash mkdir -p reports @@ -466,16 +474,39 @@ nsys profile \ python appendix/nsys_example.py --profile-once ``` -`--trace=cuda,nvtx` records CUDA APIs, GPU activity, and NVTX ranges. Disabling CPU sampling and -context-switch tracing keeps the first report focused on the CUDA timeline. If the report exposes a -long period with no GPU work, collect a separate report with the relevant host-scheduling or OS -runtime tracing enabled. +`--capture-range=cudaProfilerApi` restricts collection to the interval between +`cudaProfilerStart()` and `cudaProfilerStop()`. `--trace=cuda,nvtx` records CUDA APIs, GPU activity, +and NVTX ranges. CPU sampling and context-switch tracing are disabled here to keep the report focused +on the CUDA timeline. If that timeline contains a long GPU gap, collect a separate report with the +relevant host-scheduling or OS runtime tracing enabled. -Open the report in the GUI, or print the five most useful tables for this capture: +Once the report has been generated, open its timeline in the Nsight Systems GUI: ```bash nsys-ui reports/target-timeline.nsys-rep +``` + +### Copy, Queue, and Execution Time in the Timeline + +This example uses PyTorch to construct a copy-plus-multiple-kernel workload with little setup. The +same timeline-reading method applies to a TIRx operation. +The following real report illustrates how to interpret Nsight Systems results. It was collected on an +NVIDIA B200 with NVIDIA driver 595.58.03, CUDA 13.0, PyTorch 2.12.0+cu130, and Nsight Systems 2025.6.3. + +![A measured Nsight Systems timeline: CPU NVTX ranges and CUDA APIs appear above the H2D copy, GEMM, and ReLU on one GPU stream](../img/nsys_b200_timeline.svg) + +*This figure was redrawn from an actual capture. Each bar is scaled to the measured duration of the +corresponding activity.* + +The `7` in `GPU stream 7` is the stream identifier shown by Nsight Systems for this capture. It does +not mean the seventh execution stage and may differ in another run. The H2D copy, GEMM, and ReLU share +that stream and therefore execute in submission order. + +In addition to viewing the timeline in the GUI, use `nsys stats` to extract the timing data used +below from the same report: + +```bash nsys stats \ --format=column \ --timeunit=us \ @@ -487,16 +518,12 @@ nsys stats \ reports/target-timeline.nsys-rep ``` -### Interpret a Real Report - -The following values come from one real capture on an NVIDIA B200 with driver 595.58.03, CUDA 13.0, -PyTorch 2.12.0+cu130, and Nsight Systems 2025.6.3. They demonstrate how to read a report and should -not be treated as reference performance for this workload. - -![A measured Nsight Systems timeline: CPU NVTX ranges and CUDA APIs appear above the H2D copy, GEMM, and ReLU on one GPU stream](../img/nsys_b200_timeline.svg) - -*This figure was redrawn from the timestamps in `cuda_api_trace`, `nvtx_pushpop_trace`, and -`cuda_gpu_trace`. Bar lengths are proportional to the measured durations.* +The names following `--report` refer to summaries built into Nsight Systems. +`nvtx_gpu_proj_sum` and `nvtx_pushpop_trace` report the GPU projection of an NVTX range and its +host-side range records, respectively. `cuda_gpu_sum` summarizes kernels and CUDA memory operations; +`cuda_kern_exec_trace` correlates host launch APIs with their GPU kernels; and `cuda_api_sum` +summarizes host-side CUDA API calls. Run `nsys stats --help-reports` to list the names and definitions +available in the installed version. `cuda_gpu_sum` reports the three GPU activities: @@ -560,7 +587,7 @@ available in the installed version, and record `nsys --version` with the experim CLI and GUI; the [Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html) explains API, queue, and kernel-execution intervals in more detail. -## Analyze One Kernel with Nsight Compute +## Collect an Nsight Compute Report for One Kernel Nsight Systems shows when kernels run; Nsight Compute explains why one selected kernel behaves as it does. It collects launch configuration, occupancy, compute and memory throughput, scheduler state, @@ -572,7 +599,7 @@ The timeline above showed that the H2D copy dominates the example operation. The still selects the 93.152 μs BF16 GEMM—not because it is the operation's primary bottleneck, but to show how the metrics from one launch determine what to inspect next. -### Collect One Target Kernel +### Select One Target Kernel Launch Reuse the `--profile-once` path from the Nsight Systems example. Warm-up stays outside the capture range, and only one target operation runs inside it. That operation launches a GEMM followed by @@ -639,7 +666,7 @@ ncu --import reports/bf16-gemm-basic.ncu-rep \ `header` is a compact first view. Expand the Work ID/CLC and throughput tables in the GUI, or replace `header` with `all` to print the complete details used below. -## Interpret a Real NCU Report +## Analyze an Nsight Compute Report The following values come from a real capture on the same B200 with Nsight Compute 2026.1. NCU used nine replay passes to build the `basic` report. The percentages are NCU throughput metrics relative @@ -711,7 +738,7 @@ sustained peaks: L2, shared memory, or a memory-instruction pipeline. Expand the breakdown before calling a kernel DRAM-bound. -### Expand the report before continuing +### Choose the Next Metrics from the Basic Report The `basic` report covers steps 1–3. Use its evidence to collect only the sections needed for the next question: @@ -811,7 +838,7 @@ When `TVM_KERNEL_DUMP` is set, TVM retains the generated files and passes `-line information to a compiled binary. A Python line may lower to many CUDA or SASS instructions, and an asynchronous tile primitive may be understandable only in those lower-level views. -### Advanced: NCU Changes the Experiment +### Effects of NCU Collection on Experimental Conditions NCU collection can change the execution conditions: @@ -833,7 +860,7 @@ If NCU reports `ERR_NVGPUCTRPERM`, hardware-counter access is restricted. Follow or ask the system administrator to enable the required access; do not make running every experiment as root the default solution. -## Optional Tool: IKET +## Analyze In-Kernel Stages with IKET (Optional) Nsight Systems represents one kernel as a single GPU activity, while NCU aggregates hardware metrics across the kernel. Neither produces a named timeline for load, compute, and wait phases inside a diff --git a/img/nsys_b200_timeline_zh.svg b/img/nsys_b200_timeline_zh_en_tracks.svg similarity index 98% rename from img/nsys_b200_timeline_zh.svg rename to img/nsys_b200_timeline_zh_en_tracks.svg index 3e032218..703ecd2a 100644 --- a/img/nsys_b200_timeline_zh.svg +++ b/img/nsys_b200_timeline_zh_en_tracks.svg @@ -25,9 +25,9 @@ 800 μs 900 μs -外层 NVTX +Outer NVTX -子 NVTX ranges +Child NVTX ranges CUDA APIs diff --git a/img/scripts/gen_nsys_b200_timeline.py b/img/scripts/gen_nsys_b200_timeline.py index 8af6f988..467abd46 100644 --- a/img/scripts/gen_nsys_b200_timeline.py +++ b/img/scripts/gen_nsys_b200_timeline.py @@ -27,11 +27,7 @@ def render(*, chinese: bool, output: Path) -> None: if chinese else "4096×4096 BF16: H2D copy → GEMM → ReLU" ) - rows = ( - ["外层 NVTX", "子 NVTX ranges", "CUDA APIs", "GPU stream 7"] - if chinese - else ["Outer NVTX", "Child NVTX ranges", "CUDA APIs", "GPU stream 7"] - ) + rows = ["Outer NVTX", "Child NVTX ranges", "CUDA APIs", "GPU stream 7"] note = ( "时间以外层 NVTX range 的起点为 0；横向长度来自真实采集，并非示意比例。" if chinese @@ -113,7 +109,7 @@ def rect(start: float, end: float, y: float, color: str, label: str = "", text_c def main() -> None: image_dir = Path(__file__).resolve().parents[1] render(chinese=False, output=image_dir / "nsys_b200_timeline.svg") - render(chinese=True, output=image_dir / "nsys_b200_timeline_zh.svg") + render(chinese=True, output=image_dir / "nsys_b200_timeline_zh_en_tracks.svg") if __name__ == "__main__": diff --git a/zh/appendix/benchmarking_gpu_kernels.md b/zh/appendix/benchmarking_gpu_kernels.md index c5fd1066..1f8e1813 100644 --- a/zh/appendix/benchmarking_gpu_kernels.md +++ b/zh/appendix/benchmarking_gpu_kernels.md @@ -21,7 +21,7 @@ | Nsight Compute（`ncu`） | 一个选定的 GPU kernel 主要受哪类硬件资源或等待限制？ | | IKET（可选） | 一个选定的 kernel 内部，哪些命名阶段或 warp roles 占用了时间？ | -### Profile 结果的三种常见形式 +### 三类 Profile 视图 Profile 不是一个数字，也不只有一种报告格式。本章使用的工具会生成三种互补的视图： @@ -174,7 +174,7 @@ print(f"median end-to-end time: {median(host_samples_ms):.4f} ms") 比较多个实现时，所有实现必须使用相同的计时方法和边界。如果同时报告这两种结果，应分别命名为“CUDA Event GPU 时间”和“单次端到端时间”，而不是把使用不同 timer 得到的数字都写成同一种 latency。 -### 重叠执行时如何计时 +### 重叠执行的计时方法 前面的 GEMM 只在当前 CUDA stream 上运行，因此 start 和 end events 可以直接包住全部工作。如果一个 operator 同时使用多个 streams，仅在当前 stream 记录 events 就不够了：其他 stream 上的工作可能在 start 之前已经开始，也可能在 end 之后仍未完成。 @@ -210,9 +210,9 @@ PDL（Programmatic Dependent Launch）是另一种可能产生重叠的情况。 前面的 benchmark 只告诉我们整个 operation 用了多长时间。如果它会启动多个 kernels，还需要找出时间具体花在哪些 kernels 上。Proton 可以列出每个 kernel 的调用次数、平均时间和累计时间。 -Proton 是 Triton 项目提供的 GPU profiler。它记录 CUDA kernel 活动，因此也能看到由 TVM 编译的 TIRx kernels；这些 kernels 并不是由 Triton 编译的。前面介绍的 `bench` 也支持 `timer="proton"`：它会汇总每次调用中的 kernel 执行时间，并返回多次测量的统计结果。如果想知道其中有哪些 kernels、各调用了多少次，就需要单独采集一棵 kernel 树。 +Proton 是 Triton 项目提供的 GPU profiler。它观察的是 CUDA kernel 活动，因此也可以分析由 TVM 编译的 TIRx kernels。前面介绍的 `bench(timer="proton")` 只返回汇总后的 kernel time；这里单独创建一个 Proton session，以查看每个 kernel 的调用次数和耗时。 -下面继续使用前面分配好的矩阵，并把 GEMM 和 ReLU 组成一个两-kernel operation。代码先完成 warm-up，只采集后面的 100 次调用，最后在当前目录生成 `operator.hatchet`： +下面继续使用前面分配好的矩阵，把 GEMM 和 ReLU 组成一个 operation。代码先完成 warm-up，只采集后面的 100 次调用，最后在当前目录生成 `operator.hatchet`： ```python import torch @@ -254,7 +254,13 @@ proton-viewer --metrics time/ms,count --print-sorted operator.hatchet proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet ``` -如果 viewer 提示缺少 `pandas` 或 `hatchet`，可在 profiling 环境中运行 `python -m pip install pandas llnl-hatchet`。下面是这段代码在 B200 上的一次实际结果；为了便于阅读，缩短了 kernel 名称： +如果 `proton-viewer` 报告缺少可选依赖，再安装： + +```bash +python -m pip install pandas llnl-hatchet +``` + +下面是这段代码在 B200 上的一次实际结果；为了便于阅读，缩短了 kernel 名称： ```text target_operation calls avg/us total/ms @@ -262,17 +268,19 @@ target_operation calls avg/us total/ms └── ReLU kernel 100 4.23 0.423 ``` -先确认预期的 kernels 和调用次数是否正确，再比较叶节点的平均时间与总时间。这个例子中 GEMM 的总时间最大，因此它是更值得继续使用 Nsight Compute 分析的对象。一个 kernel 即使单次很短，也可能因为调用次数很多而占用大量总时间；父 scope 的平均值则不能当作一次完整 operation 的 latency。 +首先核对预期的 kernels 是否都出现、调用次数是否正确，再比较各 kernel 的平均时间和累计时间。这个例子中 GEMM 的累计时间最大，因此下一步应优先用 Nsight Compute 分析 GEMM。也要留意调用次数很多的短 kernel：它们单次耗时不高，累计开销却可能很大。 + +Proton 只统计捕获到的 kernel 时间，不包含内存拷贝、同步和 stream 空隙；如果 kernels 发生重叠，各项 duration 的总和还会重复计算重叠区间。因此，这些数据用于定位需要继续分析的 kernel，不表示整个 operation 的 latency。 -这棵树适合寻找耗时的 kernel，不能替代前面的计时结果。存在重叠时，各 kernel 的时间会重复覆盖同一段区间；内存拷贝、同步和 stream 空隙也不一定出现在树中。此外，这段采集会复用同一组矩阵，没有采用前面 TVM timer 的 L2 驱逐策略，因此两处数字不能直接比较。完整 operation 的时间仍由前面的 CUDA Event 或 wall-clock benchmark 给出。 +这次采集反复使用同一组矩阵，而前面的 TVM timer 会在每次测量前驱逐 L2，两者的 cache 条件也不相同。完整 operation 的 latency 仍应使用 CUDA Events 或同步的 wall-clock timer 测量。 -## 使用 Nsight Systems 阅读应用时间线 +## 使用 Nsight Systems 分析应用时间线 Proton 可以汇总各 kernel 的时间，却看不到它们以什么顺序执行，也看不到 kernel 之间的空隙、数据拷贝和 host 等待。分析这些问题时，需要使用 Nsight Systems 的时间线。 -### 采集一份可复现的报告 +### 采集目标 operation 的时间线 -仓库中的 `appendix/nsys_example.py` 构造了一个简单的三阶段 operation：先把一个 $4096\times4096$ 的 BF16 matrix 从 pinned host memory 复制到 GPU，再执行 GEMM 和 ReLU。输入和输出都在采集前分配。脚本中的核心代码是： +下面用一个简单例子说明怎样限定 Nsight Systems 的采集范围。`appendix/nsys_example.py` 中的 operation 依次完成三个步骤：把一个 $4096\times4096$ 的 BF16 matrix 从 pinned host memory 复制到 GPU（host-to-device，H2D），执行 GEMM，再执行 ReLU。所需 tensors 均在采集前分配，因此报告只聚焦这三个步骤。脚本的核心代码如下： ```python import torch @@ -300,9 +308,9 @@ def run_once_for_profiler(run, *, warmup_calls): cudart.cudaProfilerStop() ``` -Warm-up 位于采集范围之外，`target operation` 则为这次调用提供一个容易识别的 NVTX 名称。同步放在这个 range 内，保证三个 GPU operations 都在 range 结束前完成。`cudaProfilerStart()` 和 `cudaProfilerStop()` 只控制采集范围，不负责性能计时。 +`run_once_for_profiler` 先在 profiler 尚未启动时完成 warm-up，并等待 GPU 上的 warm-up 工作结束。随后，`cudaProfilerStart()` 开始采集；NVTX range `target operation` 为这次 operation 添加名称，便于在时间线中定位。这个 range 内的同步确保三项 GPU 工作在 `cudaProfilerStop()` 之前完成。`cudaProfilerStart()` 和 `cudaProfilerStop()` 只用于限定采集范围，不用于计时。 -运行下面的命令会生成 `reports/target-timeline.nsys-rep`： +下面的命令运行这个脚本，并将报告写入 `reports/target-timeline.nsys-rep`： ```bash mkdir -p reports @@ -317,13 +325,29 @@ nsys profile \ python appendix/nsys_example.py --profile-once ``` -`--trace=cuda,nvtx` 记录 CUDA API、GPU activity 和 NVTX ranges。这里关闭 CPU sampling 与 context-switch tracing，让第一份报告只聚焦 CUDA 时间线。如果报告显示 GPU 长时间没有工作，再单独采集 host scheduling 或 OS runtime 信息。 +`--capture-range=cudaProfilerApi` 只采集 `cudaProfilerStart()` 与 `cudaProfilerStop()` 之间的区间。`--trace=cuda,nvtx` 记录 CUDA API、GPU activity 和 NVTX ranges。这里先关闭 CPU sampling 与 context-switch tracing，使报告集中显示 CUDA 时间线。如果时间线中出现较长的 GPU 空隙，再单独采集一份包含 host scheduling 或 OS runtime 信息的报告。 -可以在 GUI 中打开报告，也可以从命令行打印这次采集最有用的五张表： +报告生成后，可以直接用 Nsight Systems GUI 打开时间线： ```bash nsys-ui reports/target-timeline.nsys-rep +``` + +### 时间线中的拷贝、排队与执行时间 + +这个例子使用 PyTorch，是为了用较少的代码构造数据拷贝和多个 CUDA kernels；后面的时间线读法同样适用于 TIRx operation。 + +下面用一份实际采集的报告说明如何阅读 Nsight Systems 的结果。报告来自一台 NVIDIA B200，软件版本为 NVIDIA driver 595.58.03、CUDA 13.0、PyTorch 2.12.0+cu130 和 Nsight Systems 2025.6.3。 +![一次真实的 Nsight Systems 时间线：CPU 上的 NVTX ranges 和 CUDA APIs 位于上方，实际 H2D、GEMM 与 ReLU 位于同一条 GPU stream 上](../../img/nsys_b200_timeline_zh_en_tracks.svg) + +*本图根据实际采集结果重绘，横条长度与各项操作的实测时长成比例。* + +`GPU stream 7` 中的 `7` 是 Nsight Systems 在这次采集中显示的 stream 标识，不表示第七个执行阶段，换一次运行也可能不同。图中的 H2D copy、GEMM 和 ReLU 位于同一条 stream 上，因此按提交顺序执行。 + +除了在 GUI 中查看时间线，也可以让 `nsys stats` 从同一份报告中整理出下面使用的时间数据： + +```bash nsys stats \ --format=column \ --timeunit=us \ @@ -335,13 +359,11 @@ nsys stats \ reports/target-timeline.nsys-rep ``` -### 读懂一份真实报告 - -下面的数据来自一次真实采集：NVIDIA B200、driver 595.58.03、CUDA 13.0、PyTorch 2.12.0+cu130 和 Nsight Systems 2025.6.3。数值只用于演示读法，不能作为这个 workload 的性能基准。 - -![一次真实的 Nsight Systems 时间线：CPU 上的 NVTX ranges 和 CUDA APIs 位于上方，实际 H2D、GEMM 与 ReLU 位于同一条 GPU stream 上](../../img/nsys_b200_timeline_zh.svg) - -*时间线根据这份报告中的 `cuda_api_trace`、`nvtx_pushpop_trace` 和 `cuda_gpu_trace` 时间戳重绘。横条长度与实测时长成比例。* +`--report` 后面是 Nsight Systems 自带的统计名称。`nvtx_gpu_proj_sum` 和 +`nvtx_pushpop_trace` 分别给出 NVTX range 在 GPU 上覆盖的区间和 host 端的 range 记录； +`cuda_gpu_sum` 汇总 kernels 与 CUDA memory operations；`cuda_kern_exec_trace` 将 host 上的 +launch API 与对应的 GPU kernel 关联起来；`cuda_api_sum` 则汇总 host 端的 CUDA API 调用。 +运行 `nsys stats --help-reports` 可以查看当前版本支持的全部名称和定义。 `cuda_gpu_sum` 给出三项 GPU activity 的时间： @@ -379,13 +401,13 @@ python appendix/nsys_example.py --event-samples 20 Report scripts 会随 Nsight Systems 版本变化。运行 `nsys stats --help-reports` 可以查看当前版本支持的名称，并应在实验记录中保留 `nsys --version`。[Nsight Systems User Guide](https://docs.nvidia.com/nsight-systems/UserGuide/index.html) 介绍了 CLI 和 GUI，[Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html) 则进一步解释 API、queue 和 kernel execution time。 -## 使用 Nsight Compute 分析单个 kernel +## 采集单个 kernel 的 Nsight Compute 报告 Nsight Systems 告诉我们各个 kernel 在什么时候运行；选定一个 kernel 后，Nsight Compute 可以继续查看它的启动配置、occupancy、计算与访存吞吐，以及 scheduler 状态。采集这些指标时，NCU 可能多次重放同一个 kernel，因此它适合诊断原因，不应使用报告中的 `Duration` 代替正常运行时测得的 latency。 上一节的时间线表明，示例 operation 的主要时间花在 H2D copy。下面选择其中耗时 93.152 μs 的 BF16 GEMM 演示 NCU 的读法；这并不表示 GEMM 是整个 operation 最该优化的部分。 -### 只采集一次目标 kernel +### 选择一次目标 kernel launch 继续使用上一节脚本的 `--profile-once` 模式，可以把 warm-up 留在采集范围之外，并且只执行一次目标 operation。这个 operation 会依次启动 GEMM 和 ReLU；下面通过 kernel-name filter 选中 GEMM，并用 `--launch-count 1` 只采集第一个匹配的 launch。命令采用 kernel replay，因此只适合能够独立重放的 kernel。若一段工作包含跨 kernel 依赖或并发，应先查看 Nsight Systems 时间线，再决定是否需要后文介绍的其他 replay mode。 @@ -434,7 +456,7 @@ ncu --import reports/bf16-gemm-basic.ncu-rep \ `header` 适合先查看主要指标；后文使用的 Work ID/CLC 明细和完整 throughput breakdown 可在 GUI 中展开，或把命令中的 `header` 改为 `all` 后打印。 -## 读懂一份真实的 NCU 报告 +## 分析 Nsight Compute 报告 下面的数据来自同一台 B200 上的一次真实采集，使用 Nsight Compute 2026.1。NCU 用 9 个 replay passes 完成了 `basic` 报告。表中的百分比表示相应 throughput 指标占硬件子系统可持续峰值的比例，不是直接用应用 FLOPs 除以芯片标称峰值得到的利用率。 @@ -483,7 +505,7 @@ NCU 还会显示 `Est. Speedup` 等规则生成的提示。它们是在若干简 “Memory throughput 高”不等于“DRAM-bound”。限制项也可能来自 L1、L2、shared memory 或 memory-instruction pipeline。展开 breakdown 后才能判断。 -### 继续阅读前，先按需扩展报告 +### 根据 Basic 报告选择下一组指标 `basic` 报告只覆盖前 3 步。根据其中的线索，只采集回答下一个问题所需的 sections： @@ -572,7 +594,7 @@ mkdir -p "$TVM_KERNEL_DUMP" 设置 `TVM_KERNEL_DUMP` 后，TVM 会保留生成文件，并在 NVCC 编译时加入 `-lineinfo`。NCU 采集命令还要加入 `--import-source yes --source-folders "$TVM_KERNEL_DUMP"`。保存 `inspect_source("cuda")` 仍便于手工对照，但它本身不能给已经编译的 binary 补上 line information。一个 Python line 可能 lower 成多条 CUDA 或 SASS instructions，异步 tile primitive 也可能只能在这些 lower-level views 中看清。 -### 高级情况：NCU 会改变实验条件 +### NCU 采集对实验条件的影响 NCU 采集会改变执行条件： @@ -590,7 +612,7 @@ filter。不要将 NCU 的 `Duration` 直接与无 profiler 的 hot-cache CUDA E [counter permission 指南](https://developer.nvidia.com/nvidia-development-tools-solutions-err-nvgpuctrperm-nsightcompute) 配置权限，或请系统管理员开放所需访问；不应把所有实验长期使用 root 运行作为默认方案。 -## 可选工具：IKET +## 使用 IKET 分析 kernel 内部阶段（可选） 在 Nsight Systems 时间线中，一个 kernel 只显示为完整的 GPU 执行区间；NCU 给出的指标也覆盖整个 kernel。对于已经加入阶段标记的 warp-specialized TIRx kernel，可以使用 IKET（In-Kernel Event Tracing）查看不同 warp 分工在何时工作、等待或发生重叠。 From 3c993630c92f506f523271641bfaf9b178d3c3a9 Mon Sep 17 00:00:00 2001 From: tlopex <820958424@qq.com> Date: Wed, 19 Aug 2026 22:14:57 -0400 Subject: [PATCH 3/4] Improve GPU kernel profiling tutorial --- appendix/benchmarking_gpu_kernels.md | 1481 +++++++++++++++-------- appendix/iket_example.py | 73 ++ appendix/nsys_example.py | 53 +- img/nsys_b200_timeline.svg | 81 +- img/nsys_b200_timeline_zh_en_tracks.svg | 83 +- img/scripts/gen_nsys_b200_timeline.py | 43 +- zh/appendix/benchmarking_gpu_kernels.md | 857 ++++++++----- 7 files changed, 1742 insertions(+), 929 deletions(-) create mode 100644 appendix/iket_example.py diff --git a/appendix/benchmarking_gpu_kernels.md b/appendix/benchmarking_gpu_kernels.md index a1016514..4f1121a5 100644 --- a/appendix/benchmarking_gpu_kernels.md +++ b/appendix/benchmarking_gpu_kernels.md @@ -2,18 +2,12 @@ # Measuring and Analyzing GPU Kernel Performance GPU kernel optimization involves two separate questions: how fast the operation is, and where its -time goes. A benchmark answers the first question; a profile helps answer the second. Because a -profiler changes the execution environment, final claims about complete-operator or application -latency should be confirmed with an unprofiled measurement. +time goes. A benchmark answers the first question; a profile helps answer the second. One Python call does not necessarily correspond to one GPU kernel. It may launch several kernels, enqueue memory copies, or wait for GPU work to finish. Before timing, define which of those steps belong to the operation and use the same boundary for every implementation. -In practice, first verify correctness and establish an unprofiled baseline, then use a profiler to -investigate where the time goes. After changing the implementation, repeat the same measurement. An -optimization is successful only when it improves the unprofiled baseline. - The performance chapter ({ref}`chap_performance`) uses roofline analysis to reason about compute and memory limits. Here the focus shifts to experiments: defining the timing boundary, choosing warm-up and repeat budgets, and reading profiler reports. @@ -25,29 +19,15 @@ The tools used in this workflow have different jobs: | Tool | Primary question | |---|---| | CUDA Events | How much GPU-stream time elapsed across the measured region? A stream is an ordered queue of GPU work. | +| Synchronized wall-clock timer | How much wall-clock time elapsed between starting a host call and completing all GPU work required by it? | | Proton (provided by Triton) | Which GPU kernels were launched, how often, and which kernels account for most of the time? | | Nsight Systems | How do host work, streams, copies, kernels, and communication overlap on a timeline? | -| Nsight Compute (`ncu`) | Why does one selected GPU kernel spend its cycles the way it does? | -| IKET (optional) | Which named phases or warp roles consume time inside one selected kernel? | - -### Three Profile Views - -A profile is not a single number or a single report format. The tools in this chapter produce three -complementary views: - -| View | Tools | How to read it | -|---|---|---| -| Aggregation tree | Proton | Compare call count, average duration, and total duration to locate expensive kernels. | -| Timeline | Nsight Systems; IKET inside one kernel | Read time from left to right across tracks; inspect gaps, overlap, and dependencies. | -| Per-kernel metric report | Nsight Compute | Read launch configuration, utilization, scheduler state, memory traffic, and source/SASS evidence for one launch. | - -Profiles explain where time is spent; they do not replace the performance measurement. After changing -an implementation, disable profiling and measure it again with the same timing boundary used for the -baseline. +| Nsight Compute (`ncu`) | What is the selected kernel doing internally, and which resource or stall should be investigated next? | +| IKET (optional) | After adding in-kernel markers, when are the data-movement, compute, and writeback phases active, waiting, or overlapping? | ## Verify Correctness Before Timing -Verify correctness separately before collecting performance: +Verify correctness before measuring performance: 1. Construct representative inputs, including relevant boundary cases. 2. Run the implementation and synchronize so that GPU work has completed. @@ -55,9 +35,30 @@ Verify correctness separately before collecting performance: 4. If the kernel accumulates into an existing output or modifies an input in place, restore the same initial state before each correctness check. -Reference computation and result comparison stay outside performance timing. Design the benchmark -only after correctness passes; whether state reset is timed depends on the operation boundary defined -in the next section. +For a custom GEMM, let `actual` be the implementation's output. One possible reference is a PyTorch +GEMM computed in FP32 and then converted to the target output type: + +```python +import torch + + +torch.set_float32_matmul_precision("highest") +actual = my_gemm(a, b) # replace with your implementation +torch.cuda.synchronize() +expected = torch.mm(a.float(), b.float()).to(actual.dtype) +rtol = 1e-2 # example only; adjust for the output dtype, accumulation, and shape +atol = 1e-2 +torch.testing.assert_close(actual, expected, rtol=rtol, atol=atol) +``` + +`torch.set_float32_matmul_precision("highest")` prevents this CUDA FP32 reference from using +reduced-precision internal matrix multiplication. The example `rtol` and `atol` control relative and +absolute error, respectively. The value `1e-2` is only a runnable starting point; adjust it for the +output dtype, accumulation method, shape, and operator contract. Use the same reference and +tolerances throughout one comparison. + +Reference computation and result comparison stay outside performance timing. Whether state reset is +timed depends on the operation boundary defined in the next section. ## Define the Timing Boundary @@ -67,7 +68,13 @@ result. State explicitly whether compilation, input construction, allocation, or part of that operation. Implementations are directly comparable only when they perform the same work inside the measured boundary. -After fixing the scope, choose the timer: +The GEMM-plus-ReLU example later in this chapter can use three different boundaries. CUDA Events +around `torch.mm` measure GEMM GPU-stream time. Events around the complete `run()` measure +GEMM-plus-ReLU GPU-stream time. A CPU timer started before `run()` and stopped after synchronizing +measures end-to-end latency for one Python call. Matrix allocation and warm-up remain outside all +three boundaries by default. + +Once the scope is defined, choose the timer: - **CUDA Events** record timestamps when a GPU stream reaches two points. They can measure the device-timeline interval around one kernel or a complete operator. Kernels, memory copies, and idle @@ -79,22 +86,20 @@ After fixing the scope, choose the timer: For example, if the GPU executes the start event before the host submits the next launch, that idle stream time remains inside the CUDA Event interval. An Event interval is therefore not necessarily -the same as a kernel's start-to-finish execution interval in a profiler. Profilers are useful for -examining execution and overlap, but diagnostic profiles do not directly replace unprofiled timing at -the same boundary. +the same as a kernel's start-to-finish execution interval in a profiler. + +## Measure GPU Time and Single-Call Latency -## Measure GPU Time with CUDA Events +### Measure GPU Stream Time with CUDA Events -CUDA launches are normally asynchronous. Python can continue after submitting work to a CUDA stream, -before the GPU has finished. A CPU timer placed immediately around that Python call can therefore -stop too early and mostly measure host submission time. Use CUDA Events to measure elapsed time on a -GPU stream. Use the synchronized wall-clock timer shown later when the boundary runs from the Python -call through GPU completion. The +CUDA launches are normally asynchronous, so an unsynchronized CPU timer can stop before the GPU +finishes and mostly reflect host submission time. The [PyTorch CUDA semantics documentation](https://docs.pytorch.org/docs/stable/notes/cuda.html#asynchronous-execution) -describes this behavior in more detail. +describes this behavior. The following benchmark uses CUDA Events to measure elapsed time on the +current stream. -Start with a complete CUDA Event benchmark. The following runnable example allocates its matrices, -runs a warm-up, and then measures five rounds of one FP16 GEMM before reporting their median: +The following runnable CUDA Event benchmark allocates its matrices, runs a warm-up, measures an FP16 +GEMM over five rounds, and reports the median round result: ```python from statistics import median @@ -112,7 +117,7 @@ def gemm(): def measure_batch_ms(fn, calls): - """Return mean CUDA Event time per call over one consecutive batch, in ms.""" + """Return mean CUDA Event time per call for one batch of back-to-back calls, in ms.""" start = torch.cuda.Event(enable_timing=True) end = torch.cuda.Event(enable_timing=True) @@ -141,43 +146,37 @@ print(f"median CUDA Event time: {median(samples_ms):.4f} ms") ``` `measure_batch_ms` records start and end events in the current CUDA stream and divides their elapsed -time by the number of calls. The result is the mean GPU-stream time per GEMM during consecutive -execution. `end.synchronize()` only makes the CPU wait for that round of GPU work so that the Event -result can be read. - -Here `warmup_calls=500`, `repeat=100`, and `rounds=5` are invocation or measurement counts. They are -example values selected from measurements of this GEMM on a B200, not universal defaults. In a -ten-round calibration, 50 warm-up calls still produced a first-to-last decline from 0.01425 ms to -0.01296 ms. At 500 calls, the change narrowed to 0.01332 ms to 0.01301 ms. Results with `repeat=100` -were also more stable than with `repeat=10`. - -For another workload, increase `warmup_calls` until the first rounds no longer become consistently -faster or slower. Then increase `repeat` or `rounds` until the variation is acceptable for the -experiment. Larger counts are not automatically better: if longer runs systematically shift the -timing level, inspect temperature, power, and clock behavior and decide whether the experiment should -represent short bursts or sustained execution. Long-running kernels generally need smaller counts. - -This code reuses the same matrices, so later calls may find some data in cache. It therefore represents -a warm-cache workload. A published result should also record the GPU model, software versions, and -clock settings. - -When benchmarking the TIRx kernels in this book, there is no need to rewrite the warm-up, repeated -timing, and statistics loop for every kernel. TVM's +time by the number of calls. The result is the average GPU-stream time per GEMM across back-to-back +calls. `end.synchronize()` makes the CPU wait for that round of GPU work so that the Event result can +be read. + +`warmup_calls=500` and `repeat=100` count invocations; `rounds=5` requests five independent +measurement batches. These values came from stability testing on the B200: results were still falling +after 50 warm-up calls but stabilized by 500, and `repeat=100` was more stable than `repeat=10`. For +another workload, first increase +`warmup_calls` until the early rounds stop drifting; if variation remains large, increase `repeat` or +`rounds`. If longer runs instead shift the overall timing level, inspect temperature, power, and clock +behavior. + +This code always reuses the same matrices, so it represents a warm-cache workload with repeated +inputs. Whether the accesses actually hit in cache still depends on the total amount of data revisited +by this computation and the hardware cache capacity. + +This book uses TVM's [`tvm.tirx.bench.bench`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/bench.py) -already provides those steps. Pass it a function that launches the prepared implementation; inputs, -outputs, and workspace remain allocated outside the measured interval. - -The helper uses a different cache policy from the manual example. The example repeatedly reuses the -same matrices, whereas `bench` evicts L2 before each measured invocation and records an independent -CUDA Event interval. Invoke it as follows: +to handle warm-up, repeated timing, and statistics. The supplied function only launches a prepared +implementation; inputs, outputs, and workspace remain outside the timed interval. Unlike the manual +warm-cache example, `bench` writes a 256 MiB buffer before each measured invocation to reduce reuse +of data retained in L2 from the previous invocation, then records an independent CUDA Event interval: ```python from tvm.tirx.bench import bench -# run is a no-argument function that launches the operation on preallocated tensors. +# Reuse gemm from above. For a custom TIRx kernel, substitute its no-argument callable. +run = gemm result = bench( - {"tirx": run}, + {"gemm": run}, timer="event", warmup=25, repeat=100, @@ -185,41 +184,28 @@ result = bench( cooldown_s=1.0, ) -print(result["impls"]["tirx"]) # five-round mean, in us -print(result["round_samples"]["tirx"]) # result from each round +print(result["impls"]["gemm"]) # five-round mean, in us +print(result["round_samples"]["gemm"]) # result from each round ``` -Here `warmup=25` and `repeat=100` are time budgets in milliseconds, not invocation counts. The Event -timer first performs a short estimate, then converts the 25 ms warm-up budget and 100 ms measurement -budget into iteration counts. That estimate includes both L2 eviction and the measured call, so the -resulting counts are approximate. In the reported samples, the L2 eviction occurs before the start -event and only the invocation is timed. Short kernels therefore run more times than long kernels. -`rounds=5` repeats the complete measurement five times, while `cooldown_s=1.0` waits one second before -measuring an implementation in each round. `impls` contains the five-round mean and `round_samples` -retains the individual results. The 25/100 ms values are the Event timer defaults. Five rounds are the -default used by the TIRx-kernels CLI; `bench` itself defaults to one round. - -These values are starting points rather than a standard for every workload. Increase the warm-up -budget if results continue to drift between rounds. Increase the measurement budget or number of -rounds if the results remain noisy. Use the same timer, budgets, and rounds for every implementation, -and retain all round results instead of reporting only the fastest one. - -TIRx-kernels uses this helper in its `run_bench` entry points; see -[`tirx_kernels/attention/flash_attention4.py`](https://github.com/mlc-ai/tirx-kernels/blob/5be39749e7dfd2c4bdae9b4d396f8ec35af07126/tirx_kernels/attention/flash_attention4.py). -The local benchmark defaults to Proton when `timer` is omitted. Specify `timer="event"` as above when -the intended result is a CUDA Event interval. The two timers report different quantities, so a result -must identify which one was used. +`warmup=25` and `repeat=100` are millisecond budgets. The Event timer uses a short calibration run to +convert them into invocation counts. The reported Event intervals cover only the measured calls; the +256 MiB write used to reduce L2 reuse occurs before each start event. `rounds=5` runs five rounds, +and `cooldown_s=1.0` pauses before each one. `impls` stores the five-round mean, while +`round_samples` stores the individual results. Adjust budgets +and rounds by the same stability criteria used above, and use the same settings for every +implementation. -The function passed to `bench` is invoked repeatedly. If a kernel accumulates into its output or -modifies an input in place, restore equivalent state before every call or ensure that every measured -invocation receives fresh preallocated state. A reset inside the measured function belongs to the -operation boundary defined earlier. Otherwise later calls no longer represent the same workload. +The `run_bench` entry points in TIRx-kernels also use this helper. Outside distributed mode, omitting +`timer` selects Proton by default; specify `timer="event"` for a CUDA Event interval. A repeatedly +invoked in-place kernel must still follow the reset rule above. If reset occurs inside the measured +function, its cost belongs to the operation. ### Measure End-to-End Time for One Call -Use a synchronized wall-clock timer when the target is the complete interval from one Python call -until its GPU work finishes. The following code continues with the `gemm()` defined and warmed up -above: +Use a synchronized wall-clock timer when the target is the complete interval from the start of a +Python call until its GPU work finishes. The following code continues with the `gemm()` defined and +warmed up above: ```python from statistics import median @@ -246,68 +232,127 @@ print(f"median end-to-end time: {median(host_samples_ms):.4f} ms") The first synchronization keeps unfinished earlier work outside the measurement. The second ensures that this GEMM finishes before the timer stops. Each sample contains exactly one call, so the result -includes the Python call, CUDA launch, GPU execution, and the wait for completion. The CUDA Event -benchmark above instead reports mean GPU-stream time per GEMM during consecutive execution. +includes the Python call, CUDA launch, GPU execution, and the wait for completion. By comparison, the +CUDA Event benchmark above reports average GPU-stream time per GEMM across back-to-back calls. Every implementation in a comparison must use the same timer and boundary. If both results are -reported, name them separately as *CUDA Event GPU time* and *single-call end-to-end time* rather than -placing values from different timers under one generic latency label. +reported, name them separately as *CUDA Event GPU time* and *single-call end-to-end time* so that +their different boundaries remain visible. -### Timing Overlapping GPU Work +### Advanced: Timing a Multi-Stream Operation -The GEMM examples above run entirely in the current CUDA stream, so their start and end events cover -all of the work. An operator that uses several streams needs additional synchronization. Events -recorded only in the current stream do not automatically include work elsewhere, which may begin -before the start event or remain unfinished after the end event. +When one operation submits work to several CUDA streams, the timing stream must connect the start +and finish of every branch. In this example, `sin` and `cos` run on separate streams. The timing +stream waits for both branches before adding their results: -To time the complete operator, use the start event as a common starting signal. Every work stream -waits for start before beginning the measured work and signals completion when it finishes. The stream -that records the end event waits for all of those completion signals first. The resulting interval -then spans the operation from its earliest start through its final completion. +```python +import torch -Programmatic Dependent Launch (PDL) is another source of possible overlap. On GPUs with compute -capability 9.0 or newer, it can start a later kernel early in the same stream. That kernel may perform -preparation that does not depend on earlier results, then wait before consuming those results. PDL -must be enabled explicitly and follow its trigger-and-wait contract; see the -[CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html) -for the API details. -The timing rule is the same whether overlap comes from multiple streams or PDL: place CUDA Events -around the complete operation. Overlapping kernels can cover the same time interval, so adding their -profiler durations does not produce operator latency. A Nsight Systems timeline shows the actual -ordering and overlap. Because PDL overlap is opportunistic, program correctness cannot require it to -occur. +x = torch.randn(1 << 20, device="cuda") +left = torch.empty_like(x) +right = torch.empty_like(x) +output = torch.empty_like(x) + +stream_left = torch.cuda.Stream() +stream_right = torch.cuda.Stream() +timing_stream = torch.cuda.current_stream() + +start = torch.cuda.Event(enable_timing=True) +left_done = torch.cuda.Event() +right_done = torch.cuda.Event() +end = torch.cuda.Event(enable_timing=True) + + +def measure_operation_ms(): + torch.cuda.synchronize() + start.record(timing_stream) + + stream_left.wait_event(start) + with torch.cuda.stream(stream_left): + torch.sin(x, out=left) + left_done.record(stream_left) + + stream_right.wait_event(start) + with torch.cuda.stream(stream_right): + torch.cos(x, out=right) + right_done.record(stream_right) + + timing_stream.wait_event(left_done) + timing_stream.wait_event(right_done) + torch.add(left, right, out=output) + + end.record(timing_stream) + end.synchronize() + return start.elapsed_time(end) + + +elapsed_ms = measure_operation_ms() +torch.testing.assert_close(output, torch.sin(x) + torch.cos(x)) +print(f"multi-stream operation: {elapsed_ms:.4f} ms") +``` + +`start` is the common starting signal. `left_done` and `right_done` mark the ends of the two +branches. The timing stream waits for both completion events, performs the final addition, and then +records `end`. Each `wait_event` creates a GPU-side dependency while the CPU continues submitting +work; only `end.synchronize()` waits on the CPU. + +The event graph is: + +```text +timing stream: start ──────────── wait(left_done, right_done) ─ add ─ end +left stream: wait(start) ─ sin ─ left_done +right stream: wait(start) ─ cos ─ right_done +``` + +The graph permits the two branches to execute concurrently; actual overlap depends on their GPU +resource use. Confirm the realized schedule in a Nsight Systems timeline. For formal measurement, +call `measure_operation_ms()` several times for warm-up, then call it repeatedly to collect +single-call samples and report their median and variation. + +#### PDL Within One Stream + +Programmatic Dependent Launch (PDL) applies to a custom CUDA or DSL launch path that explicitly +enables it. The primary and secondary kernels are submitted to the same stream. After the primary +emits a trigger, the secondary may begin preparation that does not depend on the primary's result; +it performs the PDL dependency synchronization before consuming that result. + +```text +primary: initial work ─ trigger ─ remaining work +secondary: preamble ─ wait ─ dependent work +``` + +Record `start` before launching the primary and `end` after launching the secondary. That complete +Event interval is the GPU time of the launch sequence. The two kernels may overlap, so the sum of +their profiler durations can exceed the complete operation latency. Use a Nsight Systems timeline +to confirm the realized overlap. + +PDL creates an opportunity for concurrent execution, while the runtime may still serialize the +kernels. Kernel correctness must cover both schedules. The `torch.cuda.Stream` interface above does +not expose PDL launch attributes; a custom CUDA or DSL implementation supplies them. See the +[CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html) +for the enablement details. ## Keep Experimental Conditions Consistent -The preceding sections established the timing boundary and timer. The remaining experimental -conditions must also be held constant. The examples above start after allocation and warm-up, so they -measure subsequent repeated calls. A first call may also include CUDA initialization, JIT compilation, -autotuning, or other one-time work. If first-call latency or a complete application path is the target, -include those steps in the boundary and report the result separately from repeated-call performance. - -One measurement is not enough to establish stability. The manual CUDA Event example retains five -rounds and reports their median. `bench` reports the mean across rounds and also stores the raw values -in `round_samples`. Whichever summary is used, keep the per-round results, inspect them for trends and -outliers, and state whether the reported value is a median or mean rather than selecting only the -fastest round. When comparing implementations, repeat the experiment in a different order so that one -implementation is not always measured on a colder or hotter device. - -Cache conditions also affect the result. The manual example repeatedly uses the same matrices and -therefore measures a warm-cache workload. -The TVM 0.26 Event and Proton timers instead write a 256 MiB buffer before every measured invocation -to evict existing L2 data; that write remains outside the timed interval. Either policy can be valid. -Choose the one that represents the target application and apply it consistently to every -implementation. `torch.cuda.empty_cache()` releases unused blocks from PyTorch's caching allocator. -It does not clear GPU L2 and cannot implement a cold-L2 measurement. - -Finally, record the GPU model, driver, CUDA runtime, framework and compiler versions, and the workload -dtype and shape. Also record clock and power settings, keep unrelated processes off the device, and -watch for thermal throttling. If clocks are locked, provide the actual values and command; "fixed -clocks" alone is not enough to reproduce the experiment. - -Matching tensor shapes alone does not make two implementations comparable. Align at least three -classes of conditions: +The examples above measure repeated calls after allocation and warm-up. If first-call latency or a +complete application path is the target, include the relevant CUDA initialization, JIT compilation, +and autotuning in the timing boundary and report that result separately. + +Keep every per-round result and state whether the summary is a median or mean. A continuing trend +across rounds calls for checking warm-up, temperature, and clock state before summarizing the full +set of measurements. When comparing implementations, alternate their measurement order so that no +implementation is consistently measured on a colder or hotter device. + +Use one cache policy throughout the comparison. The manual example repeatedly reuses the same +matrices and is therefore biased toward warm-cache reuse, although the actual hit rate still depends +on the total amount of data revisited and the cache capacity. The TVM 0.26 Event and Proton timers +instead write a 256 MiB buffer before each measured invocation to reduce L2 reuse; this write is +outside the timed interval. Choose the policy that represents the target application. +`torch.cuda.empty_cache()` only releases unused blocks from PyTorch's allocator; GPU L2 contents +remain managed by the hardware cache policy. + +Also align the following conditions across implementations: - **Numerical semantics:** input and output dtype, layout, transpose convention, alignment, accumulation precision, scale, mask, epilogue, output definition, and accuracy tolerance; @@ -316,130 +361,141 @@ classes of conditions: - **Tuning conditions:** workspace limits, whether per-shape autotuning is allowed, and the search budget available to each implementation. -Every implementation should also use the same cache, clock, warm-up, sampling, and timing policies. -For a library baseline, record its version, selected algorithm, and workspace. Autotuning may run -outside the timed interval, but its search budget and final configuration remain part of the -experimental record. +Record the GPU, driver, CUDA, framework, and compiler versions, along with dtype, shape, clock, and +power settings. For a library baseline, also record the library version, selected algorithm, and +workspace. Autotuning may run outside the timed interval, but its search budget and final +configuration are still part of the experimental setup. ## Convert Latency to Throughput -Throughput is not measured directly by the timer. It is computed by dividing a defined amount of work -by the measured latency. A table that reports TFLOP/s, GB/s, or tokens/s should therefore retain the -original latency and explain how the work was counted. This book counts GEMM as $2MNK$ FLOPs. For -attention and fused kernels, specify whether the count represents the full dense problem, the elements -actually selected, or the work executed by the kernel. See {ref}`chap_performance` for the corresponding -formulas and roofline analysis. +Compute throughput by dividing a clearly defined amount of work by the measured latency. For a +multiplication of an $M\times K$ matrix by a $K\times N$ matrix with latency `t_us` in +microseconds: -## Use Proton to Find Expensive Kernels +```text +TFLOP/s = 2 × M × N × K / t_us / 10^6 +``` -The preceding benchmark tells us how long the complete operation takes. If it launches several kernels, -we still need to determine where that time is spent. Proton reports each kernel's call count, average -time, and cumulative time. +The numerator and timing boundary must describe the same work. The complete GEMM-plus-ReLU operation +below takes 105.152 μs. Dividing $2\times4096^3$ by that duration gives about 1307 TFLOP/s, but this +is only *effective throughput*: GEMM work divided by the complete operation time, whose denominator +also contains ReLU. To report the GEMM kernel's own TFLOP/s, the timed interval must cover only GEMM. -Proton is a GPU profiler provided by the Triton project. It observes CUDA kernel activity and can -therefore analyze TIRx kernels compiled by TVM. The `bench(timer="proton")` helper introduced above -returns an aggregate kernel-time result. Here we create a separate Proton session to inspect each -kernel's call count and execution time. +A table that reports TFLOP/s, GB/s, or tokens/s should retain the original latency and explain how +the work was counted. For attention and fused kernels, also state whether the numerator represents +the full dense problem, the elements actually selected, or the work executed by the kernel. See +{ref}`chap_performance` for the corresponding formulas and roofline analysis. -The following example reuses the matrices allocated above and combines GEMM with ReLU in one -operation. It finishes warm-up, collects the next 100 calls, and writes `operator.hatchet` in the -current directory: +## Find the Most Expensive Kernel with Proton + +From this point onward, the baseline, Proton, Nsight Systems, and Nsight Compute all run the same +operation from `appendix/nsys_example.py`: multiply two $4096\times4096$ BF16 matrices and then apply +ReLU to the result. The inputs, intermediate result, and output are allocated before timing or +collection. The operation in the script is: ```python -import torch -import triton.profiler as proton +def run(): + with torch.cuda.nvtx.range("BF16 GEMM"): + torch.mm(a, b, out=c) + with torch.cuda.nvtx.range("ReLU"): + torch.clamp_min(c, 0, out=output) +``` +Before entering any timing or collection mode, the script runs one operation and synchronizes. It +then computes an FP32 `torch.mm` reference, applies ReLU, converts the result to BF16, and compares it +with the output using `rtol=2e-2` and `atol=1e-2`. A mismatch stops the command before the benchmark +or profiler begins. This preflight check is outside both the baseline timing interval and the +collection range used below. -def operation(): - torch.mm(a, b, out=c) - torch.clamp_min(c, 0, out=c) +The earlier $2048\times2048$ FP16 example was a standalone demonstration of the timing APIs. The +results below all use this BF16 operation and do not mix the two workloads. +### Establish an Unprofiled Baseline -def collect_proton(run, *, warmup_calls, profile_calls): - for _ in range(warmup_calls): - run() - torch.cuda.synchronize() +Before asking where time is spent, measure the complete operation under normal execution: - session = proton.start("operator", context="shadow", data="tree") - if session is None: - raise RuntimeError("Proton session could not be created") - try: - with proton.scope("target_operation"): - for _ in range(profile_calls): - run() - torch.cuda.synchronize() - finally: - proton.finalize(session) +```bash +python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --event-samples 20 +``` +The 500 warm-up calls finish before formal timing. Each sample then uses one pair of CUDA Events +around one GEMM-plus-ReLU operation. One actual B200 run produced: -collect_proton(operation, warmup_calls=500, profile_calls=100) +```text +median=105.152 us, min=103.136 us, max=131.200 us ``` -Here `warmup_calls` and `profile_calls` are invocation counts, not the millisecond budgets used by -`bench`. Install a Triton version compatible with TVM and record that version with the experiment. +The median is the baseline to revisit after changing code. The minimum and maximum show the sample +variation. This baseline covers the complete GEMM-plus-ReLU operation; the profiler tables below +report individual kernels from separate captures. -First list the stored metrics. The next two commands print call counts, total time, and average time: +### Use Proton to Compare the Kernels in the Operation + +Proton reports each kernel's call count, average time, and cumulative time. Before running it, +confirm that the environment has a Triton installation compatible with TVM; +Proton and `proton-viewer` are provided with Triton. The viewer also requires two Python packages: ```bash -proton-viewer --list operator.hatchet -proton-viewer --metrics time/ms,count --print-sorted operator.hatchet -proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet +python -m pip install pandas llnl-hatchet ``` -If `proton-viewer` reports missing optional dependencies, install them with: +The script's `--proton-calls` mode reuses the same `run()`, completes warm-up, profiles 100 calls to +the operation, and writes `operator.hatchet`: ```bash -python -m pip install pandas llnl-hatchet +python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --proton-calls 100 ``` -The following is one real B200 result from this code; kernel names are shortened for readability: +Here `warmup-calls` and `proton-calls` are invocation counts, while the similarly named warm-up and +repeat arguments to `bench` are millisecond budgets. After generating the report, run: + +```bash +proton-viewer --list operator.hatchet +proton-viewer --metrics time/ms,count --print-sorted operator.hatchet +proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet +``` + +`proton-viewer` prints two separate tables. First use the full kernel name to confirm the call count +and cumulative time in the `count,time/ms` output. Then find the same row in the +`avg_time/us,time/ms` output to read the average time. Select the +target primarily by `time/ms`, which is the time accumulated by that kernel across 100 operations; +`avg_time/us` reports its per-call average. The display below joins the two tables and shortens the +kernel names for readability: ```text target_operation calls avg/us total/ms -├── GEMM kernel 100 14.83 1.483 -└── ReLU kernel 100 4.23 0.423 +├── GEMM kernel 100 87.00 8.700 +└── ReLU kernel 100 11.71 1.171 ``` -First confirm that all expected kernels appear with the correct call counts, then compare their -average and cumulative times. GEMM has the largest cumulative time in this example, so it should be -the first kernel examined with Nsight Compute. Also watch for short kernels that are launched often: -their individual calls may be inexpensive while their cumulative cost is substantial. +First confirm that both expected kernels appear and that each was called 100 times, then compare +cumulative time. GEMM accounts for much more time, so it becomes the target for deeper analysis. +Before profiling it with NCU, use Nsight Systems to confirm the kernel order and gaps and correlate +them with the host launch APIs for one operation. -Proton includes only the captured kernel times, not memory copies, synchronization, or stream gaps. -When kernels overlap, summing their durations also counts the overlapping interval more than once. -These data therefore identify kernels for further analysis; they are not the latency of the complete -operation. - -This capture repeatedly uses the same matrices, whereas the earlier TVM timer evicts L2 before each -measurement, so the two experiments also have different cache conditions. Measure complete-operation -latency with CUDA Events or a synchronized wall-clock timer. +This manual Proton session preserves the normal cache state, unlike `bench(timer="proton")`, which +writes a 256 MiB buffer before each measured call. Use values from the same Proton report to rank +kernels. Compare implementations with the CUDA Event baseline above, using a timing interval that +covers the full operation. ## Analyze an Application Timeline with Nsight Systems -Proton can aggregate kernel time, but it cannot show execution order, gaps between kernels, copies, -or host waits. Use the Nsight Systems timeline to examine those relationships. +Proton provides an aggregate ranking, but it does not show kernel ordering, gaps, copies, or host +waits. Nsight Systems answers those questions with a timeline. ### Capture the Target Operation Timeline -The following example shows how to restrict Nsight Systems collection to one target operation. In -`appendix/nsys_example.py`, that operation performs three steps in sequence: it copies a -$4096\times4096$ BF16 matrix from pinned host memory to the GPU in a host-to-device (H2D) copy, runs -GEMM, and applies ReLU. All required tensors are allocated before collection, so the report focuses -on these three steps. The core of the script is: +The script's `--profile-once` mode completes warm-up while the profiler is still inactive and waits +for that work to finish. It then submits exactly one GEMM-plus-ReLU operation between +`cudaProfilerStart()` and `cudaProfilerStop()`: ```python -import torch - - -def run(): - with torch.cuda.nvtx.range("H2D input"): - a.copy_(host_a, non_blocking=True) - with torch.cuda.nvtx.range("BF16 GEMM"): - torch.mm(a, b, out=c) - with torch.cuda.nvtx.range("ReLU"): - torch.clamp_min(c, 0, out=output) - - def run_once_for_profiler(run, *, warmup_calls): for _ in range(warmup_calls): run() @@ -453,11 +509,9 @@ def run_once_for_profiler(run, *, warmup_calls): cudart.cudaProfilerStop() ``` -`run_once_for_profiler` completes warm-up before the profiler starts and waits for the warm-up work on -the GPU to finish. It then calls `cudaProfilerStart()` and labels the measured invocation with the -`target operation` NVTX range, making it easy to locate in the timeline. The synchronization inside -the range ensures that all three GPU operations finish before `cudaProfilerStop()`. These profiler -APIs define the collection range; they do not measure performance. +The NVTX range makes the target operation easy to locate in the timeline. Synchronization inside the +range ensures that both kernels finish before collection stops. `cudaProfilerStart()` and +`cudaProfilerStop()` only delimit the collection range; they do not measure time. The following command runs the script and writes `reports/target-timeline.nsys-rep`: @@ -471,7 +525,10 @@ nsys profile \ --capture-range-end=stop \ --output=reports/target-timeline \ --force-overwrite=true \ - python appendix/nsys_example.py --profile-once + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once ``` `--capture-range=cudaProfilerApi` restricts collection to the interval between @@ -486,137 +543,177 @@ Once the report has been generated, open its timeline in the Nsight Systems GUI: nsys-ui reports/target-timeline.nsys-rep ``` -### Copy, Queue, and Execution Time in the Timeline - -This example uses PyTorch to construct a copy-plus-multiple-kernel workload with little setup. The -same timeline-reading method applies to a TIRx operation. +### Locate the Most Expensive Kernel in the Timeline -The following real report illustrates how to interpret Nsight Systems results. It was collected on an -NVIDIA B200 with NVIDIA driver 595.58.03, CUDA 13.0, PyTorch 2.12.0+cu130, and Nsight Systems 2025.6.3. +The following report was collected on an NVIDIA B200 with NVIDIA driver 595.58.03, CUDA 13.0, +PyTorch 2.12.0+cu130, and Nsight Systems 2025.6.3. -![A measured Nsight Systems timeline: CPU NVTX ranges and CUDA APIs appear above the H2D copy, GEMM, and ReLU on one GPU stream](../img/nsys_b200_timeline.svg) +![A measured Nsight Systems timeline: CPU NVTX ranges and CUDA APIs appear above GEMM and ReLU on one GPU stream](../img/nsys_b200_timeline.svg) -*This figure was redrawn from an actual capture. Each bar is scaled to the measured duration of the -corresponding activity.* - -The `7` in `GPU stream 7` is the stream identifier shown by Nsight Systems for this capture. It does -not mean the seventh execution stage and may differ in another run. The H2D copy, GEMM, and ReLU share -that stream and therefore execute in submission order. - -In addition to viewing the timeline in the GUI, use `nsys stats` to extract the timing data used -below from the same report: +The `7` in `GPU stream 7` is the stream identifier in this report. Both kernels are in the same +stream and therefore execute in submission order. The times used below can also be extracted from +the command line: ```bash nsys stats \ + --force-export=true \ --format=column \ --timeunit=us \ - --report nvtx_gpu_proj_sum \ - --report nvtx_pushpop_trace \ --report cuda_gpu_sum \ --report cuda_kern_exec_trace \ + --report cuda_api_trace \ + --report nvtx_gpu_proj_sum \ + --report nvtx_pushpop_trace \ --report cuda_api_sum \ reports/target-timeline.nsys-rep ``` -The names following `--report` refer to summaries built into Nsight Systems. -`nvtx_gpu_proj_sum` and `nvtx_pushpop_trace` report the GPU projection of an NVTX range and its -host-side range records, respectively. `cuda_gpu_sum` summarizes kernels and CUDA memory operations; -`cuda_kern_exec_trace` correlates host launch APIs with their GPU kernels; and `cuda_api_sum` -summarizes host-side CUDA API calls. Run `nsys stats --help-reports` to list the names and definitions -available in the installed version. +`--force-export=true` regenerates the SQLite data from the current `.nsys-rep`, preventing stale +SQLite data with the same name from being reused. `cuda_gpu_sum` summarizes GPU activities; +`cuda_kern_exec_trace` correlates each host launch API with its GPU kernel and gives the kernel start +and duration; `cuda_api_trace` gives the start and duration of every CUDA API call. The two NVTX +reports provide the range's GPU projection and host-side record, while `cuda_api_sum` summarizes +host CUDA APIs. Run +`nsys stats --help-reports` for the complete definitions in the installed version. -`cuda_gpu_sum` reports the three GPU activities: +The command prints several tables in sequence. Extract values in this order: + +1. Compare cumulative durations in `cuda_gpu_sum` to find the most expensive GPU kernel. +2. In `cuda_kern_exec_trace`, read each kernel's start and duration and identify its host launch API. + Compute its end as `end = start + duration`; use the same calculation for CUDA APIs. +3. Compute the inter-kernel gap as `ReLU start - (GEMM start + GEMM duration)`. This uses the trace + timestamp and duration columns. +4. Finally, use the `nvtx_*` reports, `cuda_api_trace`, and `cuda_api_sum` to interpret the host + range and synchronization APIs. To determine whether a synchronization call actually waited for + the GPU, compare its start time in `cuda_api_trace` with the end of the final kernel. + +`cuda_gpu_sum` reports the execution time of both kernels: | GPU activity | Count | GPU duration | Share of listed GPU time | |---|---:|---:|---:| -| 32 MiB H2D copy | 1 | 607.230 μs | 85.4% | -| BF16 GEMM | 1 | 93.152 μs | 13.1% | -| ReLU | 1 | 11.072 μs | 1.6% | +| BF16 GEMM | 1 | 92.608 μs | 89.4% | +| ReLU | 1 | 10.944 μs | 10.6% | `cuda_kern_exec_trace` correlates each kernel with its launch API and reports API time, positive queue time, and GPU execution separately. Positive queue time is the interval from API return to a later -kernel start; it has no positive value when the kernel starts earlier. +kernel start. If the kernel starts before the API returns, the report shows no positive queue time. | Kernel | API time | Positive queue time | GPU execution | |---|---:|---:|---:| -| BF16 GEMM | 34.270 μs | 403.214 μs | 93.152 μs | -| ReLU | 11.064 μs | 442.166 μs | 11.072 μs | - -This report supports four concrete observations: - -1. **The H2D copy dominates this region.** It accounts for 85.4% of the three GPU-duration sum. - Looking only at `cuda_gpu_kern_sum` would omit the copy entirely, so this example uses - `cuda_gpu_sum`, which includes both kernels and memory operations. -2. **Queue time is not launch overhead.** GEMM and ReLU wait behind earlier work on the same stream. - Their queue intervals are long because they follow the H2D copy and GEMM, not because their launch - APIs took hundreds of microseconds. -3. **A long synchronization interval usually means that the host is waiting for the GPU.** - `cudaDeviceSynchronize` occupied 384.821 μs on the CPU because GPU work remained unfinished when - the host called it. That number is neither one kernel's duration nor the complete operation - latency. -4. **Intervals at different scopes cannot be added.** The three GPU durations sum to 711.454 μs. - `nvtx_gpu_proj_sum` measures from the first enclosed GPU operation's start to the last one's end, - producing 715.582 μs; the roughly 4.1 μs difference is gaps between activities. The original - `target operation` range on the CPU is 870.561 μs because it also includes dispatch and the final - synchronization wait. - -In a separate unprofiled run, 20 CUDA Event samples of the same operation had a median of 722.816 μs -and a range of 718.400–726.336 μs. The same measurement method can be rerun with: +| BF16 GEMM | 50.717 μs | — | 92.608 μs | +| ReLU | 13.474 μs | 5.074 μs | 10.944 μs | -```bash -python appendix/nsys_example.py --event-samples 20 -``` +Read the results in this order: -That unprofiled result is the appropriate performance number to report. The single Nsight Systems -timeline explains where time went; the two numbers need not match exactly. +1. **Select the target first.** GEMM accounts for 89.4% of the two kernels' total execution time, so + GEMM becomes the NCU analysis target. +2. **Distinguish API time from GPU time.** The GEMM launch API took 50.717 μs on the host, while the + GPU executed GEMM for 92.608 μs; these are different intervals. ReLU's 5.074 μs positive queue + time means that it waited briefly after the launch API returned before starting on the GPU. +3. **Check the gap between kernels.** Their durations sum to 103.552 μs. The GPU span from the start + of GEMM through the end of ReLU is 103.776 μs, leaving only a 0.224 μs gap. +4. **Interpret the host range according to its timing boundary.** The `target operation` host range + in the figure covers Python/PyTorch dispatch, both launches, and the synchronization API. + Because `cudaDeviceSynchronize` begins after ReLU ends, its measured duration mostly reflects + host-side API overhead; GPU execution was already complete. -These numbers also show why the timing boundary matters. If the production operation truly includes -the H2D copy, reducing or overlapping the transfer is the first place to look. If the application -already holds its input on the GPU, the copy does not belong inside the measured region. Do not -assume that GEMM is the first target merely because it is the main compute kernel. +When inspecting another Nsight Systems timeline, follow the same order: confirm the capture range, +inspect GPU kernels, copies, gaps, and overlap, and then correlate them with host launch or +synchronization APIs. See the +[Nsight Systems Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html) for +the interval definitions and additional UI details. -Apply the same reading order to other reports: first confirm that the NVTX range excludes warm-up and -initialization; inspect kernels, copies, gaps, and overlap on the GPU streams; then follow correlation -back to launch or synchronization APIs on the host. Durations on different streams cannot simply be -added, and visible overlap proves only that it occurred in this capture. Verify any claimed latency -benefit with the same unprofiled boundary. +## Use Nsight Compute to Analyze a Single Kernel -Report scripts vary across Nsight Systems releases. Run `nsys stats --help-reports` to list those -available in the installed version, and record `nsys --version` with the experiment. The -[Nsight Systems User Guide](https://docs.nvidia.com/nsight-systems/UserGuide/index.html) covers the -CLI and GUI; the [Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html) -explains API, queue, and kernel-execution intervals in more detail. +Use Nsight Systems to select the target kernel and NCU to explain its hardware behavior. This section +begins with a reusable workflow, applies it to a BF16 GEMM on a B200, and then explains the supporting +metrics and calculations. -## Collect an Nsight Compute Report for One Kernel +### How to Read an NCU Report -Nsight Systems shows when kernels run; Nsight Compute explains why one selected kernel behaves as it -does. It collects launch configuration, occupancy, compute and memory throughput, scheduler state, -and other hardware metrics. NCU can replay the kernel several times while collecting those metrics, -so it is a diagnostic tool: do not substitute the report's `Duration` for latency measured during a -normal run. +An SM (streaming multiprocessor) is a GPU compute unit that hosts thread blocks and executes their +instructions. A warp contains 32 threads and is the basic group that a scheduler selects when issuing +an instruction. A block or warp is resident after it has been assigned to an SM and before it finishes. -The timeline above showed that the H2D copy dominates the example operation. The NCU walkthrough -still selects the 93.152 μs BF16 GEMM—not because it is the operation's primary bottleneck, but to -show how the metrics from one launch determine what to inspect next. +Analyze a new NCU report in this order: -### Select One Target Kernel Launch - -Reuse the `--profile-once` path from the Nsight Systems example. Warm-up stays outside the capture -range, and only one target operation runs inside it. That operation launches a GEMM followed by -ReLU. The kernel-name filter below selects the GEMM, and `--launch-count 1` collects only the first -matching launch. Kernel replay is appropriate only when the selected kernel can be replayed in -isolation. Inspect a dependent or concurrent multi-kernel region in Nsight Systems before choosing a -different replay mode. - -Start with the `basic` section set: +| Current question | Where to look first | Purpose | +|---|---|---| +| Which kernel does this report describe? | Kernel, device, and grid/block in the report header, plus Warnings/Errors | Confirm the filter and collection result | +| Does the launch expose enough device-wide parallelism? | `Grid Size` and `Waves Per SM` in `LaunchStats` | Determine whether the grid provides enough device-wide parallelism | +| How much work can reside on each SM at once? | `Block Limit` fields and theoretical/achieved occupancy in `Occupancy` | Quantify theoretical and achieved residency, and identify which resources set the theoretical limit | +| Which part of the hardware should be investigated first? | Compute, Memory, and DRAM throughput in `SpeedOfLight` | Choose the compute, memory, or scheduling path | +| Can the scheduler consistently find an instruction to issue? | `Scheduler Statistics` → `Warps Per Scheduler`; when issue-ready work is scarce, continue to `Warp State Statistics` → `Warp State (All Cycles)` | Compare resident, ready, and issued warps; inspect the main wait states when ready work is scarce | + +`Grid Size` is the number of blocks submitted by the launch. `Waves Per SM` divides that count by the +number of blocks that could theoretically reside across the whole GPU at once. A value of 1 means +the two counts are equal; 3.46 means that the grid contains 3.46 times the device-wide theoretical +resident-block capacity. This ratio helps determine whether the grid contains enough blocks to cover +the GPU. It does not record the actual block-scheduling order. + +Each `Block Limit` reports how many blocks one SM could host if registers, shared memory, threads, or +another listed resource were the only constraint. The smallest value sets the theoretical block +limit. Theoretical occupancy is the maximum resident-warp count allowed by those limits as a fraction +of the hardware capacity. Achieved occupancy is the average active-warp count observed during +collection, expressed against the same capacity. These metrics describe resident concurrency; +scheduler metrics show whether that concurrency affects instruction issue. + +In `SpeedOfLight`, Compute represents the busiest SM compute path, Memory the busiest memory-side +path, and DRAM only the external-memory interface. Each uses its own sustainable peak as the +denominator. On a B200, HBM provides external device memory, L2 is shared across the GPU, and L1TEX is +the SM-side path that handles memory requests. A high Memory value means that some memory-side path +is busy; DRAM shows whether the external HBM interface is near saturation. In +`ComputeWorkloadAnalysis`, active cycles show when a pipeline still has work in flight, while +`Issue Slots Busy` reports how many scheduler issue opportunities were used. + +`SchedulerStats` distinguishes several warp states. An active warp is resident and unfinished. An +eligible warp has a decoded next instruction whose dependencies are ready and whose required +execution unit is available. An issued warp issued an instruction in the current cycle. After +confirming the target launch and examining the grid, residency, and `SpeedOfLight`, select the next +metrics: + +- **Compute is closer to its peak:** In `ComputeWorkloadAnalysis`, open + `Pipe Utilization (Elapsed Cycles)` → `Pipe Utilization (% of elapsed cycles)` and find the compute + path with the most active cycles. Then inspect `Issue Slots Busy` in the same section's summary. + If one path keeps processing work while most issue slots remain empty, continue to + `Scheduler Statistics` → `Warps Per Scheduler` to explain the low instruction-issue rate. +- **Memory is closer to its peak:** Collect `MemoryWorkloadAnalysis`, + `MemoryWorkloadAnalysis_Chart`, and `MemoryWorkloadAnalysis_Tables`. First follow the data path + through DRAM, L2, and L1/TEX under `Memory Workload Analysis Chart` → `Memory Chart`. Then read + throughput, read/write bytes, hit rate, and the shared- and local-memory fields under + `Memory Workload Analysis Tables` → `Memory Tables`. If DRAM is also near its peak, investigate + external HBM traffic first. If DRAM is low, shift attention to L2, L1/TEX, shared memory, or local + memory. Local memory is a per-thread private address space; its physical traffic traverses the + L1/L2 caches and external device memory. +- **Compute and Memory are both low:** Use the grid size and waves to determine whether the launch + provides enough blocks to cover the GPU. If it does not, inspect the grid, block, and cluster + configuration or divide the work into more blocks. If it does, inspect `SchedulerStats`. If active + warps exist but eligible warps are scarce, use `WarpStateStats` to see whether they are waiting on + data, synchronization, or another dependency. +- **Compute and Memory are both high:** Expand both sides, identify one specific bottleneck candidate + on each, and change one factor at a time to determine which candidate affects kernel time. + +`MemoryWorkloadAnalysis` summarizes traffic and cache behavior for the whole kernel across DRAM, L2, +L1TEX, and other memory paths. `SourceCounters` then maps sampled stalls and instruction activity to +SASS (GPU machine instructions) or source locations to help trace a specific load dependency. + +Judge “high” and “low” in the context of the current GPU, workload, and the metrics in the same +report. Once the report points to code that can be changed and predicts how its metrics and latency +should respond, the current pass has produced a testable hypothesis. If the scope is still too broad, +collect the next section along the relevant path above. + +### Worked Example: A B200 BF16 GEMM + +#### 1. Collect the First `basic` Report + +Continue using the script's `--profile-once` mode. The application submits one GEMM and then one ReLU +inside the collection range; NCU waits for the range to begin and profiles only the GEMM: ```bash mkdir -p reports ncu \ --config-file off \ - --target-processes application-only \ --profile-from-start off \ - --kernel-name-base function \ --kernel-name 'regex:.*nvjet_sm100.*' \ --launch-count 1 \ --set basic \ @@ -626,27 +723,30 @@ ncu \ --pipeline-boost-state stable \ --export reports/bf16-gemm-basic \ --force-overwrite \ - python appendix/nsys_example.py --profile-once + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once ``` -`--profile-from-start off` makes NCU wait for the profiler API range in the script. The regular -expression then selects the GEMM whose name contains `nvjet_sm100`. Generated kernel names can change -with PyTorch and CUDA versions, so copy the actual name from Nsight Systems before writing a narrower -filter for another program. - -`--set basic` collects launch, occupancy, workload-distribution, and high-level throughput sections. -Cache and clock controls are explicit because they change the profiling conditions. In particular, -`--cache-control all` flushes the GPU caches that NCU can control before every replay iteration. That -helps stabilize counter collection but does not reproduce the hot-cache policy of the headline -benchmark. Section sets and defaults can change between releases, so record `ncu --version` and -inspect `ncu --config-file off --list-sets` on the collection machine. - -If the script cannot use profiler start/stop, rerun the command with `--profile-from-start on` (or -remove `--profile-from-start off`) and use `--launch-skip N --launch-count 1` to select an invocation -after warm-up. `--launch-skip` counts matching kernel launches, so a changed filter or launch order -can select a different invocation. The -[Nsight Compute CLI documentation](https://docs.nvidia.com/nsight-compute/NsightComputeCli/) -describes the kernel and launch filters in detail. +The key options control the capture range, target, metrics, and collection conditions: + +- `--profile-from-start off` makes NCU wait until the script calls `cudaProfilerStart()`. +- `--kernel-name` selects kernels in that range whose name contains `nvjet_sm100`, and + `--launch-count 1` produces a report for the first matching launch. The expression comes from the + preceding timeline; replace it with the name from your own program. +- `--set basic` collects the launch, occupancy, and high-level throughput sections needed for this + first pass. +- `--replay-mode kernel` allows NCU to replay the selected GEMM while collecting hardware counters. + `--cache-control all` flushes controllable caches before replay, while the remaining options control + clocks and pipeline boost state during collection. Kernel replay is appropriate here because this + GEMM can be replayed independently. Workloads with cross-kernel dependencies or concurrency need a + replay mode that preserves the required application or range state. + +Here, “one GEMM” means one launch submitted by the application. NCU can still replay that GEMM +internally to collect all requested hardware counters. The 500 warm-up calls stay outside the +collection range, avoiding initialization and lazy-loading work. NCU's cache control then changes the +normal warm-cache conditions. Open the report in the GUI: @@ -654,244 +754,577 @@ Open the report in the GUI: ncu-ui reports/bf16-gemm-basic.ncu-rep ``` -or inspect it in the terminal: +Without a GUI, print the Details page in the terminal: ```bash ncu --import reports/bf16-gemm-basic.ncu-rep \ --page details \ - --print-details header \ + --print-details all \ --print-metric-name label-name ``` -`header` is a compact first view. Expand the Work ID/CLC and throughput tables in the GUI, or replace -`header` with `all` to print the complete details used below. +See the [Nsight Compute CLI documentation](https://docs.nvidia.com/nsight-compute/NsightComputeCli/) +for other filters and collection options. If NCU reports `ERR_NVGPUCTRPERM`, follow NVIDIA's +[counter-permission guidance](https://developer.nvidia.com/nvidia-development-tools-solutions-err-nvgpuctrperm-nsightcompute) +or ask the system administrator to enable access. + +#### 2. Use the `basic` Report to Choose What to Investigate + +First confirm the device, kernel name, grid and block dimensions, and warnings in the report header. +Once they match the intended launch, use these three observations to choose what to inspect next: -## Analyze an Nsight Compute Report +| Observation in `basic` | Next step in this example | +|---|---| +| `Grid Size = 512 blocks`; `Waves Per SM = 3.46` | The grid supplies enough blocks to occupy the whole GPU; next check how much work can reside on each SM | +| The register and shared-memory `Block Limit` values are both 1; theoretical/achieved occupancy is 12.50%/8.97% | Each SM can theoretically host at most eight resident warps, while the observed average is lower; use scheduler metrics to inspect instruction readiness | +| Compute 77.74%, Memory 38.71%, DRAM 12.88% | Compute is closest to its own peak, so expand the compute side first; aggregate HBM throughput across the device still has ample headroom | + +Start with the wave calculation. The current resource limits allow one resident block per SM, and +this B200 has 148 SMs, so the whole GPU's theoretical simultaneous capacity is 148 blocks. The grid +contains 512 blocks, and $512 / 148 = 3.46$: its block count is 3.46 times that theoretical capacity. +This calculation only explains the capacity ratio reported by NCU. Because this example uses +thread-block clusters, use the reported `Waves Per SM = 3.46` as the authoritative value. The ratio +confirms that the grid contains enough blocks to cover the device. + +Next consider occupancy. A B200 SM can hold at most 2,048 threads. At 32 threads per warp, that is a +hardware limit of 64 warps. This launch configuration permits at most one 256-thread block per SM, or +eight resident warps, so its theoretical occupancy is $8 / 64 = 12.50\%$. This theoretical resident +concurrency is low relative to the hardware capacity. `Achieved Occupancy = 8.97%` is the average +number of active warps observed during collection as a fraction of the same hardware capacity, below +the 12.50% theoretical maximum. The theoretical value establishes that all schedulers on one SM +have at most eight resident warps in total; `SchedulerStats` then reports how many warps assigned to +each scheduler are ready on average. + +Finally, compare throughput. Compute is closer to its peak than Memory, so investigate the compute +side first. The term `compute-bound` makes a stronger claim: further speedup is ultimately constrained +by the throughput limit of the compute units. The `basic` report only selects an entry point. The +follow-up report shows that issue opportunities are rarely used and most scheduler cycles have no +eligible (ready) warp; calling the kernel `compute-bound` at this point would hide that key clue. + +The grid can occupy the whole device, but each SM can host at most eight resident warps, and the +compute side is closest to its peak. The next report expands the compute pipelines and collects +scheduler metrics. + +#### 3. Collect Follow-Up Metrics Along the Compute Path + +The next command collects all three sections in a single run. Read them in the order +`ComputeWorkloadAnalysis` → `SchedulerStats` → `WarpStateStats`. For a new kernel, decide whether to +add each section after reading the preceding one. -The following values come from a real capture on the same B200 with Nsight Compute 2026.1. NCU used -nine replay passes to build the `basic` report. The percentages are NCU throughput metrics relative -to the sustained peak of the corresponding hardware subsystem; they are not application FLOPs -divided by the chip's advertised peak. +```bash +ncu \ + --config-file off \ + --profile-from-start off \ + --kernel-name 'regex:.*nvjet_sm100.*' \ + --launch-count 1 \ + --section ComputeWorkloadAnalysis \ + --section SchedulerStats \ + --section WarpStateStats \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --export reports/bf16-gemm-followup \ + --force-overwrite \ + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once +``` -| Metric | Measured value | +This is an independent NCU run. Percentages can vary slightly between the reports—for example, +77.74% becomes 78.39%—without indicating a performance change. The earlier `basic` report selected +the investigation path; the follow-up report uses compute, scheduler, and warp-state metrics from +one collection to narrow the cause. + +#### 4. Read the Three Sections in Order + +Start at `ComputeWorkloadAnalysis` → `Pipe Utilization (Elapsed Cycles)` → +`Pipe Utilization (% of elapsed cycles)`, the active-cycle view. `Tensor (FP)` is the floating-point +tensor-compute path. `TMEM (Tensor Memory)` is an on-chip memory path that serves tensor operations; +it is distinct from external HBM/DRAM and from TMA, the asynchronous data-movement engine. These +paths are active during about 78% of clock cycles, while the same section's summary reports +`Issue Slots Busy = 3.20%`. A multi-cycle operation can issue once and keep a pipeline active, so +frequent pipeline activity and infrequent new instruction issue can occur together. + +Next, open `Scheduler Statistics` → `Warps Per Scheduler` to see why instructions issue so +infrequently. Each scheduler averages 1.44 active warps that have not finished, but only 0.04 +eligible (ready) warps; 0.04 is an average warp count, not 4%. The denominator for `No Eligible` +includes only cycles in which the scheduler's SM subpartition has at least one warp in flight. Its +value of 96.11% means that no eligible warp is available in 96.11% of those cycles. This explains +the low 3.20% issue rate. Work resides on the SM, but most of the time no warp can continue. + +Finally, open `Warp State Statistics` → `Warp State (All Cycles)` to see what those warps are waiting +for. One warp spending one clock cycle in a state contributes one warp-cycle. Each issued warp +instruction corresponds to 37.00 warp-cycles, of which 32.11—about 87%—are attributed to +`Long Scoreboard`. This 87% is the share of warp-state cycles after normalization by issued warp +instructions, not a share of kernel execution time. A scoreboard is the hardware +dependency table that records whether results of earlier operations are ready. `Long Scoreboard` +means that the next instruction is still waiting for a memory operation handled by L1TEX. L1TEX is +the SM-side path for global, local, surface, and texture memory requests; the requested data may +ultimately come from L1, L2, or DRAM, so this field alone cannot identify the serving memory level. +`MemoryWorkloadAnalysis` characterizes aggregate L1, L2, and DRAM behavior for the whole kernel; use +the SASS/source view in `SourceCounters` to continue locating the specific load. + +The three sections now form one reading path: + +```text +Tensor/TMEM paths are often active +→ the scheduler issues few new instructions +→ most cycles have no eligible (ready) warp +→ Long Scoreboard is the largest wait category +``` + +Start with the L1TEX-related data dependency. Each SM can host at most eight resident warps, which may +make data-access latency harder to cover with other work. In the earlier `basic` report, +`DRAM Throughput = 12.88%` shows that aggregate HBM bandwidth is far from saturation; individual +requests can still reach DRAM and incur long latency. Investigate the data dependency first, then see +whether Tensor Core throughput becomes the next limit. The `nvjet` GEMM in this example is a library +implementation, so the next report can collect `MemoryWorkloadAnalysis` and compare aggregate L1, +L2, and DRAM traffic, throughput, and cache behavior for the whole kernel. + +For a kernel you control, the report suggests two directions that can be tested separately: + +- **Increase resident warps:** Adjust the tile, block, or resource use so that more work can reside + on each SM. +- **Shorten the data-dependency wait:** Hold residency constant and move the load or prefetch + earlier, or shorten the dependency chain. + +The later section “Test the Hypothesis with a Code Change” gives the concrete modifications and +validation steps. + +### Metric Calculations, Units, and Boundaries + +The main walkthrough already established the reading order and the conclusion for this kernel. The +following sections retain only complete field lists, calculations, units, and boundaries that are +easy to misread; consult them as needed for another kernel. + +#### `LaunchStats` and `Occupancy` + +The complete launch fields used by this example appear in `LaunchStats`: + +| Field | Value in this report | +|---|---:| +| `Grid Size` | 512 blocks | +| `Block Size` | 256 threads | +| `Cluster Size` | 4 blocks | +| `Waves Per SM` | 3.46 | + +- `Block Size = 256` means that each block has 256 threads, or eight warps. This value enters the + occupancy calculation below. +- `Cluster Size = 4` groups every four blocks into one thread-block cluster, giving 128 clusters in + the grid. The cluster's blocks are scheduled together according to a layout supported by the + hardware. +- For cluster launches, including this one, use the `Waves Per SM` value reported by NCU; use + `cudaOccupancyMaxActiveClusters` to calculate resident clusters programmatically. + +Per-block resource use appears in `LaunchStats`: + +| Field | Value in this report | |---|---:| -| Kernel duration | 95.87 μs | -| Grid / block size | 512 blocks / 256 threads | -| Cluster size | 4 blocks | -| Registers | 255 / thread | -| Dynamic shared memory | 213.28 KB / block | -| Waves per SM | 3.46 | -| Theoretical / achieved occupancy | 12.50% / 8.98% | -| SM compute-throughput metric | 77.34% | -| Memory-throughput metric | 38.51% | -| DRAM / L2 / L1-TEX throughput metrics | 20.42% / 34.60% / 46.93% | +| `Registers Per Thread` | 255 | +| `Dynamic Shared Memory Per Block` | 213.28 KB | -NCU's 95.87 μs differs slightly from the 93.152 μs captured by Nsight Systems above. The values come -from separate profiling runs, and NCU also changes cache, clock, and replay conditions. This is why -the report's `Duration` cannot replace the headline benchmark. +The `Occupancy` section shows the residency permitted by those resources: -Read the table in three passes: verify the selected launch, see how its work covers the GPU, and only -then choose which throughput breakdown to expand. +| Field | Value in this report | +|---|---:| +| `Block Limit Registers` | 1 block / SM | +| `Block Limit Shared Mem` | 1 block / SM | +| `Theoretical Occupancy` | 12.50% | +| `Achieved Occupancy` | 8.97% | + +When a block becomes resident, the SM reserves its registers and shared memory. The +[NVIDIA Blackwell Tuning Guide](https://docs.nvidia.com/cuda/blackwell-tuning-guide/index.html#occupancy) +specifies 65,536 32-bit registers, 228 KB of shared memory, and at most 2,048 resident threads per +B200 SM. + +- `Registers Per Thread = 255`. With 256 threads, one block needs about + $255 \times 256 = 65{,}280$ registers, nearly the entire register file. Two blocks would need + 130,560 registers, exceeding 65,536. +- `Dynamic Shared Memory Per Block = 213.28 KB`. One block already uses most of the 228 KB available; + two blocks would need at least 426.56 KB and cannot reside together. +- `Block Limit Registers = 1` and `Block Limit Shared Mem = 1` are the result of those two resource + calculations. Either resource alone permits only one resident block per SM. NCU's exact calculation + also accounts for allocation granularity and driver-reserved shared memory. +- `Achieved Occupancy = 8.97%` is the average active-warp ratio observed during collection, not + “percent of peak kernel performance.” It is below 12.50%, so execution did not sustain the + theoretical residency limit throughout the launch. The two percentages measure residency, not + closeness to peak performance. + +#### `SpeedOfLight` Denominators and `Duration` + +The `SpeedOfLight` section reports these four fields: + +| Field | Value in this report | +|---|---:| +| `Duration` | 95.30 μs | +| `Compute (SM) Throughput` | 77.74% | +| `Memory Throughput` | 38.71% | +| `DRAM Throughput` | 12.88% | -### 1. Launch Statistics and Workload Distribution +Treat `Duration` as specific to its profiler run. This NCU capture reports 95.30 μs, while a separate +Nsight Systems capture reports 92.608 μs. Compare implementations with the unprofiled CUDA Event +baseline. NCU also controls clocks, flushes caches, and may replay or serialize kernels; see +[Nsight Compute's workload-duration guidance](https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#workload-durations). -The captured name is `nvjet_sm100_tst_128x256_64x6_2x2_2cta_h_bz_NNT`, which matches the GEMM in the -timeline above. It launches 512 blocks of 256 threads and groups four blocks into each cluster. -`Waves Per SM = 3.46` means that the grid requires three full waves and one partial wave. It describes -how the grid covers the GPU over time; it is not occupancy. +The three throughput percentages use separate hardware peaks as their denominators. They cannot be +added and do not represent fractions of execution time. The +[Nsight Compute Profiling Guide](https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-structure) +defines throughput metrics and their constituents. -This report also carries a Work ID/Cluster Launch Control warning. Although the nominal launch has -512 CTAs, only 380 were granted. When this warning appears, treat metrics derived from block, warp, -or thread counts cautiously rather than assuming that the nominal launch count is the executed -count. +#### Compute Pipelines, the Scheduler, and Warp States -### 2. Occupancy +##### `Compute Throughput Breakdown` Fields -This kernel uses 255 registers per thread and 213.28 KB of dynamic shared memory per block. Both the -register and shared-memory limits allow only one block to reside on an SM. The resulting theoretical -occupancy is 12.50%, and the measured achieved occupancy is 8.98%. +The report location is `SpeedOfLight` → `GPU Throughput Breakdown` → +`Compute Throughput Breakdown`: -Those values show that few warps are resident; they do not establish occupancy as the bottleneck. -This GEMM deliberately uses four-CTA clusters and an asynchronous pipeline. Reducing registers or -shared memory merely to raise occupancy can introduce spills or sacrifice tile reuse and make the -kernel slower. +| Field | Value in this report | +|---|---:| +| `SM: Mem Tensor Cycles Active` | 77.74% | +| `SM: Pipe Tc Cycles Active` | 77.48% | +| `SM: Pipe Tensor Cycles Active` | 77.42% | +| `SM: Pipe Alu Cycles Active` | 1.36% | +| `SM: Pipe Tma Cycles Active` | 1.10% | +| `SM: Pipe Fma Cycles Active` | 0.61% | + +- `Mem Tensor` is an on-chip path related to Blackwell tensor memory. External data is stored in + DRAM/HBM, while TMA handles asynchronous multidimensional transfers; these three names refer to + distinct hardware paths. At 77.74%, `Mem Tensor` is the most heavily utilized compute-side path in + this report. +- `Pipe Tc` and `Pipe Tensor` are two distinct pipeline counters in NCU. Both are near 77%, + consistent with a BF16 GEMM performing substantial tensor-MMA-related work. They may cover + overlapping hardware activity, so read them separately; summing them would double-count activity. +- `Pipe Alu` primarily corresponds to general integer and logic operations, while `Pipe Fma` covers + ordinary FP32 arithmetic and some integer multiply-add operations. At 1.36% and 0.61%, respectively, + this GEMM does not approach the peak of either path. +- `Pipe Tma` corresponds to Tensor Memory Accelerator's asynchronous data-movement path. Its 1.10% + value shows that TMA is far from its own peak. Data supply also involves TMEM, caches, shared + memory, and dependency latency, each described by other metrics. + +##### The Two `Pipe Utilization` Denominators + +The `ComputeWorkloadAnalysis` summary reports `Issue Slots Busy = 3.20%`. The two full view names are +`Pipe Utilization (% of elapsed cycles)` and +`Pipe Utilization (% of peak instructions executed over elapsed cycles)`. + +Putting each pipeline on one row makes the contrast easier to see: + +| Pipeline field | Active-cycle view | Instruction-rate view | +|---|---:|---:| +| `TMEM (Tensor Memory)` | 78.39% | 0.04% | +| `TC` | 78.12% | 0.38% | +| `Tensor (FP)` | 78.07% | 0.61% | + +The views use different denominators to compare occupied pipeline cycles with instruction execution +rate; do not add or subtract their values. + +##### Complete `SchedulerStats` Fields + +The report location is `Scheduler Statistics` → `Warps Per Scheduler`: + +| Field | Value in this report | +|---|---:| +| `GPU Maximum Warps Per Scheduler` | 16 | +| `Theoretical Warps Per Scheduler` | 2.00 | +| `Active Warps Per Scheduler` | 1.44 | +| `Eligible Warps Per Scheduler` | 0.04 | +| `Issued Warp Per Scheduler` | 0.04 | -NCU also emits rule-based `Est. Speedup` suggestions. They are local upper bounds under simplified -assumptions and are useful as investigation prompts, not as expected speedups from changing the -kernel. +The summary in the same section also reports `No Eligible = 96.11%`. -### 3. Speed of Light +`GPU Maximum = 16` is the hardware capacity of one scheduler. `Theoretical = 2.00` comes from this +kernel's theoretical maximum of eight resident warps per SM divided among four schedulers. -The basic report shows an SM compute-throughput metric of 77.34% and a memory-throughput metric of -38.51%, with DRAM at only 20.42%. The evidence therefore does not support calling the kernel -DRAM-bound; expanding the compute pipelines is the more useful next step. +##### `WarpStateStats` Normalization -For another kernel, compare the high-level compute and memory metrics against their respective -sustained peaks: +One warp spending one cycle in a state contributes one warp-cycle; four warps spending the same cycle +in a state contribute four warp-cycles. NCU then normalizes those cycles by the number of issued warp +instructions. The report summary gives the first field below, and `Warp State (All Cycles)` gives the +second: -- high compute and lower memory throughput suggests a compute-pipeline limit; -- high memory and lower compute throughput suggests investigating the memory hierarchy; -- both low suggests underfill, dependency latency, synchronization, imbalance, or too few eligible - warps before it suggests a peak-throughput limit. +| Field | Value in this report | +|---|---:| +| `Warp Cycles Per Issued Instruction` | 37.00 warp-cycles / issued instruction | +| `Stall Long Scoreboard` | 32.11 warp-cycles / issued instruction | -"Memory throughput" is not synonymous with DRAM throughput. Its limiting contributor can be L1, -L2, shared memory, or a memory-instruction pipeline. Expand the breakdown before calling a kernel -DRAM-bound. +The report averages 37.00 warp-cycles per issued warp instruction. Of those, 32.11, or about 86.8%, +are assigned to `Long Scoreboard`. These values are normalized across all warps and use +warp-cycles / issued instruction, a different unit from the scheduler's average issue rate. -### Choose the Next Metrics from the Basic Report +The `Est. Speedup` shown next to an NCU rule is a model-based estimate of the potential reduction in +workload time for that rule. Use it to prioritize the investigation; only a benchmark of the modified +kernel can establish the actual speedup. -The `basic` report covers steps 1–3. Use its evidence to collect only the sections needed for the -next question: +##### Other Common Warp States -| Evidence from the basic report | Add next | -|---|---| -| Register, shared-memory, or resident-block limit | `LaunchStats`, `Occupancy` are already in `basic`; inspect their limit tables before collecting more | -| Compute path appears dominant | `ComputeWorkloadAnalysis` | -| Memory hierarchy appears dominant | `MemoryWorkloadAnalysis`; add `_Chart` for the visual breakdown or `_Tables` for detailed requests and sectors | -| Too few eligible warps or unexplained issue gaps | `SchedulerStats`, then `WarpStateStats` | -| A source or instruction location is required | `SourceCounters` | +| Warp state | Direct meaning | What to inspect next | +|---|---|---| +| `Short Scoreboard` | Usually waiting for shared memory or another on-chip unit to produce a result | Inspect shared-memory accesses and the corresponding source | +| `Barrier` | Waiting for other warps to reach a synchronization point | Compare the work assigned to different warps and when they arrive | +| `Not Selected` | The warp is ready, but another warp was selected in this cycle | Check whether many ready warps are competing for issue opportunities | + +#### Locate SASS or Source with `SourceCounters` + +`WarpStateStats` shows what the kernel spends time waiting for as a whole. `SourceCounters` takes the +next step by placing sampled waits and execution counts beside individual SASS instructions. The +Source page then reveals where those waits are concentrated. If the binary contains line information +and NCU can find the source file, the instructions are also correlated with CUDA source lines. The +data comes from periodic sampling of warp stall reasons together with instruction counts and selected +memory-access metrics. -The compute metric is higher in this report, so add `ComputeWorkloadAnalysis` for the same isolated -launch: +The `nvjet` GEMM is a library implementation, so this tutorial has no CUDA source file to import for +it. The SASS view is still available. Reuse the earlier filter and collection conditions: ```bash -mkdir -p reports ncu \ --config-file off \ - --target-processes application-only \ --profile-from-start off \ - --kernel-name-base function \ --kernel-name 'regex:.*nvjet_sm100.*' \ --launch-count 1 \ - --section ComputeWorkloadAnalysis \ + --section SourceCounters \ --replay-mode kernel \ --cache-control all \ --clock-control boost \ --pipeline-boost-state stable \ - --export reports/bf16-gemm-compute \ + --export reports/bf16-gemm-source \ --force-overwrite \ - python appendix/nsys_example.py --profile-once + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once ``` -The follow-up report contains these values: +Select SASS on the GUI's Source page, or print the same view in the terminal: -| Pipeline | Throughput metric | -|---|---:| -| TMEM | 77.23% | -| Tensor Core | 77.04% | -| Tensor FP | 76.90% | -| ALU / TMA / FMA | all below 2% | +```bash +ncu --import reports/bf16-gemm-source.ncu-rep \ + --page source \ + --print-source sass +``` -That evidence resolves the basic report's aggregate 77.34% compute metric to the Tensor Core and -Tensor Memory path. For another hypothesis, keep the command shape and replace the `--section` lines -rather than accumulating every section in one report. This keeps the report smaller and reduces -replay overhead. +Start with `Warp Stall Sampling (Not-issued Samples)` and `Instructions Executed`. The first counts +sampling observations taken when the warp scheduler issued no instruction; the second counts +executions of the corresponding SASS instruction per warp. If `WarpStateStats` was dominated by +`Long Scoreboard` and the Source page concentrates corresponding samples near one load, that +instruction becomes a candidate for closer inspection. These values come from periodic sampling and +identify hotspot locations. Determining whether the data came from L1, L2, or DRAM still requires +the kernel-wide evidence in `MemoryWorkloadAnalysis` together with the code's access relationships. -### 4. Compute and Memory Workload Analysis +For a TIRx kernel that you compile yourself, the SASS can also be correlated with generated CUDA. +First select NVCC, retain the generated source, and enable line information: -Compute Workload Analysis identifies which execution pipelines are active. Check the Tensor Core, -FMA, ALU, special-function, and relevant asynchronous pipelines rather than inferring Tensor Core -utilization from one aggregate compute percentage. +```bash +export TVM_CUDA_COMPILE_MODE=nvcc +export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels" +mkdir -p "$TVM_KERNEL_DUMP" +``` -Memory Workload Analysis separates DRAM, L2, L1/TEX, shared memory, and local-memory effects. Read -traffic volume together with bandwidth, cache hit rate, and local-memory spill. Detailed sector and -request tables require `MemoryWorkloadAnalysis_Tables`; source-level coalescing and shared-memory -conflict evidence can require `SourceCounters`. A high cache hit rate alone says little when the -traffic volume is small. +After setting the variables, restart the workload so that the target kernel is recompiled in that +process. The following is a collection-command template; replace `YOUR_KERNEL_NAME` and the program +path on the final line: + +```bash +ncu \ + --config-file off \ + --kernel-name 'regex:.*YOUR_KERNEL_NAME.*' \ + --launch-count 1 \ + --section SourceCounters \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --import-source yes \ + --source-folders "$TVM_KERNEL_DUMP" \ + --export reports/tirx-source \ + --force-overwrite \ + python path/to/your_workload.py +``` -### 5. Scheduler and Warp States +Open `ncu-ui reports/tirx-source.ncu-rep` and use the CUDA/SASS correlation on the Source page. The +terminal equivalent is +`ncu --import reports/tirx-source.ncu-rep --page source --print-source cuda,sass`. +If `executable` is your TIRx compilation result, +`executable.mod.imports[0].inspect_source("cuda")` prints the generated source for manual inspection. +NCU's line-by-line correlation depends on line information embedded in the binary during this +recompilation. + +#### Test the Hypothesis with a Code Change + +Because this example calls the library-provided `torch.mm`, the kernel itself is not editable here. +The following steps give concrete modifications and validation checks for a custom TIRx kernel or +another DSL kernel. + +To test the hypothesis that the kernel has too few resident warps to hide L1TEX latency, adjust the +tile, block, or pipeline stages to reduce `Registers Per Thread` and +`Dynamic Shared Memory Per Block`. Then read both block limits again. A second block can reside only +if both `Block Limit Registers` and `Block Limit Shared Mem` rise from one to at least two; improving +only one is insufficient. Reducing register use can cause spills into local memory, and reducing +shared-memory use may reduce data reuse. Also confirm that every other block limit is at least two, +then inspect NCU's recalculated theoretical active blocks and occupancy. Use the latency measurement +to determine whether the trade-off helps overall. + +For a separate experiment, hold residency constant while moving a load or prefetch earlier or +shortening the dependency chain. If `Long Scoreboard` and latency fall together, that supports the +explanation that warps now spend less time waiting for data. Because this metric is normalized per +issued instruction, a decrease in this value alone does not establish a speedup. Change one key +factor at a time, then check three things: + +1. **Correctness:** Compare both outputs with the same reference on the same inputs and with the same + tolerance. The current script already performs an FP32 reference check before timing or collection + begins; use the same reference and tolerance after substituting your own implementation. +2. **Predicted metrics:** Collect the same NCU sections again and inspect the metrics relevant to the + prediction: for a residency experiment, inspect theoretical/active and eligible/issued warps; + for a dependency-chain experiment, inspect `Long Scoreboard`. Higher occupancy shows that the + residency experiment reached its resource target; the next check determines whether that change + also improves latency. +3. **Actual latency:** Disable Proton, Nsight Systems, and NCU. Measure both implementations with + exactly the same shape, dtype, input policy, warm-up, CUDA Event boundary, and sample count used + at the beginning. -Scheduler Statistics shows active, eligible, and issued warps. First determine whether schedulers -often have no eligible instruction to issue. Only then use Warp State Statistics to investigate why. -The NCU guide explicitly warns that stalls are not all avoidable and do not automatically limit -performance. +```bash +python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --event-samples 20 +``` -Common states should be interpreted as clues: +Compare the unprofiled medians before and after the change, and inspect sample variation as well. If +the correctness check passes, the metrics change as predicted, and CUDA Event latency decreases +consistently, the hypothesis is supported. If the NCU metrics move as expected but latency does not +improve, check whether another bottleneck has emerged or the original hypothesis was incomplete. -| State | Useful interpretation | Do not conclude from it alone | -|---|---|---| -| Long Scoreboard | Waiting on a dependency associated with the L1TEX path | Every wait reached DRAM | -| Short Scoreboard | Waiting on an MIO-path dependency, often involving shared memory | A bank conflict definitely exists | -| Barrier | Waiting for a synchronization dependency | The barrier is unnecessary | -| Not Selected | The warp was eligible but another warp issued | The scheduler is starved | -| Math/MIO Throttle | A pipeline or queue is under pressure | Removing arbitrary instructions will improve runtime | +## Use IKET to Inspect Phases Inside a DSL Kernel -For warp-specialized kernels, aggregate stall percentages also combine roles with intentionally -different behavior. Relate the result to the producer, MMA, softmax, or writeback role before changing -synchronization. +IKET (In-Kernel Event Tracing) adds an internal timeline to a warp-specialized TIRx kernel. Nsight +Systems shows the beginning and end of the whole kernel, NCU aggregates hardware metrics across the +launch, and IKET records when each warp role executes producer, wait, consumer, or other marked +regions. -### 6. Source and SASS Correlation +### Run a Complete Example -The SASS view and instruction attribution do not require CUDA line information. Correlating generated -CUDA source back to SASS does: the binary needs line information, and NCU must be able to find the -source file. For a TIRx module compiled through NVCC, dump the generated source before compilation: +TVM 0.26 uses a version-pinned `cutlass-4.6.0` profiling profile. For the CUDA 13 environment used in +this chapter, install the matching dependencies and confirm that `run-iket` is available: ```bash -export TVM_CUDA_COMPILE_MODE=nvcc -export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels" -mkdir -p "$TVM_KERNEL_DUMP" +python -m pip install \ + 'nvidia-cutlass-dsl[cu13]==4.6.0' \ + 'nvidia-cuda-nvdisasm==13.3.73' \ + 'nvidia-cuda-nvrtc==13.2.78' +run-iket --help ``` -When `TVM_KERNEL_DUMP` is set, TVM retains the generated files and passes `-lineinfo` to NVCC. Add -`--import-source yes --source-folders "$TVM_KERNEL_DUMP"` to the NCU collection command. Saving -`inspect_source("cuda")` is still useful for manual comparison, but by itself it cannot add line -information to a compiled binary. A Python line may lower to many CUDA or SASS instructions, and an -asynchronous tile primitive may be understandable only in those lower-level views. +The complete script below is available as `appendix/iket_example.py`. One CTA contains two warps. +Warp 0 moves 256 elements from global to shared memory, both warps meet at a CTA barrier, and warp 1 +reads shared memory, computes the result, and writes it out. Three `range_push()` / `range_pop()` +pairs mark the producer, wait, and consumer regions: -### Effects of NCU Collection on Experimental Conditions +```python +"""Minimal TIRx workload with IKET ranges for two warp roles.""" -NCU collection can change the execution conditions: +from pathlib import Path -- it can replay a kernel to collect counter groups; -- its default cache control can flush GPU caches between replay iterations; -- it can control GPU clocks; -- replay can serialize or otherwise alter concurrent work; -- application replay reruns the entire program and requires deterministic execution and launch - matching; it is not a remedy for a nondeterministic launch order; -- a dependent multi-kernel region may require range replay rather than replaying one kernel in - isolation. +import numpy as np -Record the NCU version, replay mode, cache control, clock control, selected sections, and kernel -filter. Do not compare NCU's `Duration` column directly with an unprofiled hot-cache CUDA Event result. -Do not run Proton and NCU in the same profiling process. +import tvm +from tvm.script import tirx as T +from tvm.tirx.cuda import iket -If NCU reports `ERR_NVGPUCTRPERM`, hardware-counter access is restricted. Follow NVIDIA's -[counter-permission guidance](https://developer.nvidia.com/nvidia-development-tools-solutions-err-nvgpuctrperm-nsightcompute) -or ask the system administrator to enable the required access; do not make running every experiment as -root the default solution. -## Analyze In-Kernel Stages with IKET (Optional) +N = 256 +ELEMS_PER_LANE = 8 -Nsight Systems represents one kernel as a single GPU activity, while NCU aggregates hardware metrics -across the kernel. Neither produces a named timeline for load, compute, and wait phases inside a -warp-specialized kernel. A TIRx kernel with phase annotations can use IKET (In-Kernel Event Tracing) -to show when different warp roles work, wait, or overlap. -TVM 0.26 integrates IKET for SM90-or-newer CUDA targets and validates a strict set of CUTLASS DSL, -NVRTC, and related tool versions. The instrumentation changes the generated kernel, so its timing is -useful for understanding phase relationships, not for reporting latency. See -[`python/tvm/backend/cuda/iket.py`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/backend/cuda/iket.py) -and the -[NVIDIA IKET guide](https://github.com/NVIDIA/cutlass/blob/v4.6.0/media/docs/pythonDSL/cute_dsl_general/iket_profiling.rst) -for the required versions, annotations, and Perfetto trace workflow. +@T.prim_func +def warp_role_example(inp: T.Buffer((N,), "float32"), out: T.Buffer((N,), "float32")): + T.device_entry() + profiler = iket.IketProfiler() + warp_id = T.warp_id([2]) + lane = T.lane_id([32]) + shared = T.alloc_buffer((N,), "float32", scope="shared") -## Benchmark Checklist + profiler.mark("kernel_start", warp_id) + if warp_id == 0: + profiler.range_push("producer_load") + for i in T.serial(ELEMS_PER_LANE, unroll=False): + index = i * 32 + lane + shared[index] = inp[index] + profiler.range_pop() -Before publishing a benchmark table or pull request, use this checklist to ensure that another person -can reconstruct the measurement: + profiler.range_push("wait_for_data") + T.cuda.cta_sync() + profiler.range_pop() -| Category | Record | -|---|---| -| Hardware | Exact GPU, number of devices, topology when relevant, clock and power policy | -| Software | Driver, CUDA, framework, compiler, library versions, and source commit | -| Workload | Shapes, dtype, layouts, mask, scale, epilogue, input distribution, batch/sequence details, state and reset policy | -| Correctness | Reference, tolerance, accumulation and output dtype, exceptional-input policy | -| Timing | Timer type, kernel/operator/end-to-end boundary, stream policy, CUDA Graph use, values and units for `warmup`/`repeat`, `rounds`, and raw per-round results | -| Cache | Reused inputs, rotating inputs, explicit flush policy, and whether the policy models the application | -| Statistics | Raw latency unit, median or mean, spread, independent runs, implementation order | -| Baseline | Library and algorithm, workspace, tuning budget, selected configuration | -| Profiling | Proton/IKET/Nsight Systems/NCU versions, kernel filters, IKET ranges and trace format, Nsight Systems capture options and trace, NCU sections and replay/cache/clock controls | - -The final workflow is deliberately circular. A benchmark establishes that a change matters; a profile -suggests why; the next unprofiled benchmark determines whether the explanation led to a real -improvement. + if warp_id == 1: + profiler.range_push("consumer_compute") + for i in T.serial(ELEMS_PER_LANE, unroll=False): + index = i * 32 + lane + out[index] = shared[index] * T.float32(2) + T.float32(1) + profiler.range_pop() + + +def profile_workload(): + target = tvm.target.Target({"kind": "cuda", "arch": "sm_100a"}) + executable = tvm.compile(warp_role_example, target=target, tir_pipeline="tirx") + module = executable.jit() + + input_numpy = np.arange(N, dtype=np.float32) + inp = tvm.runtime.tensor(input_numpy, device=tvm.cuda()) + out = tvm.runtime.empty((N,), "float32", device=tvm.cuda()) + module.main(inp, out) + tvm.cuda().sync() + + expected = input_numpy * 2 + 1 + np.testing.assert_array_equal(out.numpy(), expected) + + +def main(): + result = iket.run( + profile_workload, + output_dir=Path("reports/iket-warp-roles"), + postprocess="all", + clobber=True, + timeout=600.0, + ) + print(f"IKET output directory: {result.output_dir}") + for path in (*result.json_traces, *result.perfetto_traces, *result.html_reports): + print(f"IKET artifact: {path}") + + +if __name__ == "__main__": + main() +``` + +Run it directly on a B200: + +```bash +python appendix/iket_example.py +``` + +`iket.run` restarts the current script inside the IKET collection process and calls +`profile_workload()`. Keeping `tvm.compile()` and `.jit()` inside that function ensures that the +kernel is compiled and loaded while IKET recording is active. The script also verifies that the +output equals `input * 2 + 1`. + +With `postprocess="all"`, the `reports/iket-warp-roles` directory receives JSON, `*.pftrace`, and HTML +artifacts. Load the `*.pftrace` file in Perfetto to inspect `producer_load`, `wait_for_data`, and +`consumer_compute`. Warp 1 reaches the barrier before warp 0 and therefore usually has a longer +`wait_for_data` region. For an H100, change `sm_100a` in the script to `sm_90a`. + +### Move the Annotations into Your Kernel + +Create an `IketProfiler` inside the `PrimFunc`. Use `mark()` for an instantaneous event, and use +matched `range_push()` / `range_pop()` or `range_start()` / `range_end()` calls around a phase. Every +warp's actual control-flow path must keep ranges balanced. Mark waiting explicitly, as the example +does with `wait_for_data`. + +Keep compilation and the first JIT load inside the function passed to `iket.run`. IKET supports +Hopper and newer architectures and validates the CUTLASS DSL packages, NVRTC, `nvdisasm`, and related +binaries against the pinned profile. The recording code inserted by IKET changes the generated +kernel and adds overhead, so use the IKET trace to study phases and warp roles. Measure formal +latency with the uninstrumented CUDA Event benchmark. See +[`python/tvm/backend/cuda/iket.py`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/backend/cuda/iket.py) +and the +[NVIDIA IKET guide](https://github.com/NVIDIA/cutlass/blob/v4.6.0/media/docs/pythonDSL/cute_dsl_general/iket_profiling.rst) +for the complete API and trace options. diff --git a/appendix/iket_example.py b/appendix/iket_example.py new file mode 100644 index 00000000..99203986 --- /dev/null +++ b/appendix/iket_example.py @@ -0,0 +1,73 @@ +"""Minimal TIRx workload with IKET ranges for two warp roles.""" + +from pathlib import Path + +import numpy as np + +import tvm +from tvm.script import tirx as T +from tvm.tirx.cuda import iket + + +N = 256 +ELEMS_PER_LANE = 8 + + +@T.prim_func +def warp_role_example(inp: T.Buffer((N,), "float32"), out: T.Buffer((N,), "float32")): + T.device_entry() + profiler = iket.IketProfiler() + warp_id = T.warp_id([2]) + lane = T.lane_id([32]) + shared = T.alloc_buffer((N,), "float32", scope="shared") + + profiler.mark("kernel_start", warp_id) + if warp_id == 0: + profiler.range_push("producer_load") + for i in T.serial(ELEMS_PER_LANE, unroll=False): + index = i * 32 + lane + shared[index] = inp[index] + profiler.range_pop() + + profiler.range_push("wait_for_data") + T.cuda.cta_sync() + profiler.range_pop() + + if warp_id == 1: + profiler.range_push("consumer_compute") + for i in T.serial(ELEMS_PER_LANE, unroll=False): + index = i * 32 + lane + out[index] = shared[index] * T.float32(2) + T.float32(1) + profiler.range_pop() + + +def profile_workload(): + target = tvm.target.Target({"kind": "cuda", "arch": "sm_100a"}) + executable = tvm.compile(warp_role_example, target=target, tir_pipeline="tirx") + module = executable.jit() + + input_numpy = np.arange(N, dtype=np.float32) + inp = tvm.runtime.tensor(input_numpy, device=tvm.cuda()) + out = tvm.runtime.empty((N,), "float32", device=tvm.cuda()) + module.main(inp, out) + tvm.cuda().sync() + + expected = input_numpy * 2 + 1 + np.testing.assert_array_equal(out.numpy(), expected) + + +def main(): + result = iket.run( + profile_workload, + output_dir=Path("reports/iket-warp-roles"), + postprocess="all", + clobber=True, + timeout=600.0, + ) + print(f"IKET output directory: {result.output_dir}") + for path in (*result.json_traces, *result.perfetto_traces, *result.html_reports): + print(f"IKET artifact: {path}") + + +if __name__ == "__main__": + main() diff --git a/appendix/nsys_example.py b/appendix/nsys_example.py index 9fe981fb..9a82e97f 100644 --- a/appendix/nsys_example.py +++ b/appendix/nsys_example.py @@ -1,4 +1,4 @@ -"""Small multi-stage CUDA workload for the Nsight Systems appendix example.""" +"""Reusable CUDA workload for the benchmarking and profiling appendix.""" import argparse from statistics import median @@ -7,21 +7,23 @@ def make_workload(size: int): - host_a = torch.randn((size, size), dtype=torch.bfloat16, pin_memory=True) - a = torch.empty_like(host_a, device="cuda") + a = torch.randn((size, size), dtype=torch.bfloat16, device="cuda") b = torch.randn((size, size), dtype=torch.bfloat16, device="cuda") c = torch.empty((size, size), dtype=torch.bfloat16, device="cuda") output = torch.empty_like(c) def run(): - with torch.cuda.nvtx.range("H2D input"): - a.copy_(host_a, non_blocking=True) with torch.cuda.nvtx.range("BF16 GEMM"): torch.mm(a, b, out=c) with torch.cuda.nvtx.range("ReLU"): torch.clamp_min(c, 0, out=output) - return run, output + def validate(): + torch.set_float32_matmul_precision("highest") + expected = torch.clamp_min(torch.mm(a.float(), b.float()), 0).to(output.dtype) + torch.testing.assert_close(output, expected, rtol=2e-2, atol=1e-2) + + return run, validate def run_once_for_profiler(run, *, warmup_calls: int): @@ -54,16 +56,41 @@ def measure_event_us(run, *, warmup_calls: int, samples: int): return values +def collect_proton(run, *, warmup_calls: int, profile_calls: int, output: str): + import triton.profiler as proton + + for _ in range(warmup_calls): + run() + torch.cuda.synchronize() + + session = proton.start(output, context="shadow", data="tree") + if session is None: + raise RuntimeError("Proton session could not be created") + try: + with proton.scope("target_operation"): + for _ in range(profile_calls): + run() + torch.cuda.synchronize() + finally: + proton.finalize(session) + + def main(): parser = argparse.ArgumentParser() mode = parser.add_mutually_exclusive_group() mode.add_argument("--profile-once", action="store_true") mode.add_argument("--event-samples", type=int, default=0) + mode.add_argument("--proton-calls", type=int, default=0) parser.add_argument("--size", type=int, default=4096) - parser.add_argument("--warmup-calls", type=int, default=5) + parser.add_argument("--warmup-calls", type=int, default=500) + parser.add_argument("--proton-output", default="operator") args = parser.parse_args() - run, output = make_workload(args.size) + run, validate = make_workload(args.size) + run() + torch.cuda.synchronize() + validate() + if args.profile_once: run_once_for_profiler(run, warmup_calls=args.warmup_calls) elif args.event_samples: @@ -76,13 +103,17 @@ def main(): f"median={median(values):.3f} us, " f"min={min(values):.3f} us, max={max(values):.3f} us" ) + elif args.proton_calls: + collect_proton( + run, + warmup_calls=args.warmup_calls, + profile_calls=args.proton_calls, + output=args.proton_output, + ) else: run() torch.cuda.synchronize() - if not torch.isfinite(output).all().item(): - raise RuntimeError("workload produced a non-finite output") - if __name__ == "__main__": main() diff --git a/img/nsys_b200_timeline.svg b/img/nsys_b200_timeline.svg index 91310f96..8eec35da 100644 --- a/img/nsys_b200_timeline.svg +++ b/img/nsys_b200_timeline.svg @@ -3,28 +3,26 @@ One Real Nsight Systems Capture on B200 -4096×4096 BF16: H2D copy → GEMM → ReLU +4096×4096 BF16: GEMM → ReLU 0 μs - -100 μs - -200 μs - -300 μs - -400 μs - -500 μs - -600 μs - -700 μs - -800 μs + +50 μs + +100 μs + +150 μs + +200 μs + +250 μs + +300 μs + +350 μs -900 μs +400 μs Outer NVTX Child NVTX ranges @@ -33,32 +31,25 @@ GPU stream 7 - -target operation · host NVTX 870.6 μs - -H2D input - -BF16 GEMM - -ReLU - - - - -cudaDeviceSynchronize · 384.8 μs - -memcpy API - -GEMM launch - -ReLU launch - -H2D copy · 607.2 μs - -GEMM -93.2 μs - - -ReLU · 11.1 μs + +target operation · host NVTX 376.0 μs + +BF16 GEMM + +ReLU + + + + +GEMM launch · 50.7 μs + +ReLU launch · 13.5 μs + +sync · 20.7 μs + +GEMM · 92.6 μs + + +ReLU · 10.9 μs Time is relative to the outer NVTX-range start; horizontal lengths come from the measured capture. \ No newline at end of file diff --git a/img/nsys_b200_timeline_zh_en_tracks.svg b/img/nsys_b200_timeline_zh_en_tracks.svg index 703ecd2a..3d9eef8e 100644 --- a/img/nsys_b200_timeline_zh_en_tracks.svg +++ b/img/nsys_b200_timeline_zh_en_tracks.svg @@ -3,28 +3,26 @@ B200 上的一次真实 Nsight Systems 采集 -4096×4096 BF16：H2D copy → GEMM → ReLU +4096×4096 BF16：GEMM → ReLU 0 μs - -100 μs - -200 μs - -300 μs - -400 μs - -500 μs - -600 μs - -700 μs - -800 μs + +50 μs + +100 μs + +150 μs + +200 μs + +250 μs + +300 μs + +350 μs -900 μs +400 μs Outer NVTX Child NVTX ranges @@ -33,32 +31,25 @@ GPU stream 7 - -target operation · host NVTX 870.6 μs - -H2D input - -BF16 GEMM - -ReLU - - - - -cudaDeviceSynchronize · 384.8 μs - -memcpy API - -GEMM launch - -ReLU launch - -H2D copy · 607.2 μs - -GEMM -93.2 μs - - -ReLU · 11.1 μs -时间以外层 NVTX range 的起点为 0；横向长度来自真实采集，并非示意比例。 + +target operation · host NVTX 376.0 μs + +BF16 GEMM + +ReLU + + + + +GEMM launch · 50.7 μs + +ReLU launch · 13.5 μs + +sync · 20.7 μs + +GEMM · 92.6 μs + + +ReLU · 10.9 μs +时间以外层 NVTX range 的起点为 0；横向长度按真实采集比例绘制。 \ No newline at end of file diff --git a/img/scripts/gen_nsys_b200_timeline.py b/img/scripts/gen_nsys_b200_timeline.py index 467abd46..04d0a192 100644 --- a/img/scripts/gen_nsys_b200_timeline.py +++ b/img/scripts/gen_nsys_b200_timeline.py @@ -8,7 +8,7 @@ HEIGHT = 500 LEFT = 235 RIGHT = 1440 -T_MAX = 900.0 +T_MAX = 400.0 BAR_HEIGHT = 34 @@ -23,13 +23,13 @@ def render(*, chinese: bool, output: Path) -> None: else "One Real Nsight Systems Capture on B200" ) subtitle = ( - "4096×4096 BF16：H2D copy → GEMM → ReLU" + "4096×4096 BF16：GEMM → ReLU" if chinese - else "4096×4096 BF16: H2D copy → GEMM → ReLU" + else "4096×4096 BF16: GEMM → ReLU" ) rows = ["Outer NVTX", "Child NVTX ranges", "CUDA APIs", "GPU stream 7"] note = ( - "时间以外层 NVTX range 的起点为 0；横向长度来自真实采集，并非示意比例。" + "时间以外层 NVTX range 的起点为 0；横向长度按真实采集比例绘制。" if chinese else "Time is relative to the outer NVTX-range start; horizontal lengths come from the measured capture." ) @@ -45,7 +45,7 @@ def render(*, chinese: bool, output: Path) -> None: axis_y = 88 parts.append(f'') - for tick in range(0, 901, 100): + for tick in range(0, 401, 50): x = x_pos(float(tick)) parts.append(f'') parts.append(f'{tick} μs') @@ -69,37 +69,32 @@ def rect(start: float, end: float, y: float, color: str, label: str = "", text_c ) # Outer and child NVTX ranges, relative to target-operation start. - rect(0.0, 870.561, row_y[0], "#365f9d", "target operation · host NVTX 870.6 μs") - rect(12.553, 174.087, row_y[1], "#72a7d8", "H2D input") - rect(190.849, 359.104, row_y[1], "#598bc2", "BF16 GEMM") - rect(364.257, 413.265, row_y[1], "#86b6df", "ReLU", "#243247") + rect(0.0, 375.992, row_y[0], "#365f9d", "target operation · host NVTX 376.0 μs") + rect(5.927, 201.012, row_y[1], "#598bc2", "BF16 GEMM") + rect(213.209, 272.981, row_y[1], "#86b6df", "ReLU", "#243247") # Host CUDA API intervals. - rect(116.267, 150.436, row_y[2], "#e39c45") - rect(319.529, 353.799, row_y[2], "#d7832f") - rect(397.159, 408.223, row_y[2], "#c76d25") - rect(482.407, 867.228, row_y[2], "#b85f46", "cudaDeviceSynchronize · 384.8 μs") + rect(135.682, 186.399, row_y[2], "#d7832f") + rect(255.070, 268.544, row_y[2], "#c76d25") + rect(352.196, 372.905, row_y[2], "#b85f46") for time_us, label, anchor in [ - (133.351, "memcpy API", "middle"), - (336.664, "GEMM launch", "middle"), - (402.691, "ReLU launch", "start"), + (161.041, "GEMM launch · 50.7 μs", "middle"), + (261.807, "ReLU launch · 13.5 μs", "middle"), + (362.551, "sync · 20.7 μs", "end"), ]: x = x_pos(time_us) - dx = 12 if anchor == "start" else 0 + dx = 12 if anchor == "start" else (-12 if anchor == "end" else 0) parts.append(f'') parts.append( f'{escape(label)}' ) # GPU activity on the default stream. - rect(145.879, 753.109, row_y[3], "#3d9b72", "H2D copy · 607.2 μs") - rect(757.013, 850.165, row_y[3], "#6f55b5", "GEMM") - gemm_x = (x_pos(757.013) + x_pos(850.165)) / 2 - parts.append(f'93.2 μs') - rect(850.389, 861.461, row_y[3], "#a574d1") - relu_x = x_pos(855.925) + rect(180.786, 273.394, row_y[3], "#6f55b5", "GEMM · 92.6 μs") + rect(273.618, 284.562, row_y[3], "#a574d1") + relu_x = x_pos(279.090) parts.append(f'') - parts.append(f'ReLU · 11.1 μs') + parts.append(f'ReLU · 10.9 μs') parts.append(f'{escape(note)}') parts.append('') diff --git a/zh/appendix/benchmarking_gpu_kernels.md b/zh/appendix/benchmarking_gpu_kernels.md index 1f8e1813..50eae3e8 100644 --- a/zh/appendix/benchmarking_gpu_kernels.md +++ b/zh/appendix/benchmarking_gpu_kernels.md @@ -1,12 +1,10 @@ (chap_benchmarking)= # GPU Kernel 性能测量与分析 -优化 GPU kernel 时，需要分别回答两个问题：运行一次要多久，时间主要花在哪里。Benchmark 负责测前者，profile 用来分析后者。Profiler 会改变程序的执行条件，因此涉及整个 operator 或应用路径的最终性能数字应回到关闭 profiler 后的测量中确认。 +优化 GPU kernel 时，需要分别回答两个问题：运行一次要多久，时间主要花在哪里。Benchmark 负责测前者，profile 用来分析后者。 一次 Python 调用不一定只对应一个 GPU kernel。它可能启动多个 kernels、提交内存拷贝，或者等待 GPU 完成工作。计时前要先确定被测 operation 包含哪些步骤，并让所有实现采用相同的边界。 -实际操作时，先验证结果并测出无 profiler 的基线，再用 profiler 查找耗时的原因。修改实现后，使用相同的计时方法重新测量；只有基线时间确实缩短，才能说明优化有效。 - {ref}`chap_performance` 介绍了如何用 roofline 判断性能受计算吞吐还是内存带宽限制。接下来讨论实验方法：如何确定计时范围、选择 warm-up 和 repeat，以及解读 profiler 报告。 ## 区分性能测量与性能诊断 @@ -16,22 +14,11 @@ | 工具 | 主要回答的问题 | |---|---| | CUDA Events | 被测区间在 GPU stream 上经过了多长时间？Stream 是按提交顺序执行 GPU 工作的队列。 | +| 同步的 wall-clock timer | 从 host 发起一次调用到这次调用所需的 GPU 工作完成，经过了多长时间？ | | Proton（Triton 提供的 profiler） | 启动了哪些 GPU kernels、各被调用多少次，哪些 kernel 占用了主要时间？ | | Nsight Systems | Host、streams、拷贝、kernels 和通信如何在时间线上重叠？ | -| Nsight Compute（`ncu`） | 一个选定的 GPU kernel 主要受哪类硬件资源或等待限制？ | -| IKET（可选） | 一个选定的 kernel 内部，哪些命名阶段或 warp roles 占用了时间？ | - -### 三类 Profile 视图 - -Profile 不是一个数字，也不只有一种报告格式。本章使用的工具会生成三种互补的视图： - -| 视图 | 工具 | 阅读重点 | -|---|---|---| -| 聚合树 | Proton | 比较调用次数、平均时间和总时间，找出主要耗时的 kernels。 | -| 时间线 | Nsight Systems；单个 kernel 内部使用 IKET | 沿横轴阅读不同 tracks 上的事件，检查空隙、重叠和依赖关系。 | -| 单 kernel 指标报告 | Nsight Compute | 查看一次 launch 的配置、利用率、scheduler 状态、内存流量以及 source/SASS 证据。 | - -这些 profile 用来解释时间花在哪里，不能替代正式的性能测量。修改实现后，应关闭 profiler，并使用与基线相同的计时边界重新测量。 +| Nsight Compute（`ncu`） | 选定的 kernel 内部在做什么，下一步应调查哪类硬件资源或等待？ | +| IKET（可选） | 加入 kernel 内标记后，数据搬运、计算和写回等阶段何时 active、等待或重叠？ | ## 计时前先验证正确性 @@ -42,24 +29,45 @@ Profile 不是一个数字，也不只有一种报告格式。本章使用的工 3. 使用明确的 tolerance，将结果与 reference 比较。 4. 如果 kernel 会在已有 output 上累加或原地修改输入，每次验证前都恢复相同的初始状态。 -Reference 计算和结果比较不属于性能计时。正确性通过后，再开始设计 benchmark；是否把状态重置计入时间，由下一节定义的 operation 边界决定。 +以自写 GEMM 为例，`actual` 是被测实现的输出，`expected` 可以先用 PyTorch 的 FP32 GEMM 计算，再转换成目标输出类型： + +```python +import torch + + +torch.set_float32_matmul_precision("highest") +actual = my_gemm(a, b) # 换成自己的实现 +torch.cuda.synchronize() +expected = torch.mm(a.float(), b.float()).to(actual.dtype) +rtol = 1e-2 # 示例值；根据输出 dtype、累加方式和 shape 调整 +atol = 1e-2 +torch.testing.assert_close(actual, expected, rtol=rtol, atol=atol) +``` + +`torch.set_float32_matmul_precision("highest")` 避免这个 CUDA FP32 reference 使用较低的内部计算精度。示例中的 `rtol` 和 `atol` 分别控制相对误差与绝对误差；`1e-2` 只是可运行的起点，实际值要根据输出 dtype、累加方式、shape 和算子约定调整。同一组比较应使用相同的 reference 和 tolerance。 + +Reference 计算和结果比较不计入性能时间；状态重置是否计时，由下一节定义的 operation 边界决定。 ## 明确计时边界 计时前，先写清楚一次被测 operation 包含哪些工作。它可以只包含一个 kernel，也可以包含得到完整结果所需的全部 kernels、内存拷贝和状态重置。编译、输入构造、内存分配或数据格式转换是否属于这次 operation，也要明确说明。不同实现只有在测量相同工作时才能直接比较。 +后文的 GEMM + ReLU 例子可以采用三种边界：只用 CUDA Events 包住 `torch.mm`，得到 GEMM 的 GPU stream 时间；用 Events 包住整个 `run()`，得到 GEMM + ReLU 的 GPU stream 时间；在调用 `run()` 前启动 CPU timer，并在调用后同步，得到一次 Python 调用的端到端时间。矩阵分配和 warm-up 默认位于这三种边界之外。 + 范围确定后，再选择计时方法： - **CUDA Events** 记录 GPU stream 执行到两个位置时的时间戳，适合测量一个 kernel 或整个 operator 在 device 时间线上的区间。区间内的 kernels、内存拷贝和 stream 空隙都会被计入。对于多 stream operation，所有参与计算的 streams 都必须在 start event 之后开始被测工作，并在记录 end event 前完成汇合。 - **同步的 wall-clock timer** 从 host 发起调用前开始计时，在调用所需的 GPU 工作全部完成后停止。它还会计入 Python dispatch、CUDA launch 和等待 GPU 完成的时间，适合测量一次调用的端到端 latency。 -例如，GPU 已经执行 start event，但 host 还没有提交下一个 launch 时，这段 stream 空闲时间仍会落在 CUDA Event 区间内。因此，CUDA Event 区间不一定等于 profiler 中某个 kernel 从开始到结束的执行区间。诊断型 profiler 适合观察 kernel 执行和重叠关系，但其结果不能直接替代相同边界下的无 profiler 计时。 +例如，GPU 已经执行 start event，但 host 还没有提交下一个 launch 时，这段 stream 空闲时间仍会落在 CUDA Event 区间内。因此，CUDA Event 区间不一定等于 profiler 中某个 kernel 从开始到结束的执行区间。 + +## 测量 GPU 时间与单次调用延迟 -## 使用 CUDA Events 测量 GPU 时间 +### 使用 CUDA Events 测量 GPU stream 时间 -CUDA launch 通常是异步的：Python 把工作提交到 CUDA stream 后就可以继续运行，此时 GPU 不一定已经完成。如果只用 CPU 时钟记录这次 Python 调用前后的时间，计时可能在 GPU 完成前就已经停止，得到的主要是 host 提交耗时。测量 GPU stream 上经过的时间时，应使用 CUDA Events；测量从 Python 发起调用到 GPU 完成的完整时间时，则使用后面介绍的同步 wall-clock timer。[PyTorch CUDA semantics 文档](https://docs.pytorch.org/docs/stable/notes/cuda.html#asynchronous-execution)也说明了这种异步行为。 +CUDA launch 通常是异步的；未经同步的 CPU 计时可能在 GPU 完成前就已停止，主要反映 host 提交时间。[PyTorch CUDA semantics 文档](https://docs.pytorch.org/docs/stable/notes/cuda.html#asynchronous-execution) 也说明了这种行为。下面直接用 CUDA Events 测量当前 stream 上经过的时间。 -先看一份可以直接运行的 CUDA Event benchmark。它在计时前分配矩阵，先执行 warm-up，再测量五轮 FP16 GEMM 并报告中位数： +先看一份可以直接运行的 CUDA Event benchmark。它在计时前分配矩阵，先执行 warm-up，再分五轮测量 FP16 GEMM，并报告五轮结果的中位数： ```python from statistics import median @@ -107,23 +115,20 @@ print(f"median CUDA Event time: {median(samples_ms):.4f} ms") `measure_batch_ms` 在当前 CUDA stream 中记录 start 和 end events，并用两者之间的时间除以调用次数。这样得到的是连续执行时每次 GEMM 的平均 GPU stream 时间。`end.synchronize()` 只是让 CPU 等到这轮 GPU 工作完成，以便读取 Event 结果。 -这里的 `warmup_calls=500`、`repeat=100` 和 `rounds=5` 都表示调用或测量次数。它们不是通用标准，而是根据这个 GEMM 在 B200 上的实测结果选出的示例值。在一次 10 轮校准中，warm-up 50 次时，第一轮到最后一轮仍从 0.01425 ms 降到 0.01296 ms；增加到 500 次后，变化缩小为 0.01332 ms 到 0.01301 ms。`repeat=100` 的结果也比 `repeat=10` 更稳定。 - -选择其他 workload 的参数时，可以逐步增加 `warmup_calls`，直到前几轮不再持续变快或变慢；再增加 `repeat` 或 `rounds`，直到波动已经满足实验需要。次数也不是越多越好：如果延长实验后整体时间系统性变化，应检查温度、功耗和时钟频率，并先确定实验要表示短时运行还是持续运行。耗时较长的 kernel 通常可以使用更小的次数。 - -这段代码始终复用同一组矩阵，因此后续调用可能从 cache 中读取部分数据，测得的是 warm-cache 场景。发布结果时，还应记录 GPU 型号、软件版本和时钟设置。 +`warmup_calls=500` 和 `repeat=100` 表示调用次数，`rounds=5` 表示五轮独立测量。这组值来自 B200 上的稳定性测试：50 次 warm-up 后结果仍持续下降，增加到 500 次后才趋于稳定；`repeat=100` 也比 `repeat=10` 稳定。测量其他 workload 时，先增加 `warmup_calls`，直到前几轮不再持续变化；若结果仍有较大波动，再增加 `repeat` 或 `rounds`。如果更长的运行反而使整体时间系统性变化，就要检查温度、功耗和时钟频率。 -实际测量本书中的 TIRx kernels 时，不需要为每个 kernel 重新编写 warm-up、重复计时和统计逻辑。TVM 的 [`tvm.tirx.bench.bench`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/bench.py) 已经封装了这些步骤。只需传入一个负责启动被测实现的函数，输入、输出和 workspace 仍在计时前分配。 +这段代码始终复用同一组矩阵，因此结果偏向有数据复用的 warm-cache 场景；是否真的命中 cache，还取决于本次计算反复访问的数据总量与硬件 cache 容量。 -这个 helper 与上例采用不同的 cache 策略：上例连续复用同一组矩阵，而 `bench` 会在每次正式调用前驱逐 L2 cache，再用一对独立的 CUDA Events 测量被测实现。调用方式如下： +本书使用 TVM 的 [`tvm.tirx.bench.bench`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/bench.py) 统一处理 warm-up、重复计时和统计。传入的函数只启动已经准备好的实现，输入、输出和 workspace 均在计时前分配。与上面的 warm-cache 示例相比，`bench` 会在每次正式调用前写入一个 256 MiB buffer，尽量减少前一次调用留下的 L2 复用，再用独立的 CUDA Events 计时： ```python from tvm.tirx.bench import bench -# run 是无参数函数；它只使用已经分配好的 tensors 启动被测实现。 +# 这里复用上文的 gemm；分析自己的 TIRx kernel 时，换成对应的无参数 callable。 +run = gemm result = bench( - {"tirx": run}, + {"gemm": run}, timer="event", warmup=25, repeat=100, @@ -131,17 +136,13 @@ result = bench( cooldown_s=1.0, ) -print(result["impls"]["tirx"]) # 五轮平均值，单位为 us -print(result["round_samples"]["tirx"]) # 每轮结果 +print(result["impls"]["gemm"]) # 五轮平均值，单位为 us +print(result["round_samples"]["gemm"]) # 每轮结果 ``` -这里的 `warmup=25` 和 `repeat=100` 表示毫秒预算，不是固定的调用次数。Event timer 会先做一轮短测，再把 25 ms 的 warm-up 预算和 100 ms 的正式测量预算换算成实际次数。短测包含 L2 驱逐和被测调用，因此换算出的次数只是近似值；正式报告的 Event 时间只覆盖被测调用，L2 驱逐发生在 start event 之前。短 kernel 会自动执行更多次，长 kernel 则执行较少次。`rounds=5` 表示完整测量五轮，`cooldown_s=1.0` 表示每轮测量一个实现前暂停一秒。最终结果是五轮的平均值，每轮结果仍保存在 `round_samples` 中。25/100 ms 是 Event timer 的默认预算；五轮测量则是 TIRx-kernels 命令行工具采用的默认设置，`bench` 函数本身默认只运行一轮。 +`warmup=25` 和 `repeat=100` 是毫秒预算，Event timer 会根据短测结果换算调用次数。正式 Event 只覆盖被测调用；用于减少 L2 复用的 256 MiB 写入发生在 start event 之前。`rounds=5` 测量五轮，`cooldown_s=1.0` 在每轮前暂停一秒；`impls` 保存五轮平均值，`round_samples` 保存逐轮结果。预算和轮数仍按上面的稳定性标准调整，并对所有实现使用相同设置。 -这些数值只是默认起点。若各轮结果仍持续漂移，应增加 warm-up 预算；若各轮波动很大，应增加正式测量预算或轮数。所有实现必须使用相同的 timer、预算和轮数，并保留每轮结果，而不是只报告最快的一次。 - -TIRx-kernels 中的 `run_bench` 也调用这个 helper，例如 [`tirx_kernels/attention/flash_attention4.py`](https://github.com/mlc-ai/tirx-kernels/blob/5be39749e7dfd2c4bdae9b4d396f8ec35af07126/tirx_kernels/attention/flash_attention4.py)。省略 `timer` 时，本地 benchmark 默认使用 Proton；需要 CUDA Event 区间时，应像上面一样显式指定 `timer="event"`。两种 timer 的结果含义不同，报告时必须注明。 - -传给 `bench` 的函数会被反复调用。如果 kernel 会累加 output 或原地修改输入，就要在每次调用前恢复相同状态，或者保证每次正式测量都使用一份尚未修改的预分配输入。若恢复操作放在被测函数中，它的时间也属于前面定义的 operation 边界。否则后一次调用面对的已经不是同一个 workload。 +TIRx-kernels 的 `run_bench` 也使用这个 helper。未使用 distributed 模式时，省略 `timer` 会默认使用 Proton；需要 CUDA Event 区间时，应显式指定 `timer="event"`。对于会原地修改状态的 kernel，重复调用时仍要遵守前面的重置规则；若重置写在被测函数内，其时间也属于 operation。 ### 测量一次调用的端到端时间 @@ -172,129 +173,193 @@ print(f"median end-to-end time: {median(host_samples_ms):.4f} ms") 第一个同步确保计时开始前没有更早的 GPU 工作残留；第二个同步保证计时停止前，这次 GEMM 已经完成。每个 sample 只包含一次调用，因此结果包括 Python 调用、CUDA launch、GPU 执行以及等待完成的开销。前面的 CUDA Event benchmark 测量的则是连续调用时每次 GEMM 的平均 GPU stream 时间。 -比较多个实现时，所有实现必须使用相同的计时方法和边界。如果同时报告这两种结果，应分别命名为“CUDA Event GPU 时间”和“单次端到端时间”，而不是把使用不同 timer 得到的数字都写成同一种 latency。 +比较多个实现时，所有实现使用相同的计时方法和边界。如果同时报告这两种结果，可以分别命名为“CUDA Event GPU 时间”和“单次端到端时间”，让读者直接看出两个数字覆盖的范围。 + +### 进阶：测量多 stream operation -### 重叠执行的计时方法 +一个 operation 把工作提交到多条 CUDA streams 时，计时 stream 需要连接每条分支的起点和终点。下面的 `sin` 和 `cos` 分别在两条 streams 上运行，等两条分支都完成后，再回到计时 stream 相加： -前面的 GEMM 只在当前 CUDA stream 上运行，因此 start 和 end events 可以直接包住全部工作。如果一个 operator 同时使用多个 streams，仅在当前 stream 记录 events 就不够了：其他 stream 上的工作可能在 start 之前已经开始，也可能在 end 之后仍未完成。 +```python +import torch -要测量整个 operator，可以把 start event 作为所有工作 streams 的共同起点：每个 stream 先等待 start，再开始被测工作，并在完成后各自记录一个 event。最后，当前 stream 等待这些完成 events，再记录 end。这样得到的区间才覆盖从最早开始到全部完成的整个 operation。 -PDL（Programmatic Dependent Launch）是另一种可能产生重叠的情况。它允许 compute capability 9.0 或更新的 GPU 在同一 stream 中提前启动后一个 kernel：后一个 kernel 可以先完成不依赖前序结果的准备工作，在真正读取这些结果前再等待。这个过程需要显式启用，并遵守相应的 trigger 和 wait 约定；具体 API 见 [CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html)。 +x = torch.randn(1 << 20, device="cuda") +left = torch.empty_like(x) +right = torch.empty_like(x) +output = torch.empty_like(x) -无论重叠来自多个 streams 还是 PDL，计时原则都相同：用 CUDA Events 包住完整 operation。重叠的 kernels 可能覆盖同一段时间，因此不能把 profiler 中各 kernel 的 duration 直接相加作为 operator latency。实际的执行顺序和重叠情况可以在 Nsight Systems 时间线中查看。PDL 是否产生重叠由运行时决定，程序正确性不能依赖它一定发生。 +stream_left = torch.cuda.Stream() +stream_right = torch.cuda.Stream() +timing_stream = torch.cuda.current_stream() -## 固定实验条件 +start = torch.cuda.Event(enable_timing=True) +left_done = torch.cuda.Event() +right_done = torch.cuda.Event() +end = torch.cuda.Event(enable_timing=True) -前面确定了计时边界和计时器，接下来还要固定会影响结果的实验条件。前面的示例都在输入分配和 warm-up 完成后开始计时，测量的是后续重复调用的性能。首次调用则可能包含 CUDA 初始化、JIT、autotuning 或其他只发生一次的工作。如果研究目标是首次调用或完整应用路径，就应把相应步骤纳入计时边界并单独报告，不能与重复调用的结果混在一起。 -一次测量不足以说明结果是否稳定。手写 CUDA Event 示例保留五轮结果并报告中位数；`bench` 则报告各轮的平均值，同时把原始结果保存在 `round_samples` 中。无论采用哪种汇总方式，都应保留每轮结果、检查是否存在趋势或异常波动，并明确报告使用的是中位数还是平均值，而不是只挑最快的一轮。比较多个实现时，还可以更换测量顺序后再运行一次，避免某个实现总是在设备较冷或较热时被测量。 +def measure_operation_ms(): + torch.cuda.synchronize() + start.record(timing_stream) -缓存状态也会改变结果。前面的手写示例反复使用同一组矩阵，属于 warm-cache 测量。TVM 0.26 的 Event 和 Proton timers 则会在每次正式调用前写入一个 256 MiB buffer，以驱逐 L2 中已有的数据；这次写入发生在计时区间之外。两种策略都可以使用，关键是选择符合目标应用的一种，并让所有实现保持一致。`torch.cuda.empty_cache()` 只会释放 PyTorch caching allocator 中未使用的 blocks，不会清空 GPU L2 cache，因此不能用它实现 cold-L2 测量。 + stream_left.wait_event(start) + with torch.cuda.stream(stream_left): + torch.sin(x, out=left) + left_done.record(stream_left) -最后，记录 GPU 型号、driver、CUDA runtime、framework 和 compiler 版本，以及被测 workload 的 dtype 与 shape。还要记录时钟和功耗设置，避免其他进程占用设备，并留意热降频。若锁定时钟，应给出具体数值与命令；只写“fixed clocks”不足以复现实验。 + stream_right.wait_event(start) + with torch.cuda.stream(stream_right): + torch.cos(x, out=right) + right_done.record(stream_right) + + timing_stream.wait_event(left_done) + timing_stream.wait_event(right_done) + torch.add(left, right, out=output) + + end.record(timing_stream) + end.synchronize() + return start.elapsed_time(end) -相同的 tensor shape 并不代表两个实现可以直接比较。至少要对齐三类条件： + +elapsed_ms = measure_operation_ms() +torch.testing.assert_close(output, torch.sin(x) + torch.cos(x)) +print(f"multi-stream operation: {elapsed_ms:.4f} ms") +``` + +`start` 是两条分支的共同起点，`left_done` 和 `right_done` 分别标记两个分支的终点。计时 stream 等待这两个 completion events，执行最后的加法，再记录 `end`。`wait_event` 在 GPU 上建立依赖，CPU 可以继续提交后续工作；`end.synchronize()` 才让 CPU 等待这次测量完成。 + +事件关系可以写成： + +```text +timing stream: start ──────────── wait(left_done, right_done) ─ add ─ end +left stream: wait(start) ─ sin ─ left_done +right stream: wait(start) ─ cos ─ right_done +``` + +这段代码允许两条分支并发执行，实际重叠程度取决于 GPU 资源占用。Nsight Systems 时间线可以确认真实的执行关系。正式测量时，先调用若干次 `measure_operation_ms()` 完成 warm-up，再重复调用它收集多个单次样本，最后报告 median 和样本波动。 + +#### 同一 stream 内的 PDL + +Programmatic Dependent Launch（PDL）用于显式启用该机制的自定义 CUDA 或 DSL launch path。Primary 和 secondary kernels 仍提交到同一条 stream；primary 发出 trigger 后，secondary 可以提前执行与 primary 结果无关的准备阶段，并在读取 primary 的结果前完成 PDL 规定的依赖同步。 + +```text +primary: initial work ─ trigger ─ remaining work +secondary: preamble ─ wait ─ dependent work +``` + +计时时把 `start` 放在 primary launch 前，把 `end` 放在 secondary launch 后，完整 Event 区间就是这组 launches 的 GPU 时间。两条 kernel 时间可能部分重叠，因此 profiler 中两个 duration 的和可能大于完整 operation latency；实际重叠关系由 Nsight Systems 时间线确认。 + +PDL 提供并发执行的机会，运行时也可以选择串行执行，kernel 的正确性需要覆盖两种情况。上面的 `torch.cuda.Stream` 接口本身没有 PDL launch 参数；具体启用方式由自定义 CUDA/DSL 实现提供，参见 [CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html)。 + +## 固定实验条件 + +前面的示例测量的是输入已经分配、warm-up 已经完成后的重复调用。若研究首次调用或完整应用路径，应把 CUDA 初始化、JIT、autotuning 等相应步骤纳入计时边界，并与重复调用的结果分开报告。 + +每轮原始结果都应保留，并注明最终报告的是中位数还是平均值。如果结果随轮次持续变化，应继续检查 warm-up、温度和时钟状态。比较多个实现时，也可以交替测量顺序，让每个实现经历相近的设备温度与时钟状态。 + +缓存策略同样需要统一。手写示例反复使用同一组矩阵，结果偏向 warm-cache 和数据复用，但实际命中率仍取决于反复访问的数据总量与 cache 容量；TVM 0.26 的 Event 和 Proton timers 会在每次正式调用前写入一个 256 MiB buffer，以减少 L2 复用，且这次写入不计时。应选择符合目标应用的策略，并让所有实现保持一致。`torch.cuda.empty_cache()` 释放的是 PyTorch allocator 中未使用的 blocks，GPU L2 的内容仍由硬件 cache 策略管理。 + +比较不同实现时，还要对齐： - **数值语义：** 输入与输出的数据类型、布局、转置方式、对齐要求、累加精度、缩放、mask、epilogue、输出定义和误差阈值； - **被测范围：** 是否包含 allocation、数据转换、状态重置、辅助 kernels、通信和同步； - **调优条件：** workspace 上限、是否允许针对每个 shape 自动调优，以及各实现可使用的搜索预算。 -所有实现还应采用相同的 cache、时钟、warm-up、采样和计时策略。使用库实现作为 baseline 时，需要记录版本、所选算法和 workspace；自动调优可以放在计时区间之外，但搜索预算与最终配置仍应写入实验记录。 +此外还应记录 GPU、driver、CUDA、framework 与 compiler 版本，以及 dtype、shape、时钟和功耗设置。使用库实现作为 baseline 时，要记录库版本、所选算法和 workspace；自动调优可以放在计时区间之外，但搜索预算和最终配置仍属于实验条件。 ## 由延迟换算吞吐率 -吞吐率不是计时器直接测出来的，而是用约定的工作量除以延迟得到的。因此，性能表在给出 TFLOP/s、GB/s 或 tokens/s 时，也应保留原始延迟，并说明工作量如何计算。本书将 GEMM 的工作量记为 $2MNK$ FLOPs；对于 attention 和 fused kernels，则需注明统计的是完整的稠密问题、实际选中的元素，还是 kernel 真正执行的工作。相关公式和 roofline 分析见 {ref}`chap_performance`。 +吞吐率等于约定的工作量除以延迟；计时器提供公式中的延迟。对于一个 $M\times K$ 与 $K\times N$ 的 GEMM，若延迟为 `t_us` 微秒，则： -## 使用 Proton 找出耗时的 kernel +```text +TFLOP/s = 2 × M × N × K / t_us / 10^6 +``` -前面的 benchmark 只告诉我们整个 operation 用了多长时间。如果它会启动多个 kernels，还需要找出时间具体花在哪些 kernels 上。Proton 可以列出每个 kernel 的调用次数、平均时间和累计时间。 +分子和计时边界必须对应同一份工作。例如，后文完整 GEMM + ReLU operation 的延迟是 105.152 μs；用 $2\times4096^3$ 除以这个时间会得到约 1307 TFLOP/s，但它只能称为“按 GEMM 工作量计算的有效吞吐率（effective throughput）”，因为分母还包含 ReLU。要报告 GEMM kernel 自身的 TFLOP/s，计时区间也要只覆盖 GEMM。 -Proton 是 Triton 项目提供的 GPU profiler。它观察的是 CUDA kernel 活动，因此也可以分析由 TVM 编译的 TIRx kernels。前面介绍的 `bench(timer="proton")` 只返回汇总后的 kernel time；这里单独创建一个 Proton session，以查看每个 kernel 的调用次数和耗时。 +性能表在给出 TFLOP/s、GB/s 或 tokens/s 时，也应保留原始延迟，并说明工作量如何计算。对于 attention 和 fused kernels，还需注明分子统计的是完整稠密问题、实际选中的元素，还是 kernel 真正执行的工作。相关公式和 roofline 分析见 {ref}`chap_performance`。 -下面继续使用前面分配好的矩阵，把 GEMM 和 ReLU 组成一个 operation。代码先完成 warm-up，只采集后面的 100 次调用，最后在当前目录生成 `operator.hatchet`： +## 用 Proton 找出最耗时的 kernel + +从这里开始，baseline、Proton、Nsight Systems 和 Nsight Compute 都运行 `appendix/nsys_example.py` 中的同一个 operation：两个 $4096\times4096$ BF16 矩阵先做 GEMM，再对结果做 ReLU。输入、中间结果和输出都在计时或采集前分配。脚本中的 operation 是： ```python -import torch -import triton.profiler as proton +def run(): + with torch.cuda.nvtx.range("BF16 GEMM"): + torch.mm(a, b, out=c) + with torch.cuda.nvtx.range("ReLU"): + torch.clamp_min(c, 0, out=output) +``` +进入任何计时或采集模式前，脚本会先执行一次 operation 并同步，再用 FP32 `torch.mm` 计算 reference、执行 ReLU、转换为 BF16，最后以 `rtol=2e-2`、`atol=1e-2` 比较实际输出。比较失败时命令会在开始 benchmark 或 profiler 之前报错；这次 preflight check 位于后面的 baseline 和采集范围之外。 -def operation(): - torch.mm(a, b, out=c) - torch.clamp_min(c, 0, out=c) +前面的 2048×2048 FP16 代码用于讲解计时 API；下面换成这份 BF16 operation 后，不再混用两组 workload 的结果。 +### 先记录无 profiler baseline -def collect_proton(run, *, warmup_calls, profile_calls): - for _ in range(warmup_calls): - run() - torch.cuda.synchronize() +在分析“时间花在哪里”之前，先测出正常运行时的完整 operation 时间： - session = proton.start("operator", context="shadow", data="tree") - if session is None: - raise RuntimeError("Proton session could not be created") - try: - with proton.scope("target_operation"): - for _ in range(profile_calls): - run() - torch.cuda.synchronize() - finally: - proton.finalize(session) +```bash +python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --event-samples 20 +``` +500 次 warm-up 发生在正式计时前；之后，每个 sample 用一对 CUDA Events 包住一次 GEMM + ReLU。B200 上的一次实际输出为： -collect_proton(operation, warmup_calls=500, profile_calls=100) +```text +median=105.152 us, min=103.136 us, max=131.200 us ``` -这里的 `warmup_calls` 和 `profile_calls` 都是调用次数，不是 `bench` 使用的毫秒预算。运行代码需要安装与 TVM 兼容的 Triton，并在实验记录中注明 Triton 版本。 +这里的 median 是后面判断代码修改是否真的变快时要回到的 baseline，min 和 max 用于观察样本波动。这份 baseline 覆盖完整的 GEMM + ReLU operation；后面的 profiler 表格来自独立采集，并分别列出单个 kernel。 + +### 用 Proton 比较 operation 中的 kernels -先查看文件中有哪些 metrics，再用后两条命令分别打印调用次数、总时间和平均时间： +Proton 可以查看每个 kernel 的调用次数、平均时间和累计时间。运行前先确认环境中已经安装与 TVM 兼容的 Triton；Proton 和 `proton-viewer` 随 Triton 提供。Viewer 还需要下面两个 Python 依赖： ```bash -proton-viewer --list operator.hatchet -proton-viewer --metrics time/ms,count --print-sorted operator.hatchet -proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet +python -m pip install pandas llnl-hatchet ``` -如果 `proton-viewer` 报告缺少可选依赖，再安装： +脚本的 `--proton-calls` 模式复用同一个 `run()`，先 warm-up，再采集 100 次 operation，并生成 `operator.hatchet`： ```bash -python -m pip install pandas llnl-hatchet +python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --proton-calls 100 +``` + +这里的 `warmup-calls` 和 `proton-calls` 都以调用次数为单位；`bench` 中同名的 warm-up/repeat 参数则使用毫秒预算。报告生成后运行： + +```bash +proton-viewer --list operator.hatchet +proton-viewer --metrics time/ms,count --print-sorted operator.hatchet +proton-viewer --metrics avg_time/us,time/ms --print-sorted operator.hatchet ``` -下面是这段代码在 B200 上的一次实际结果；为了便于阅读，缩短了 kernel 名称： +`proton-viewer` 会分别打印两张表。先在 `count,time/ms` 输出中用完整 kernel 名称确认调用次数和累计时间，再到 `avg_time/us,time/ms` 输出中找到同一行读取平均时间。选择目标时主要看 `time/ms`，因为它表示 100 次 operation 中该 kernel 累计占用的时间；`avg_time/us` 给出单次平均值。下面为了排版合并两张表并缩短 kernel 名称： ```text target_operation calls avg/us total/ms -├── GEMM kernel 100 14.83 1.483 -└── ReLU kernel 100 4.23 0.423 +├── GEMM kernel 100 87.00 8.700 +└── ReLU kernel 100 11.71 1.171 ``` -首先核对预期的 kernels 是否都出现、调用次数是否正确，再比较各 kernel 的平均时间和累计时间。这个例子中 GEMM 的累计时间最大，因此下一步应优先用 Nsight Compute 分析 GEMM。也要留意调用次数很多的短 kernel：它们单次耗时不高，累计开销却可能很大。 +先检查预期的两个 kernels 是否都出现、调用次数是否为 100，再比较累计时间。GEMM 占用的时间明显更多，因此选它作为后续深入分析的目标。交给 NCU 之前，先用 Nsight Systems 确认单次 operation 中的执行顺序和空隙，并关联对应的 host launch APIs。 -Proton 只统计捕获到的 kernel 时间，不包含内存拷贝、同步和 stream 空隙；如果 kernels 发生重叠，各项 duration 的总和还会重复计算重叠区间。因此，这些数据用于定位需要继续分析的 kernel，不表示整个 operation 的 latency。 - -这次采集反复使用同一组矩阵，而前面的 TVM timer 会在每次测量前驱逐 L2，两者的 cache 条件也不相同。完整 operation 的 latency 仍应使用 CUDA Events 或同步的 wall-clock timer 测量。 +这次手动 Proton session 保留正常的 cache 状态，与 `bench(timer="proton")` 在每次正式调用前写入 256 MiB buffer 的策略不同。这里用同一份 Proton 报告内的数值给 kernels 排序。实现之间的快慢仍看上面的 CUDA Event baseline；完整 operation latency 则取自包住整个 operation 的计时区间。 ## 使用 Nsight Systems 分析应用时间线 -Proton 可以汇总各 kernel 的时间，却看不到它们以什么顺序执行，也看不到 kernel 之间的空隙、数据拷贝和 host 等待。分析这些问题时，需要使用 Nsight Systems 的时间线。 +Proton 给出了汇总排名，但没有显示 kernels 的先后关系、空隙、拷贝或 host 等待。Nsight Systems 用时间线回答这些问题。 ### 采集目标 operation 的时间线 -下面用一个简单例子说明怎样限定 Nsight Systems 的采集范围。`appendix/nsys_example.py` 中的 operation 依次完成三个步骤：把一个 $4096\times4096$ 的 BF16 matrix 从 pinned host memory 复制到 GPU（host-to-device，H2D），执行 GEMM，再执行 ReLU。所需 tensors 均在采集前分配，因此报告只聚焦这三个步骤。脚本的核心代码如下： +脚本的 `--profile-once` 模式先在 profiler 尚未启动时完成 warm-up，并等待这些工作结束；然后只在 `cudaProfilerStart()` 和 `cudaProfilerStop()` 之间提交一次 GEMM + ReLU： ```python -import torch - - -def run(): - with torch.cuda.nvtx.range("H2D input"): - a.copy_(host_a, non_blocking=True) - with torch.cuda.nvtx.range("BF16 GEMM"): - torch.mm(a, b, out=c) - with torch.cuda.nvtx.range("ReLU"): - torch.clamp_min(c, 0, out=output) - - def run_once_for_profiler(run, *, warmup_calls): for _ in range(warmup_calls): run() @@ -308,7 +373,7 @@ def run_once_for_profiler(run, *, warmup_calls): cudart.cudaProfilerStop() ``` -`run_once_for_profiler` 先在 profiler 尚未启动时完成 warm-up，并等待 GPU 上的 warm-up 工作结束。随后，`cudaProfilerStart()` 开始采集；NVTX range `target operation` 为这次 operation 添加名称，便于在时间线中定位。这个 range 内的同步确保三项 GPU 工作在 `cudaProfilerStop()` 之前完成。`cudaProfilerStart()` 和 `cudaProfilerStop()` 只用于限定采集范围，不用于计时。 +这里的 NVTX range 用于在时间线中定位目标 operation。Range 内的同步确保两个 kernels 在停止采集前完成；`cudaProfilerStart()` 和 `cudaProfilerStop()` 只限定采集范围，不负责计时。 下面的命令运行这个脚本，并将报告写入 `reports/target-timeline.nsys-rep`： @@ -322,10 +387,13 @@ nsys profile \ --capture-range-end=stop \ --output=reports/target-timeline \ --force-overwrite=true \ - python appendix/nsys_example.py --profile-once + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once ``` -`--capture-range=cudaProfilerApi` 只采集 `cudaProfilerStart()` 与 `cudaProfilerStop()` 之间的区间。`--trace=cuda,nvtx` 记录 CUDA API、GPU activity 和 NVTX ranges。这里先关闭 CPU sampling 与 context-switch tracing，使报告集中显示 CUDA 时间线。如果时间线中出现较长的 GPU 空隙，再单独采集一份包含 host scheduling 或 OS runtime 信息的报告。 +`--capture-range=cudaProfilerApi` 只采集 profiler API 之间的区间，`--trace=cuda,nvtx` 记录 CUDA API、GPU activity 和 NVTX ranges。这里关闭 CPU sampling 与 context-switch tracing，让第一份报告先集中显示 CUDA 时间线；如果发现较长的 GPU 空隙，再采一份包含 host scheduling 或 OS runtime 信息的报告。 报告生成后，可以直接用 Nsight Systems GUI 打开时间线： @@ -333,93 +401,106 @@ nsys profile \ nsys-ui reports/target-timeline.nsys-rep ``` -### 时间线中的拷贝、排队与执行时间 - -这个例子使用 PyTorch，是为了用较少的代码构造数据拷贝和多个 CUDA kernels；后面的时间线读法同样适用于 TIRx operation。 +### 从时间线定位最耗时的 kernel -下面用一份实际采集的报告说明如何阅读 Nsight Systems 的结果。报告来自一台 NVIDIA B200，软件版本为 NVIDIA driver 595.58.03、CUDA 13.0、PyTorch 2.12.0+cu130 和 Nsight Systems 2025.6.3。 +下面的报告来自 NVIDIA B200，软件版本为 NVIDIA driver 595.58.03、CUDA 13.0、PyTorch 2.12.0+cu130 和 Nsight Systems 2025.6.3。 -![一次真实的 Nsight Systems 时间线：CPU 上的 NVTX ranges 和 CUDA APIs 位于上方，实际 H2D、GEMM 与 ReLU 位于同一条 GPU stream 上](../../img/nsys_b200_timeline_zh_en_tracks.svg) +![一次真实的 Nsight Systems 时间线：CPU 上的 NVTX ranges 和 CUDA APIs 位于上方，GEMM 与 ReLU 位于同一条 GPU stream 上](../../img/nsys_b200_timeline_zh_en_tracks.svg) -*本图根据实际采集结果重绘，横条长度与各项操作的实测时长成比例。* - -`GPU stream 7` 中的 `7` 是 Nsight Systems 在这次采集中显示的 stream 标识，不表示第七个执行阶段，换一次运行也可能不同。图中的 H2D copy、GEMM 和 ReLU 位于同一条 stream 上，因此按提交顺序执行。 - -除了在 GUI 中查看时间线，也可以让 `nsys stats` 从同一份报告中整理出下面使用的时间数据： +`GPU stream 7` 中的 `7` 是这次报告里的 stream 标识。两个 kernels 位于同一条 stream 上，所以按提交顺序执行。也可以从命令行提取下文使用的时间： ```bash nsys stats \ + --force-export=true \ --format=column \ --timeunit=us \ - --report nvtx_gpu_proj_sum \ - --report nvtx_pushpop_trace \ --report cuda_gpu_sum \ --report cuda_kern_exec_trace \ + --report cuda_api_trace \ + --report nvtx_gpu_proj_sum \ + --report nvtx_pushpop_trace \ --report cuda_api_sum \ reports/target-timeline.nsys-rep ``` -`--report` 后面是 Nsight Systems 自带的统计名称。`nvtx_gpu_proj_sum` 和 -`nvtx_pushpop_trace` 分别给出 NVTX range 在 GPU 上覆盖的区间和 host 端的 range 记录； -`cuda_gpu_sum` 汇总 kernels 与 CUDA memory operations；`cuda_kern_exec_trace` 将 host 上的 -launch API 与对应的 GPU kernel 关联起来；`cuda_api_sum` 则汇总 host 端的 CUDA API 调用。 -运行 `nsys stats --help-reports` 可以查看当前版本支持的全部名称和定义。 +`--force-export=true` 会从当前 `.nsys-rep` 重新生成 SQLite 数据，避免误读同名旧文件。`cuda_gpu_sum` 汇总 GPU activities；`cuda_kern_exec_trace` 把 host launch API 与对应的 GPU kernel 关联起来，并给出 kernel 的 start 和 duration；`cuda_api_trace` 给出每次 CUDA API 的 start 和 duration。两个 NVTX reports 分别给出 range 在 GPU 上覆盖的区间和 host 端记录；`cuda_api_sum` 汇总 host CUDA API。运行 `nsys stats --help-reports` 可以查看当前版本的完整定义。 + +这条命令会依次打印多张表，按下面的顺序取数： -`cuda_gpu_sum` 给出三项 GPU activity 的时间： +1. 在 `cuda_gpu_sum` 中比较各 GPU activity 的累计 duration，找出最耗时的 kernel。 +2. 在 `cuda_kern_exec_trace` 中读取每个 kernel 的 start 和 duration，并查看它对应的 host launch API。用 `end = start + duration` 算出结束时间；CUDA API 也采用同样的计算。 +3. 用 `ReLU start - (GEMM start + GEMM duration)` 计算两个 kernels 之间的空隙；这里使用 trace 时间戳和 duration 两列。 +4. 最后查看 `nvtx_*`、`cuda_api_trace` 和 `cuda_api_sum`，解释 host range 和同步 API。判断一次同步是否真的在等待 GPU，要比较 `cuda_api_trace` 中同步 API 的开始时间与最后一个 kernel 的结束时间。 + +先看两个 kernels 的 GPU 执行时间： | GPU activity | 次数 | GPU duration | 占所列 GPU 时长总和的比例 | |---|---:|---:|---:| -| 32 MiB H2D copy | 1 | 607.230 μs | 85.4% | -| BF16 GEMM | 1 | 93.152 μs | 13.1% | -| ReLU | 1 | 11.072 μs | 1.6% | +| BF16 GEMM | 1 | 92.608 μs | 89.4% | +| ReLU | 1 | 10.944 μs | 10.6% | -`cuda_kern_exec_trace` 把每个 kernel 与对应的 launch API 关联起来，并分别给出 API time、positive queue time 和 GPU execution。这里的 positive queue time 是 launch API 返回后到 kernel 开始前的等待时间；如果 kernel 更早开始，这一项就没有正值。 +再把 GPU execution 与 host 上的 launch API 对应起来。Positive queue time 指 launch API 返回后，到 kernel 稍后才开始之间的时间；kernel 在 API 返回前已经开始时，该字段为空。 | Kernel | API time | Positive queue time | GPU execution | |---|---:|---:|---:| -| BF16 GEMM | 34.270 μs | 403.214 μs | 93.152 μs | -| ReLU | 11.064 μs | 442.166 μs | 11.072 μs | +| BF16 GEMM | 50.717 μs | — | 92.608 μs | +| ReLU | 13.474 μs | 5.074 μs | 10.944 μs | -这份报告可以读出以下结论： +按下面的顺序读： -1. **主要时间花在 H2D copy。** 它占三项 GPU duration 之和的 85.4%。若只看 `cuda_gpu_kern_sum`，这次 copy 会被完全漏掉；因此这里使用同时包含 kernels 和 memory operations 的 `cuda_gpu_sum`。 -2. **Queue time 不是 launch overhead。** GEMM 与 ReLU 都在同一条 stream 上等待更早提交的工作。它们分别排在 H2D copy 和 GEMM 后面，所以 queue time 远大于各自的 API time。 -3. **较长的同步 API 区间通常表示 host 正在等待 GPU。** `cudaDeviceSynchronize` 在 CPU 上持续了 384.821 μs，说明调用它时仍有 GPU 工作没有完成；这个数字不是某个 kernel 的时间，也不是整个 operation 的 latency。 -4. **不同范围的时间不能相加。** 三项 GPU duration 相加是 711.454 μs。`nvtx_gpu_proj_sum` 以该 NVTX range 中第一项 GPU work 的开始为起点、最后一项的结束为终点，得到 715.582 μs；约 4.1 μs 的差值是 activities 之间的空隙。CPU 上的原始 `target operation` NVTX range 则是 870.561 μs，其中还包含 dispatch 和最后的同步等待。 +1. **先选目标。** GEMM 占两个 kernel 总执行时间的 89.4%，因此 NCU 的分析目标选为 GEMM。 +2. **区分 API time 和 GPU time。** GEMM launch API 在 host 上用了 50.717 μs，GPU 执行 GEMM 用了 92.608 μs；这是两个不同区间。ReLU 的 5.074 μs positive queue time 表示 launch API 返回后，它还等了一小段时间才在 GPU 上开始。 +3. **检查 kernels 之间的空隙。** 两个 duration 相加为 103.552 μs；从 GEMM 开始到 ReLU 结束的 GPU span 为 103.776 μs，因此中间只有 0.224 μs 空隙。 +4. **按计时边界解释 host range。** 图中的 `target operation` NVTX range 覆盖 Python/PyTorch dispatch、两次 launch 和同步 API。`cudaDeviceSynchronize` 开始时 ReLU 已经结束，因此它的 duration 主要来自 host 侧的 API 开销；此时 GPU 执行已经完成。 -在另一轮关闭 profiler 的测量中，同一 operation 的 20 个 CUDA Event samples 得到 722.816 μs 的 median，范围为 718.400–726.336 μs。可以用下面的命令复现相同的测量方法： +分析其他报告时，也先确认采集范围，再看 GPU kernels、copies、空隙和重叠，最后关联到 host launch 或同步 API。更多定义见 [Nsight Systems Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html)。 -```bash -python appendix/nsys_example.py --event-samples 20 -``` +## 使用 Nsight Compute 分析单个 kernel + +Nsight Systems 用来选择目标 kernel；NCU 用来解释这个 kernel 的硬件行为。下面先给出一套适用于不同 kernel 的阅读顺序，再用 B200 上的 BF16 GEMM 演示完整过程，最后解释相关字段和计算。 + +### 如何阅读一份 NCU 报告 + +SM（Streaming Multiprocessor）是 GPU 上接收 thread blocks 并执行指令的计算单元。一个 warp 由 32 个 threads 组成，是 scheduler 选择和发射指令时使用的基本线程组；block 或 warp 已经分配到某个 SM、尚未执行结束时，称为驻留（resident）。 + +分析一份新的 NCU 报告时，按下面的顺序分析： -这个无 profiler 结果才适合报告性能；Nsight Systems 的一次时间线用于解释时间花在哪里，两者不要求数值完全相同。 +| 当前问题 | 最先查看的位置 | 这一步的作用 | +|---|---|---| +| 报告属于哪个 kernel？ | 报告头中的 kernel、device、grid/block，以及 Warnings/Errors | 确认筛选和采集结果 | +| Grid 是否提供了足够多的 blocks？ | `LaunchStats` 的 `Grid Size`、`Waves Per SM` | 判断 grid 是否提供足够的整机并行度 | +| 每个 SM 同时能驻留多少 blocks 和 warps？ | `Occupancy` 的各项 `Block Limit`、theoretical/achieved occupancy | 计算理论与实测驻留量，并找出决定理论上限的资源 | +| 下一步先查哪一侧？ | `SpeedOfLight` 的 Compute、Memory、DRAM throughput | 选择计算、访存或调度分支 | +| Scheduler 能否持续找到可发射的指令？ | `Scheduler Statistics` → `Warps Per Scheduler`；可发射工作较少时再看 `Warp State Statistics` → `Warp State (All Cycles)` | 比较已驻留、已就绪和实际发射的 warps；就绪工作较少时再查看主要等待状态 | + +`Grid Size` 是这次 launch 提交的 block 总数。`Waves Per SM` 用这个总数除以整块 GPU 理论上可同时驻留的 block 数；1 表示两者相等，3.46 表示 grid 的 block 总数是全卡理论同时容纳量的 3.46 倍。这个容量比例用于判断 grid 是否有足够多的 blocks 覆盖整块 GPU，实际的 block 调度顺序不包含在该字段中。 -这组数据也说明了为什么必须先确定计时边界。如果实际应用中的 operation 确实包含 H2D copy，优化重点应首先考虑减少或重叠传输；如果输入早已位于 GPU，这次 copy 就不应放进被测范围。不能因为 GEMM 是主要的计算 kernel，就默认它是首要优化对象。 +各项 `Block Limit` 分别给出 registers、shared memory、threads 等单项资源允许每个 SM 驻留的 block 上限，其中最小值决定理论 block 上限。Theoretical occupancy 是按这些上限算出的最大驻留 warp 数占硬件容量的比例；achieved occupancy 是采集期间实际平均活跃 warp 数的比例。它们描述并发驻留量，是否影响执行速度还要结合 scheduler 指标。 -分析其他报告时也采用同样的顺序：先确认 NVTX range 中没有 warm-up 和初始化，再查看 GPU streams 上的 kernels、copies、空隙与重叠，随后沿 correlation 回到 host 上的 launch 或同步 API。不同 streams 上的 duration 不能直接相加；看见重叠也只说明它在这次采集中发生，是否降低了 latency 仍要通过相同边界的无 profiler 测量验证。 +`SpeedOfLight` 中的 Compute 表示最忙的 SM 计算路径，Memory 表示最忙的内存侧路径，DRAM 只看外部显存接口；三者都以各自的可持续峰值为分母。B200 的外部显存是 HBM，L2 是全 GPU 共享的 cache，L1/TEX 是 SM 一侧处理内存请求的路径。Memory 较高只说明某条内存侧路径繁忙，外部 HBM 是否接近饱和由 DRAM 字段判断。`ComputeWorkloadAnalysis` 中的 active cycles 表示流水线仍在处理工作的周期，`Issue Slots Busy` 表示 scheduler 实际使用了多少指令发射机会。 -Report scripts 会随 Nsight Systems 版本变化。运行 `nsys stats --help-reports` 可以查看当前版本支持的名称，并应在实验记录中保留 `nsys --version`。[Nsight Systems User Guide](https://docs.nvidia.com/nsight-systems/UserGuide/index.html) 介绍了 CLI 和 GUI，[Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html) 则进一步解释 API、queue 和 kernel execution time。 +`SchedulerStats` 把 warps 分成几种状态：active warp 已经驻留且尚未结束；eligible（已就绪）warp 的下一条指令已经解码、依赖已经就绪，而且所需执行单元可用；issued warp 在当前周期实际发出了指令。确认目标 launch，并依次看完 grid、驻留量和 `SpeedOfLight` 后，再按结果选择下一组指标： -## 采集单个 kernel 的 Nsight Compute 报告 +- **Compute 更接近峰值：** 在 `ComputeWorkloadAnalysis` 的 `Pipe Utilization (Elapsed Cycles)` 中打开 `Pipe Utilization (% of elapsed cycles)`，找出 active cycles 最高的计算路径，再看同一 section 摘要中的 `Issue Slots Busy`。一条路径长时间有操作在执行、发射槽却大多为空时，继续打开 `Scheduler Statistics` → `Warps Per Scheduler`，解释较低的指令发射率。 +- **Memory 更接近峰值：** 采集 `MemoryWorkloadAnalysis`、`MemoryWorkloadAnalysis_Chart` 和 `MemoryWorkloadAnalysis_Tables`。先在 `Memory Workload Analysis Chart` → `Memory Chart` 中沿 DRAM、L2 和 L1/TEX 查看数据流，再在 `Memory Workload Analysis Tables` → `Memory Tables` 中读取 throughput、read/write bytes、hit rate，以及 shared/local memory 字段。DRAM 也接近峰值时先调查外部 HBM 流量；DRAM 较低时则把重点移到 L2、L1/TEX、shared memory 或 local memory。这里的 local memory 是每个 thread 私有的地址空间，其物理流量由 L1/L2 cache 和外部显存层级承载。 +- **Compute 和 Memory 都低：** 先用 grid 和 waves 判断这次 launch 是否提供了足以覆盖 GPU 的 blocks。Blocks 不足时，检查 grid、block、cluster 配置，或把工作拆成更多 blocks；blocks 足够时，再查看 `SchedulerStats`。有 active warps 却几乎没有 eligible（已就绪）warps 时，再用 `WarpStateStats` 查看它们在等数据、同步还是其他依赖。 +- **Compute 和 Memory 都高：** 分别展开两侧，各自找出一个具体候选，再通过只改变一个因素的实验判断哪一项真正影响 kernel 时间。 -Nsight Systems 告诉我们各个 kernel 在什么时候运行；选定一个 kernel 后，Nsight Compute 可以继续查看它的启动配置、occupancy、计算与访存吞吐，以及 scheduler 状态。采集这些指标时，NCU 可能多次重放同一个 kernel，因此它适合诊断原因，不应使用报告中的 `Duration` 代替正常运行时测得的 latency。 +`MemoryWorkloadAnalysis` 汇总整条 kernel 在 DRAM、L2、L1/TEX 和其他内存路径上的流量与 cache 行为。某一条 load 的具体依赖还要用 `SourceCounters` 把采样到的 stall 和指令活动映射到 SASS（GPU 机器指令）或源码位置。 -上一节的时间线表明，示例 operation 的主要时间花在 H2D copy。下面选择其中耗时 93.152 μs 的 BF16 GEMM 演示 NCU 的读法；这并不表示 GEMM 是整个 operation 最该优化的部分。 +“高”和“低”要结合当前 GPU、workload 和同一份报告判断。当报告已经指向一处可修改的代码，并且能够写出修改后预期变化的指标和 latency 时，这一轮分析就形成了可检验的假设。范围仍然过宽时，再沿上面的分支采集下一组 section。 -### 选择一次目标 kernel launch +### 完整示例：分析 B200 BF16 GEMM -继续使用上一节脚本的 `--profile-once` 模式，可以把 warm-up 留在采集范围之外，并且只执行一次目标 operation。这个 operation 会依次启动 GEMM 和 ReLU；下面通过 kernel-name filter 选中 GEMM，并用 `--launch-count 1` 只采集第一个匹配的 launch。命令采用 kernel replay，因此只适合能够独立重放的 kernel。若一段工作包含跨 kernel 依赖或并发，应先查看 Nsight Systems 时间线，再决定是否需要后文介绍的其他 replay mode。 +#### 1. 采集第一份 `basic` 报告 -先收集 `basic` section set： +继续使用脚本的 `--profile-once` 模式。应用在采集范围内提交一次 GEMM，再提交一次 ReLU；NCU 等待这个范围开始，并只筛选 GEMM： ```bash mkdir -p reports ncu \ --config-file off \ - --target-processes application-only \ --profile-from-start off \ - --kernel-name-base function \ --kernel-name 'regex:.*nvjet_sm100.*' \ --launch-count 1 \ --set basic \ @@ -429,15 +510,20 @@ ncu \ --pipeline-boost-state stable \ --export reports/bf16-gemm-basic \ --force-overwrite \ - python appendix/nsys_example.py --profile-once + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once ``` -`--profile-from-start off` 让 NCU 等待脚本中的 profiler API；正则表达式再从该范围中选出名称包含 `nvjet_sm100` 的 GEMM。生成的完整 kernel 名称会随 PyTorch 和 CUDA 版本变化，因此分析其他程序时，应先从 Nsight Systems 抄下实际名称，再编写更精确的 filter。 +几个关键选项分别控制采集范围、目标、指标和采集条件： -`--set basic` 采集 launch、occupancy、workload distribution 和高层 throughput sections。Cache 与 clock controls 会改变 profiling 条件，因此命令中将它们显式写出。`--cache-control all` 会在每次 kernel replay 前清理 NCU 能够控制的 GPU caches；这有助于稳定 counter 采集，却不等同于正式 benchmark 的 hot-cache 条件。不同 NCU release 的 section sets 和 defaults 可能变化；采集时应记录 `ncu --version`，并在目标机器上运行 `ncu --config-file off --list-sets` 查看实际配置。 +- `--profile-from-start off` 让 NCU 等待脚本调用 `cudaProfilerStart()`。 +- `--kernel-name` 筛选名称包含 `nvjet_sm100` 的 kernel；`--launch-count 1` 为第一个匹配的 launch 生成一份结果。筛选表达式来自前面的时间线，分析其他程序时要换成实际名称。 +- `--set basic` 收集启动配置、occupancy 和高层吞吐率等第一轮所需指标。 +- `--replay-mode kernel` 允许 NCU 为收集硬件计数器而重放选中的 GEMM；`--cache-control all` 在 replay 前清理 NCU 可控制的 cache；其余两个选项控制采集期间的时钟和 pipeline boost 状态。本例的 GEMM 可以独立重放，因此适合 kernel replay；涉及跨 kernel 依赖或并发时，应选择能够保留所需 application 或 range 状态的 replay mode。 -如果脚本不能调用 profiler start/stop，应把命令改为 `--profile-from-start on`（或删除 `--profile-from-start off`），再用 `--launch-skip N --launch-count 1` 选择 warm-up 后的某次调用。`--launch-skip` 只统计匹配 kernel 的 launches；filter 或 launch 顺序变化时，它可能选中另一个实例。[Nsight Compute CLI 文档](https://docs.nvidia.com/nsight-compute/NsightComputeCli/) -详细说明了 kernel 与 launch filters。 +这里的“一次 GEMM”指应用提交一次 launch。为了收集所需的硬件计数器，NCU 仍可能在内部重放这次 GEMM。500 次 warm-up 位于采集范围外，可避开初始化和 lazy loading；NCU 的 cache control 随后会改变正常的 warm-cache 条件。 在 GUI 中打开报告： @@ -445,194 +531,407 @@ ncu \ ncu-ui reports/bf16-gemm-basic.ncu-rep ``` -也可以直接在 terminal 中查看： +终端中也可以查看 Details 页： ```bash ncu --import reports/bf16-gemm-basic.ncu-rep \ --page details \ - --print-details header \ + --print-details all \ --print-metric-name label-name ``` -`header` 适合先查看主要指标；后文使用的 Work ID/CLC 明细和完整 throughput breakdown 可在 GUI 中展开，或把命令中的 `header` 改为 `all` 后打印。 +其他筛选条件和采集选项见 [Nsight Compute CLI 文档](https://docs.nvidia.com/nsight-compute/NsightComputeCli/)。如果命令报告 `ERR_NVGPUCTRPERM`，请按照 NVIDIA 的 [counter permission 指南](https://developer.nvidia.com/nvidia-development-tools-solutions-err-nvgpuctrperm-nsightcompute) 配置权限，或请系统管理员开放访问。 + +#### 2. 用 `basic` 报告确定下一步查什么 + +先确认报告头中的设备、kernel 名称、grid/block 和 warnings。它们与目标 launch 一致后，再看下面三组字段： + +| `basic` 中的观察 | 本例的下一步 | +|---|---| +| `Grid Size = 512 blocks`，`Waves Per SM = 3.46` | Grid 足以覆盖整块 GPU，继续看单个 SM 上的并发工作量 | +| registers 和 shared memory 的 `Block Limit` 都是 1；theoretical/achieved occupancy 为 12.50%/8.97% | 每个 SM 理论上最多驻留 8 个 warps；实际平均活跃 warp 数更低，后续用 scheduler 指标查看指令就绪情况 | +| Compute 77.74%，Memory 38.71%，DRAM 12.88% | Compute 最接近自身峰值，先展开计算侧；整机聚合 HBM throughput 仍有较大余量 | + +先解释 wave。当前资源限制只允许每个 SM 驻留 1 个 block，这台 B200 有 148 个 SM，所以全卡的理论同时容纳量是 148 个 blocks。整张 grid 有 512 个 blocks，$512 / 148 = 3.46$，也就是 grid 的工作量为这份理论容量的 3.46 倍。这个计算只用于解释 NCU 给出的容量比例；本例使用 thread-block clusters，分析时以报告中的 `Waves Per SM = 3.46` 为准。由此可以确认 grid 中有足够多的 blocks 覆盖全卡。 -## 分析 Nsight Compute 报告 +再看 occupancy。B200 每个 SM 最多驻留 2,048 个 threads；一个 warp 包含 32 个 threads，所以硬件上限是 64 个 warps。本例的启动配置每个 SM 理论上最多驻留一个 256-thread block，即 8 个 warps，因此 theoretical occupancy 为 $8 / 64 = 12.50\%$。这是相对于硬件容量较低的理论驻留并发度。`Achieved Occupancy = 8.97%` 是采集期间实际平均活跃 warp 数占同一硬件容量的比例，低于 12.50% 的理论上限。本轮先用 theoretical occupancy 确认一个 SM 的多个 schedulers 合计最多有 8 个 resident warps，再用 `SchedulerStats` 查看平均分到每个 scheduler 的 warps 中有多少已经就绪。 -下面的数据来自同一台 B200 上的一次真实采集，使用 Nsight Compute 2026.1。NCU 用 9 个 replay passes 完成了 `basic` 报告。表中的百分比表示相应 throughput 指标占硬件子系统可持续峰值的比例，不是直接用应用 FLOPs 除以芯片标称峰值得到的利用率。 +最后看吞吐率。Compute 比 Memory 更接近各自峰值，所以先进入计算分支。`compute-bound` 的含义更强：继续提速最终受到计算单元吞吐上限约束。`basic` 在这里给出初始方向；后续报告显示指令发射机会很少，而且绝大多数调度周期没有 eligible（已就绪）warp，因此此时直接贴上 `compute-bound` 标签会漏掉关键线索。 -| 指标 | 实测值 | +综合三组字段，本例的 grid 足以覆盖全卡，每个 SM 理论上可供调度的 warps 较少，计算侧又最接近自身峰值。下一份报告同时展开计算流水线并采集 scheduler 指标。 + +#### 3. 沿计算分支采集后续指标 + +下面在一次 run 中同时采集三个 sections；阅读顺序是 `ComputeWorkloadAnalysis` → `SchedulerStats` → `WarpStateStats`。分析新的 kernel 时，可以根据前一组结果决定是否追加下一组。 + +```bash +ncu \ + --config-file off \ + --profile-from-start off \ + --kernel-name 'regex:.*nvjet_sm100.*' \ + --launch-count 1 \ + --section ComputeWorkloadAnalysis \ + --section SchedulerStats \ + --section WarpStateStats \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --export reports/bf16-gemm-followup \ + --force-overwrite \ + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once +``` + +这是一次独立的 NCU run。两份报告中的百分比可能有小幅波动，例如 77.74% 变成 78.39%，这个差值不表示性能发生了变化。前一份 `basic` 报告用于选择调查入口，下面的 follow-up 报告则用同一次采集中的计算、scheduler 和 warp-state 指标收窄原因。 + +#### 4. 按顺序读取三组指标 + +先在 `ComputeWorkloadAnalysis` → `Pipe Utilization (Elapsed Cycles)` → `Pipe Utilization (% of elapsed cycles)` 中看 active-cycle 视图。`Tensor (FP)` 对应浮点 tensor 运算路径，`TMEM (Tensor Memory)` 是为 tensor 运算服务的片上内存路径；TMEM 与外部 HBM/DRAM、负责异步搬运的 TMA 都是不同硬件。它们在约 78% 的时钟周期里处于 active 状态，而同一 section 摘要中的 `Issue Slots Busy` 只有 3.20%。多周期操作发射后可以让流水线持续 active，因此计算路径经常被占用和新指令发得很少可以同时出现。 + +接着打开 `Scheduler Statistics` → `Warps Per Scheduler`，找出新指令为什么发得少。每个 scheduler 平均有 1.44 个尚未结束的 active warps，其中只有 0.04 个 eligible（已就绪）warp；这里的 0.04 是平均 warp 数。`No Eligible` 的分母只包括这个 scheduler 所属 SM subpartition 中至少有一个 warp 尚未结束的周期；`No Eligible = 96.11%` 表示在这些周期中，96.11% 的周期找不到 eligible warp。到这一步，3.20% 的低发射率已经有了解释：工作虽然驻留在 SM 上，大部分时间却没有可以继续执行的 warp。 + +最后打开 `Warp State Statistics` → `Warp State (All Cycles)`，查看这些 warps 在等什么。一个 warp 在某种状态停留一个时钟周期，记作 1 个 warp-cycle。本例每条 issued warp instruction 对应 37.00 个 warp-cycles，其中 32.11、约 87%，归入 `Long Scoreboard`。这个 87% 是按已发射 warp 指令归一化后的 warp-state 周期占比，并非 kernel 执行时间占比。Scoreboard 是硬件记录前序操作结果是否就绪的依赖表；`Long Scoreboard` 表示下一条指令仍在等待由 L1TEX 处理的某项内存操作完成。L1TEX 位于 SM 一侧，负责处理 global、local、surface 和 texture memory 请求；请求最终可能由 L1、L2 或 DRAM 提供数据，因此这个字段无法单独确定等待发生在哪一级存储。`MemoryWorkloadAnalysis` 用于查看整条 kernel 的 L1、L2 和 DRAM 聚合行为；具体 load 则在 `SourceCounters` 的 SASS/源码视图中继续定位。 + +把三份 section 接起来，得到下面这条证据链： + +```text +tensor/TMEM 路径经常 active +→ scheduler 很少发出新指令 +→ 多数周期没有 eligible（已就绪）warp +→ 最主要的等待来自 Long Scoreboard +``` + +所以先查 L1TEX 相关的数据依赖。每个 SM 理论上最多驻留 8 个 warps，这可能让数据等待更难被其他工作覆盖。上一份 `basic` 报告中的 `DRAM Throughput = 12.88%` 表明整机聚合 HBM 带宽没有接近饱和；它仍允许个别请求访问 DRAM 并产生较长延迟。当前调查顺序是先看数据依赖，再看 Tensor Core 吞吐率是否构成下一层限制。当前的 `nvjet` GEMM 来自库实现，下一步可以采集 `MemoryWorkloadAnalysis`，比较整条 kernel 在 L1、L2 和 DRAM 上的聚合流量、throughput 与 cache 行为。 + +换成自己编写的 kernel 后，这份报告给出两个可以分别尝试的方向： + +- **增加驻留 warps：** 调整 tile、block 或资源用量，让一个 SM 能同时驻留更多工作。 +- **缩短数据依赖等待：** 保持驻留数量不变，把 load 或预取提前，或者缩短依赖链。 + +具体修改方法和验证步骤放在后面的“用代码修改检验假设”中。 + +### 字段计算、单位与边界 + +主线已经给出了阅读顺序和本例结论。下面只保留完整字段、计算过程、指标单位和容易误读的边界，分析其他 kernel 时可以按需查阅。 + +#### `LaunchStats` 与 `Occupancy` + +`LaunchStats` 中与本例有关的完整启动字段如下： + +| 字段 | 本次报告中的值 | +|---|---:| +| `Grid Size` | 512 blocks | +| `Block Size` | 256 threads | +| `Cluster Size` | 4 blocks | +| `Waves Per SM` | 3.46 | + +- `Block Size = 256` 表示每个 block 有 256 个 threads，即 8 个 warps；稍后用它计算 occupancy。 +- `Cluster Size = 4` 表示每 4 个 blocks 组成一个 thread-block cluster，整张 grid 有 128 个 clusters；cluster 的 blocks 会按硬件支持的布局共同调度。 +- 对于包含本例在内的 cluster launch，直接读取 NCU 的 `Waves Per SM`；需要在代码中计算可驻留 clusters 时，使用 `cudaOccupancyMaxActiveClusters`。 + +每个 block 的资源用量位于 `LaunchStats`： + +| 字段 | 本次报告中的值 | |---|---:| -| Kernel duration | 95.87 μs | -| Grid / block size | 512 blocks / 256 threads | -| Cluster size | 4 blocks | -| Registers | 255 / thread | -| Dynamic shared memory | 213.28 KB / block | -| Waves per SM | 3.46 | -| Theoretical / achieved occupancy | 12.50% / 8.98% | -| SM compute throughput 指标 | 77.34% | -| Memory throughput 指标 | 38.51% | -| DRAM / L2 / L1-TEX throughput 指标 | 20.42% / 34.60% / 46.93% | +| `Registers Per Thread` | 255 | +| `Dynamic Shared Memory Per Block` | 213.28 KB | -NCU 中的 95.87 μs 与上一节 Nsight Systems 采集到的 93.152 μs 不完全相同。两者来自不同的 profiling run，而且 NCU 还改变了 cache、clock 和 replay 条件；这种差异正说明 profile 中的 `Duration` 不能替代正式 benchmark。 +这些资源最终允许的驻留数量位于 `Occupancy`： -阅读这张表时，先确认捕获对象，再看工作怎样铺满 GPU，最后才判断应展开哪类指标。 +| 字段 | 本次报告中的值 | +|---|---:| +| `Block Limit Registers` | 1 block / SM | +| `Block Limit Shared Mem` | 1 block / SM | +| `Theoretical Occupancy` | 12.50% | +| `Achieved Occupancy` | 8.97% | -### 1. Launch Statistics 与 Workload Distribution +一个 block 驻留到 SM 时，硬件会为它保留所需的 registers 和 shared memory。根据 [NVIDIA Blackwell Tuning Guide](https://docs.nvidia.com/cuda/blackwell-tuning-guide/index.html#occupancy)，B200 的每个 SM 有 65,536 个 32-bit registers、228 KB shared memory，最多驻留 2,048 个 threads。 -报告中的名称是 `nvjet_sm100_tst_128x256_64x6_2x2_2cta_h_bz_NNT`，与上一节时间线中的 GEMM 一致。它以 512 个 blocks、每 block 256 个 threads 启动，并把 4 个 blocks 组成一个 cluster。`Waves Per SM = 3.46` 表示整张 grid 需要三个完整 waves，再加一个不完整 wave 才能执行完；它描述的是 grid 在时间上覆盖 GPU 的方式，不是 occupancy。 +- `Registers Per Thread = 255`。一个 block 有 256 个 threads，因此约需 $255 \times 256 = 65{,}280$ 个 registers，几乎占满一个 SM 的 register file。两个 blocks 需要 130,560 个，已经超过 65,536。 +- `Dynamic Shared Memory Per Block = 213.28 KB`。一个 block 已经使用了 228 KB 上限中的绝大部分；两个 blocks 至少需要 426.56 KB，也无法同时驻留。 +- `Block Limit Registers = 1` 和 `Block Limit Shared Mem = 1` 正是前两项计算的结果：只看 registers 或只看 shared memory，任意一项资源都只允许每个 SM 驻留一个 block。NCU 的精确计算还会考虑资源分配粒度和驱动占用的 shared memory。 +- `Achieved Occupancy = 8.97%`。这是采集期间实际平均活跃的 warps 占硬件上限的比例。它低于 12.50%，说明实际执行期间未始终保持理论驻留上限；它衡量驻留并发度，与“达到峰值性能的百分比”采用不同定义。 -这次报告还给出了 Work ID/Cluster Launch Control 警告：名义上启动 512 个 CTAs，报告只记录到 380 个获准执行的 CTAs。只要出现这种警告，依赖 block、warp 或 thread 数量的指标都要谨慎解释，不能把名义 launch 数量直接当作实际执行数量。 +#### `SpeedOfLight` 的分母与 `Duration` -### 2. Occupancy +`SpeedOfLight` 指标组给出下面四个字段： -这个 kernel 每个 thread 使用 255 个 registers，每个 block 使用 213.28 KB dynamic shared memory。报告中的 register limit 和 shared-memory limit 都只允许每个 SM 驻留一个 block，因此 theoretical occupancy 为 12.50%，实际采集到 8.98%。 +| 字段 | 本次报告中的值 | +|---|---:| +| `Duration` | 95.30 μs | +| `Compute (SM) Throughput` | 77.74% | +| `Memory Throughput` | 38.71% | +| `DRAM Throughput` | 12.88% | -这只能说明同时驻留的 warps 较少，不能单独证明 occupancy 是性能瓶颈。这个 GEMM 本来就采用 4-CTA clusters 和异步流水线；如果只为提高 occupancy 而减少 registers 或 shared memory，可能引入 spills、减少 tile reuse，反而变慢。 +`Duration` 只在各自的 profiler run 中解读：这次 NCU 采集为 95.30 μs，另一轮 Nsight Systems 采集为 92.608 μs。实现之间的性能比较采用无 profiler 的 CUDA Event baseline。NCU 还会控制时钟、清理 cache，并可能 replay 或串行化 kernel；详见 [NCU 的 workload-duration 说明](https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#workload-durations)。 -NCU 还会显示 `Est. Speedup` 等规则生成的提示。它们是在若干简化假设下估算的局部上限，用来提示值得调查的方向，不是修改 kernel 后可以期待的实际加速比。 +三个 throughput 百分比分别使用各自的硬件峰值作为分母，彼此不能相加，也不能当作执行时间占比。[Nsight Compute Profiling Guide](https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-structure) 给出了 throughput metric 的组成规则。 -### 3. Speed of Light +#### 计算流水线、scheduler 与 warp 状态 -这次 basic 报告中的 SM compute throughput 指标为 77.34%，memory throughput 指标为 38.51%，其中 DRAM 只有 20.42%。因此，现有证据不支持把它称为 DRAM-bound；下一步更合理的是展开 compute pipelines。 +##### `Compute Throughput Breakdown` 字段 -对其他 kernel，仍可先比较 compute 和 memory throughput 相对于各自 sustained peak 的比例： +报告位置是 `SpeedOfLight` → `GPU Throughput Breakdown` → `Compute Throughput Breakdown`： -- Compute 高而 memory 较低，说明应继续检查 compute pipelines； -- Memory 高而 compute 较低，说明应继续检查 memory hierarchy； -- 两者都低，通常应先检查 underfill、dependency latency、synchronization、imbalance 或缺少 - eligible warps，而不是直接归因于 peak throughput。 +| 字段 | 本次报告中的值 | +|---|---:| +| `SM: Mem Tensor Cycles Active` | 77.74% | +| `SM: Pipe Tc Cycles Active` | 77.48% | +| `SM: Pipe Tensor Cycles Active` | 77.42% | +| `SM: Pipe Alu Cycles Active` | 1.36% | +| `SM: Pipe Tma Cycles Active` | 1.10% | +| `SM: Pipe Fma Cycles Active` | 0.61% | -“Memory throughput 高”不等于“DRAM-bound”。限制项也可能来自 L1、L2、shared memory 或 -memory-instruction pipeline。展开 breakdown 后才能判断。 +- `Mem Tensor` 是 Blackwell 的 tensor-memory 相关片上路径。外部数据存放在 DRAM/HBM，异步多维搬运由 TMA 负责；这三个名称指向不同的硬件路径。77.74% 表明 `Mem Tensor` 是这份报告中最忙的计算侧组成项。 +- `Pipe Tc` 和 `Pipe Tensor` 是 NCU 报告中的两条不同流水线计数器。它们都在约 77%，与 BF16 GEMM 大量执行 tensor MMA 相关工作相符；两项可能覆盖重叠的硬件活动，需要分别解读，直接相加会重复计算。 +- `Pipe Alu` 主要对应通用整数与逻辑运算，`Pipe Fma` 覆盖普通 FP32 算术和部分整数乘加操作。它们分别只有 1.36% 和 0.61%，说明这个 GEMM 没有接近这些路径各自的峰值。 +- `Pipe Tma` 对应 Tensor Memory Accelerator 的异步数据搬运路径。1.10% 表示 TMA 距离自身峰值较远；数据供应还涉及 TMEM、cache、shared memory 和依赖延迟，需要结合相应指标判断。 -### 根据 Basic 报告选择下一组指标 +##### `Pipe Utilization` 的两个分母 -`basic` 报告只覆盖前 3 步。根据其中的线索，只采集回答下一个问题所需的 sections： +`ComputeWorkloadAnalysis` 摘要中的 `Issue Slots Busy` 为 3.20%。`Pipe Utilization` 的两个完整视图名称分别是 `Pipe Utilization (% of elapsed cycles)` 和 `Pipe Utilization (% of peak instructions executed over elapsed cycles)`。 -| `basic` 报告中的线索 | 下一步加入 | -|---|---| -| Registers、shared memory 或 resident blocks 构成限制 | `LaunchStats` 和 `Occupancy` 已包含在 `basic` 中；先阅读其中的 limit tables | -| Compute path 更可能构成限制 | `ComputeWorkloadAnalysis` | -| Memory hierarchy 更可能构成限制 | `MemoryWorkloadAnalysis`；用 `_Chart` 查看图形 breakdown，用 `_Tables` 查看详细 requests 与 sectors | -| Eligible warps 太少，或指令发射存在无法解释的空隙 | `SchedulerStats`，再根据结果加入 `WarpStateStats` | -| 需要定位到 source 或 instruction | `SourceCounters` | +把同一条流水线放在一行后，两种视图的差异会更直观： + +| 流水线字段 | Active-cycle 视图 | Instruction-rate 视图 | +|---|---:|---:| +| `TMEM (Tensor Memory)` | 78.39% | 0.04% | +| `TC` | 78.12% | 0.38% | +| `Tensor (FP)` | 78.07% | 0.61% | -这次报告中的 compute 指标更高，因此对同一次隔离 launch 加入 `ComputeWorkloadAnalysis`： +两个视图采用不同分母，用于对照流水线占用周期和指令执行率；它们之间不做加减。 + +##### `SchedulerStats` 的完整字段 + +报告位置是 `Scheduler Statistics` → `Warps Per Scheduler`： + +| 字段 | 本次报告中的值 | +|---|---:| +| `GPU Maximum Warps Per Scheduler` | 16 | +| `Theoretical Warps Per Scheduler` | 2.00 | +| `Active Warps Per Scheduler` | 1.44 | +| `Eligible Warps Per Scheduler` | 0.04 | +| `Issued Warp Per Scheduler` | 0.04 | + +同一 section 的摘要还给出 `No Eligible = 96.11%`。 + +`GPU Maximum = 16` 是每个 scheduler 的硬件容量上限；`Theoretical = 2.00` 来自本 kernel 每个 SM 理论上最多驻留的 8 个 warps 除以 4 个 schedulers。 + +##### `WarpStateStats` 的归一化单位 + +一个 warp 在某种状态中经历一个周期记作一个 warp-cycle；4 个 warps 同时经历一个周期，就记作 4 个 warp-cycles。NCU 再用已发射的 warp 指令数对这些周期归一化。报告摘要给出第一项，`Warp State (All Cycles)` 表给出第二项： + +| 字段 | 本次报告中的值 | +|---|---:| +| `Warp Cycles Per Issued Instruction` | 37.00 warp-cycles / issued instruction | +| `Stall Long Scoreboard` | 32.11 warp-cycles / issued instruction | + +本例每发射一条 warp 指令，对应 37.00 个 warp-cycles，其中 32.11、约 86.8%，归入 `Long Scoreboard`。这里按所有 warps 的状态周期归一化，单位是 warp-cycles / issued instruction，与 scheduler 的平均发射率不同。 + +报告 rule 旁的 `Est. Speedup` 是该规则按自身模型估算的潜在 workload-time 降幅，用于安排调查优先级；实际加速仍由修改后的 benchmark 给出。 + +##### 其他常见的 warp states + +| Warp state | 直接含义 | 下一步检查 | +|---|---|---| +| `Short Scoreboard` | 通常在等待 shared-memory 或其他片上单元产生结果 | 查看 shared-memory 访问和对应源码 | +| `Barrier` | 等待其他 warps 到达同步点 | 检查不同 warps 的工作量和到达时间 | +| `Not Selected` | warp 已经 eligible（已就绪），但本周期选择了另一个 warp | 检查是否有许多 eligible warps 竞争发射机会 | + +#### 用 `SourceCounters` 定位 SASS 或源码 + +前面的 `WarpStateStats` 只能看到整条 kernel 主要在等什么；`SourceCounters` 再把采样到的等待和执行次数标到一条条 SASS 指令旁边。打开 NCU 的 Source 页后，可以看到等待集中在哪些指令附近。Binary 带有 line information 且 NCU 能找到源文件时，这些指令还会关联到 CUDA 源码行。它的数据来自对 warp stall reason 的周期性采样，以及指令数和部分访存指标。 + +当前 `nvjet` GEMM 来自库实现，没有可供本书导入的 CUDA 源文件，但 SASS 视图仍然可用。沿用前面的筛选和采集条件，运行： ```bash -mkdir -p reports ncu \ --config-file off \ - --target-processes application-only \ --profile-from-start off \ - --kernel-name-base function \ --kernel-name 'regex:.*nvjet_sm100.*' \ --launch-count 1 \ - --section ComputeWorkloadAnalysis \ + --section SourceCounters \ --replay-mode kernel \ --cache-control all \ --clock-control boost \ --pipeline-boost-state stable \ - --export reports/bf16-gemm-compute \ + --export reports/bf16-gemm-source \ --force-overwrite \ - python appendix/nsys_example.py --profile-once + python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --profile-once ``` -这份进一步采集的报告给出： +在 GUI 的 Source 页选择 SASS，或者从终端打印同一视图： -| Pipeline | Throughput 指标 | -|---|---:| -| TMEM | 77.23% | -| Tensor Core | 77.04% | -| Tensor FP | 76.90% | -| ALU / TMA / FMA | 均低于 2% | +```bash +ncu --import reports/bf16-gemm-source.ncu-rep \ + --page source \ + --print-source sass +``` -这些数据才把 basic 报告中笼统的 77.34% compute 指标落实到 Tensor Core 和 Tensor Memory 路径上。分析其他假设时沿用同一命令结构，并替换其中的 `--section` 行，不要把所有 sections 逐次累加到一份报告中。这样可以缩小报告并减少 replay overhead。 +先看 `Warp Stall Sampling (Not-issued Samples)` 和 `Instructions Executed`。前者记录采样时 warp scheduler 没有发出指令的观测次数，后者是对应 SASS 指令按 warp 统计的执行次数。如果前面的 `WarpStateStats` 以 `Long Scoreboard` 为主，而 Source 页又把相应 samples 集中到某条 load 附近，这条指令就是下一步检查的候选。这里使用的是周期性采样，结果表示热点位置；数据最终来自 L1、L2 还是 DRAM，仍要结合 `MemoryWorkloadAnalysis` 的整 kernel 聚合指标和代码中的访问关系判断。 -### 4. Compute 与 Memory Workload Analysis +对于自己编译的 TIRx kernel，可以把 SASS 继续关联到生成的 CUDA。先让 TVM 使用 NVCC、保留源码并写入 line information： -Compute Workload Analysis 会显示哪些 execution pipelines 正在工作。应分别检查 Tensor Core、 -FMA、ALU、special-function，以及相关的 asynchronous pipelines，不要从一个 aggregate compute -百分比推断 Tensor Core utilization。 +```bash +export TVM_CUDA_COMPILE_MODE=nvcc +export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels" +mkdir -p "$TVM_KERNEL_DUMP" +``` -Memory Workload Analysis 会区分 DRAM、L2、L1/TEX、shared memory 和 local-memory effects。将实际 traffic volume 与 bandwidth、cache hit rate 和 local-memory spill 放在一起看。更详细的 -sector 与 request tables 需要 `MemoryWorkloadAnalysis_Tables`；source-level coalescing 和 -shared-memory conflict 证据可能需要 `SourceCounters`。若 traffic 很小，单独一个高 cache hit -rate 并不能说明性能良好。 +设置环境变量后，重新启动 workload，让目标 kernel 在这个进程中重新编译。下面是采集命令模板；把 `YOUR_KERNEL_NAME` 和最后一行的程序路径换成自己的值： -### 5. Scheduler 与 Warp States +```bash +ncu \ + --config-file off \ + --kernel-name 'regex:.*YOUR_KERNEL_NAME.*' \ + --launch-count 1 \ + --section SourceCounters \ + --replay-mode kernel \ + --cache-control all \ + --clock-control boost \ + --pipeline-boost-state stable \ + --import-source yes \ + --source-folders "$TVM_KERNEL_DUMP" \ + --export reports/tirx-source \ + --force-overwrite \ + python path/to/your_workload.py +``` -Scheduler Statistics 展示 active、eligible 和 issued warps。首先判断 scheduler 是否经常没有 -eligible instruction 可以发出；只有这时，才需要使用 Warp State Statistics 深挖原因。NCU -文档明确提醒，并非所有 stalls 都可避免，它们也不会自动成为性能瓶颈。 +`ncu-ui reports/tirx-source.ncu-rep` 的 Source 页可以在 CUDA/SASS 关联视图中逐行查看指标。终端对应命令是 `ncu --import reports/tirx-source.ncu-rep --page source --print-source cuda,sass`。如果 `executable` 是自己的 TIRx 编译结果，`executable.mod.imports[0].inspect_source("cuda")` 可以直接打印生成代码；NCU 的逐行关联则依赖这次重新编译写入 binary 的 line information。 -常见状态只应作为线索： +#### 用代码修改检验假设 -| 状态 | 可以帮助判断 | 不能单独推出 | -|---|---|---| -| Long Scoreboard | 正在等待与 L1TEX 路径相关的数据依赖 | 每次等待都访问了 DRAM | -| Short Scoreboard | 正在等待 MIO 路径依赖，常见于 shared memory | 一定存在 bank conflict | -| Barrier | Warp 正在等待 synchronization dependency | 这个 barrier 没有必要 | -| Not Selected | Warp 已 eligible，但 scheduler 发出了另一个 warp | Scheduler 缺少可执行工作 | -| Math/MIO Throttle | 某个 pipeline 或 queue 承受较高压力 | 随意删掉几条指令就会变快 | +本例调用库提供的 `torch.mm`，无法直接修改 kernel 实现。下面补充自写 TIRx 或其他 DSL kernel 的具体修改和验证步骤。 -对于 warp-specialized kernels，aggregate stall percentage 还混合了行为刻意不同的 roles。修改 -synchronization 前，应先将结果对应到 producer、MMA、softmax 或 writeback role。 +例如，要检验“驻留 warps 太少，难以隐藏 L1TEX 等待”，可以调整 tile、block 或 pipeline stages，减少 `Registers Per Thread` 和 `Dynamic Shared Memory Per Block`。修改后重新看两项 block limit：`Block Limit Registers` 和 `Block Limit Shared Mem` 都从 1 提高到至少 2 时，第二个 block 才具备同时驻留的资源条件。随后确认其他 block limits 也不低于 2，并查看 NCU 重新计算的 theoretical active blocks 和 occupancy。减少 registers 可能造成 register spill（寄存器不足时溢出到 local memory），减少 shared memory 也可能损失数据复用，所以还要用 latency 判断总体收益。 -### 6. Source 与 SASS 对应关系 +另一种实验保持驻留数量不变，只把 load 或预取提前，或者缩短依赖链。若 `Long Scoreboard` 和 latency 一起下降，就支持“warps 等待数据的时间减少了”这个判断。这个值按已发射指令归一化，因此还要结合 latency 判断是否产生实际提速。一次只改变一个关键因素，然后检查三个对象： -SASS 视图和 instruction attribution 不依赖 CUDA line information；要把生成的 CUDA source -对应到 SASS，则 binary 必须包含 line information，而且 NCU 必须能找到 source 文件。通过 NVCC -编译 TIRx module 时，可以在编译前设置： +1. **正确性：** 对相同输入比较两份输出与 reference，使用相同 tolerance。本例脚本已经在计时和采集开始前使用 FP32 reference；替换成自己的实现后沿用同一份 reference 和 tolerance。 +2. **预测的指标：** 重新采集同一组 NCU sections，检查与本次预测相关的指标：驻留实验看 theoretical/active 和 eligible（已就绪）/issued warps，依赖链实验看 `Long Scoreboard` 等待。Occupancy 变高说明驻留实验达到了资源目标；性能收益仍看下一项 latency。 +3. **实际 latency：** 关闭 Proton、Nsight Systems 和 NCU，用与开头完全相同的 shape、dtype、输入策略、warm-up、CUDA Event 边界和 samples 分别测量两份实现。 ```bash -export TVM_CUDA_COMPILE_MODE=nvcc -export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels" -mkdir -p "$TVM_KERNEL_DUMP" +python appendix/nsys_example.py \ + --size 4096 \ + --warmup-calls 500 \ + --event-samples 20 ``` -设置 `TVM_KERNEL_DUMP` 后，TVM 会保留生成文件，并在 NVCC 编译时加入 `-lineinfo`。NCU 采集命令还要加入 `--import-source yes --source-folders "$TVM_KERNEL_DUMP"`。保存 -`inspect_source("cuda")` 仍便于手工对照，但它本身不能给已经编译的 binary 补上 line information。一个 Python line 可能 lower 成多条 CUDA 或 SASS instructions，异步 tile primitive 也可能只能在这些 lower-level views 中看清。 +比较修改前后的无 profiler median，并同时查看样本波动。正确性通过、指标按预测变化，而且 CUDA Event latency 稳定下降时，这份假设得到了支持。若只有 NCU 指标变化，说明代码已经改变了预期硬件行为，但这次修改尚未带来实际提速；接下来检查性能限制是否转移，或原先的判断是否还缺一环。 -### NCU 采集对实验条件的影响 +## 用 IKET 查看 DSL kernel 内部阶段 -NCU 采集会改变执行条件： +编写 warp-specialized TIRx kernel 时，IKET（In-Kernel Event Tracing）可以把 kernel 内部的阶段画成时间线。Nsight Systems 显示整个 kernel 的起止区间，NCU 汇总整个 launch 的硬件指标，IKET 则记录各个 warp role 何时执行 producer、等待、consumer 等代码段。 -- 它可能重放 kernel，以收集不同 counter groups； -- 默认 cache control 可能在 replay iterations 之间清理 GPU caches； -- 它可以控制 GPU clocks； -- Replay 可能串行化或改变原本并发的工作； -- Application replay 会重新运行整个程序，并要求各次运行的执行过程和 launch matching 具有确定性；它不能解决不确定的 launch 顺序； -- 具有跨 kernel 依赖的一段工作可能需要 range replay，而不适合只隔离重放一个 kernel。 +### 运行一个完整示例 -报告中应记录 NCU version、replay mode、cache control、clock control、所选 sections 与 kernel -filter。不要将 NCU 的 `Duration` 直接与无 profiler 的 hot-cache CUDA Event 结果比较，也不要在同一个 profiling 进程中同时启动 Proton 和 NCU。 +TVM 0.26 使用版本锁定的 `cutlass-4.6.0` profiling profile。本章的 CUDA 13 环境可以安装对应依赖，并先确认 `run-iket` 命令可用： -如果 NCU 报告 `ERR_NVGPUCTRPERM`，说明 hardware-counter access 受到限制。应按照 NVIDIA 的 -[counter permission 指南](https://developer.nvidia.com/nvidia-development-tools-solutions-err-nvgpuctrperm-nsightcompute) -配置权限，或请系统管理员开放所需访问；不应把所有实验长期使用 root 运行作为默认方案。 +```bash +python -m pip install \ + 'nvidia-cutlass-dsl[cu13]==4.6.0' \ + 'nvidia-cuda-nvdisasm==13.3.73' \ + 'nvidia-cuda-nvrtc==13.2.78' +run-iket --help +``` -## 使用 IKET 分析 kernel 内部阶段（可选） +下面的完整脚本位于 `appendix/iket_example.py`。一个 CTA 包含两个 warps：warp 0 把 256 个元素从 global memory 搬到 shared memory，两个 warps 在 CTA barrier 处汇合，warp 1 再读取 shared memory、执行计算并写回结果。三个 `range_push()` / `range_pop()` 区间分别标出 producer、等待和 consumer： -在 Nsight Systems 时间线中，一个 kernel 只显示为完整的 GPU 执行区间；NCU 给出的指标也覆盖整个 kernel。对于已经加入阶段标记的 warp-specialized TIRx kernel，可以使用 IKET（In-Kernel Event Tracing）查看不同 warp 分工在何时工作、等待或发生重叠。 +```python +"""Minimal TIRx workload with IKET ranges for two warp roles.""" -TVM 0.26 已为 TIRx 接入 IKET，目前要求 SM90 或更新的 CUDA target，并对 CUTLASS DSL、NVRTC 等工具版本有严格要求。IKET 会在 kernel 中加入记录代码，因此得到的时间只适合分析阶段关系，不能作为正式 latency。具体的版本要求、标记方法和 Perfetto trace 生成流程见 [`python/tvm/backend/cuda/iket.py`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/backend/cuda/iket.py) 和 [NVIDIA IKET 文档](https://github.com/NVIDIA/cutlass/blob/v4.6.0/media/docs/pythonDSL/cute_dsl_general/iket_profiling.rst)。 +from pathlib import Path -## 性能实验检查清单 +import numpy as np -发布 benchmark table 或 pull request 前，可以用下面的清单确认其他人能够重建这次测量： +import tvm +from tvm.script import tirx as T +from tvm.tirx.cuda import iket -| 类别 | 需要记录的内容 | -|---|---| -| Hardware | 准确的 GPU、设备数量、相关 topology、clock 与 power policy | -| Software | Driver、CUDA、framework、compiler、library versions 与 source commit | -| Workload | Shapes、dtype、layouts、mask、scale、epilogue、输入分布、batch/sequence 信息、状态与重置策略 | -| Correctness | Reference、tolerance、accumulation/output dtype、异常输入策略 | -| Timing | Timer 类型、kernel/operator/end-to-end 边界、stream policy、CUDA Graph、`warmup`/`repeat` 的数值与单位、`rounds` 以及每轮原始结果 | -| Cache | 是否复用输入、是否轮换输入、显式 flush policy，以及该策略是否符合应用 | -| Statistics | 原始 latency 单位、median/mean、spread、独立 runs、实现顺序 | -| Baseline | Library 与 algorithm、workspace、tuning budget、最终 configuration | -| Profiling | Proton/IKET/Nsight Systems/NCU versions、kernel filters、IKET ranges 与 trace 格式、Nsight Systems 采集选项与 trace、NCU sections 与 replay/cache/clock controls | - -整套流程是一个循环：benchmark 证明变化确实影响性能；profile 提供原因线索；下一次无 profiler -benchmark 再判断这个解释是否带来了真实改进。 + +N = 256 +ELEMS_PER_LANE = 8 + + +@T.prim_func +def warp_role_example(inp: T.Buffer((N,), "float32"), out: T.Buffer((N,), "float32")): + T.device_entry() + profiler = iket.IketProfiler() + warp_id = T.warp_id([2]) + lane = T.lane_id([32]) + shared = T.alloc_buffer((N,), "float32", scope="shared") + + profiler.mark("kernel_start", warp_id) + if warp_id == 0: + profiler.range_push("producer_load") + for i in T.serial(ELEMS_PER_LANE, unroll=False): + index = i * 32 + lane + shared[index] = inp[index] + profiler.range_pop() + + profiler.range_push("wait_for_data") + T.cuda.cta_sync() + profiler.range_pop() + + if warp_id == 1: + profiler.range_push("consumer_compute") + for i in T.serial(ELEMS_PER_LANE, unroll=False): + index = i * 32 + lane + out[index] = shared[index] * T.float32(2) + T.float32(1) + profiler.range_pop() + + +def profile_workload(): + target = tvm.target.Target({"kind": "cuda", "arch": "sm_100a"}) + executable = tvm.compile(warp_role_example, target=target, tir_pipeline="tirx") + module = executable.jit() + + input_numpy = np.arange(N, dtype=np.float32) + inp = tvm.runtime.tensor(input_numpy, device=tvm.cuda()) + out = tvm.runtime.empty((N,), "float32", device=tvm.cuda()) + module.main(inp, out) + tvm.cuda().sync() + + expected = input_numpy * 2 + 1 + np.testing.assert_array_equal(out.numpy(), expected) + + +def main(): + result = iket.run( + profile_workload, + output_dir=Path("reports/iket-warp-roles"), + postprocess="all", + clobber=True, + timeout=600.0, + ) + print(f"IKET output directory: {result.output_dir}") + for path in (*result.json_traces, *result.perfetto_traces, *result.html_reports): + print(f"IKET artifact: {path}") + + +if __name__ == "__main__": + main() +``` + +在 B200 上直接运行： + +```bash +python appendix/iket_example.py +``` + +`iket.run` 会在 IKET 采集进程中重新启动当前脚本，并调用 `profile_workload()`。把 `tvm.compile()` 和 `.jit()` 写在这个函数里，可以保证 kernel 在 IKET 已启用时重新编译和加载。脚本还会检查输出是否等于 `input * 2 + 1`。 + +`postprocess="all"` 在 `reports/iket-warp-roles` 下生成 JSON、`*.pftrace` 和 HTML。把 `*.pftrace` 加载到 Perfetto 后，可以分别查看 `producer_load`、`wait_for_data` 和 `consumer_compute`；warp 1 会比 warp 0 更早到达 barrier，因此它的 `wait_for_data` 区间通常更长。目标为 H100 时，把脚本中的 `sm_100a` 改成 `sm_90a`。 + +### 把标记迁移到自己的 kernel + +在自己的 `PrimFunc` 中创建 `IketProfiler`，用 `mark()` 记录瞬时事件，用成对的 `range_push()` / `range_pop()` 或 `range_start()` / `range_end()` 包住阶段。每个 warp 实际经过的控制流都要保持 range 成对；等待时间也需要像示例中的 `wait_for_data` 一样显式包住。 + +编译和第一次 JIT load 继续放在 `iket.run` 调用的函数内。IKET 支持 Hopper 或更新架构，并会检查 CUTLASS DSL packages、NVRTC、`nvdisasm` 和相关 binary 是否与锁定 profile 一致。IKET 插入的记录代码会改变生成的 kernel 并带来额外开销，因此 trace 用于解释阶段和 warp roles；正式 latency 仍由未插桩版本的 CUDA Event benchmark 给出。完整 API 和 trace 选项见 [`python/tvm/backend/cuda/iket.py`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/backend/cuda/iket.py) 和 [NVIDIA IKET 文档](https://github.com/NVIDIA/cutlass/blob/v4.6.0/media/docs/pythonDSL/cute_dsl_general/iket_profiling.rst)。 From 2e54fe3156c382c5976d596a7d47a782db703436 Mon Sep 17 00:00:00 2001 From: tlopex <820958424@qq.com> Date: Wed, 19 Aug 2026 23:07:14 -0400 Subject: [PATCH 4/4] Polish English GPU profiling tutorial --- appendix/benchmarking_gpu_kernels.md | 386 ++++++++++++++------------- 1 file changed, 197 insertions(+), 189 deletions(-) diff --git a/appendix/benchmarking_gpu_kernels.md b/appendix/benchmarking_gpu_kernels.md index 4f1121a5..572cb1bf 100644 --- a/appendix/benchmarking_gpu_kernels.md +++ b/appendix/benchmarking_gpu_kernels.md @@ -18,12 +18,12 @@ The tools used in this workflow have different jobs: | Tool | Primary question | |---|---| -| CUDA Events | How much GPU-stream time elapsed across the measured region? A stream is an ordered queue of GPU work. | +| CUDA events | How much GPU-stream time elapsed across the measured region? A stream is an ordered queue of GPU work. | | Synchronized wall-clock timer | How much wall-clock time elapsed between starting a host call and completing all GPU work required by it? | | Proton (provided by Triton) | Which GPU kernels were launched, how often, and which kernels account for most of the time? | | Nsight Systems | How do host work, streams, copies, kernels, and communication overlap on a timeline? | | Nsight Compute (`ncu`) | What is the selected kernel doing internally, and which resource or stall should be investigated next? | -| IKET (optional) | After adding in-kernel markers, when are the data-movement, compute, and writeback phases active, waiting, or overlapping? | +| IKET (optional) | After adding in-kernel markers, when do marked phases run, and where do warp roles wait or overlap? | ## Verify Correctness Before Timing @@ -51,14 +51,14 @@ atol = 1e-2 torch.testing.assert_close(actual, expected, rtol=rtol, atol=atol) ``` -`torch.set_float32_matmul_precision("highest")` prevents this CUDA FP32 reference from using -reduced-precision internal matrix multiplication. The example `rtol` and `atol` control relative and -absolute error, respectively. The value `1e-2` is only a runnable starting point; adjust it for the +`torch.set_float32_matmul_precision("highest")` prevents PyTorch from using reduced-precision +internal computation for this CUDA FP32 reference. The example `rtol` and `atol` specify the relative +and absolute error tolerances. The value `1e-2` is only a practical starting point; adjust it for the output dtype, accumulation method, shape, and operator contract. Use the same reference and tolerances throughout one comparison. -Reference computation and result comparison stay outside performance timing. Whether state reset is -timed depends on the operation boundary defined in the next section. +Keep the reference computation and result comparison outside the timed region. Whether state reset +is timed depends on the operation boundary defined in the next section. ## Define the Timing Boundary @@ -68,7 +68,7 @@ result. State explicitly whether compilation, input construction, allocation, or part of that operation. Implementations are directly comparable only when they perform the same work inside the measured boundary. -The GEMM-plus-ReLU example later in this chapter can use three different boundaries. CUDA Events +The GEMM-plus-ReLU example later in this chapter can use three different boundaries. CUDA events around `torch.mm` measure GEMM GPU-stream time. Events around the complete `run()` measure GEMM-plus-ReLU GPU-stream time. A CPU timer started before `run()` and stopped after synchronizing measures end-to-end latency for one Python call. Matrix allocation and warm-up remain outside all @@ -76,16 +76,17 @@ three boundaries by default. Once the scope is defined, choose the timer: -- **CUDA Events** record timestamps when a GPU stream reaches two points. They can measure the - device-timeline interval around one kernel or a complete operator. Kernels, memory copies, and idle - stream gaps inside that interval all count. For a multi-stream operation, measured work on every - participating stream must begin after the start event and join before the end event is recorded. +- **CUDA events** record timestamps when a GPU stream reaches two points. They can measure an + interval on the device timeline around one kernel or a complete operator. Kernels, memory copies, + and idle stream gaps inside that interval all count. For a multi-stream operation, work on every + participating stream must be ordered after the start event and complete before the end event is + recorded. - A **synchronized wall-clock timer** starts before the host call and stops after all GPU work required by that call has completed. It also includes Python dispatch, CUDA launch, and the wait for GPU completion, so it is appropriate for end-to-end call latency. -For example, if the GPU executes the start event before the host submits the next launch, that idle -stream time remains inside the CUDA Event interval. An Event interval is therefore not necessarily +For example, if the stream reaches the start event before the host submits the next launch, that idle +stream time remains inside the CUDA event interval. A CUDA event interval is therefore not necessarily the same as a kernel's start-to-finish execution interval in a profiler. ## Measure GPU Time and Single-Call Latency @@ -95,10 +96,10 @@ the same as a kernel's start-to-finish execution interval in a profiler. CUDA launches are normally asynchronous, so an unsynchronized CPU timer can stop before the GPU finishes and mostly reflect host submission time. The [PyTorch CUDA semantics documentation](https://docs.pytorch.org/docs/stable/notes/cuda.html#asynchronous-execution) -describes this behavior. The following benchmark uses CUDA Events to measure elapsed time on the +describes this behavior. The following benchmark uses CUDA events to measure elapsed time on the current stream. -The following runnable CUDA Event benchmark allocates its matrices, runs a warm-up, measures an FP16 +This runnable CUDA event benchmark allocates its matrices, runs a warm-up, measures an FP16 GEMM over five rounds, and reports the median round result: ```python @@ -117,7 +118,7 @@ def gemm(): def measure_batch_ms(fn, calls): - """Return mean CUDA Event time per call for one batch of back-to-back calls, in ms.""" + """Return mean CUDA event time per call for one batch of back-to-back calls, in ms.""" start = torch.cuda.Event(enable_timing=True) end = torch.cuda.Event(enable_timing=True) @@ -142,21 +143,21 @@ samples_ms = [measure_batch_ms(gemm, repeat) for _ in range(rounds)] print(f"device: {torch.cuda.get_device_name()}") print(f"calls per round: {repeat}") print("round samples (ms):", [round(x, 4) for x in samples_ms]) -print(f"median CUDA Event time: {median(samples_ms):.4f} ms") +print(f"median CUDA event time: {median(samples_ms):.4f} ms") ``` -`measure_batch_ms` records start and end events in the current CUDA stream and divides their elapsed +`measure_batch_ms` records start and end events on the current CUDA stream and divides their elapsed time by the number of calls. The result is the average GPU-stream time per GEMM across back-to-back -calls. `end.synchronize()` makes the CPU wait for that round of GPU work so that the Event result can -be read. +calls. `end.synchronize()` makes the CPU wait for that round of GPU work so that the event timing +result can be read. -`warmup_calls=500` and `repeat=100` count invocations; `rounds=5` requests five independent +`warmup_calls=500` and `repeat=100` are invocation counts; `rounds=5` requests five independent measurement batches. These values came from stability testing on the B200: results were still falling after 50 warm-up calls but stabilized by 500, and `repeat=100` was more stable than `repeat=10`. For -another workload, first increase -`warmup_calls` until the early rounds stop drifting; if variation remains large, increase `repeat` or -`rounds`. If longer runs instead shift the overall timing level, inspect temperature, power, and clock -behavior. +another workload, first increase `warmup_calls` until the early rounds stop drifting; if variation +remains large, increase `repeat` or +`rounds`. If timings shift systematically as the run gets longer, inspect temperature, power, and +clock behavior. This code always reuses the same matrices, so it represents a warm-cache workload with repeated inputs. Whether the accesses actually hit in cache still depends on the total amount of data revisited @@ -167,7 +168,7 @@ This book uses TVM's to handle warm-up, repeated timing, and statistics. The supplied function only launches a prepared implementation; inputs, outputs, and workspace remain outside the timed interval. Unlike the manual warm-cache example, `bench` writes a 256 MiB buffer before each measured invocation to reduce reuse -of data retained in L2 from the previous invocation, then records an independent CUDA Event interval: +of data left in L2 by the previous invocation, then records an independent CUDA event interval: ```python from tvm.tirx.bench import bench @@ -188,8 +189,8 @@ print(result["impls"]["gemm"]) # five-round mean, in us print(result["round_samples"]["gemm"]) # result from each round ``` -`warmup=25` and `repeat=100` are millisecond budgets. The Event timer uses a short calibration run to -convert them into invocation counts. The reported Event intervals cover only the measured calls; the +`warmup=25` and `repeat=100` are millisecond budgets. The event timer uses a short calibration run to +convert them into invocation counts. The reported event intervals cover only the measured calls; the 256 MiB write used to reduce L2 reuse occurs before each start event. `rounds=5` runs five rounds, and `cooldown_s=1.0` pauses before each one. `impls` stores the five-round mean, while `round_samples` stores the individual results. Adjust budgets @@ -197,15 +198,14 @@ and rounds by the same stability criteria used above, and use the same settings implementation. The `run_bench` entry points in TIRx-kernels also use this helper. Outside distributed mode, omitting -`timer` selects Proton by default; specify `timer="event"` for a CUDA Event interval. A repeatedly +`timer` selects Proton by default; specify `timer="event"` for a CUDA event interval. A repeatedly invoked in-place kernel must still follow the reset rule above. If reset occurs inside the measured function, its cost belongs to the operation. ### Measure End-to-End Time for One Call -Use a synchronized wall-clock timer when the target is the complete interval from the start of a -Python call until its GPU work finishes. The following code continues with the `gemm()` defined and -warmed up above: +Use a synchronized wall-clock timer to measure the interval from the start of a Python call until its +GPU work finishes. The following code continues with the `gemm()` defined and warmed up above: ```python from statistics import median @@ -233,17 +233,18 @@ print(f"median end-to-end time: {median(host_samples_ms):.4f} ms") The first synchronization keeps unfinished earlier work outside the measurement. The second ensures that this GEMM finishes before the timer stops. Each sample contains exactly one call, so the result includes the Python call, CUDA launch, GPU execution, and the wait for completion. By comparison, the -CUDA Event benchmark above reports average GPU-stream time per GEMM across back-to-back calls. +CUDA event benchmark above reports average GPU-stream time per GEMM across back-to-back calls. Every implementation in a comparison must use the same timer and boundary. If both results are -reported, name them separately as *CUDA Event GPU time* and *single-call end-to-end time* so that +reported, name them separately as *CUDA event GPU time* and *single-call end-to-end time* so that their different boundaries remain visible. ### Advanced: Timing a Multi-Stream Operation -When one operation submits work to several CUDA streams, the timing stream must connect the start -and finish of every branch. In this example, `sin` and `cos` run on separate streams. The timing -stream waits for both branches before adding their results: +When one operation submits work to several CUDA streams, the event dependencies must ensure that +every branch starts after the start event and finishes before the end event is recorded. In this +example, `sin` and `cos` run on separate streams. The timing stream waits for both branches before +adding their results: ```python import torch @@ -305,48 +306,49 @@ left stream: wait(start) ─ sin ─ left_done right stream: wait(start) ─ cos ─ right_done ``` -The graph permits the two branches to execute concurrently; actual overlap depends on their GPU -resource use. Confirm the realized schedule in a Nsight Systems timeline. For formal measurement, -call `measure_operation_ms()` several times for warm-up, then call it repeatedly to collect -single-call samples and report their median and variation. +This dependency structure allows the two branches to execute concurrently; actual overlap depends on +their GPU resource usage. Confirm the actual execution schedule in a Nsight Systems timeline. For +reported measurements, call `measure_operation_ms()` several times for warm-up, then call it +repeatedly to collect single-call samples and report their median and variation. #### PDL Within One Stream -Programmatic Dependent Launch (PDL) applies to a custom CUDA or DSL launch path that explicitly -enables it. The primary and secondary kernels are submitted to the same stream. After the primary -emits a trigger, the secondary may begin preparation that does not depend on the primary's result; -it performs the PDL dependency synchronization before consuming that result. +Programmatic Dependent Launch (PDL) is available to custom CUDA and DSL launch paths that explicitly +enable it. The primary and secondary kernels are submitted to the same stream. After the primary +emits a trigger, the secondary may begin preparation that does not depend on the primary's result. +Before consuming that result, the secondary waits on the PDL dependency. ```text primary: initial work ─ trigger ─ remaining work secondary: preamble ─ wait ─ dependent work ``` -Record `start` before launching the primary and `end` after launching the secondary. That complete -Event interval is the GPU time of the launch sequence. The two kernels may overlap, so the sum of -their profiler durations can exceed the complete operation latency. Use a Nsight Systems timeline -to confirm the realized overlap. +Record `start` before launching the primary and `end` after launching the secondary. The resulting +CUDA event interval measures GPU-stream elapsed time for the complete launch sequence. The two +kernels may overlap, so the sum of their profiler durations can exceed the complete operation +latency. Use a Nsight Systems timeline to confirm the observed overlap. PDL creates an opportunity for concurrent execution, while the runtime may still serialize the -kernels. Kernel correctness must cover both schedules. The `torch.cuda.Stream` interface above does -not expose PDL launch attributes; a custom CUDA or DSL implementation supplies them. See the +kernels. Validate correctness both when the kernels overlap and when the runtime serializes them. The +`torch.cuda.Stream` interface above does not expose PDL launch attributes; a custom CUDA or DSL +implementation supplies them. See the [CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html) for the enablement details. -## Keep Experimental Conditions Consistent +## Keep Benchmark Conditions Consistent The examples above measure repeated calls after allocation and warm-up. If first-call latency or a complete application path is the target, include the relevant CUDA initialization, JIT compilation, and autotuning in the timing boundary and report that result separately. -Keep every per-round result and state whether the summary is a median or mean. A continuing trend -across rounds calls for checking warm-up, temperature, and clock state before summarizing the full -set of measurements. When comparing implementations, alternate their measurement order so that no +Keep every per-round result and state whether the summary is a median or mean. If results continue to +trend across rounds, check warm-up, temperature, and clock state before summarizing the full set of +measurements. When comparing implementations, alternate their measurement order so that no implementation is consistently measured on a colder or hotter device. Use one cache policy throughout the comparison. The manual example repeatedly reuses the same -matrices and is therefore biased toward warm-cache reuse, although the actual hit rate still depends -on the total amount of data revisited and the cache capacity. The TVM 0.26 Event and Proton timers +matrices and therefore models a warm-cache workload, although the actual hit rate still depends +on the total amount of data revisited and the cache capacity. The TVM 0.26 CUDA event and Proton timers instead write a 256 MiB buffer before each measured invocation to reduce L2 reuse; this write is outside the timed interval. Choose the policy that represents the target application. `torch.cuda.empty_cache()` only releases unused blocks from PyTorch's allocator; GPU L2 contents @@ -378,13 +380,14 @@ TFLOP/s = 2 × M × N × K / t_us / 10^6 The numerator and timing boundary must describe the same work. The complete GEMM-plus-ReLU operation below takes 105.152 μs. Dividing $2\times4096^3$ by that duration gives about 1307 TFLOP/s, but this -is only *effective throughput*: GEMM work divided by the complete operation time, whose denominator -also contains ReLU. To report the GEMM kernel's own TFLOP/s, the timed interval must cover only GEMM. +is only *effective throughput*: GEMM work divided by the complete operation time. The denominator is +the full operation time and therefore includes ReLU. To report the GEMM kernel's own TFLOP/s, the +timed interval must cover only GEMM. -A table that reports TFLOP/s, GB/s, or tokens/s should retain the original latency and explain how -the work was counted. For attention and fused kernels, also state whether the numerator represents -the full dense problem, the elements actually selected, or the work executed by the kernel. See -{ref}`chap_performance` for the corresponding formulas and roofline analysis. +A table that reports TFLOP/s, GB/s, or tokens/s should include the underlying latency measurement and +explain how the work was counted. For attention and fused kernels, also state whether the numerator +represents the full dense problem, the elements actually selected, or the work executed by the +kernel. See {ref}`chap_performance` for the corresponding formulas and roofline analysis. ## Find the Most Expensive Kernel with Proton @@ -401,8 +404,8 @@ def run(): torch.clamp_min(c, 0, out=output) ``` -Before entering any timing or collection mode, the script runs one operation and synchronizes. It -then computes an FP32 `torch.mm` reference, applies ReLU, converts the result to BF16, and compares it +Before any timed or profiled run, the script runs one operation and synchronizes. It then computes an +FP32 `torch.mm` reference, applies ReLU, converts the result to BF16, and compares it with the output using `rtol=2e-2` and `atol=1e-2`. A mismatch stops the command before the benchmark or profiler begins. This preflight check is outside both the baseline timing interval and the collection range used below. @@ -421,16 +424,16 @@ python appendix/nsys_example.py \ --event-samples 20 ``` -The 500 warm-up calls finish before formal timing. Each sample then uses one pair of CUDA Events -around one GEMM-plus-ReLU operation. One actual B200 run produced: +The 500 warm-up calls complete before measurement begins. Each sample then uses one pair of CUDA +events around one GEMM-plus-ReLU operation. A representative run on a B200 produced: ```text median=105.152 us, min=103.136 us, max=131.200 us ``` -The median is the baseline to revisit after changing code. The minimum and maximum show the sample -variation. This baseline covers the complete GEMM-plus-ReLU operation; the profiler tables below -report individual kernels from separate captures. +Use the median as the reference point when evaluating subsequent code changes. The minimum and +maximum show the sample variation. This baseline covers the complete GEMM-plus-ReLU operation; the +profiler tables below report individual kernels from separate captures. ### Use Proton to Compare the Kernels in the Operation @@ -442,7 +445,7 @@ Proton and `proton-viewer` are provided with Triton. The viewer also requires tw python -m pip install pandas llnl-hatchet ``` -The script's `--proton-calls` mode reuses the same `run()`, completes warm-up, profiles 100 calls to +The script's `--proton-calls` mode reuses the same `run()`, runs the warm-up, profiles 100 calls to the operation, and writes `operator.hatchet`: ```bash @@ -476,12 +479,12 @@ target_operation calls avg/us total/ms First confirm that both expected kernels appear and that each was called 100 times, then compare cumulative time. GEMM accounts for much more time, so it becomes the target for deeper analysis. -Before profiling it with NCU, use Nsight Systems to confirm the kernel order and gaps and correlate -them with the host launch APIs for one operation. +Before profiling it with NCU, use Nsight Systems to confirm the kernel order and gaps, then correlate +each kernel with its host launch API for one operation. This manual Proton session preserves the normal cache state, unlike `bench(timer="proton")`, which writes a 256 MiB buffer before each measured call. Use values from the same Proton report to rank -kernels. Compare implementations with the CUDA Event baseline above, using a timing interval that +kernels. Compare implementations with the CUDA event baseline above, using a timing interval that covers the full operation. ## Analyze an Application Timeline with Nsight Systems @@ -489,9 +492,9 @@ covers the full operation. Proton provides an aggregate ranking, but it does not show kernel ordering, gaps, copies, or host waits. Nsight Systems answers those questions with a timeline. -### Capture the Target Operation Timeline +### Capture a Timeline for the Target Operation -The script's `--profile-once` mode completes warm-up while the profiler is still inactive and waits +The script's `--profile-once` mode runs the warm-up while the profiler is still inactive and waits for that work to finish. It then submits exactly one GEMM-plus-ReLU operation between `cudaProfilerStart()` and `cudaProfilerStop()`: @@ -595,7 +598,7 @@ The command prints several tables in sequence. Extract values in this order: | ReLU | 1 | 10.944 μs | 10.6% | `cuda_kern_exec_trace` correlates each kernel with its launch API and reports API time, positive queue -time, and GPU execution separately. Positive queue time is the interval from API return to a later +time, and GPU execution time separately. Positive queue time is the interval from API return to a later kernel start. If the kernel starts before the API returns, the report shows no positive queue time. | Kernel | API time | Positive queue time | GPU execution | @@ -633,7 +636,7 @@ metrics and calculations. An SM (streaming multiprocessor) is a GPU compute unit that hosts thread blocks and executes their instructions. A warp contains 32 threads and is the basic group that a scheduler selects when issuing -an instruction. A block or warp is resident after it has been assigned to an SM and before it finishes. +an instruction. A block or warp is resident from the time it is assigned to an SM until it finishes. Analyze a new NCU report in this order: @@ -643,34 +646,34 @@ Analyze a new NCU report in this order: | Does the launch expose enough device-wide parallelism? | `Grid Size` and `Waves Per SM` in `LaunchStats` | Determine whether the grid provides enough device-wide parallelism | | How much work can reside on each SM at once? | `Block Limit` fields and theoretical/achieved occupancy in `Occupancy` | Quantify theoretical and achieved residency, and identify which resources set the theoretical limit | | Which part of the hardware should be investigated first? | Compute, Memory, and DRAM throughput in `SpeedOfLight` | Choose the compute, memory, or scheduling path | -| Can the scheduler consistently find an instruction to issue? | `Scheduler Statistics` → `Warps Per Scheduler`; when issue-ready work is scarce, continue to `Warp State Statistics` → `Warp State (All Cycles)` | Compare resident, ready, and issued warps; inspect the main wait states when ready work is scarce | +| Can the scheduler consistently find an instruction to issue? | `Scheduler Statistics` → `Warps Per Scheduler`; when few warps are ready to issue an instruction, continue to `Warp State Statistics` → `Warp State (All Cycles)` | Compare resident, ready, and issued warps; inspect the main wait states when ready work is scarce | -`Grid Size` is the number of blocks submitted by the launch. `Waves Per SM` divides that count by the -number of blocks that could theoretically reside across the whole GPU at once. A value of 1 means -the two counts are equal; 3.46 means that the grid contains 3.46 times the device-wide theoretical +`Grid Size` is the number of blocks submitted by the launch. `Waves Per SM` is the grid size divided +by the number of blocks that could theoretically reside across the whole GPU at once. A value of 1 +means the two counts are equal; 3.46 means that the grid contains 3.46 times the device-wide theoretical resident-block capacity. This ratio helps determine whether the grid contains enough blocks to cover the GPU. It does not record the actual block-scheduling order. Each `Block Limit` reports how many blocks one SM could host if registers, shared memory, threads, or another listed resource were the only constraint. The smallest value sets the theoretical block -limit. Theoretical occupancy is the maximum resident-warp count allowed by those limits as a fraction -of the hardware capacity. Achieved occupancy is the average active-warp count observed during -collection, expressed against the same capacity. These metrics describe resident concurrency; -scheduler metrics show whether that concurrency affects instruction issue. - -In `SpeedOfLight`, Compute represents the busiest SM compute path, Memory the busiest memory-side -path, and DRAM only the external-memory interface. Each uses its own sustainable peak as the -denominator. On a B200, HBM provides external device memory, L2 is shared across the GPU, and L1TEX is -the SM-side path that handles memory requests. A high Memory value means that some memory-side path -is busy; DRAM shows whether the external HBM interface is near saturation. In -`ComputeWorkloadAnalysis`, active cycles show when a pipeline still has work in flight, while +limit. Theoretical occupancy is the maximum resident-warp count allowed by those limits, expressed as +a fraction of the hardware capacity. Achieved occupancy is the average active-warp count observed +during collection, expressed as a fraction of the same capacity. These metrics describe resident +concurrency; scheduler metrics show whether that concurrency affects instruction issue. + +In `SpeedOfLight`, Compute represents the busiest SM compute path, while Memory represents the +busiest memory-side path. DRAM specifically measures the external-memory interface. Each uses its own +sustainable peak as the denominator. On a B200, HBM provides external device memory, L2 is shared +across the GPU, and L1TEX is the SM-side path that handles memory requests. A high Memory value means +that some memory-side path is busy; DRAM shows whether the external HBM interface is near saturation. +In `ComputeWorkloadAnalysis`, active cycles show when a pipeline still has work in flight, while `Issue Slots Busy` reports how many scheduler issue opportunities were used. `SchedulerStats` distinguishes several warp states. An active warp is resident and unfinished. An eligible warp has a decoded next instruction whose dependencies are ready and whose required -execution unit is available. An issued warp issued an instruction in the current cycle. After -confirming the target launch and examining the grid, residency, and `SpeedOfLight`, select the next -metrics: +execution unit is available. An issued warp has issued an instruction in the current cycle. +After confirming the target launch and examining the grid, residency, and `SpeedOfLight`, select the +next metrics: - **Compute is closer to its peak:** In `ComputeWorkloadAnalysis`, open `Pipe Utilization (Elapsed Cycles)` → `Pipe Utilization (% of elapsed cycles)` and find the compute @@ -691,16 +694,16 @@ metrics: warps exist but eligible warps are scarce, use `WarpStateStats` to see whether they are waiting on data, synchronization, or another dependency. - **Compute and Memory are both high:** Expand both sides, identify one specific bottleneck candidate - on each, and change one factor at a time to determine which candidate affects kernel time. + on each side, and change one factor at a time to determine which candidate affects kernel time. `MemoryWorkloadAnalysis` summarizes traffic and cache behavior for the whole kernel across DRAM, L2, L1TEX, and other memory paths. `SourceCounters` then maps sampled stalls and instruction activity to SASS (GPU machine instructions) or source locations to help trace a specific load dependency. -Judge “high” and “low” in the context of the current GPU, workload, and the metrics in the same -report. Once the report points to code that can be changed and predicts how its metrics and latency -should respond, the current pass has produced a testable hypothesis. If the scope is still too broad, -collect the next section along the relevant path above. +Interpret “high” and “low” in the context of the current GPU and workload, and compare values from the +same report. Once the evidence points to a specific code change and supports a prediction about how +the relevant metrics and latency should respond, the current pass has produced a testable hypothesis. +If the scope is still too broad, collect the next section along the relevant path above. ### Worked Example: A B200 BF16 GEMM @@ -732,7 +735,7 @@ ncu \ The key options control the capture range, target, metrics, and collection conditions: - `--profile-from-start off` makes NCU wait until the script calls `cudaProfilerStart()`. -- `--kernel-name` selects kernels in that range whose name contains `nvjet_sm100`, and +- `--kernel-name` selects kernels in that range whose names contain `nvjet_sm100`, and `--launch-count 1` produces a report for the first matching launch. The expression comes from the preceding timeline; replace it with the name from your own program. - `--set basic` collects the launch, occupancy, and high-level throughput sections needed for this @@ -775,35 +778,36 @@ Once they match the intended launch, use these three observations to choose what | Observation in `basic` | Next step in this example | |---|---| -| `Grid Size = 512 blocks`; `Waves Per SM = 3.46` | The grid supplies enough blocks to occupy the whole GPU; next check how much work can reside on each SM | +| `Grid Size = 512 blocks`; `Waves Per SM = 3.46` | The grid supplies enough blocks to cover all SMs; next check how much work can reside on each SM | | The register and shared-memory `Block Limit` values are both 1; theoretical/achieved occupancy is 12.50%/8.97% | Each SM can theoretically host at most eight resident warps, while the observed average is lower; use scheduler metrics to inspect instruction readiness | | Compute 77.74%, Memory 38.71%, DRAM 12.88% | Compute is closest to its own peak, so expand the compute side first; aggregate HBM throughput across the device still has ample headroom | -Start with the wave calculation. The current resource limits allow one resident block per SM, and -this B200 has 148 SMs, so the whole GPU's theoretical simultaneous capacity is 148 blocks. The grid -contains 512 blocks, and $512 / 148 = 3.46$: its block count is 3.46 times that theoretical capacity. +Start with `Waves Per SM`. The current resource limits allow one resident block per SM, and this B200 +has 148 SMs, so the theoretical device-wide residency capacity is 148 blocks. The grid contains 512 +blocks, and $512 / 148 = 3.46$: its block count is 3.46 times that theoretical capacity. This calculation only explains the capacity ratio reported by NCU. Because this example uses thread-block clusters, use the reported `Waves Per SM = 3.46` as the authoritative value. The ratio confirms that the grid contains enough blocks to cover the device. Next consider occupancy. A B200 SM can hold at most 2,048 threads. At 32 threads per warp, that is a hardware limit of 64 warps. This launch configuration permits at most one 256-thread block per SM, or -eight resident warps, so its theoretical occupancy is $8 / 64 = 12.50\%$. This theoretical resident -concurrency is low relative to the hardware capacity. `Achieved Occupancy = 8.97%` is the average -number of active warps observed during collection as a fraction of the same hardware capacity, below -the 12.50% theoretical maximum. The theoretical value establishes that all schedulers on one SM +eight resident warps, so its theoretical occupancy is $8 / 64 = 12.50\%$. The launch can therefore +use at most 12.50% of the hardware's warp-residency capacity. `Achieved Occupancy = 8.97%` is the +observed average active-warp count expressed as a fraction of that same capacity, below the 12.50% +theoretical maximum. The theoretical value establishes that all schedulers on one SM have at most eight resident warps in total; `SchedulerStats` then reports how many warps assigned to each scheduler are ready on average. Finally, compare throughput. Compute is closer to its peak than Memory, so investigate the compute side first. The term `compute-bound` makes a stronger claim: further speedup is ultimately constrained -by the throughput limit of the compute units. The `basic` report only selects an entry point. The -follow-up report shows that issue opportunities are rarely used and most scheduler cycles have no -eligible (ready) warp; calling the kernel `compute-bound` at this point would hide that key clue. +by the throughput limit of the compute units. The `basic` report only identifies where to begin the +investigation. The follow-up report shows that issue opportunities are rarely used and most scheduler +cycles have no eligible (ready) warp; calling the kernel `compute-bound` at this point would hide that +key clue. -The grid can occupy the whole device, but each SM can host at most eight resident warps, and the -compute side is closest to its peak. The next report expands the compute pipelines and collects -scheduler metrics. +The grid is large enough to cover the whole device, but each SM can host at most eight resident +warps, and the compute side is closest to its peak. The next report expands the compute pipelines and +collects scheduler metrics. #### 3. Collect Follow-Up Metrics Along the Compute Path @@ -835,24 +839,25 @@ ncu \ This is an independent NCU run. Percentages can vary slightly between the reports—for example, 77.74% becomes 78.39%—without indicating a performance change. The earlier `basic` report selected the investigation path; the follow-up report uses compute, scheduler, and warp-state metrics from -one collection to narrow the cause. +one collection to narrow the investigation. #### 4. Read the Three Sections in Order -Start at `ComputeWorkloadAnalysis` → `Pipe Utilization (Elapsed Cycles)` → -`Pipe Utilization (% of elapsed cycles)`, the active-cycle view. `Tensor (FP)` is the floating-point -tensor-compute path. `TMEM (Tensor Memory)` is an on-chip memory path that serves tensor operations; -it is distinct from external HBM/DRAM and from TMA, the asynchronous data-movement engine. These -paths are active during about 78% of clock cycles, while the same section's summary reports +Start with `ComputeWorkloadAnalysis` → `Pipe Utilization (Elapsed Cycles)` → +`Pipe Utilization (% of elapsed cycles)`. This is the active-cycle view. `Tensor (FP)` is the +floating-point tensor-compute path. `TMEM (Tensor Memory)` is an on-chip memory path that serves +tensor operations; it is distinct from external HBM/DRAM and from TMA, the asynchronous data-movement +engine. These paths are active for about 78% of clock cycles, while the same section's summary reports `Issue Slots Busy = 3.20%`. A multi-cycle operation can issue once and keep a pipeline active, so -frequent pipeline activity and infrequent new instruction issue can occur together. +high active-cycle utilization can coexist with a low instruction issue rate. -Next, open `Scheduler Statistics` → `Warps Per Scheduler` to see why instructions issue so -infrequently. Each scheduler averages 1.44 active warps that have not finished, but only 0.04 -eligible (ready) warps; 0.04 is an average warp count, not 4%. The denominator for `No Eligible` -includes only cycles in which the scheduler's SM subpartition has at least one warp in flight. Its -value of 96.11% means that no eligible warp is available in 96.11% of those cycles. This explains -the low 3.20% issue rate. Work resides on the SM, but most of the time no warp can continue. +Next, open `Scheduler Statistics` → `Warps Per Scheduler` to see why the scheduler issues new +instructions so infrequently. Each scheduler has an average of 1.44 active warps that have not +finished, but only 0.04 eligible (ready) warps; 0.04 is an average warp count, not 4%. The denominator +for `No Eligible` includes only cycles in which the scheduler's SM subpartition has at least one warp +in flight. A value of 96.11% means that no eligible warp is available in 96.11% of those cycles. This +explains the low 3.20% issue rate. The SM has resident work, but most of the time no warp is ready to +continue. Finally, open `Warp State Statistics` → `Warp State (All Cycles)` to see what those warps are waiting for. One warp spending one clock cycle in a state contributes one warp-cycle. Each issued warp @@ -862,11 +867,12 @@ instructions, not a share of kernel execution time. A scoreboard is the hardware dependency table that records whether results of earlier operations are ready. `Long Scoreboard` means that the next instruction is still waiting for a memory operation handled by L1TEX. L1TEX is the SM-side path for global, local, surface, and texture memory requests; the requested data may -ultimately come from L1, L2, or DRAM, so this field alone cannot identify the serving memory level. -`MemoryWorkloadAnalysis` characterizes aggregate L1, L2, and DRAM behavior for the whole kernel; use -the SASS/source view in `SourceCounters` to continue locating the specific load. +ultimately come from L1, L2, or DRAM, so this field alone cannot identify which memory level served +the request. `MemoryWorkloadAnalysis` characterizes aggregate L1, L2, and DRAM behavior for the whole +kernel; use the SASS/source view in `SourceCounters` to narrow the investigation to specific load +instructions. -The three sections now form one reading path: +Together, the three sections tell this story: ```text Tensor/TMEM paths are often active @@ -875,9 +881,9 @@ Tensor/TMEM paths are often active → Long Scoreboard is the largest wait category ``` -Start with the L1TEX-related data dependency. Each SM can host at most eight resident warps, which may -make data-access latency harder to cover with other work. In the earlier `basic` report, -`DRAM Throughput = 12.88%` shows that aggregate HBM bandwidth is far from saturation; individual +Start with the L1TEX-related data dependency. Each SM can host at most eight resident warps, leaving +fewer independent warps available to run while another warp waits for data. In the earlier `basic` +report, `DRAM Throughput = 12.88%` shows that aggregate HBM bandwidth is far from saturation; individual requests can still reach DRAM and incur long latency. Investigate the data dependency first, then see whether Tensor Core throughput becomes the next limit. The `nvjet` GEMM in this example is a library implementation, so the next report can collect `MemoryWorkloadAnalysis` and compare aggregate L1, @@ -885,7 +891,7 @@ L2, and DRAM traffic, throughput, and cache behavior for the whole kernel. For a kernel you control, the report suggests two directions that can be tested separately: -- **Increase resident warps:** Adjust the tile, block, or resource use so that more work can reside +- **Increase resident warps:** Adjust the tile, block, or resource usage so that more work can reside on each SM. - **Shorten the data-dependency wait:** Hold residency constant and move the load or prefetch earlier, or shorten the dependency chain. @@ -896,12 +902,12 @@ validation steps. ### Metric Calculations, Units, and Boundaries The main walkthrough already established the reading order and the conclusion for this kernel. The -following sections retain only complete field lists, calculations, units, and boundaries that are -easy to misread; consult them as needed for another kernel. +following sections serve as a reference for complete field lists, calculations, units, and boundaries +that are easy to misread; consult them as needed for another kernel. #### `LaunchStats` and `Occupancy` -The complete launch fields used by this example appear in `LaunchStats`: +The `LaunchStats` fields used in this example are: | Field | Value in this report | |---|---:| @@ -918,7 +924,7 @@ The complete launch fields used by this example appear in `LaunchStats`: - For cluster launches, including this one, use the `Waves Per SM` value reported by NCU; use `cudaOccupancyMaxActiveClusters` to calculate resident clusters programmatically. -Per-block resource use appears in `LaunchStats`: +Per-block resource usage appears in `LaunchStats`: | Field | Value in this report | |---|---:| @@ -963,9 +969,10 @@ The `SpeedOfLight` section reports these four fields: | `Memory Throughput` | 38.71% | | `DRAM Throughput` | 12.88% | -Treat `Duration` as specific to its profiler run. This NCU capture reports 95.30 μs, while a separate -Nsight Systems capture reports 92.608 μs. Compare implementations with the unprofiled CUDA Event -baseline. NCU also controls clocks, flushes caches, and may replay or serialize kernels; see +Interpret `Duration` only within the profiling run that produced it. This NCU capture reports +95.30 μs, while a separate Nsight Systems capture reports 92.608 μs. Compare implementations with the +unprofiled CUDA event baseline. NCU also controls clocks, flushes caches, and may replay or serialize +kernels; see [Nsight Compute's workload-duration guidance](https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#workload-durations). The three throughput percentages use separate hardware peaks as their denominators. They cannot be @@ -994,7 +1001,7 @@ The report location is `SpeedOfLight` → `GPU Throughput Breakdown` → distinct hardware paths. At 77.74%, `Mem Tensor` is the most heavily utilized compute-side path in this report. - `Pipe Tc` and `Pipe Tensor` are two distinct pipeline counters in NCU. Both are near 77%, - consistent with a BF16 GEMM performing substantial tensor-MMA-related work. They may cover + consistent with a BF16 GEMM performing substantial tensor-core MMA work. They may cover overlapping hardware activity, so read them separately; summing them would double-count activity. - `Pipe Alu` primarily corresponds to general integer and logic operations, while `Pipe Fma` covers ordinary FP32 arithmetic and some integer multiply-add operations. At 1.36% and 0.61%, respectively, @@ -1005,8 +1012,8 @@ The report location is `SpeedOfLight` → `GPU Throughput Breakdown` → ##### The Two `Pipe Utilization` Denominators -The `ComputeWorkloadAnalysis` summary reports `Issue Slots Busy = 3.20%`. The two full view names are -`Pipe Utilization (% of elapsed cycles)` and +The `ComputeWorkloadAnalysis` summary reports `Issue Slots Busy = 3.20%`. The report labels the two +views `Pipe Utilization (% of elapsed cycles)` and `Pipe Utilization (% of peak instructions executed over elapsed cycles)`. Putting each pipeline on one row makes the contrast easier to see: @@ -1017,8 +1024,8 @@ Putting each pipeline on one row makes the contrast easier to see: | `TC` | 78.12% | 0.38% | | `Tensor (FP)` | 78.07% | 0.61% | -The views use different denominators to compare occupied pipeline cycles with instruction execution -rate; do not add or subtract their values. +The two views use different denominators: one reports active pipeline cycles, while the other reports +instruction execution rate relative to peak. Their values cannot be added or subtracted. ##### Complete `SchedulerStats` Fields @@ -1049,8 +1056,8 @@ second: | `Warp Cycles Per Issued Instruction` | 37.00 warp-cycles / issued instruction | | `Stall Long Scoreboard` | 32.11 warp-cycles / issued instruction | -The report averages 37.00 warp-cycles per issued warp instruction. Of those, 32.11, or about 86.8%, -are assigned to `Long Scoreboard`. These values are normalized across all warps and use +The report shows an average of 37.00 warp-cycles per issued warp instruction. Of those, 32.11, or +about 86.8%, are assigned to `Long Scoreboard`. These values are normalized across all warps and use warp-cycles / issued instruction, a different unit from the scheduler's average issue rate. The `Est. Speedup` shown next to an NCU rule is a model-based estimate of the potential reduction in @@ -1065,17 +1072,17 @@ kernel can establish the actual speedup. | `Barrier` | Waiting for other warps to reach a synchronization point | Compare the work assigned to different warps and when they arrive | | `Not Selected` | The warp is ready, but another warp was selected in this cycle | Check whether many ready warps are competing for issue opportunities | -#### Locate SASS or Source with `SourceCounters` +#### Correlate Metrics with SASS and Source Code Using `SourceCounters` -`WarpStateStats` shows what the kernel spends time waiting for as a whole. `SourceCounters` takes the -next step by placing sampled waits and execution counts beside individual SASS instructions. The -Source page then reveals where those waits are concentrated. If the binary contains line information -and NCU can find the source file, the instructions are also correlated with CUDA source lines. The -data comes from periodic sampling of warp stall reasons together with instruction counts and selected -memory-access metrics. +`WarpStateStats` summarizes wait states across the entire kernel. `SourceCounters` narrows them down +by placing stall samples and execution counts beside individual SASS instructions. The Source page +then reveals where those waits are concentrated. If the binary contains line information and NCU can +find the source file, the instructions are also correlated with CUDA source lines. The data comes from +periodic sampling of warp stall reasons together with instruction counts and selected memory-access +metrics. -The `nvjet` GEMM is a library implementation, so this tutorial has no CUDA source file to import for -it. The SASS view is still available. Reuse the earlier filter and collection conditions: +Because the `nvjet` GEMM is a library implementation, this tutorial does not have access to its CUDA +source. The SASS view is still available. Reuse the earlier filter and collection conditions: ```bash ncu \ @@ -1107,13 +1114,13 @@ ncu --import reports/bf16-gemm-source.ncu-rep \ Start with `Warp Stall Sampling (Not-issued Samples)` and `Instructions Executed`. The first counts sampling observations taken when the warp scheduler issued no instruction; the second counts executions of the corresponding SASS instruction per warp. If `WarpStateStats` was dominated by -`Long Scoreboard` and the Source page concentrates corresponding samples near one load, that -instruction becomes a candidate for closer inspection. These values come from periodic sampling and +`Long Scoreboard` and the corresponding samples cluster around a particular load, that instruction +becomes a candidate for closer inspection. These values come from periodic sampling and identify hotspot locations. Determining whether the data came from L1, L2, or DRAM still requires -the kernel-wide evidence in `MemoryWorkloadAnalysis` together with the code's access relationships. +the kernel-wide evidence in `MemoryWorkloadAnalysis` together with the code's memory-access pattern. -For a TIRx kernel that you compile yourself, the SASS can also be correlated with generated CUDA. -First select NVCC, retain the generated source, and enable line information: +For a TIRx kernel that you compile yourself, the SASS can also be correlated with the generated CUDA +source. First select NVCC, retain the generated source, and enable line information: ```bash export TVM_CUDA_COMPILE_MODE=nvcc @@ -1121,9 +1128,9 @@ export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels" mkdir -p "$TVM_KERNEL_DUMP" ``` -After setting the variables, restart the workload so that the target kernel is recompiled in that -process. The following is a collection-command template; replace `YOUR_KERNEL_NAME` and the program -path on the final line: +After setting the variables, rerun the workload in a fresh process so that the target kernel is +recompiled with these settings. The following template collects the report; replace +`YOUR_KERNEL_NAME` and the program path on the final line: ```bash ncu \ @@ -1152,13 +1159,13 @@ recompilation. #### Test the Hypothesis with a Code Change -Because this example calls the library-provided `torch.mm`, the kernel itself is not editable here. -The following steps give concrete modifications and validation checks for a custom TIRx kernel or -another DSL kernel. +Because this example calls the library-provided `torch.mm`, we cannot modify the kernel implementation +here. The following steps give concrete modifications and validation checks for a custom TIRx kernel +or another DSL kernel. To test the hypothesis that the kernel has too few resident warps to hide L1TEX latency, adjust the tile, block, or pipeline stages to reduce `Registers Per Thread` and -`Dynamic Shared Memory Per Block`. Then read both block limits again. A second block can reside only +`Dynamic Shared Memory Per Block`. Then recheck both block limits. A second block can reside only if both `Block Limit Registers` and `Block Limit Shared Mem` rise from one to at least two; improving only one is insufficient. Reducing register use can cause spills into local memory, and reducing shared-memory use may reduce data reuse. Also confirm that every other block limit is at least two, @@ -1180,7 +1187,7 @@ factor at a time, then check three things: residency experiment reached its resource target; the next check determines whether that change also improves latency. 3. **Actual latency:** Disable Proton, Nsight Systems, and NCU. Measure both implementations with - exactly the same shape, dtype, input policy, warm-up, CUDA Event boundary, and sample count used + exactly the same shape, dtype, input policy, warm-up, CUDA event boundary, and sample count used at the beginning. ```bash @@ -1191,7 +1198,7 @@ python appendix/nsys_example.py \ ``` Compare the unprofiled medians before and after the change, and inspect sample variation as well. If -the correctness check passes, the metrics change as predicted, and CUDA Event latency decreases +the correctness check passes, the metrics change as predicted, and CUDA event time decreases consistently, the hypothesis is supported. If the NCU metrics move as expected but latency does not improve, check whether another bottleneck has emerged or the original hypothesis was incomplete. @@ -1199,13 +1206,14 @@ improve, check whether another bottleneck has emerged or the original hypothesis IKET (In-Kernel Event Tracing) adds an internal timeline to a warp-specialized TIRx kernel. Nsight Systems shows the beginning and end of the whole kernel, NCU aggregates hardware metrics across the -launch, and IKET records when each warp role executes producer, wait, consumer, or other marked -regions. +launch, and IKET records when each warp role is active in marked regions such as producer, wait, and +consumer. ### Run a Complete Example -TVM 0.26 uses a version-pinned `cutlass-4.6.0` profiling profile. For the CUDA 13 environment used in -this chapter, install the matching dependencies and confirm that `run-iket` is available: +TVM 0.26 uses the `cutlass-4.6.0` IKET profile, which pins the profiling dependencies to specific +versions. For the CUDA 13 environment used in this chapter, install the matching dependencies and +confirm that `run-iket` is available: ```bash python -m pip install \ @@ -1307,23 +1315,23 @@ python appendix/iket_example.py kernel is compiled and loaded while IKET recording is active. The script also verifies that the output equals `input * 2 + 1`. -With `postprocess="all"`, the `reports/iket-warp-roles` directory receives JSON, `*.pftrace`, and HTML -artifacts. Load the `*.pftrace` file in Perfetto to inspect `producer_load`, `wait_for_data`, and -`consumer_compute`. Warp 1 reaches the barrier before warp 0 and therefore usually has a longer -`wait_for_data` region. For an H100, change `sm_100a` in the script to `sm_90a`. +With `postprocess="all"`, `iket.run` writes JSON, `*.pftrace`, and HTML artifacts to +`reports/iket-warp-roles`. Load the `*.pftrace` file in Perfetto to inspect `producer_load`, +`wait_for_data`, and `consumer_compute`. Warp 1 usually reaches the barrier before warp 0, giving it +a longer `wait_for_data` region. For an H100, change `sm_100a` in the script to `sm_90a`. -### Move the Annotations into Your Kernel +### Add IKET Annotations to Your Own Kernel Create an `IketProfiler` inside the `PrimFunc`. Use `mark()` for an instantaneous event, and use -matched `range_push()` / `range_pop()` or `range_start()` / `range_end()` calls around a phase. Every -warp's actual control-flow path must keep ranges balanced. Mark waiting explicitly, as the example -does with `wait_for_data`. +matched `range_push()` / `range_pop()` or `range_start()` / `range_end()` calls around a phase. Keep +IKET ranges balanced along every control-flow path that a warp may take. Mark waiting explicitly, as +the example does with `wait_for_data`. Keep compilation and the first JIT load inside the function passed to `iket.run`. IKET supports Hopper and newer architectures and validates the CUTLASS DSL packages, NVRTC, `nvdisasm`, and related binaries against the pinned profile. The recording code inserted by IKET changes the generated -kernel and adds overhead, so use the IKET trace to study phases and warp roles. Measure formal -latency with the uninstrumented CUDA Event benchmark. See +kernel and adds overhead, so use the IKET trace to study phases and warp roles. For reported latency, +use the uninstrumented CUDA event benchmark. See [`python/tvm/backend/cuda/iket.py`](https://github.com/apache/tvm/blob/v0.26.0/python/tvm/backend/cuda/iket.py) and the [NVIDIA IKET guide](https://github.com/NVIDIA/cutlass/blob/v4.6.0/media/docs/pythonDSL/cute_dsl_general/iket_profiling.rst)