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..572cb1bf
--- /dev/null
+++ b/appendix/benchmarking_gpu_kernels.md
@@ -0,0 +1,1338 @@
+(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.
+
+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.
+
+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. |
+| 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 do marked phases run, and where do warp roles wait or overlap? |
+
+## Verify Correctness Before Timing
+
+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.
+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.
+
+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 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.
+
+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
+
+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.
+
+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 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 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
+
+### Measure GPU Stream Time with CUDA Events
+
+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
+current stream.
+
+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
+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 for one batch of back-to-back calls, 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 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 timing
+result can be read.
+
+`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 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
+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)
+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 left in L2 by the previous invocation, then records an independent CUDA event interval:
+
+```python
+from tvm.tirx.bench import bench
+
+
+# Reuse gemm from above. For a custom TIRx kernel, substitute its no-argument callable.
+run = gemm
+result = bench(
+ {"gemm": run},
+ timer="event",
+ warmup=25,
+ repeat=100,
+ rounds=5,
+ cooldown_s=1.0,
+)
+
+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
+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 `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 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
+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. 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* so that
+their different boundaries remain visible.
+
+### Advanced: Timing a Multi-Stream Operation
+
+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
+
+
+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
+```
+
+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) 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. 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. 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 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. 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 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
+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;
+- **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.
+
+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
+
+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:
+
+```text
+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. 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 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
+
+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
+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 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.
+
+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
+
+Before asking where time is spent, measure the complete operation under normal execution:
+
+```bash
+python appendix/nsys_example.py \
+ --size 4096 \
+ --warmup-calls 500 \
+ --event-samples 20
+```
+
+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
+```
+
+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
+
+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
+python -m pip install pandas llnl-hatchet
+```
+
+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
+python appendix/nsys_example.py \
+ --size 4096 \
+ --warmup-calls 500 \
+ --proton-calls 100
+```
+
+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 87.00 8.700
+└── ReLU kernel 100 11.71 1.171
+```
+
+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, 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
+covers the full operation.
+
+## Analyze an Application Timeline with Nsight Systems
+
+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 a Timeline for the Target Operation
+
+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()`:
+
+```python
+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()
+```
+
+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`:
+
+```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 \
+ --size 4096 \
+ --warmup-calls 500 \
+ --profile-once
+```
+
+`--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.
+
+Once the report has been generated, open its timeline in the Nsight Systems GUI:
+
+```bash
+nsys-ui reports/target-timeline.nsys-rep
+```
+
+### Locate the Most Expensive Kernel in the Timeline
+
+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.
+
+
+
+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 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
+```
+
+`--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.
+
+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 |
+|---|---:|---:|---:|
+| 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 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 |
+|---|---:|---:|---:|
+| BF16 GEMM | 50.717 μs | — | 92.608 μs |
+| ReLU | 13.474 μs | 5.074 μs | 10.944 μs |
+
+Read the results in this order:
+
+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.
+
+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.
+
+## Use Nsight Compute to Analyze a Single Kernel
+
+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.
+
+### How to Read an NCU Report
+
+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 from the time it is assigned to an SM until it finishes.
+
+Analyze a new NCU report in this order:
+
+| 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 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` 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, 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 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
+ 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 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.
+
+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
+
+#### 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 \
+ --profile-from-start off \
+ --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 \
+ --size 4096 \
+ --warmup-calls 500 \
+ --profile-once
+```
+
+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 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
+ 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:
+
+```bash
+ncu-ui reports/bf16-gemm-basic.ncu-rep
+```
+
+Without a GUI, print the Details page in the terminal:
+
+```bash
+ncu --import reports/bf16-gemm-basic.ncu-rep \
+ --page details \
+ --print-details all \
+ --print-metric-name label-name
+```
+
+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:
+
+| Observation in `basic` | Next step in this example |
+|---|---|
+| `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 `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\%$. 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 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 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
+
+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.
+
+```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
+```
+
+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 investigation.
+
+#### 4. Read the Three Sections in Order
+
+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
+high active-cycle utilization can coexist with a low instruction issue rate.
+
+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
+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 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.
+
+Together, the three sections tell this story:
+
+```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, 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,
+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 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.
+
+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 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 `LaunchStats` fields used in this example are:
+
+| 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 usage appears in `LaunchStats`:
+
+| Field | Value in this report |
+|---|---:|
+| `Registers Per Thread` | 255 |
+| `Dynamic Shared Memory Per Block` | 213.28 KB |
+
+The `Occupancy` section shows the residency permitted by those resources:
+
+| 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% |
+
+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
+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.
+
+#### Compute Pipelines, the Scheduler, and Warp States
+
+##### `Compute Throughput Breakdown` Fields
+
+The report location is `SpeedOfLight` → `GPU Throughput Breakdown` →
+`Compute Throughput Breakdown`:
+
+| 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-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,
+ 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 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:
+
+| 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 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
+
+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 |
+
+The summary in the same section also reports `No Eligible = 96.11%`.
+
+`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.
+
+##### `WarpStateStats` Normalization
+
+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:
+
+| 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 |
+
+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
+workload time for that rule. Use it to prioritize the investigation; only a benchmark of the modified
+kernel can establish the actual speedup.
+
+##### Other Common Warp States
+
+| 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 |
+
+#### Correlate Metrics with SASS and Source Code Using `SourceCounters`
+
+`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.
+
+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 \
+ --config-file off \
+ --profile-from-start off \
+ --kernel-name 'regex:.*nvjet_sm100.*' \
+ --launch-count 1 \
+ --section SourceCounters \
+ --replay-mode kernel \
+ --cache-control all \
+ --clock-control boost \
+ --pipeline-boost-state stable \
+ --export reports/bf16-gemm-source \
+ --force-overwrite \
+ python appendix/nsys_example.py \
+ --size 4096 \
+ --warmup-calls 500 \
+ --profile-once
+```
+
+Select SASS on the GUI's Source page, or print the same view in the terminal:
+
+```bash
+ncu --import reports/bf16-gemm-source.ncu-rep \
+ --page source \
+ --print-source sass
+```
+
+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 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 memory-access pattern.
+
+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
+export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels"
+mkdir -p "$TVM_KERNEL_DUMP"
+```
+
+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 \
+ --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
+```
+
+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`, 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 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,
+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.
+
+```bash
+python appendix/nsys_example.py \
+ --size 4096 \
+ --warmup-calls 500 \
+ --event-samples 20
+```
+
+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 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.
+
+## Use IKET to Inspect Phases Inside a DSL Kernel
+
+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 is active in marked regions such as producer, wait, and
+consumer.
+
+### Run a Complete Example
+
+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 \
+ 'nvidia-cutlass-dsl[cu13]==4.6.0' \
+ 'nvidia-cuda-nvdisasm==13.3.73' \
+ 'nvidia-cuda-nvrtc==13.2.78'
+run-iket --help
+```
+
+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:
+
+```python
+"""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()
+```
+
+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"`, `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`.
+
+### 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. 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. 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)
+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/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..9a82e97f
--- /dev/null
+++ b/appendix/nsys_example.py
@@ -0,0 +1,119 @@
+"""Reusable CUDA workload for the benchmarking and profiling appendix."""
+
+import argparse
+from statistics import median
+
+import torch
+
+
+def make_workload(size: int):
+ 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("BF16 GEMM"):
+ torch.mm(a, b, out=c)
+ with torch.cuda.nvtx.range("ReLU"):
+ torch.clamp_min(c, 0, out=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):
+ 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 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=500)
+ parser.add_argument("--proton-output", default="operator")
+ args = parser.parse_args()
+
+ 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:
+ 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"
+ )
+ 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 __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..8eec35da
--- /dev/null
+++ b/img/nsys_b200_timeline.svg
@@ -0,0 +1,55 @@
+
+
\ 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
new file mode 100644
index 00000000..3d9eef8e
--- /dev/null
+++ b/img/nsys_b200_timeline_zh_en_tracks.svg
@@ -0,0 +1,55 @@
+
+
\ 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..04d0a192
--- /dev/null
+++ b/img/scripts/gen_nsys_b200_timeline.py
@@ -0,0 +1,111 @@
+"""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 = 400.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：GEMM → ReLU"
+ if chinese
+ else "4096×4096 BF16: GEMM → ReLU"
+ )
+ rows = ["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'')
+ 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_en_tracks.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..50eae3e8
--- /dev/null
+++ b/zh/appendix/benchmarking_gpu_kernels.md
@@ -0,0 +1,937 @@
+(chap_benchmarking)=
+# GPU Kernel 性能测量与分析
+
+优化 GPU kernel 时，需要分别回答两个问题：运行一次要多久，时间主要花在哪里。Benchmark 负责测前者，profile 用来分析后者。
+
+一次 Python 调用不一定只对应一个 GPU kernel。它可能启动多个 kernels、提交内存拷贝，或者等待 GPU 完成工作。计时前要先确定被测 operation 包含哪些步骤，并让所有实现采用相同的边界。
+
+{ref}`chap_performance` 介绍了如何用 roofline 判断性能受计算吞吐还是内存带宽限制。接下来讨论实验方法：如何确定计时范围、选择 warm-up 和 repeat，以及解读 profiler 报告。
+
+## 区分性能测量与性能诊断
+
+这套流程中的工具各自回答不同问题：
+
+| 工具 | 主要回答的问题 |
+|---|---|
+| 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`） | 选定的 kernel 内部在做什么，下一步应调查哪类硬件资源或等待？ |
+| IKET（可选） | 加入 kernel 内标记后，数据搬运、计算和写回等阶段何时 active、等待或重叠？ |
+
+## 计时前先验证正确性
+
+性能计时前，先单独验证正确性：
+
+1. 构造有代表性的输入，并覆盖相关的边界情况。
+2. 运行被测实现并同步，确保 GPU 已经完成计算。
+3. 使用明确的 tolerance，将结果与 reference 比较。
+4. 如果 kernel 会在已有 output 上累加或原地修改输入，每次验证前都恢复相同的初始状态。
+
+以自写 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 从开始到结束的执行区间。
+
+## 测量 GPU 时间与单次调用延迟
+
+### 使用 CUDA Events 测量 GPU stream 时间
+
+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，并报告五轮结果的中位数：
+
+```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` 表示五轮独立测量。这组值来自 B200 上的稳定性测试：50 次 warm-up 后结果仍持续下降，增加到 500 次后才趋于稳定；`repeat=100` 也比 `repeat=10` 稳定。测量其他 workload 时，先增加 `warmup_calls`，直到前几轮不再持续变化；若结果仍有较大波动，再增加 `repeat` 或 `rounds`。如果更长的运行反而使整体时间系统性变化，就要检查温度、功耗和时钟频率。
+
+这段代码始终复用同一组矩阵，因此结果偏向有数据复用的 warm-cache 场景；是否真的命中 cache，还取决于本次计算反复访问的数据总量与硬件 cache 容量。
+
+本书使用 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
+
+
+# 这里复用上文的 gemm；分析自己的 TIRx kernel 时，换成对应的无参数 callable。
+run = gemm
+result = bench(
+ {"gemm": run},
+ timer="event",
+ warmup=25,
+ repeat=100,
+ rounds=5,
+ cooldown_s=1.0,
+)
+
+print(result["impls"]["gemm"]) # 五轮平均值，单位为 us
+print(result["round_samples"]["gemm"]) # 每轮结果
+```
+
+`warmup=25` 和 `repeat=100` 是毫秒预算，Event timer 会根据短测结果换算调用次数。正式 Event 只覆盖被测调用；用于减少 L2 复用的 256 MiB 写入发生在 start event 之前。`rounds=5` 测量五轮，`cooldown_s=1.0` 在每轮前暂停一秒；`impls` 保存五轮平均值，`round_samples` 保存逐轮结果。预算和轮数仍按上面的稳定性标准调整，并对所有实现使用相同设置。
+
+TIRx-kernels 的 `run_bench` 也使用这个 helper。未使用 distributed 模式时，省略 `timer` 会默认使用 Proton；需要 CUDA Event 区间时，应显式指定 `timer="event"`。对于会原地修改状态的 kernel，重复调用时仍要遵守前面的重置规则；若重置写在被测函数内，其时间也属于 operation。
+
+### 测量一次调用的端到端时间
+
+如果关心的是从 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 时间”和“单次端到端时间”，让读者直接看出两个数字覆盖的范围。
+
+### 进阶：测量多 stream operation
+
+一个 operation 把工作提交到多条 CUDA streams 时，计时 stream 需要连接每条分支的起点和终点。下面的 `sin` 和 `cos` 分别在两条 streams 上运行，等两条分支都完成后，再回到计时 stream 相加：
+
+```python
+import torch
+
+
+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` 是两条分支的共同起点，`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 自动调优，以及各实现可使用的搜索预算。
+
+此外还应记录 GPU、driver、CUDA、framework 与 compiler 版本，以及 dtype、shape、时钟和功耗设置。使用库实现作为 baseline 时，要记录库版本、所选算法和 workspace；自动调优可以放在计时区间之外，但搜索预算和最终配置仍属于实验条件。
+
+## 由延迟换算吞吐率
+
+吞吐率等于约定的工作量除以延迟；计时器提供公式中的延迟。对于一个 $M\times K$ 与 $K\times N$ 的 GEMM，若延迟为 `t_us` 微秒，则：
+
+```text
+TFLOP/s = 2 × M × N × K / t_us / 10^6
+```
+
+分子和计时边界必须对应同一份工作。例如，后文完整 GEMM + ReLU operation 的延迟是 105.152 μs；用 $2\times4096^3$ 除以这个时间会得到约 1307 TFLOP/s，但它只能称为“按 GEMM 工作量计算的有效吞吐率（effective throughput）”，因为分母还包含 ReLU。要报告 GEMM kernel 自身的 TFLOP/s，计时区间也要只覆盖 GEMM。
+
+性能表在给出 TFLOP/s、GB/s 或 tokens/s 时，也应保留原始延迟，并说明工作量如何计算。对于 attention 和 fused kernels，还需注明分子统计的是完整稠密问题、实际选中的元素，还是 kernel 真正执行的工作。相关公式和 roofline 分析见 {ref}`chap_performance`。
+
+## 用 Proton 找出最耗时的 kernel
+
+从这里开始，baseline、Proton、Nsight Systems 和 Nsight Compute 都运行 `appendix/nsys_example.py` 中的同一个 operation：两个 $4096\times4096$ BF16 矩阵先做 GEMM，再对结果做 ReLU。输入、中间结果和输出都在计时或采集前分配。脚本中的 operation 是：
+
+```python
+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 和采集范围之外。
+
+前面的 2048×2048 FP16 代码用于讲解计时 API；下面换成这份 BF16 operation 后，不再混用两组 workload 的结果。
+
+### 先记录无 profiler baseline
+
+在分析“时间花在哪里”之前，先测出正常运行时的完整 operation 时间：
+
+```bash
+python appendix/nsys_example.py \
+ --size 4096 \
+ --warmup-calls 500 \
+ --event-samples 20
+```
+
+500 次 warm-up 发生在正式计时前；之后，每个 sample 用一对 CUDA Events 包住一次 GEMM + ReLU。B200 上的一次实际输出为：
+
+```text
+median=105.152 us, min=103.136 us, max=131.200 us
+```
+
+这里的 median 是后面判断代码修改是否真的变快时要回到的 baseline，min 和 max 用于观察样本波动。这份 baseline 覆盖完整的 GEMM + ReLU operation；后面的 profiler 表格来自独立采集，并分别列出单个 kernel。
+
+### 用 Proton 比较 operation 中的 kernels
+
+Proton 可以查看每个 kernel 的调用次数、平均时间和累计时间。运行前先确认环境中已经安装与 TVM 兼容的 Triton；Proton 和 `proton-viewer` 随 Triton 提供。Viewer 还需要下面两个 Python 依赖：
+
+```bash
+python -m pip install pandas llnl-hatchet
+```
+
+脚本的 `--proton-calls` 模式复用同一个 `run()`，先 warm-up，再采集 100 次 operation，并生成 `operator.hatchet`：
+
+```bash
+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
+```
+
+`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 87.00 8.700
+└── ReLU kernel 100 11.71 1.171
+```
+
+先检查预期的两个 kernels 是否都出现、调用次数是否为 100，再比较累计时间。GEMM 占用的时间明显更多，因此选它作为后续深入分析的目标。交给 NCU 之前，先用 Nsight Systems 确认单次 operation 中的执行顺序和空隙，并关联对应的 host launch APIs。
+
+这次手动 Proton session 保留正常的 cache 状态，与 `bench(timer="proton")` 在每次正式调用前写入 256 MiB buffer 的策略不同。这里用同一份 Proton 报告内的数值给 kernels 排序。实现之间的快慢仍看上面的 CUDA Event baseline；完整 operation latency 则取自包住整个 operation 的计时区间。
+
+## 使用 Nsight Systems 分析应用时间线
+
+Proton 给出了汇总排名，但没有显示 kernels 的先后关系、空隙、拷贝或 host 等待。Nsight Systems 用时间线回答这些问题。
+
+### 采集目标 operation 的时间线
+
+脚本的 `--profile-once` 模式先在 profiler 尚未启动时完成 warm-up，并等待这些工作结束；然后只在 `cudaProfilerStart()` 和 `cudaProfilerStop()` 之间提交一次 GEMM + ReLU：
+
+```python
+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()
+```
+
+这里的 NVTX range 用于在时间线中定位目标 operation。Range 内的同步确保两个 kernels 在停止采集前完成；`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 \
+ --size 4096 \
+ --warmup-calls 500 \
+ --profile-once
+```
+
+`--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 打开时间线：
+
+```bash
+nsys-ui reports/target-timeline.nsys-rep
+```
+
+### 从时间线定位最耗时的 kernel
+
+下面的报告来自 NVIDIA B200，软件版本为 NVIDIA driver 595.58.03、CUDA 13.0、PyTorch 2.12.0+cu130 和 Nsight Systems 2025.6.3。
+
+
+
+`GPU stream 7` 中的 `7` 是这次报告里的 stream 标识。两个 kernels 位于同一条 stream 上，所以按提交顺序执行。也可以从命令行提取下文使用的时间：
+
+```bash
+nsys stats \
+ --force-export=true \
+ --format=column \
+ --timeunit=us \
+ --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
+```
+
+`--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` 可以查看当前版本的完整定义。
+
+这条命令会依次打印多张表，按下面的顺序取数：
+
+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 时长总和的比例 |
+|---|---:|---:|---:|
+| BF16 GEMM | 1 | 92.608 μs | 89.4% |
+| ReLU | 1 | 10.944 μs | 10.6% |
+
+再把 GPU execution 与 host 上的 launch API 对应起来。Positive queue time 指 launch API 返回后，到 kernel 稍后才开始之间的时间；kernel 在 API 返回前已经开始时，该字段为空。
+
+| Kernel | API time | Positive queue time | GPU execution |
+|---|---:|---:|---:|
+| BF16 GEMM | 50.717 μs | — | 92.608 μs |
+| ReLU | 13.474 μs | 5.074 μs | 10.944 μs |
+
+按下面的顺序读：
+
+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 执行已经完成。
+
+分析其他报告时，也先确认采集范围，再看 GPU kernels、copies、空隙和重叠，最后关联到 host launch 或同步 API。更多定义见 [Nsight Systems Analysis Guide](https://docs.nvidia.com/nsight-systems/AnalysisGuide/index.html)。
+
+## 使用 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 报告时，按下面的顺序分析：
+
+| 当前问题 | 最先查看的位置 | 这一步的作用 |
+|---|---|---|
+| 报告属于哪个 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 调度顺序不包含在该字段中。
+
+各项 `Block Limit` 分别给出 registers、shared memory、threads 等单项资源允许每个 SM 驻留的 block 上限，其中最小值决定理论 block 上限。Theoretical occupancy 是按这些上限算出的最大驻留 warp 数占硬件容量的比例；achieved occupancy 是采集期间实际平均活跃 warp 数的比例。它们描述并发驻留量，是否影响执行速度还要结合 scheduler 指标。
+
+`SpeedOfLight` 中的 Compute 表示最忙的 SM 计算路径，Memory 表示最忙的内存侧路径，DRAM 只看外部显存接口；三者都以各自的可持续峰值为分母。B200 的外部显存是 HBM，L2 是全 GPU 共享的 cache，L1/TEX 是 SM 一侧处理内存请求的路径。Memory 较高只说明某条内存侧路径繁忙，外部 HBM 是否接近饱和由 DRAM 字段判断。`ComputeWorkloadAnalysis` 中的 active cycles 表示流水线仍在处理工作的周期，`Issue Slots Busy` 表示 scheduler 实际使用了多少指令发射机会。
+
+`SchedulerStats` 把 warps 分成几种状态：active warp 已经驻留且尚未结束；eligible（已就绪）warp 的下一条指令已经解码、依赖已经就绪，而且所需执行单元可用；issued warp 在当前周期实际发出了指令。确认目标 launch，并依次看完 grid、驻留量和 `SpeedOfLight` 后，再按结果选择下一组指标：
+
+- **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 时间。
+
+`MemoryWorkloadAnalysis` 汇总整条 kernel 在 DRAM、L2、L1/TEX 和其他内存路径上的流量与 cache 行为。某一条 load 的具体依赖还要用 `SourceCounters` 把采样到的 stall 和指令活动映射到 SASS（GPU 机器指令）或源码位置。
+
+“高”和“低”要结合当前 GPU、workload 和同一份报告判断。当报告已经指向一处可修改的代码，并且能够写出修改后预期变化的指标和 latency 时，这一轮分析就形成了可检验的假设。范围仍然过宽时，再沿上面的分支采集下一组 section。
+
+### 完整示例：分析 B200 BF16 GEMM
+
+#### 1. 采集第一份 `basic` 报告
+
+继续使用脚本的 `--profile-once` 模式。应用在采集范围内提交一次 GEMM，再提交一次 ReLU；NCU 等待这个范围开始，并只筛选 GEMM：
+
+```bash
+mkdir -p reports
+ncu \
+ --config-file off \
+ --profile-from-start off \
+ --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 \
+ --size 4096 \
+ --warmup-calls 500 \
+ --profile-once
+```
+
+几个关键选项分别控制采集范围、目标、指标和采集条件：
+
+- `--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。
+
+这里的“一次 GEMM”指应用提交一次 launch。为了收集所需的硬件计数器，NCU 仍可能在内部重放这次 GEMM。500 次 warm-up 位于采集范围外，可避开初始化和 lazy loading；NCU 的 cache control 随后会改变正常的 warm-cache 条件。
+
+在 GUI 中打开报告：
+
+```bash
+ncu-ui reports/bf16-gemm-basic.ncu-rep
+```
+
+终端中也可以查看 Details 页：
+
+```bash
+ncu --import reports/bf16-gemm-basic.ncu-rep \
+ --page details \
+ --print-details all \
+ --print-metric-name label-name
+```
+
+其他筛选条件和采集选项见 [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 覆盖全卡。
+
+再看 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 中有多少已经就绪。
+
+最后看吞吐率。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`：
+
+| 字段 | 本次报告中的值 |
+|---|---:|
+| `Registers Per Thread` | 255 |
+| `Dynamic Shared Memory Per Block` | 213.28 KB |
+
+这些资源最终允许的驻留数量位于 `Occupancy`：
+
+| 字段 | 本次报告中的值 |
+|---|---:|
+| `Block Limit Registers` | 1 block / SM |
+| `Block Limit Shared Mem` | 1 block / SM |
+| `Theoretical Occupancy` | 12.50% |
+| `Achieved Occupancy` | 8.97% |
+
+一个 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。
+
+- `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%，说明实际执行期间未始终保持理论驻留上限；它衡量驻留并发度，与“达到峰值性能的百分比”采用不同定义。
+
+#### `SpeedOfLight` 的分母与 `Duration`
+
+`SpeedOfLight` 指标组给出下面四个字段：
+
+| 字段 | 本次报告中的值 |
+|---|---:|
+| `Duration` | 95.30 μs |
+| `Compute (SM) Throughput` | 77.74% |
+| `Memory Throughput` | 38.71% |
+| `DRAM Throughput` | 12.88% |
+
+`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)。
+
+三个 throughput 百分比分别使用各自的硬件峰值作为分母，彼此不能相加，也不能当作执行时间占比。[Nsight Compute Profiling Guide](https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-structure) 给出了 throughput metric 的组成规则。
+
+#### 计算流水线、scheduler 与 warp 状态
+
+##### `Compute Throughput Breakdown` 字段
+
+报告位置是 `SpeedOfLight` → `GPU Throughput Breakdown` → `Compute Throughput Breakdown`：
+
+| 字段 | 本次报告中的值 |
+|---|---:|
+| `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` 是 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 和依赖延迟，需要结合相应指标判断。
+
+##### `Pipe Utilization` 的两个分母
+
+`ComputeWorkloadAnalysis` 摘要中的 `Issue Slots Busy` 为 3.20%。`Pipe Utilization` 的两个完整视图名称分别是 `Pipe Utilization (% of elapsed cycles)` 和 `Pipe Utilization (% of peak instructions executed over elapsed cycles)`。
+
+把同一条流水线放在一行后，两种视图的差异会更直观：
+
+| 流水线字段 | Active-cycle 视图 | Instruction-rate 视图 |
+|---|---:|---:|
+| `TMEM (Tensor Memory)` | 78.39% | 0.04% |
+| `TC` | 78.12% | 0.38% |
+| `Tensor (FP)` | 78.07% | 0.61% |
+
+两个视图采用不同分母，用于对照流水线占用周期和指令执行率；它们之间不做加减。
+
+##### `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
+ncu \
+ --config-file off \
+ --profile-from-start off \
+ --kernel-name 'regex:.*nvjet_sm100.*' \
+ --launch-count 1 \
+ --section SourceCounters \
+ --replay-mode kernel \
+ --cache-control all \
+ --clock-control boost \
+ --pipeline-boost-state stable \
+ --export reports/bf16-gemm-source \
+ --force-overwrite \
+ python appendix/nsys_example.py \
+ --size 4096 \
+ --warmup-calls 500 \
+ --profile-once
+```
+
+在 GUI 的 Source 页选择 SASS，或者从终端打印同一视图：
+
+```bash
+ncu --import reports/bf16-gemm-source.ncu-rep \
+ --page source \
+ --print-source sass
+```
+
+先看 `Warp Stall Sampling (Not-issued Samples)` 和 `Instructions Executed`。前者记录采样时 warp scheduler 没有发出指令的观测次数，后者是对应 SASS 指令按 warp 统计的执行次数。如果前面的 `WarpStateStats` 以 `Long Scoreboard` 为主，而 Source 页又把相应 samples 集中到某条 load 附近，这条指令就是下一步检查的候选。这里使用的是周期性采样，结果表示热点位置；数据最终来自 L1、L2 还是 DRAM，仍要结合 `MemoryWorkloadAnalysis` 的整 kernel 聚合指标和代码中的访问关系判断。
+
+对于自己编译的 TIRx kernel，可以把 SASS 继续关联到生成的 CUDA。先让 TVM 使用 NVCC、保留源码并写入 line information：
+
+```bash
+export TVM_CUDA_COMPILE_MODE=nvcc
+export TVM_KERNEL_DUMP="$PWD/reports/tvm-kernels"
+mkdir -p "$TVM_KERNEL_DUMP"
+```
+
+设置环境变量后，重新启动 workload，让目标 kernel 在这个进程中重新编译。下面是采集命令模板；把 `YOUR_KERNEL_NAME` 和最后一行的程序路径换成自己的值：
+
+```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
+```
+
+`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。
+
+#### 用代码修改检验假设
+
+本例调用库提供的 `torch.mm`，无法直接修改 kernel 实现。下面补充自写 TIRx 或其他 DSL kernel 的具体修改和验证步骤。
+
+例如，要检验“驻留 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 判断总体收益。
+
+另一种实验保持驻留数量不变，只把 load 或预取提前，或者缩短依赖链。若 `Long Scoreboard` 和 latency 一起下降，就支持“warps 等待数据的时间减少了”这个判断。这个值按已发射指令归一化，因此还要结合 latency 判断是否产生实际提速。一次只改变一个关键因素，然后检查三个对象：
+
+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
+python appendix/nsys_example.py \
+ --size 4096 \
+ --warmup-calls 500 \
+ --event-samples 20
+```
+
+比较修改前后的无 profiler median，并同时查看样本波动。正确性通过、指标按预测变化，而且 CUDA Event latency 稳定下降时，这份假设得到了支持。若只有 NCU 指标变化，说明代码已经改变了预期硬件行为，但这次修改尚未带来实际提速；接下来检查性能限制是否转移，或原先的判断是否还缺一环。
+
+## 用 IKET 查看 DSL kernel 内部阶段
+
+编写 warp-specialized TIRx kernel 时，IKET（In-Kernel Event Tracing）可以把 kernel 内部的阶段画成时间线。Nsight Systems 显示整个 kernel 的起止区间，NCU 汇总整个 launch 的硬件指标，IKET 则记录各个 warp role 何时执行 producer、等待、consumer 等代码段。
+
+### 运行一个完整示例
+
+TVM 0.26 使用版本锁定的 `cutlass-4.6.0` profiling profile。本章的 CUDA 13 环境可以安装对应依赖，并先确认 `run-iket` 命令可用：
+
+```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
+```
+
+下面的完整脚本位于 `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：
+
+```python
+"""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()
+```
+
+在 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)。
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
```