Skip to content

Commit 8fa3ebc

Browse files
acollins3acollins3
authored
Add guidance for coding agents (#84)
Add AGENTS.md with repository guidance for coding agents
1 parent f80a480 commit 8fa3ebc

1 file changed

Lines changed: 346 additions & 0 deletions

File tree

‎AGENTS.md‎

Lines changed: 346 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,346 @@
1+
# Agent instructions
2+
3+
Nsight Python (`nsight-python` on PyPI, imported as `nsight`) profiles GPU
4+
kernels from Python. You decorate a benchmark function with
5+
`@nsight.analyze.kernel`, mark kernel launches with `nsight.annotate`, and get
6+
back a `ProfileResults` whose `.to_dataframe()` holds aggregated metrics.
7+
Collection is done by NVIDIA Nsight Compute (`ncu`, default) or, for kernel
8+
time only, by the experimental CUPTI backend.
9+
10+
This file has two parts:
11+
12+
- **Part 1: Using Nsight Python.** How the profiler behaves and how to write
13+
and read benchmarks. Read this whenever you write or run profiling code,
14+
including tests and examples in this repository.
15+
- **Part 2: Working on this repository.** Layout, commands and rules for
16+
changing the library.
17+
18+
## Sources of truth
19+
20+
This file records rules and pitfalls. Facts that change between releases live
21+
elsewhere. Check them there rather than relying on memory or on this file. If
22+
this file disagrees with them, they win, and this file should be fixed.
23+
24+
| Topic | Where to look |
25+
| --- | --- |
26+
| Required `ncu`, `ncu-report` and CUPTI versions | `docs/source/installation/runtime_requirements.rst`; enforced by `MIN_NCU_VERSION` in `nsight/collection/ncu.py` |
27+
| Decorator and annotation parameters | Docstrings of `kernel` in `nsight/analyze.py` and `annotate` in `nsight/annotation.py` |
28+
| Result columns and their meaning | `docs/source/overview/profiling_results.rst`, `ProfileResults.to_dataframe` |
29+
| Output files and CSV columns | `kernel` docstring in `nsight/analyze.py` |
30+
| Backends, CUPTI limitations | `docs/source/tools.rst` |
31+
| Known issues | `docs/source/release_notes/known_issues.rst` |
32+
| Metric names | `ncu --query-metrics --chip <chip>` (for example `ga102`, `gh100`) |
33+
| Dev commands, lint and test setup | `pyproject.toml` (hatch envs), `conftest.py` |
34+
| CI jobs and GPU matrix | `.github/workflows/ci.yaml` |
35+
| Commit and PR rules | `Contributing.md` |
36+
37+
Do not invent parameters, result columns, metric names, or units.
38+
39+
# Part 1: Using Nsight Python
40+
41+
## Before running anything
42+
43+
- Check `nvidia-smi` and `ncu --version` against the runtime requirements.
44+
Many machines have an older `ncu` from a CUDA toolkit first on `PATH`. List
45+
every `ncu` on `PATH` with `which -a ncu` (Linux) or `where.exe ncu`
46+
(Windows).
47+
- An older or missing `ncu` does not fail at `import nsight`. The error is
48+
stored and raised on the first call to a decorated function.
49+
- Non-admin profiling needs access to GPU performance counters. If they are
50+
restricted to administrators, `ncu` fails with `ERR_NVGPUCTRPERM`. On Linux
51+
you can check beforehand: `/proc/driver/nvidia/params` shows
52+
`RmProfilingAdminOnly: 0` when access is open. On Windows the setting is in
53+
the NVIDIA Control Panel (Developer > Manage GPU Performance Counters). Do
54+
not change driver settings or the Control Panel setting, add capabilities,
55+
or run as root or administrator to get around this. Report it instead.
56+
- On Windows, set `$env:MPLBACKEND = "Agg"` (PowerShell) before using
57+
`@nsight.analyze.plot` in a non-interactive run. matplotlib defaults to the
58+
TkAgg GUI backend there, and plotting fails where Tk is unusable, as on
59+
the CI runners. On Linux, matplotlib falls back to `Agg` by itself when
60+
there is no display.
61+
- On Windows, a Python process that loaded NCU can abort at exit with code
62+
`0xC0000409` after all its work has finished. Judge success from the
63+
program's own output (for example pytest's summary line), not only from
64+
the exit code. CI works around this in `.github/workflows/ci.yaml`.
65+
- `UserWarning: Nsight Python Thermovision is not supported on this machine`
66+
is harmless. Thermal throttling protection is simply off on that GPU.
67+
68+
## How the profiler runs your code
69+
70+
These facts decide how to write benchmark functions correctly:
71+
72+
1. **NCU must attach before CUDA is initialized.** `import nsight` activates
73+
NCU by default (`NSPY_NCU_INIT_AT_IMPORT=1`), so import `nsight` before
74+
anything that touches CUDA. Only one tool can be active per process, and
75+
NCU can never be deactivated.
76+
2. **The decorated function is called `len(configs) * runs` times.** Its
77+
return value is discarded (a `RuntimeWarning` is emitted if it returns
78+
something). The call returns `ProfileResults`. Keep anything whose result
79+
you need, such as correctness checks, in an undecorated function.
80+
3. **Only kernels inside `nsight.annotate(...)` are measured.** Setup such as
81+
`torch.randn` outside annotations is not profiled. Annotation names must be
82+
unique per run, and annotations cannot be nested.
83+
4. **One kernel per annotation by default.** More than one raises
84+
`RuntimeError: More than one ... kernel is launched`. Options are
85+
`combine_kernel_metrics=lambda a, b: a + b`, `ignore_kernel_list=[...]`, or
86+
`replay_mode="range"`. Every annotation must launch the **same number of
87+
kernels in every run and config**. The check is only
88+
`total % (configs * runs) == 0`, so varying counts can pass and give wrong
89+
numbers.
90+
5. **NCU replays each kernel** with caches flushed (`cache_control="all"`,
91+
the default; `"none"` keeps caches warm) and measures it in isolation.
92+
`gpu__time_duration.sum` is per-kernel GPU time in ns. It does not
93+
include launch overhead or host time, and it is not end-to-end
94+
wall-clock.
95+
6. Output goes to a fresh `nspy_*` directory in the system temp directory
96+
unless `output_prefix` is set.
97+
98+
## Warm up first-call work
99+
100+
Code that does extra work on its first call breaks the
101+
one-kernel-per-annotation rule and makes runs non-uniform. Examples are Triton
102+
`@triton.autotune` (which benchmarks every config, launching dozens of
103+
kernels), `torch.compile`, and library heuristics that probe on first use.
104+
105+
- Before the decorated call, run the code once for **every shape** in
106+
`configs`, outside profiling. The kernel-count errors suggest this too.
107+
Follow that suggestion rather than reaching for `combine_kernel_metrics`,
108+
range replay or `ignore_kernel_list`, which would hide the warm-up kernels
109+
in the result instead of removing them.
110+
- Triton JIT compilation launches no kernels and is safe inside an
111+
annotation, but warming up outside keeps runs uniform.
112+
113+
## Kernel names
114+
115+
- The `Kernel` column holds the launched kernel's name, for example the
116+
`@triton.jit` function name. Variants of one kernel (different
117+
`num_warps`, block sizes or template arguments) can share a name, so encode
118+
the variant in the annotation name (`f"triton_bm{bm}_w{w}"`).
119+
- Names are **demangled** under NCU (`vectorized_elementwise_kernel`) but
120+
**mangled** under CUPTI (`_ZN2at6native29vectorized_elementwise_kernel...`).
121+
`ignore_kernel_list` matches exact names, so write it for the backend in use.
122+
123+
## Failures inside annotations
124+
125+
- By default, an exception inside an annotation stops the whole profiling
126+
call, so one failing candidate aborts a sweep.
127+
- With `nsight.annotate(name, ignore_failures=True)`, the failed run gets
128+
`NaN`, so the row can show `AvgValue=NaN`, `NumRuns=0` and
129+
`Kernel="dummy_kernel_failure"`. The exception itself is reported only as a
130+
`RuntimeWarning` ("Ignored failure in annotation ..."), so capture warnings
131+
if you need to know why a candidate failed.
132+
133+
## CUPTI backend (experimental)
134+
135+
Opt in by setting `NSPY_NCU_INIT_AT_IMPORT=0` before `import nsight` (in the
136+
shell, or with `os.environ` at the top of the script), then call
137+
`nsight.activate(nsight.Tool.CUPTI)`. It has lower overhead than NCU for
138+
timing-only loops. Check `docs/source/tools.rst` for supported platforms
139+
(currently not Windows), metrics and limitations. Rules that are not written
140+
down there:
141+
142+
- **Activate CUPTI before importing torch.** If `torch` is imported before
143+
`nsight.activate(nsight.Tool.CUPTI)` *and* a CUDA context is created before
144+
the first decorated call, no kernel launch records arrive. Any torch op,
145+
tensor allocation, Triton launch or `torch.cuda.init()` creates a context.
146+
Every profiled call in that process then raises `RuntimeError: No kernels
147+
were profiled`, even though the annotated code did launch kernels. Warm-up
148+
and correctness checks can run anywhere once CUPTI is active. Importing
149+
Triton or `cuda.core` first is fine. (Observed with cupti-python 13.4 and
150+
PyTorch 2.12.)
151+
- CUPTI ignores `cache_control`, so caches are not flushed between runs and
152+
short kernels vary much more than under NCU. Check `RelativeStdDevPct`
153+
before trusting a difference.
154+
155+
## Reading results
156+
157+
`ProfileResults.to_dataframe()` returns one row per
158+
(annotation, config, metric). Inspect `df.columns` and the results docs
159+
rather than assuming columns. Interpretation rules:
160+
161+
- **`NaN` means failed or unavailable, not zero.** `NumRuns` counts valid
162+
samples only.
163+
- **With `normalize_against`, only `AvgValue` is normalized.** `Unit` still
164+
shows the original unit, for example `ns`, and the spread columns stay in
165+
raw units. For comparisons, prefer computing ratios yourself from
166+
un-normalized data. If you do use it on a time metric, speedup is
167+
`1 / AvgValue`.
168+
- `derive_metric(m1, ..., mk, *config_args)` receives metrics in the order
169+
given in `metrics=`, followed by the function's parameters. Return
170+
`(value, unit)` or `{name: (value, unit)}`. Bare scalars get `Unit=NaN` and
171+
a warning.
172+
- With `runs=1` (the default), there is no spread estimate. Use `runs >= 3`
173+
before claiming a difference, and compare it against `RelativeStdDevPct`.
174+
- Throughput metrics (`*.pct_of_peak_sustained_elapsed`, unit `%`) tell you
175+
whether a kernel is memory- or compute-bound. Collect them alongside time
176+
rather than inferring a bottleneck from time alone.
177+
- Do not read plot PNGs to make decisions. Everything `@nsight.analyze.plot`
178+
shows comes from the same DataFrame.
179+
180+
## Writing benchmarks
181+
182+
Use this shape. Warm up and validate first, outside profiling, then profile
183+
with a small, bounded configuration set:
184+
185+
```python
186+
import torch # importing torch does not initialize CUDA
187+
188+
import nsight # must be imported before the first CUDA call
189+
190+
SHAPES = [(4096, 1024), (4096, 4096)]
191+
192+
193+
def my_softmax(x: torch.Tensor) -> torch.Tensor: ... # the candidate to evaluate
194+
195+
196+
@nsight.analyze.kernel(
197+
configs=SHAPES,
198+
runs=3,
199+
metrics=[
200+
"gpu__time_duration.sum",
201+
"dram__throughput.avg.pct_of_peak_sustained_elapsed",
202+
],
203+
)
204+
def bench(rows: int, cols: int) -> None:
205+
x = torch.randn(rows, cols, device="cuda")
206+
with nsight.annotate("baseline"):
207+
torch.softmax(x, dim=1)
208+
with nsight.annotate("candidate"):
209+
my_softmax(x)
210+
211+
212+
def main() -> None:
213+
# Warm up and validate every shape outside profiling. assert_close picks
214+
# tolerances from the dtype; pass atol/rtol if the candidate needs looser.
215+
for rows, cols in SHAPES:
216+
x = torch.randn(rows, cols, device="cuda")
217+
torch.testing.assert_close(my_softmax(x), torch.softmax(x, dim=1))
218+
219+
df = bench().to_dataframe() # configs supply the arguments
220+
columns = ["Annotation", "rows", "cols", "Metric", "AvgValue", "RelativeStdDevPct"]
221+
print(df[columns].to_string(index=False))
222+
223+
224+
if __name__ == "__main__":
225+
main()
226+
```
227+
228+
- A candidate written as several PyTorch ops launches several kernels and
229+
raises the more-than-one-kernel error. To time it as a whole, pass
230+
`combine_kernel_metrics=lambda a, b: a + b`. Summing is right for times and
231+
counts, but not for percentages or rates such as
232+
`*.pct_of_peak_sustained_elapsed`, so profile those per kernel or in a
233+
separate run.
234+
- Validate every candidate with a dtype-appropriate tolerance **before**
235+
accepting its timing. Report incorrect candidates separately from profiling
236+
failures. Faster but wrong is not an improvement.
237+
- Each extra NCU metric can add replay passes. A sweep's cost is roughly
238+
configs × runs × annotations × passes. Start with one config and
239+
`runs=1` as a smoke test.
240+
- Check metric names before a long run. An invalid name only raises
241+
`ProfilerException: Invalid value ...` **after** the whole workload has run.
242+
- If no annotation launches a kernel, NCU writes no report and you get
243+
`ProfilerException: No NVIDIA Nsight Compute report found`. Check that
244+
every annotation really launches a kernel (views, no-ops and cached results
245+
do not).
246+
247+
## GPU execution rules
248+
249+
- Reading and editing code does not need a GPU.
250+
- Assume the GPU is shared. Do not kill other processes, lock clocks with
251+
`nvidia-smi`, change power limits, or start background profiling jobs. Use
252+
`clock_control="base"` (via ncu) only if asked.
253+
- Run with a time limit: your tool's command timeout, or
254+
`timeout 600 python ...` on Linux. Windows has no direct equivalent
255+
(`timeout.exe` only waits), so rely on the tool's timeout there.
256+
Smoke-test with one config before a sweep.
257+
- If the GPU or `ncu` is unavailable, run lint and any CPU-only checks, then
258+
state exactly what was not verified. Never report a GPU test as passing
259+
without having run it.
260+
261+
# Part 2: Working on this repository
262+
263+
## Repository map
264+
265+
| Path | What it is |
266+
| --- | --- |
267+
| `nsight/analyze.py` | Public decorators: `kernel`, `plot`, `ignore_failures`. Docstrings here are the API reference. |
268+
| `nsight/annotation.py` | `nsight.annotate` (context manager and decorator). |
269+
| `nsight/collection/core.py` | Config handling, the run loop (`run_profile_session`), `ProfileResults`. |
270+
| `nsight/collection/ncu.py` | NCU injection loading, `ncu --mode attach` launch, version checks. |
271+
| `nsight/collection/cupti.py` | CUPTI activity-based timing backend. |
272+
| `nsight/extraction.py` | Report to raw DataFrame: kernel-per-annotation checks, `derive_metric`. |
273+
| `nsight/transformation.py` | Aggregation into `AvgValue`/`StdDev`/CI columns, `normalize_against`, `Geomean`. |
274+
| `nsight/visualization.py` | Matplotlib plotting used by `@nsight.analyze.plot`. |
275+
| `nsight/info_collector.py`, `nsight/experimental.py` | Custom info collectors (`@nsight.experimental.collect`). |
276+
| `nsight/thermovision.py` | GPU thermal-headroom throttling guard. |
277+
| `nsight/tools_manager.py` | `nsight.activate` / `deactivate` / `get_active_tool`, `nsight.Tool`. |
278+
| `examples/` | Numbered runnable examples, indexed in `examples/README.md`, run by `examples/test_examples.py`. |
279+
| `tests/` | Pytest suite. Most tests need a GPU and `ncu`. |
280+
| `conftest.py` | Adds `--tool-init`, which selects how the profiling tool is set up. |
281+
| `docs/source/` | Sphinx docs. Release notes in `docs/source/release_notes/`. |
282+
283+
## Commands
284+
285+
The hatch environments are defined in `pyproject.toml`. CI runs the same ones.
286+
287+
Tests and examples also need pytest and a CUDA build of PyTorch, which are
288+
not package dependencies (CI installs PyTorch separately; the hatch test
289+
environments use system packages). Some examples need Triton or cuTile and
290+
are skipped without them.
291+
292+
```bash
293+
pip install -e ".[cu13]" # or [cu12]; add ,cupti for the CUPTI backend
294+
pip install pytest torch # pick the torch build matching your CUDA, see pytorch.org
295+
hatch run lint:check # black --check, isort --check, mypy --strict (whole repo, incl. examples/tests)
296+
hatch run lint:format # apply black + isort
297+
hatch run docs:build # sphinx-build -W (warnings are errors)
298+
pytest tests -v # unit/integration tests (GPU + ncu)
299+
pytest examples -v # runs every example's main()
300+
pytest tests/test_profiler.py::test_name -v # one test
301+
pytest -v --tool-init=cupti # same suites on the CUPTI backend (not on Windows)
302+
hatch run test_cu13:run # what CI runs
303+
```
304+
305+
- `mypy --strict` covers `examples/` and `tests/` too. Annotate every
306+
function. Use the existing `# type: ignore[untyped-decorator]` on
307+
`@triton.jit` and `@pytest.mark.*` decorators.
308+
- `scripts/git-hooks/pre-commit` runs `hatch run lint:check`. Enable it with
309+
`git config core.hooksPath scripts/git-hooks`, unless the user has their own
310+
hooks set up.
311+
- Follow `Contributing.md` for commit sign-off and title format.
312+
313+
## Adding an example
314+
315+
- Use the next free number, `NN_snake_case.py`. Include the SPDX header,
316+
a module docstring with the concepts it introduces, and a typed `main()`
317+
under `if __name__ == "__main__":`.
318+
- Add a `test_NN_*` that imports it and calls `main()` in
319+
`examples/test_examples.py`. Use `pytest.importorskip` for optional
320+
frameworks (Triton, cuTile), and the `_skip_cupti` marker if it needs NCU
321+
metrics.
322+
- Add it to `examples/README.md`, including any extra install step.
323+
- Keep default workloads small. `pytest examples` runs every example in CI,
324+
on several GPUs and on both Linux and Windows.
325+
326+
## Changing library code
327+
328+
- Public behavior changes need a docstring update in `nsight/analyze.py` (the
329+
docs are autodoc-generated), relevant `docs/source/*.rst` pages, and a
330+
release-note entry under `docs/source/release_notes/topics/`.
331+
- New result columns or output formats must be documented in the `kernel`
332+
docstring (CSV section), `ProfileResults.to_dataframe` and
333+
`docs/source/overview/profiling_results.rst`.
334+
- New limitations go in `docs/source/release_notes/known_issues.rst` or the
335+
relevant docs page, not only in this file.
336+
- Add tests in `tests/`. Use `nsight.get_active_tool() == nsight.Tool.CUPTI`
337+
with a `skipif` for NCU-only behavior, as `tests/test_profiler.py` does.
338+
Do not weaken assertions to make a test pass.
339+
- Don't add runtime dependencies without justification. The core deps are
340+
listed in `pyproject.toml`.
341+
342+
# Completion report
343+
344+
Summarize what changed and why, which commands you ran and their outcome,
345+
whether GPU profiling actually ran (and on which GPU and `ncu` version), and
346+
anything not verified.

0 commit comments

Comments
 (0)