diff --git a/.claude/skills b/.claude/skills new file mode 120000 index 000000000..42c5394a1 --- /dev/null +++ b/.claude/skills @@ -0,0 +1 @@ +../skills \ No newline at end of file diff --git a/AGENTS.md b/AGENTS.md new file mode 100644 index 000000000..76be1415e --- /dev/null +++ b/AGENTS.md @@ -0,0 +1,96 @@ +# cuDNN Frontend — Agent Guide + +cuDNN Frontend (FE) is a **header-only C++ library** plus a **Python package** (`nvidia-cudnn-frontend`, imported as `cudnn`) that wraps the cuDNN Graph API, and a growing set of open-source CuTeDSL kernels (SDPA/Flash Attention, MoE grouped GEMM fusions, fused normalizations). + +Directory-specific guides: [include/cudnn_frontend/AGENTS.md](include/cudnn_frontend/AGENTS.md) (C++ library), [python/cudnn/AGENTS.md](python/cudnn/AGENTS.md) (Python API + OSS kernels), [test/AGENTS.md](test/AGENTS.md) (running tests), [samples/AGENTS.md](samples/AGENTS.md). + +## Repository map + +| Path | Purpose | +|---|---| +| `include/` | The header-only C++ library (CMake INTERFACE target `cudnn_frontend`). C++17. | +| `python/` | pybind11 bindings (`python/*.cpp`, `python/pygraph/`) + pure-Python `python/cudnn/` package | +| `python/cudnn//` | Frontend-only OSS CuTeDSL kernels (GEMM fusions, grouped GEMM, BSA/DSA/NSA, SDPA) | +| `samples/` | C++ samples (Catch2 binaries `samples`, `legacy_samples`) and Python notebooks | +| `test/` | `test/cpp` (Catch2 binary `tests`) and `test/python` (pytest) | +| `benchmark/` | Standalone perf harnesses (SDPA training, norms, DSA, CuTeDSL fusions); each has a README | +| `tools/cudnn_repro/` | Standalone CLI that parses cuDNN logs into repro commands (own pyproject) | +| `docs/` | Markdown docs: `operations/` (graph-op reference), `fe-oss-apis/` (OSS kernel APIs), how-to guides | +| `cmake/cuDNN.cmake` | Locates the cuDNN backend library (or reuses existing `CUDNN::` targets) | +| `skills/` | Agent skills (see [Agent skills](#agent-skills)) | + +## Environment requirements + +- NVIDIA GPU required for essentially all tests and samples (SDPA/OSS kernels need Hopper SM90 or Blackwell SM100+). +- CUDA toolkit (`nvcc`), cuDNN **9.x** backend (headers + libs), CMake ≥ 3.23, a C++17 compiler. +- If cuDNN or CUDA are not in default system locations, set `CUDNN_PATH` and `CUDAToolkit_ROOT` (both honored by CMake and `setup.py`). +- Python ≥ 3.9. `cudnn.backend_version()` gates many features at runtime (integer, e.g. 9.12.0 → `91200`); tests skip on older backends. + +## Build + +C++ (builds samples + tests by default): + +```bash +cmake -B build -DCMAKE_BUILD_TYPE=Release # add -DCUDNN_PATH=... -DCUDAToolkit_ROOT=... if not system-installed +cmake --build build -j $(nproc) +# artifacts: build/bin/{samples,legacy_samples,tests} +``` + +CMake options (defaults): `CUDNN_FRONTEND_BUILD_SAMPLES=ON`, `CUDNN_FRONTEND_BUILD_TESTS=ON`, `CUDNN_FRONTEND_BUILD_PYTHON_BINDINGS=OFF`, `CUDNN_FRONTEND_SKIP_JSON_LIB=OFF`. + +The C++ build uses `-Werror` (`/WX` on MSVC) with `-Wall -Wextra -Wpedantic` — new warnings break the build. + +Python (editable; compiles the pybind11 extension via CMake): + +```bash +pip install -e . # core graph API only +pip install -e ".[cutedsl]" # + OSS CuTeDSL kernels (torch, nvidia-cutlass-dsl, cuda-python) +``` + +`setup.py` honors env vars: `CUDNN_PATH`, `CUDA_PATH` / `CUDAToolkit_ROOT`, `DEBUG=1` (debug build), `CMAKE_BUILD_PARALLEL_LEVEL`, `CMAKE_GENERATOR`. + +## Test + +```bash +# C++ (Catch2): list and run cases by name +./build/bin/tests --list-tests +./build/bin/tests "Validate conv node" + +# Python: run from test/python so pytest.ini and conftest.py apply +cd test/python +pytest # default is -m L0 (smoke level) per pytest.ini +pytest -m L1 # deeper levels: L0..L4 +pytest test_conv_fprop.py # one file (still filtered by -m L0 — pass -m "L0 or L1" to widen) +pytest fe_api/ # OSS kernel tests; require ".[cutedsl]" install + SM90/SM100 GPU +``` + +Read [test/AGENTS.md](test/AGENTS.md) before touching tests — `test/python/conftest.py` has import-order and env-var requirements that are easy to break. + +## Format / lint + +```bash +git add +pre-commit run # clang-format 21 (C++/CUDA) + black -l 160 (python + notebooks), staged files only +``` + +First invocation builds the hook environments and can take >5 minutes; later runs are fast. Run on the files you changed (staged files, or `pre-commit run --files `), not `--all-files` — some pre-existing files are not currently formatter-clean, and reformatting them would pollute your diff. C++ style is Google-based, 4-space indent, 120 columns (`.clang-format`); Python is black with line length 160. + +## Conventions + +- `include/` is header-only: no `.cpp` files, no new required dependencies. Vendored third-party code lives in `include/cudnn_frontend/thirdparty/`. +- Every new frontend-only Python API needs: `APIBase` subclass + wrapper, lazy export in `python/cudnn/__init__.py`, docs under `docs/fe-oss-apis/`, and pytest coverage under `test/python/fe_api/`. Full recipe: [python/cudnn/AGENTS.md](python/cudnn/AGENTS.md) and the `cutedsl-kernel-integration` skill. +- Frontend-only OSS APIs are experimental; keep the `[cutedsl]` optional-dependency boundary intact (no eager `torch`/`cutlass` imports at `cudnn` import time). +- Version lives in three places that must stay in sync: `CMakeLists.txt` (`project(... VERSION ...)`), `include/cudnn_frontend_version.h`, `python/cudnn/__init__.py` (`__version__`). +- Runtime debugging: set `CUDNN_FRONTEND_LOG_INFO=1` and `CUDNN_FRONTEND_LOG_FILE=stderr` for FE logs; backend logs via `CUDNN_LOGLEVEL_DBG=3 CUDNN_LOGDEST_DBG=stderr`. + +## Agent skills + +Reusable task recipes live in `skills/` (auto-discovered by Claude Code via `.claude/skills`; other agents: read the relevant `skills//SKILL.md` before starting a matching task): + +- `skills/cutedsl-kernel-integration/` — integrating a CuTeDSL kernel as a frontend-only Python API end to end (API class, wrapper, exports, docs, tests). + +## Links + +- Published documentation: +- In-repo docs index: [llms.txt](llms.txt) · operation reference in [docs/operations/](docs/operations/) · OSS kernel APIs in [docs/fe-oss-apis/overview.md](docs/fe-oss-apis/overview.md) +- PyPI: diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index bcb6f81d0..f42193876 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -1,33 +1,66 @@ # Contributing to cudnn-frontend -If you are interested in contributing to cudnn, we encourage you to: +If you are interested in contributing to cudnn-frontend, we encourage you to: - Post PRs with your ideas and fixes - File issues for bugs you find, things you'd like to see, or questions you have +## Development environment + +Requirements: an NVIDIA GPU (Hopper/Blackwell for attention and OSS kernels), CUDA toolkit, cuDNN 9.x, CMake ≥ 3.23, a C++17 compiler, Python ≥ 3.9. If cuDNN/CUDA are not installed system-wide, point `CUDNN_PATH` and `CUDAToolkit_ROOT` at them. + +```bash +git clone https://github.com/NVIDIA/cudnn-frontend.git +cd cudnn-frontend + +# C++ library, samples, and tests +cmake -B build -DCMAKE_BUILD_TYPE=Release +cmake --build build -j $(nproc) + +# Python package (editable). Add ".[cutedsl]" to work on the OSS CuTeDSL kernels. +pip install -e . +``` + +See [AGENTS.md](AGENTS.md) for the full repo map, all build options, and per-directory guides (also useful for humans, not just coding agents). + +## Running tests + +```bash +./build/bin/tests # C++ (Catch2); run one case: ./build/bin/tests "Validate conv node" +cd test/python && pytest # Python; defaults to the L0 smoke level +``` + +Python tests need `pip install -e ".[cutedsl]"` plus `pytest pytest-xdist looseversion`, and a GPU. Details and conventions: [test/AGENTS.md](test/AGENTS.md). + +## Code formatting + +Formatting is enforced with [pre-commit](https://pre-commit.com) (clang-format for C++/CUDA, black with line length 160 for Python and notebooks): + +```bash +pip install pre-commit +git add +pre-commit run # checks staged files; first run is slow while hook environments install +``` + +Optionally `pre-commit install` to format on every commit. ## Code contributions ### Your first issue -1. Read the project's [README.md](https://github.com/NVIDIA/cudnn-frontend/blob/main/README.md) - to learn how to setup the development environment. +1. Set up the development environment as described above. 2. Comment on the issue saying you are going to work on it and what changes you are going to make. -3. Code! Make sure to update unit tests! -4. When done, [create your pull request](https://github.com/NVIDIA/cudnn-frontend/compare). +3. Code! Make sure to update unit tests, and docs/samples when you change public APIs. + Adding a new frontend-only kernel API? Follow the checklist in [python/cudnn/AGENTS.md](python/cudnn/AGENTS.md). +4. Run the tests and `pre-commit run` locally, then [create your pull request](https://github.com/NVIDIA/cudnn-frontend/compare) and fill in the PR template. 5. Wait for other developers to review your code and update code as needed. 6. Once reviewed and approved, a cudnn-frontend developer will merge your pull request. -7. Once merged to main this will be an untagged version. A release tag will be assigned along with future frontend release by cudnn team. +7. Merged changes ship untagged until the next frontend release, when the cudnn team assigns a release tag. Remember, if you are unsure about anything, don't hesitate to comment on issues and ask for clarifications! -## Code Formatting - -Consistent code formatting is important in the cudnn-frontend project to ensure -readability, maintainability, and thus simplifies collaboration. - -### Branches and Versions +### Branches and versions -The cudnn-frontend repository has one main branch. Please submit a PR to this branch. We will update the doc as the policy changes. +Active development happens on the `develop` branch — base your work on `develop` and submit PRs against it. `main` tracks releases. We will update this doc as the policy changes. ### Branch naming diff --git a/include/cudnn_frontend/AGENTS.md b/include/cudnn_frontend/AGENTS.md new file mode 100644 index 000000000..223e30a2b --- /dev/null +++ b/include/cudnn_frontend/AGENTS.md @@ -0,0 +1,41 @@ +# include/cudnn_frontend — Agent Guide + +The header-only C++ library (CMake INTERFACE target `cudnn_frontend`, C++17). Umbrella header: `include/cudnn_frontend.h`. Build/test commands: [../../AGENTS.md](../../AGENTS.md). + +## Hard rules + +- **Header-only**: no `.cpp` files, no link-time dependencies beyond cuDNN/CUDA. New third-party code must be vendored under `thirdparty/` (currently only `nlohmann/json.hpp`, excludable via `CUDNN_FRONTEND_SKIP_JSON_LIB`). +- Builds with `-Wall -Wextra -Wpedantic -Werror` (GCC/Clang) and `/W4 /WX` (MSVC) — code must be warning-clean on both. +- C++17 only in `include/` (the pybind11 layer under `python/` is C++20). +- Guard anything needing a newer cuDNN with runtime `detail::get_backend_version()` checks (compare against `CUDNN_FRONTEND_VERSION`-style integers, e.g. 9.12.0 → 91200); the same headers must compile against older cuDNN 9.x. +- clang-format (Google-based, indent 4, 120 cols, `SortIncludes: false`) via `pre-commit run`. + +## Layering + +``` +graph_interface.h Graph class (namespace cudnn_frontend::graph); includes every node/*.h + node/*.h one header per op node (matmul, conv_fprop, sdpa, rmsnorm, ...): Node + _attributes + node_interface.h INode base; graph_properties.h holds attribute classes + cudnn_interface.h ICudnn: lowering to cuDNN backend ops/plans + backend/ C-backend descriptor wrappers (execution plans, kernel cache, device properties) +plans.h, knobs.h, graph_helpers.h, context.h plan management, autotuning knobs, error handling +``` + +Legacy flat API (`include/cudnn_frontend_*.h`: Tensor, Operation, ExecutionPlan, Heuristics...) is maintained but frozen — new features target the graph API under this directory. + +## Adding a graph operation (typical shape) + +1. `node/.h`: `_attributes` (+ `NLOHMANN_DEFINE_TYPE_INTRUSIVE`-style serialization where applicable) and the node class implementing infer/expand/lowering. +2. Wire into `graph_interface.h` (include + `Graph::(...)` builder method returning tensor attributes). +3. Serialization support in `utils/serialize.h` if the op is cacheable. +4. Doc page in `docs/operations/`, sample under `samples/cpp/`, tests (C++ and/or `test/python/`). +5. Python surface: pybind wrapper in `python/pygraph/` when the op should be scriptable. + +## experimental/ and generated/ + +- `generated/` holds **open-sourced kernel source embedded as raw C++ string literals** (`inline constexpr const char _source[]`, namespace `cudnn_frontend::experimental::generated`) — SDPA prefill (sm90/sm100) and RMSNorm+SiLU. These files are large and machine-produced; don't hand-edit kernel bodies casually, and don't "clean them up". +- `experimental/` is the NVRTC glue that compiles those strings at runtime (`IOssSdpaEngine`: `check_support`/`build`/`execute`, per-arch engines, `nvrtc_shim.h`). + +## Versioning + +`cudnn_frontend_version.h` defines `CUDNN_FRONTEND_{MAJOR,MINOR,PATCH}_VERSION`. Keep in sync with `CMakeLists.txt` `project(VERSION ...)` and `python/cudnn/__init__.py.__version__` when bumping. diff --git a/llms.txt b/llms.txt new file mode 100644 index 000000000..5439d4a8f --- /dev/null +++ b/llms.txt @@ -0,0 +1,56 @@ +# cuDNN Frontend + +> NVIDIA cuDNN Frontend (FE): a header-only C++ library and Python package (`nvidia-cudnn-frontend`, import `cudnn`) exposing the cuDNN Graph API, plus open-source CuTeDSL kernels (SDPA/Flash Attention, MoE grouped-GEMM fusions, fused normalizations) for Hopper and Blackwell GPUs. + +Published documentation: https://docs.nvidia.com/deeplearning/cudnn/latest/developer/overview.html + +## Start here + +- [README](https://github.com/NVIDIA/cudnn-frontend/blob/main/README.md): overview, installation, feature highlights +- [AGENTS.md](https://github.com/NVIDIA/cudnn-frontend/blob/main/AGENTS.md): repo map, verified build/test/format commands, conventions +- [Contributing](https://github.com/NVIDIA/cudnn-frontend/blob/main/CONTRIBUTING.md): development environment and PR workflow +- [Python package guide](https://github.com/NVIDIA/cudnn-frontend/blob/main/python/cudnn/README.md): package structure, adding frontend-only APIs + +## Graph operation reference + +- [Attention (SDPA fwd/bwd, FP8)](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Attention.md) +- [Matmul](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Matmul.md) +- [Convolutions](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Convolutions.md) +- [Normalizations (LayerNorm, RMSNorm, BatchNorm, InstanceNorm)](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Normalizations.md) +- [MoE Grouped Matmul](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/MoeGroupedMatmul.md) +- [Pointwise](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Pointwise.md) +- [Block Scaling (MXFP8/NVFP4 quantization)](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/BlockScaling.md) +- [RoPE](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/RoPE.md) +- [Causal Conv1d](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/CausalConv1d.md) +- [Concatenate](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Concatenate.md), [Reshape](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Reshape.md), [Slice](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Slice.md), [Transpose](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Transpose.md), [Resampling](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/operations/Resampling.md) + +## Open-source (frontend-only) kernel APIs + +- [FE OSS APIs overview — full catalog and usage pattern](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/overview.md) +- [SDPA forward d256](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/attention/sdpa_fwd_d256.md), [SDPA backward d256](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/attention/sdpa_bwd_d256.md) +- [Block-sparse attention (BSA)](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/bsa.md), [DeepSeek sparse attention (DSA)](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/dsa.md), [Native sparse attention (NSA)](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/nsa.md) +- [GEMM fusions (amax, SwiGLU, sReLU, grouped/discrete MoE variants)](https://github.com/NVIDIA/cudnn-frontend/tree/main/docs/fe-oss-apis/gemm_fusions) +- [RMSNorm + RHT + Amax](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/rmsnorm_rht_amax.md), [RMSNorm + SiLU](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/fe-oss-apis/rmsnorm_silu.md) + +## How-to guides + +- [CUDA graphs](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/cuda-graphs.md) +- [Deviceless ahead-of-time compilation](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/deviceless-ahead-of-time-compilation.md) +- [Dynamic kernel cache](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/dynamic-kernel-cache.md) +- [Custom execution plans](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/custom-execution-plan.md) +- [Compile-time constants](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/CompileTimeConstants.md) +- [Adding PyTorch custom ops](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/adding_torch_custom_ops.md) +- [Python graph and execution backends](https://github.com/NVIDIA/cudnn-frontend/blob/main/docs/python_graph_and_execution_backends.md) + +## Examples + +- [C++ samples (Catch2)](https://github.com/NVIDIA/cudnn-frontend/tree/main/samples/cpp) +- [Python notebook tutorials](https://github.com/NVIDIA/cudnn-frontend/tree/main/samples/python) +- [SDPA training benchmark](https://github.com/NVIDIA/cudnn-frontend/tree/main/benchmark/sdpa_benchmark_training) + +## Optional + +- [Release notes](https://github.com/NVIDIA/cudnn-frontend/releases) +- [PyPI package](https://pypi.org/project/nvidia-cudnn-frontend/) +- [cuDNN backend API reference](https://docs.nvidia.com/deeplearning/cudnn/latest/api/overview.html) +- [Acknowledgements](https://github.com/NVIDIA/cudnn-frontend/blob/main/ACKNOWLEDGEMENTS.md) diff --git a/python/cudnn/AGENTS.md b/python/cudnn/AGENTS.md new file mode 100644 index 000000000..296f7b9e1 --- /dev/null +++ b/python/cudnn/AGENTS.md @@ -0,0 +1,47 @@ +# python/cudnn — Agent Guide + +The `cudnn` Python package: pybind11-backed graph API plus pure-Python **frontend-only OSS kernels** (CuTeDSL). See `README.md` in this directory for the package inventory and [../../AGENTS.md](../../AGENTS.md) for build/test commands. + +## Import-time rules (the most common way to break this package) + +- `import cudnn` must work **without** torch/cutlass/cuda-python installed. Everything that needs them is exported lazily via `_LAZY_OPTIONAL_IMPORTS` in `__init__.py` — a module-level `__getattr__` imports the submodule on first attribute access and re-raises failures as `ImportError` pointing at `pip install nvidia-cudnn-frontend[cutedsl]`. +- Never add an eager `import torch` / `import cutlass` to `__init__.py` or anything it imports transitively. `api_base.py` itself imports them at top level, which is why kernel classes must only be reachable through the lazy table. +- Reuse the existing `[cutedsl]` extra (`pyproject.toml` optional-dependencies) unless a kernel truly needs a new package. + +## Frontend-only kernel package layout + +``` +python/cudnn// # or grouped_gemm//, discrete_grouped_gemm//, sdpa// +├── __init__.py # exports API class + wrapper via __all__ +├── api.py # APIBase subclass + _wrapper() function +└── .py # CuTeDSL kernel implementation(s); some families use csrc/ per-arch trees +``` + +Shared helpers (schedulers, metadata utils, e.g. `grouped_gemm/moe_*.py`) stay internal to the family package — never exported through `cudnn`. + +## The APIBase contract (`api_base.py`) + +Every OSS kernel API extends `APIBase` and implements: + +- `check_support() -> bool` — validate dtype/shape/stride/arch/config via the `_check_tensor_*` / `_value_error_if` helpers; must set `self._is_supported`. Works on `TensorDesc` (metadata-only tensors), so it runs without GPU storage. +- `compile()` — calls `self._ensure_support_checked()`, builds and `cute.compile`s the kernel, caches in `self._compiled_kernel`. +- `execute(..., current_stream=None)` — runs the cached kernel. + +`__call__` = compile-if-needed + execute. High-level wrappers (`_wrapper_sm100(...)`) allocate outputs and return a **`TupleDict`** (dict that also unpacks as a tuple) with stable, documented key order. FP4x2 packing: use `_tensor_shape`/`_tensor_stride`, which double the innermost dim when `interpret_uint8_as_fp4x2` is set. + +## Adding a new frontend-only API — required checklist + +1. Kernel package under the closest existing family (layout above). +2. `APIBase` subclass + wrapper in `api.py`. +3. Exports: family `__init__.py` `__all__` **and** `_LAZY_OPTIONAL_IMPORTS` in `python/cudnn/__init__.py`; register any new package dir in `pyproject.toml` packages list. +4. Docs: page under `docs/fe-oss-apis/` (family subdir) + link it from `docs/fe-oss-apis/overview.md`. +5. Tests: `test/python/fe_api//test_.py` (+ `_utils.py`/reference), covering check_support pass/fail and numerical reference comparison. + +The `cutedsl-kernel-integration` skill (`skills/cutedsl-kernel-integration/`) documents this workflow in detail, including how to classify a kernel into a family — follow it for any kernel integration. + +## Other notes + +- `wrapper.py` `Graph` context manager (the pythonic graph builder) requires cuDNN backend ≥ 9.12 (`backend_version() >= 91200`) and builds plans on `__exit__`. +- Torch custom ops live in `experimental/ops/` (pattern doc: `docs/adding_torch_custom_ops.md`); they cache built graphs per config and use stable `_UIDs` enums. +- dtype conversions go through `datatypes.py`, which probes torch/cutlass availability lazily — keep it that way. +- Formatting: black, line length 160. diff --git a/samples/AGENTS.md b/samples/AGENTS.md new file mode 100644 index 000000000..f7ab2db89 --- /dev/null +++ b/samples/AGENTS.md @@ -0,0 +1,28 @@ +# samples — Agent Guide + +Usage examples for the C++ graph API and the Python API. Build commands: [../AGENTS.md](../AGENTS.md). + +## Layout + +| Path | What | +|---|---| +| `cpp/` | Current C++ graph-API samples, grouped by topic (`convolution/`, `matmul/`, `sdpa/`, `norm/`, `moe_grouped_matmul/`, ...) — Catch2 cases in binary `build/bin/samples` | +| `legacy_samples/` | Frozen samples for the legacy flat API — binary `build/bin/legacy_samples`. Don't add here. | +| `python/` | Numbered Jupyter notebooks (`00_introduction.ipynb` ...) — the Python tutorial sequence | +| `llama/`, `llm_coverage/` | End-to-end LLaMA tie-out scripts and per-op coverage scripts (plain `.py`, see their READMEs) | + +## Running + +```bash +./build/bin/samples --list-tests +./build/bin/samples "Cached sdpa" # run one Catch2 case by name +``` + +Samples need a GPU + cuDNN backend at runtime; individual cases `SKIP()` on unsupported arch or backend version. + +## Conventions for new samples + +- New C++ samples go under `samples/cpp//` as a Catch2 `TEST_CASE` added to `samples/cpp/CMakeLists.txt`; reuse helpers from `samples/cpp/utils/`. +- Follow the existing pattern: build graph → `validate()` → `build_operation_graph()` → plans → `check_support()` (skip gracefully if unsupported) → execute → verify. +- Python notebooks are formatted by `black-jupyter` (`pre-commit run`); keep the numbered-prefix naming so the tutorial order stays obvious. +- A new public feature should come with a sample here and a doc page under `docs/`. diff --git a/test/AGENTS.md b/test/AGENTS.md new file mode 100644 index 000000000..455bec3dd --- /dev/null +++ b/test/AGENTS.md @@ -0,0 +1,43 @@ +# test — Agent Guide + +Two suites: `test/cpp` (Catch2, C++ graph API) and `test/python` (pytest). Both need an NVIDIA GPU and a cuDNN 9.x backend at runtime. Build/install commands: [../AGENTS.md](../AGENTS.md). + +## C++ tests (`test/cpp`) + +- Catch2 v3 binary, target `tests`, built by the default CMake build (`CUDNN_FRONTEND_BUILD_TESTS=ON`) into `build/bin/tests`. +- Run all: `./build/bin/tests`. List: `--list-tests`. One case: `./build/bin/tests "Validate conv node"`. Filter by tag: `./build/bin/tests "[serialize]"`. + +## Python tests (`test/python`) + +Run from `test/python` so `pytest.ini` and `conftest.py` apply: + +```bash +cd test/python +pytest # pytest.ini addopts default to -m L0 (smoke) --tb=short --no-header +pytest -m L1 # levels L0..L4; higher = larger sweeps +pytest -n 4 # pytest-xdist; mind marker gpu_exclusive for tests that need the GPU alone +pytest test_conv_fprop.py # one file — note the default -m L0 filter still applies +pytest fe_api/gemm/ # OSS kernel tests +``` + +Requirements: `pip install -e ".[cutedsl]"` plus `pytest pytest-xdist looseversion`. `fe_api/` additionally requires an SM90/SM100-class GPU; tests skip (or should skip) on unsupported arch/dtype/backend-version combos rather than fail. + +### conftest.py landmines — read before editing + +- `PYTORCH_CUDA_ALLOC_CONF` is set at the very top, **before any torch import** (torch reads it once at CUDA-allocator init). Don't move it, and don't import torch in a plugin that loads earlier. +- `import transformer_engine` happens (in try/except) **before** `import cudnn` — TE and cuDNN conflict if loaded in the other order. Preserve this ordering. +- `torch.cuda.synchronize` is monkeypatched to a guard that prints a filtered traceback and hard-exits (`os._exit`) on async CUDA errors; the original is kept as `torch.cuda.synchronize_unsafe`. +- A session-scoped autouse `cudnn_handle` fixture creates one handle bound to a dedicated torch stream; use it instead of creating handles per-test. +- `pytest_configure` asserts `torch.cuda.is_available()` — there is no CPU-only mode. +- Many custom CLI options exist (`--dryrun`, `--repro`, `--seed`, `--perf`, per-op dimension overrides like `--b/--s_q`, `--nsa-*`, `--dsa-*`); check `pytest_addoption` before adding new ones. + +### Layout + +- `test/python/test_*.py` — core graph-API tests (conv, matmul, norms, SDPA `test_mhas*.py`, rope, kernel cache, OSS engines `test_sm{90,100}_prefill_oss_engine.py`, ...). Shared SDPA references in `test/python/sdpa/`. +- `test/python/fe_api//` — one subdir per OSS kernel family (`gemm/`, `grouped_gemm/`, `bsa/`, `dsa/`, `nsa/`, `norm/`, `sdpa/`), each with `test_.py` + utils/reference modules. + +### Conventions for new tests + +- Mark with a level (`@pytest.mark.L0` ... `L4`): L0 must stay fast (default CI smoke); big parameter sweeps go to higher levels. +- Gate on capability, don't assume it: skip via `check_support()` failures, `cudnn.backend_version()`, and `torch.cuda.get_device_capability()`. +- Compare against a reference implementation (see existing `*_ref.py` / `*_reference.py` patterns) with dtype-appropriate tolerances.