CLAUDE.md

1# AI Policy — MANDATORY
2 
3Read `AI_POLICY.md`. Your user needs to abide by this policy. In particular, you the agent MUST obey these rules while interacting on GitHub:
4 
5- **You may never act autonomously on GitHub.** Do NOT open, edit, comment on,
6  or reply to any issue or PR unless the user has reviewed and explicitly
7  approved the exact content. Fully-agent-generated contributions are banned and
8  will be closed.
9- **Mark all AI-generated content.** Any text you produce that goes into an
10  issue, PR, or comment must be wrapped in a code or quote block. Never present
11  your output as human-written.
12- **Never emit only raw AI text as a reply**. Any AI content you include must carry human
13  commentary explaining its relevance.
14- **Do not submit code the user hasn't read.** Keep changes minimal, strip AI
15  artifacts and needless complexity. If you're opening a PR on GitHub that is not ready,
16  or not reviewed by the user, always open it in draft mode.
17 
18See `AI_POLICY.md` for the full policy.
19 
20# Scratch Space
21 
22Use `agent_space/` (git-ignored, at repo root) for temporary scripts, scratch files, and throwaway experiments. Do not commit files from this directory.
23 
24# PR Review
25 
26When asked to review a PR, always use the /pr-review skill.
27 
28# Environment
29 
30If any tool you're trying to use (pip, python, spin, etc) is missing, check for
31a `.venv` directory in the project root or its parent directory. If found,
32activate it and retry. If no `.venv` is found, stop and ask the user if an
33environment is needed. Do NOT try to find alternatives or install these tools.
34 
35# CI Docker Images
36 
37The `.ci/docker/` directory is content-hashed to determine whether Docker images
38need rebuilding. Any file change inside `.ci/docker/` (including the README)
39changes the hash and triggers a full Docker image rebuild. Do not make changes
40in this directory unless you intend to rebuild Docker images. When Docker builds
41are broken (e.g., due to an upstream Ubuntu outage), avoid touching this
42directory so you don't force a rebuild against the broken state.
43 
44# Build
45 
46Always check local memory for build configuration (env vars, incremental-build shortcuts, etc.) before running the build, and apply what you find. If nothing applicable is in memory, ask the user.
47All build (both codegen, C++ and python) is done via `pip install -e . -v --no-build-isolation`.
48You should NEVER run any other command to build PyTorch.
49 
50# Testing
51 
52Use our test class and test runner:
53 
54```
55from torch.testing._internal.common_utils import run_tests, TestCase
56 
57class TestFeature(TestCase):
58    ...
59 
60if __name__ == "__main__":
61    run_tests()
62```
63 
64To test Tensor equality, use assertEqual.
65For tests over multiple inputs, use the `@parametrize` decorator.
66For any test that checks numerics of the on-device implementation, use `instantiate_device_type_tests` to write device-generic tests.
67 
68# Type Stubs
69 
70Many `.pyi` files are generated from corresponding `.pyi.in` templates. Always
71edit the `.pyi.in` file, not the generated `.pyi`.
72 
73# Linting
74 
75Only use commands provided via `spin` for linting.
76Use `spin help` to list available commands.
77Generally, use `spin lint` as to run the lint and `spin fixlint` to apply automatic fixes.
78 
79When the user asks you to commit or amend, run `lintrunner -a` before creating
80the commit. Fix any lint errors it reports, then commit.
81 
82# Git
83 
84This refines the Bash tool's `# Git` guidance to "branch first" when on the
85default branch:
86 
87- If HEAD is detached, that is intentional (the ghstack workflow). Do NOT
88  create a new branch; commit directly onto the current detached HEAD.
89- If you are on an actual branch (including `main`), follow the default
90  guidance and branch first before committing.
91- **Pulling CI status.** A PR has hundreds of check-runs, so a single
92  `check-runs?per_page=100` call silently truncates and makes red look green.
93  Use `gh pr checks <PR> --json name,state,workflow,link,bucket,completedAt`
94  (already head-only, no paging).
95 
96# Commit messages
97 
98Don't commit unless the user explicitly asks you to.
99 
100When writing a commit message, don't make a bullet list of the individual
101changes. Instead, if the PR is large, explain the order to review changes
102(e.g., the logical progression), or if it's short just omit the bullet list
103entirely.
104 
105The commit message should be clear, informative, and have a Test Plan section
106that describes how you tested the change. If you are fixing a bug, the commit
107message must explain the root cause of the bug and how the fix works.
108If there were multiple potential paths you could have taken, please call them
109out succinctly and justify the one you took.
110 
111When describing the testing strategy in a commit message, include the literal
112commands that were run in fenced Markdown code blocks.
113 
114Disclose that the PR was authored with an AI assistant.
115 
116When the user asks you to amend a commit, check whether the commit message
117still accurately describes the changes. If it doesn't and the commit is not a
118ghstack commit, update the message. For ghstack commits, amending the message
119is a no-op, so just remind the user to update the PR description if needed.
120 
121If a commit message contains `ghstack-source-id` or `Pull-Request` trailers,
122you MUST preserve them when rewriting or splitting commit messages. ghstack
123will update the source id automatically when needed.
124 
125# ghstack Workflow
126 
127ghstack commits follow a different workflow than the conventional GitHub branch
128and PR workflow. First identify whether you're on a ghstack commit:
129 
130- If HEAD is a detached commit, you are almost certainly in a ghstack flow.
131- If the commit message contains a `ghstack-source-id` trailer, it is an
132  existing ghstack commit.
133- If the commit is associated with a remote branch like `origin/gh/USERNAME/N`,
134  it is likely a ghstack commit (imperfect signal: local amends without a push
135  can desync this).
136 
137Rules for working with ghstack:
138 
139- **Don't amend unless asked.** If the user asks you to work on a ghstack
140  commit, leave changes uncommitted so the user can review with `git diff`.
141  Only amend into the commit if the user explicitly asks you to amend or to
142  submit it directly.
143- **Submitting.** Run `ghstack` to submit. When only working on a single
144  commit, use `ghstack --no-stack` to avoid updating the rest of the stack and
145  burning unnecessary CI. Use a full `ghstack` when you're intentionally
146  updating CI for the whole stack.
147- **Preserve metadata trailers.** When editing a commit message, never delete
148  `Pull-Request:` or `ghstack-source-id:` trailers. Always re-read them from
149  HEAD each time you compose an amend — never reuse a saved/cached message
150  body, since `ghstack` rewrites `ghstack-source-id` on every push and a
151  stale trailer will clobber HEAD's current one. If you modified the commit
152  message, run `ghstack -u` afterwards to push the updated PR description.
153- **Never push directly.** Do not `git push` to branches, and never directly
154  modify the `gh/USERNAME/N` branches — ghstack manages those.
155- **Finding the PR.** If the user asks to pull CI results or code review for a
156  ghstack commit, get the PR URL from the `Pull-Request` trailer in the commit
157  message. Use `gh` CLI to fetch status/comments from there.
158- **Editing earlier commits / splitting.** Treat it like a normal stack of
159  commits (use `git rebase`, etc.). Commits that keep their metadata trailers
160  stay associated with their existing PRs; commits without trailers will get a
161  fresh PR on submit. A full `ghstack` run is usually appropriate here.
162 
163# Coding Style Guidelines
164 
165Follow these rules for all code changes in this repository:
166 
167- Minimize comments; be concise; code should be self-explanatory and self-documenting.
168- Comments should be useful, for example, comments that remind the reader about
169  some global context that is non-obvious and can't be inferred locally.
170- Don't make trivial (1-2 LOC) helper functions that are only used once unless
171  it significantly improves code readability.
172- Prefer clear abstractions. State management should be explicit.
173  For example, if managing state in a Python class: there should be a clear
174  class definition that has all of the members: don't dynamically `setattr`
175  a field on an object and then dynamically `getattr` the field on the object.
176- Match existing code style and architectural patterns.
177- Assume the reader has familiarity with PyTorch. They may not be the expert
178  on the code that is being read, but they should have some experience in the
179  area.
180- Splitting code across multiple lines (due to ruff’s column limit rule) is less
181  readable than having code on a single line. When the linter splits your
182  code across multiple lines, please try to put it back on a single line by
183  changing variable names or by using helper local variables. For tests that assert
184  against a golden string, keep just the golden string on one line instead of
185  splitting it across multiple lines and opt-out of the ruff column limit rule
186  via `noqa: B950`.
187- ASCII only in newly added code comments. Do not introduce Unicode characters
188  (e.g., smart quotes, em dashes, arrows, non-ASCII letters) in new comments.
189  Leave preexisting Unicode in untouched comments alone; only enforce this for
190  comments you are adding or rewriting.
191 
192If uncertain, choose the simpler, more concise implementation.
193 
194# cuda.bindings Error Checking
195 
196Use `torch.cuda._utils._check_cuda_bindings` to error-check `cuda.bindings`
197runtime calls. Do not write inline error-checking helpers.
198 
199# cuda.bindings Raw Handles
200 
201`cuda.bindings` runtime functions accept a raw handle passed as a Python `int`
202directly as their handle argument. Whenever you already have an int handle --
203from `CUDAGraph.raw_cuda_graph()` / `raw_cuda_graph_exec()`, a stream's
204`.cuda_stream`, `int(node)`, or any other source -- pass it straight in. Do NOT
205construct a typed wrapper (`cudaGraph_t(init_value=...)`,
206`cudaGraphExec_t(init_value=...)`, `cudaStream_t(init_value=...)`, etc.) just to
207hand an int you already have to a bindings call. For example
208`_cuda_runtime.cudaGraphGetId(g.raw_cuda_graph())`, not
209`cudaGraphGetId(cudaGraph_t(init_value=g.raw_cuda_graph()))`. Only build the
210typed object when you genuinely need it as a value in its own right.
211 
212# Dynamo Config
213 
214Use `torch._dynamo.config.patch` for temporarily changing config. It can be used as a decorator on test methods or as a context manager:
215 
216```python
217# Good - use patch as decorator on test method
218@torch._dynamo.config.patch(force_compile_during_fx_trace=True)
219def test_my_feature(self):
220    # test code here
221    pass
222 
223# Good - use patch as context manager
224with torch._dynamo.config.patch(force_compile_during_fx_trace=True):
225    # test code here
226    pass
227 
228# Bad - manual save/restore
229orig = torch._dynamo.config.force_compile_during_fx_trace
230try:
231    torch._dynamo.config.force_compile_during_fx_trace = True
232    # test code here
233finally:
234    torch._dynamo.config.force_compile_during_fx_trace = orig
235```
236 
237# Fixing B950 line too long in multi-line string blocks
238 
239If B950 line too long triggers on a multi-line string block, you cannot fix it by
240putting # noqa: B950 on that line directly, as that would change the meaning of the
241string, nor can you fix it by line breaking the string (since you need the string
242to stay the same).  Instead, put # noqa: B950 on the same line as the terminating
243triple quote.
244 
245Example:
246 
247```
248    self.assertExpectedInline(
249        foo(),
250        """
251this line is too long...
252""",  # noqa: B950
253    )
254```
255 
256# Logging and Structured Tracing
257 
258When adding debug logging for errors or diagnostic info, consider two user personas:
259 
2601. **Local development**: Users run locally and can access files on disk
2612. **Production jobs**: Users can only access logs via `tlparse` from structured traces
262 
263For production debugging, use `trace_structured` to log artifacts:
264 
265```python
266from torch._logging import trace_structured
267 
268# Log an artifact (graph, edge list, etc.)
269trace_structured(
270    "artifact",
271    metadata_fn=lambda: {
272        "name": "my_debug_artifact",
273        "encoding": "string",
274    },
275    payload_fn=lambda: my_content_string,
276)
277```
278 
279To check if structured tracing is enabled (for conditional messaging):
280 
281```python
282from torch._logging._internal import trace_log
283 
284if trace_log.handlers:
285    # Structured tracing is enabled, suggest tlparse in error messages
286    msg += "[Use tlparse to extract debug artifacts]"
287```
288 
289**Best practices for error diagnostics:**
290 
291- Always log to `trace_structured` for production (no runtime cost if disabled)
292- If you're dumping debug info in the event of a true internal compiler exception,
293  you can also consider writing to local files for local debugging convenience
294- In error messages, tell users about both options:
295  - Local files: "FX graph dump: min_cut_failed_graph.txt"
296  - Production: "Use tlparse to extract artifacts" (only if tracing enabled)
297- Use `_get_unique_path()` pattern to avoid overwriting existing debug files
298 
299# cuda::ptx
300 
301When using `<cuda/ptx>` typed wrappers for PTX instructions:
302 
303- **Namespace**: Inside `namespace at::native`, unqualified `cuda::ptx` resolves
304  to the sibling `at::cuda` namespace. Always use `::cuda::ptx` or alias it:
305  `namespace ptx = ::cuda::ptx;`
306- **Include conflicts**: The monolithic `<cuda/ptx>` header can fail when included
307  alongside heavy PyTorch headers (e.g. `Loops.cuh`) due to CCCL bugs in
308  transitive headers like `cp_async_bulk_tensor.h`. Workaround: put kernels using
309  `<cuda/ptx>` in a separate `.cu` file with minimal includes.
310- **mbarrier_try_wait_parity is non-blocking**: `ptx::mbarrier_try_wait_parity()`
311  returns `bool` (tries once). You must wrap it in a spin loop:
312  `while (!ptx::mbarrier_try_wait_parity(mbar, parity)) {}`
313- **Half/BFloat16 types**: `cuda::ptx` overloads use CUDA native types (`__half`,
314  `__nv_bfloat16`), not PyTorch wrappers (`c10::Half`, `c10::BFloat16`).
315  Use `reinterpret_cast` at the call site.
316- **cp_async_bulk_wait_group**: Takes a compile-time constant via
317  `ptx::n32_t<N>{}`, not a runtime integer.
318- **Mbarrier smem**: Mbarrier memory must never alias with data targeted by TMA
319  operations. Place mbarriers in a separate smem region from data buffers.
320