Comparison

A · CLAUDE.md · pytorch/pytorch B · CLAUDE.md · pytorch/pytorch

What each file covers, counted
Dimension	Shared	Only in A	Only in B	Overlap
Sections	1	22	22	2%
Commands	0	3	6	0%
Section tags	3	1	6	30%

What each file covers

Sections

1 shared · 22 only in A · 22 only in B

− torch/_dynamo
− Architecture Overview
− Key Abstractions
− VariableTracker (`variables/`)
− Source (`source.py`)
− Guards (`guards.py`)
− Side Effects (`side_effects.py`)
− Other key files
− Graph Breaks
− Common patterns
− Debugging
− TORCH_LOGS
− Reproducing crashes
− Structured tracing (for production)
− Compile-time profiling
− comptime.breakpoint() (`comptime.py`)
− Bytecode Debugger (`bytecode_debugger.py`)
− C++ Runtime (`torch/csrc/dynamo/`)
− Frame Evaluation
− Cache
− Guard Evaluation Tree (`guards.cpp`)
− Other C++ files
+ AI Policy — MANDATORY
+ Scratch Space
+ PR Review
+ Environment
+ CI Docker Images
+ Build
+ Type Stubs
+ Linting
+ Git
+ Commit messages
+ ghstack Workflow
+ Coding Style Guidelines
+ cuda.bindings Error Checking
+ cuda.bindings Raw Handles
+ Dynamo Config
+ Good - use patch as decorator on test method
+ Good - use patch as context manager
+ Bad - manual save/restore
+ Fixing B950 line too long in multi-line string blocks
+ Logging and Structured Tracing
+ Log an artifact (graph, edge list, etc.)
+ cuda::ptx
Testing

Commands

0 shared · 3 only in A · 6 only in B

− python test/dynamo/test_misc.py
− python test/dynamo/test_misc.py MiscTests.test_foo
− python test/dynamo/test_misc.py -k test_foo
+ pip install -e . -v --no-build-isolation
+ gh pr checks <PR> --json name,state,workflow,link,bucket,completedAt
+ git diff
+ git push
+ gh/USERNAME/N
+ git rebase

Section tags

3 shared · 1 only in A · 6 only in B

− architecture
+ setup
+ lint-format
+ types
+ git-pr
+ do-not
+ agent-behaviour
build
test
code-style

Line diff

+249 added−235 removed71 unchanged22.2% identical

pytorch/pytorch · torch/_dynamo/CLAUDE.md

@@ −1 @@

1# torch/_dynamo

3TorchDynamo is a Python-level JIT compiler that captures PyTorch programs into

4FX graphs by symbolically executing Python bytecode. It hooks into CPython's

5PEP 523 frame evaluation API to intercept execution, traces operations into an

6FX graph, compiles the graph with a backend (e.g. Inductor), and generates new

7bytecode that calls the compiled code.

9## Architecture Overview

11The compilation pipeline, in execution order:

131. **`eval_frame.py`** — Runtime entry point. `torch.compile()` wraps a

14 function in an `OptimizedModule`. At runtime, the C extension

15 (`torch._C._dynamo.eval_frame`) intercepts Python frames via PEP 523.

162. **`convert_frame.py`** — `ConvertFrameAssert.__call__` checks caches,

17 handles recompilation limits, calls `_compile()` → `trace_frame()`.

183. **`symbolic_convert.py`** — The heart of Dynamo. `InstructionTranslator`

19 symbolically executes bytecode instruction-by-instruction. Maintains a

20 symbolic `stack` (list of `VariableTracker`s) and `symbolic_locals` (dict of

21 name → `VariableTracker`). Opcodes are dispatched via a `dispatch_table`

22 built by `BytecodeDispatchTableMeta`.

234. **`output_graph.py`** — `OutputGraph` owns the FX graph being built (via

24 `SubgraphTracer`), the `SideEffects` tracker, guards, shape environment,

25 and graph args. `compile_subgraph()` finalizes the graph, calls the backend,

26 and generates output bytecode.

275. **`codegen.py`** — `PyCodegen` emits output bytecode: loads graph inputs

28 (via `Source.reconstruct()`), calls the compiled graph, unpacks outputs, and

29 replays side effects.

306. **`resume_execution.py`** — Generates continuation functions for execution

31 after graph breaks.

33## Key Abstractions

35### VariableTracker (`variables/`)

37Every Python value encountered during tracing is wrapped in a `VariableTracker`

38subclass. Key interface: `as_python_constant()`, `as_proxy()`,

39`call_function()`, `call_method()`, `getattro_impl()`, `reconstruct()`.

41Key fields: `source` (where the value came from, for guards) and

42`mutation_type` (whether/how mutations are tracked).

44**Factory**: `VariableTracker.build(tx, value, source=...)` dispatches to

45`VariableBuilder` (sourced values needing guards) or `SourcelessBuilder`

46(values created during tracing).

48Key subclass families in `variables/`: `TensorVariable` / `SymNodeVariable`

49(tensor.py), `ConstantVariable` (constant.py), `ListVariable` /

50`TupleVariable` (lists.py), `ConstDictVariable` (dicts.py), `SetVariable` (sets.py),

51`UserFunctionVariable` (functions.py), `BuiltinVariable` (builtin.py),

52`NNModuleVariable` (nn_module.py), `UserDefinedObjectVariable`

53(user_defined.py), `TorchHigherOrderOperatorVariable` (higher_order_ops.py),

54`LazyVariableTracker` (lazy.py). `VariableBuilder` and `SourcelessBuilder` are

55in builder.py.

57### Source (`source.py`)

59Tracks value provenance — how to access a value at runtime. Used for guard

60generation (`source.make_guard(GuardBuilder.XXX)`) and bytecode reconstruction

61(`source.reconstruct(codegen)`). Root sources: `LocalSource`, `GlobalSource`.

62Chained sources: `AttrSource`, `GetItemSource`, `NNModuleSource`, etc.

64### Guards (`guards.py`)

66Runtime conditions that must hold for cached compiled code to be reused.

67Install via `install_guard(source.make_guard(GuardBuilder.TYPE_MATCH))`.

68Common types: `TYPE_MATCH`, `ID_MATCH`, `EQUALS_MATCH`, `TENSOR_MATCH`,

69`SEQUENCE_LENGTH`. At finalization, `CheckFunctionManager` builds a tree of

70C++ `GuardManager` objects for fast runtime checking.

72### Side Effects (`side_effects.py`)

74Tracks mutations during tracing (attribute stores, list mutations, cell

75variable updates, tensor hooks) and replays them as bytecode after graph

76execution. The `MutationType` system (`variables/base.py`) controls what

77mutations are allowed: `ValueMutationNew/Existing`,

78`AttributeMutationNew/Existing`, or `None` (immutable). The `scope` field

79prevents cross-scope mutations inside higher-order operators.

81### Other key files

83- `trace_rules.py` — inline/skip/graph-break decisions per function

84- `exc.py` — exception hierarchy: `Unsupported` (graph break), `RestartAnalysis`

85 (restart tracing), `ObservedException` (user exceptions during tracing),

86 `BackendCompilerFailed`

87- `config.py` — configuration flags, supports `config.patch()` context

88 manager/decorator

89- `bytecode_transformation.py` / `bytecode_analysis.py` — low-level bytecode

90 manipulation, liveness analysis

91- `pgo.py` — profile-guided optimization for dynamic shapes

92- `polyfills/` — traceable replacements for stdlib functions

93- `repro/` — reproduction/minification tools

95## Graph Breaks

97Call `unimplemented()` (from `exc.py`) to trigger a graph break:

99```python

100from torch._dynamo.exc import unimplemented

101from torch._dynamo import graph_break_hints

102

103unimplemented(

104 gb_type="short_category_name",

105 context=f"dynamic details: {value}",

106 explanation="Human-readable explanation of why this breaks the graph.",

107 hints=[*graph_break_hints.SUPPORTABLE],

108)

109```

110

111- `gb_type`: Context-free category (no dynamic strings).

112- `context`: Developer-facing details (can be dynamic).

113- `explanation`: User-facing explanation (can be dynamic).

114- `hints`: Use constants from `graph_break_hints.py`: `SUPPORTABLE`,

115 `FUNDAMENTAL`, `DIFFICULT`, `DYNAMO_BUG`, `USER_ERROR`,

116 `CAUSED_BY_EARLIER_GRAPH_BREAK`.

117

118The `break_graph_if_unsupported` decorator on instruction handlers catches

119`Unsupported`, logs the graph break, updates the `SpeculationLog`, and restarts

120analysis. On the second pass, the partial graph is compiled at the break point

121and a resume function handles the rest.

122

123## Testing

124

125Tests live in `test/dynamo/`. Use `torch._dynamo.test_case.TestCase` as base

126class — it calls `torch._dynamo.reset()` in setUp/tearDown and patches config

127for strict error checking.

128

129```bash

130python test/dynamo/test_misc.py # whole file

131python test/dynamo/test_misc.py MiscTests.test_foo # single test

132python test/dynamo/test_misc.py -k test_foo # pattern match

133```

134

135### Common patterns

136

137The default backend to `torch.compile()` is `backend="eager"`.

138

139**CompileCounter** — count compilations and graph ops:

140```python

141cnt = torch._dynamo.testing.CompileCounter()

142

143@torch.compile(backend=cnt)

144def fn(x):

145 return x + 1

146

147fn(torch.randn(10))

148self.assertEqual(cnt.frame_count, 1)

149self.assertEqual(cnt.op_count, 1)

150```

151

152**fullgraph=True** — assert no graph breaks:

153```python

154torch.compile(fn, backend="eager", fullgraph=True)(x)

155```

156

157**EagerAndRecordGraphs** — inspect captured FX graphs:

158```python

159backend = torch._dynamo.testing.EagerAndRecordGraphs()

160torch.compile(fn, backend=backend)(x)

161graph = backend.graphs[0]

162```

163

164**normalize_gm + assertExpectedInline** — snapshot test graph output:

165```python

166from torch._dynamo.testing import normalize_gm

167self.assertExpectedInline(

168 normalize_gm(backend.graphs[0].print_readable(False)),

169 """\

170expected output here

171""",

172)

173```

174

175Call `torch._dynamo.reset()` within a test when testing multiple compilation

176scenarios in a single test method. The base class handles setUp/tearDown reset

177automatically.

178

179## Debugging

180

181### TORCH_LOGS

182

183```bash

184TORCH_LOGS="graph_breaks" python script.py # see graph breaks

185TORCH_LOGS="guards,recompiles" python script.py # see guards and recompilation reasons

186TORCH_LOGS="graph_code" python script.py # see captured FX graph code

187TORCH_LOGS="+dynamo" python script.py # full debug logging

188TORCH_LOGS="bytecode" python script.py # see bytecode transformations

189```

190

191### Reproducing crashes

192

193When writing a minimal repro for a Dynamo crash:

194

1951. **Capture `TORCH_LOGS="+dynamo"` for the known-failing case.** Look for

196 the frame ID (e.g. `[3/0_1]`), INLINING/FAILED/COMPILING/Restart events,

197 speculation behavior, and graph break reasons.

1982. **Capture the same logs for your repro attempt** and diff against the

199 failing case. The divergence point tells you what condition you're missing.

2003. **Match each condition from the crash traceback:**

201 - Does the code need a speculation checkpoint? Add tensor ops before the

202 failing code so `compile_subgraph` creates one.

203 - Does the function need to be inlined on retry (not skipped)? The graph

204 break type matters — `step_unsupported` keeps inlining on retry while

205 `unimplemented` may skip the function.

206 - Does the crash happen in a resume function? Add a `graph_break()` earlier

207 so PEP 523 compiles the resume as a fresh frame.

208

209### Structured tracing (for production)

210

211```bash

212TORCH_TRACE=/path/to/dir python script.py # explicit trace directory

213```

214

215Analyze with `tlparse`.

216

217### Compile-time profiling

218

219`TORCH_COMPILE_DYNAMO_PROFILER=1` prints per-function cumtime/tottime

220(cProfile-style) showing where Dynamo spends time during tracing. Set to a

221file path instead to save a profile loadable by `snakeviz`.

222

223### comptime.breakpoint() (`comptime.py`)

224

225Drops into pdb during **compilation** to inspect Dynamo state. Call

226`comptime.breakpoint()` in user code; in the pdb session use `ctx`

227(`ComptimeContext`) to call `print_locals()`, `print_bt()`, `print_graph()`,

228or `get_local("x").as_fake()`.

229

230### Bytecode Debugger (`bytecode_debugger.py`)

231

232pdb-like debugger for stepping through Dynamo-generated bytecode. Useful for

233debugging segfaults (no Python traceback) and codegen errors.

234

235```python

236with torch._dynamo.bytecode_debugger.debug():

237 my_compiled_fn(x)

238```

239

240**Programmatic breakpoints** (no graph break): call

241`torch._dynamo.bytecode_debugger.breakpoint()` in user code, or

242`codegen.extend_output(create_breakpoint())` in codegen. Auto-activates

243without an explicit `debug()` wrapper.

244

245**Segfault debugging**: `v` (verbose) then `c` (continue) — every instruction

246is printed with `flush=True` before execution, so the last line before a crash

247is the culprit. On exceptions, the debugger stops at the faulting instruction

248automatically.

249

250## C++ Runtime (`torch/csrc/dynamo/`)

251

252The C/C++ layer implements the PEP 523 frame evaluation hook, the cache, and

253the guard evaluation tree. Performance-critical runtime on every Python frame.

254

255### Frame Evaluation

256

257**`eval_frame.c`** — Installs a custom frame evaluation function via

258`_PyInterpreterState_SetEvalFrameFunc`. A thread-local callback controls

259behavior: `None` (disabled), `Py_False` (run-only / cache lookup), or a

260callable (full Dynamo).

261

262**`eval_frame_cpp.cpp`** — `dynamo__custom_eval_frame` is called for every

263frame: gets `ExtraState` from the code object, builds a `FrameLocalsMapping`

264(O(1) access to locals without dict materialization), evaluates guards via

265`run_root_guard_manager()` across all `CacheEntry`s (LRU ordered). On cache

266hit, executes compiled code via a shadow frame (`dynamo_eval_custom_code_impl`

267copies `localsplus` into a new frame with the compiled code object). On miss,

268calls the Python callback to trigger compilation.

269

270### Cache

271

272**`extra_state.cpp/.h`** — `ExtraState` is attached per code object via

273`_PyCode_SetExtra`. Contains a `cache_entry_list` (LRU linked list),

274`frame_state` (dynamic shapes detection), and `FrameExecStrategy`.

275

276**`cache_entry.cpp/.h`** — Each `CacheEntry` stores a `RootGuardManager*`

277(raw C++ pointer for fast guard eval), the compiled code object, and the

278backend.

279

280### Guard Evaluation Tree (`guards.cpp`)

281

282Guards are organized as a C++ tree (~7800 lines) mirroring the data access

283pattern. `RootGuardManager` is the root, receiving a `FrameLocalsMapping`.

284Each `GuardManager` node has leaf guards and child accessors.

285

286**LeafGuard** subclasses: `TYPE_MATCH` (Py_TYPE pointer comparison),

287`ID_MATCH`, `EQUALS_MATCH`, `TENSOR_MATCH` (dtype/device/shape/strides/dispatch

288keys in C++), `DICT_VERSION`, `GLOBAL_STATE` (grad mode, autocast, etc.).

289

290**GuardAccessor** subclasses define tree edges: `FrameLocalsGuardAccessor`

291(O(1) index), `GetAttrGuardAccessor`, `DictGetItemGuardAccessor`,

292`GlobalsGuardAccessor`, etc.

293

294Key optimizations: fail-fast accessor reordering, dict version tag matching to

295skip subtrees, `FrameLocalsMapping` avoids dict construction,

296`check_nopybind()` avoids pybind11 overhead.

297

298### Other C++ files

299

300- `framelocals_mapping.cpp` — O(1) frame locals/cells/freevars access

301- `cpython_defs.c` — copied CPython internals for frame manipulation

302- `init.cpp` — `torch._C._dynamo` module and pybind11 bindings

303- `debug_macros.h` — `DEBUG_TRACE`, `NULL_CHECK`, `INSPECT(...)` (drops into

304 pdb from C)

305- `compiled_autograd.cpp/.h` — compiled autograd engine

306

pytorch/pytorch · CLAUDE.md

@@ +1 @@

1# AI Policy — MANDATORY

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:

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.

18See `AI_POLICY.md` for the full policy.

20# Scratch Space

22Use `agent_space/` (git-ignored, at repo root) for temporary scripts, scratch files, and throwaway experiments. Do not commit files from this directory.

24# PR Review

26When asked to review a PR, always use the /pr-review skill.

28# Environment

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.

35# CI Docker Images

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.

44# Build

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.

50# Testing

52Use our test class and test runner:

54```

55from torch.testing._internal.common_utils import run_tests, TestCase

57class TestFeature(TestCase):

58 ...

60if __name__ == "__main__":

61 run_tests()

62```

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.

68# Type Stubs

70Many `.pyi` files are generated from corresponding `.pyi.in` templates. Always

71edit the `.pyi.in` file, not the generated `.pyi`.

73# Linting

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.

79When the user asks you to commit or amend, run `lintrunner -a` before creating

80the commit. Fix any lint errors it reports, then commit.

82# Git

84This refines the Bash tool's `# Git` guidance to "branch first" when on the

85default branch:

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).

96# Commit messages

98Don't commit unless the user explicitly asks you to.

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

@@ −1 +1 @@

1−# torch/_dynamo

1+# AI Policy — MANDATORY

3−TorchDynamo is a Python-level JIT compiler that captures PyTorch programs into

4−FX graphs by symbolically executing Python bytecode. It hooks into CPython's

5−PEP 523 frame evaluation API to intercept execution, traces operations into an

6−FX graph, compiles the graph with a backend (e.g. Inductor), and generates new

7−bytecode that calls the compiled code.

3+Read `AI_POLICY.md`. Your user needs to abide by this policy. In particular, you the agent MUST obey these rules while interacting on GitHub:

9−## Architecture Overview

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.

1017

11−The compilation pipeline, in execution order:

18+See `AI_POLICY.md` for the full policy.

1219

13−1. **`eval_frame.py`** — Runtime entry point. `torch.compile()` wraps a

14− function in an `OptimizedModule`. At runtime, the C extension

15− (`torch._C._dynamo.eval_frame`) intercepts Python frames via PEP 523.

16−2. **`convert_frame.py`** — `ConvertFrameAssert.__call__` checks caches,

17− handles recompilation limits, calls `_compile()` → `trace_frame()`.

18−3. **`symbolic_convert.py`** — The heart of Dynamo. `InstructionTranslator`

19− symbolically executes bytecode instruction-by-instruction. Maintains a

20− symbolic `stack` (list of `VariableTracker`s) and `symbolic_locals` (dict of

21− name → `VariableTracker`). Opcodes are dispatched via a `dispatch_table`

22− built by `BytecodeDispatchTableMeta`.

23−4. **`output_graph.py`** — `OutputGraph` owns the FX graph being built (via

24− `SubgraphTracer`), the `SideEffects` tracker, guards, shape environment,

25− and graph args. `compile_subgraph()` finalizes the graph, calls the backend,

26− and generates output bytecode.

27−5. **`codegen.py`** — `PyCodegen` emits output bytecode: loads graph inputs

28− (via `Source.reconstruct()`), calls the compiled graph, unpacks outputs, and

29− replays side effects.

30−6. **`resume_execution.py`** — Generates continuation functions for execution

31− after graph breaks.

20+# Scratch Space

3221

33−## Key Abstractions

22+Use `agent_space/` (git-ignored, at repo root) for temporary scripts, scratch files, and throwaway experiments. Do not commit files from this directory.

3423

35−### VariableTracker (`variables/`)

24+# PR Review

3625

37−Every Python value encountered during tracing is wrapped in a `VariableTracker`

38−subclass. Key interface: `as_python_constant()`, `as_proxy()`,

39−`call_function()`, `call_method()`, `getattro_impl()`, `reconstruct()`.

26+When asked to review a PR, always use the /pr-review skill.

4027

41−Key fields: `source` (where the value came from, for guards) and

42−`mutation_type` (whether/how mutations are tracked).

28+# Environment

4329

44−**Factory**: `VariableTracker.build(tx, value, source=...)` dispatches to

45−`VariableBuilder` (sourced values needing guards) or `SourcelessBuilder`

46−(values created during tracing).

30+If any tool you're trying to use (pip, python, spin, etc) is missing, check for

31+a `.venv` directory in the project root or its parent directory. If found,

32+activate it and retry. If no `.venv` is found, stop and ask the user if an

33+environment is needed. Do NOT try to find alternatives or install these tools.

4734

48−Key subclass families in `variables/`: `TensorVariable` / `SymNodeVariable`

49−(tensor.py), `ConstantVariable` (constant.py), `ListVariable` /

50−`TupleVariable` (lists.py), `ConstDictVariable` (dicts.py), `SetVariable` (sets.py),

51−`UserFunctionVariable` (functions.py), `BuiltinVariable` (builtin.py),

52−`NNModuleVariable` (nn_module.py), `UserDefinedObjectVariable`

53−(user_defined.py), `TorchHigherOrderOperatorVariable` (higher_order_ops.py),

54−`LazyVariableTracker` (lazy.py). `VariableBuilder` and `SourcelessBuilder` are

55−in builder.py.

35+# CI Docker Images

5636

57−### Source (`source.py`)

37+The `.ci/docker/` directory is content-hashed to determine whether Docker images

38+need rebuilding. Any file change inside `.ci/docker/` (including the README)

39+changes the hash and triggers a full Docker image rebuild. Do not make changes

40+in this directory unless you intend to rebuild Docker images. When Docker builds

41+are broken (e.g., due to an upstream Ubuntu outage), avoid touching this

42+directory so you don't force a rebuild against the broken state.

5843

59−Tracks value provenance — how to access a value at runtime. Used for guard

60−generation (`source.make_guard(GuardBuilder.XXX)`) and bytecode reconstruction

61−(`source.reconstruct(codegen)`). Root sources: `LocalSource`, `GlobalSource`.

62−Chained sources: `AttrSource`, `GetItemSource`, `NNModuleSource`, etc.

44+# Build

6345

64−### Guards (`guards.py`)

46+Always 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.

47+All build (both codegen, C++ and python) is done via `pip install -e . -v --no-build-isolation`.

48+You should NEVER run any other command to build PyTorch.

6549

66−Runtime conditions that must hold for cached compiled code to be reused.

67−Install via `install_guard(source.make_guard(GuardBuilder.TYPE_MATCH))`.

68−Common types: `TYPE_MATCH`, `ID_MATCH`, `EQUALS_MATCH`, `TENSOR_MATCH`,

69−`SEQUENCE_LENGTH`. At finalization, `CheckFunctionManager` builds a tree of

70−C++ `GuardManager` objects for fast runtime checking.

50+# Testing

7151

72−### Side Effects (`side_effects.py`)

52+Use our test class and test runner:

7353

74−Tracks mutations during tracing (attribute stores, list mutations, cell

75−variable updates, tensor hooks) and replays them as bytecode after graph

76−execution. The `MutationType` system (`variables/base.py`) controls what

77−mutations are allowed: `ValueMutationNew/Existing`,

78−`AttributeMutationNew/Existing`, or `None` (immutable). The `scope` field

79−prevents cross-scope mutations inside higher-order operators.

54+```

55+from torch.testing._internal.common_utils import run_tests, TestCase

8056

81−### Other key files

57+class TestFeature(TestCase):

58+ ...

8259

83−- `trace_rules.py` — inline/skip/graph-break decisions per function

84−- `exc.py` — exception hierarchy: `Unsupported` (graph break), `RestartAnalysis`

85− (restart tracing), `ObservedException` (user exceptions during tracing),

86− `BackendCompilerFailed`

87−- `config.py` — configuration flags, supports `config.patch()` context

88− manager/decorator

89−- `bytecode_transformation.py` / `bytecode_analysis.py` — low-level bytecode

90− manipulation, liveness analysis

91−- `pgo.py` — profile-guided optimization for dynamic shapes

92−- `polyfills/` — traceable replacements for stdlib functions

93−- `repro/` — reproduction/minification tools

60+if __name__ == "__main__":

61+ run_tests()

62+```

9463

95−## Graph Breaks

64+To test Tensor equality, use assertEqual.

65+For tests over multiple inputs, use the `@parametrize` decorator.

66+For any test that checks numerics of the on-device implementation, use `instantiate_device_type_tests` to write device-generic tests.

9667

97−Call `unimplemented()` (from `exc.py`) to trigger a graph break:

68+# Type Stubs

9869

99−```python

100−from torch._dynamo.exc import unimplemented

101−from torch._dynamo import graph_break_hints

70+Many `.pyi` files are generated from corresponding `.pyi.in` templates. Always

71+edit the `.pyi.in` file, not the generated `.pyi`.

10272

103−unimplemented(

104− gb_type="short_category_name",

105− context=f"dynamic details: {value}",

106− explanation="Human-readable explanation of why this breaks the graph.",

107− hints=[*graph_break_hints.SUPPORTABLE],

108−)

109−```

73+# Linting

11074

111−- `gb_type`: Context-free category (no dynamic strings).

112−- `context`: Developer-facing details (can be dynamic).

113−- `explanation`: User-facing explanation (can be dynamic).

114−- `hints`: Use constants from `graph_break_hints.py`: `SUPPORTABLE`,

115− `FUNDAMENTAL`, `DIFFICULT`, `DYNAMO_BUG`, `USER_ERROR`,

116− `CAUSED_BY_EARLIER_GRAPH_BREAK`.

75+Only use commands provided via `spin` for linting.

76+Use `spin help` to list available commands.

77+Generally, use `spin lint` as to run the lint and `spin fixlint` to apply automatic fixes.

11778

118−The `break_graph_if_unsupported` decorator on instruction handlers catches

119−`Unsupported`, logs the graph break, updates the `SpeculationLog`, and restarts

120−analysis. On the second pass, the partial graph is compiled at the break point

121−and a resume function handles the rest.

79+When the user asks you to commit or amend, run `lintrunner -a` before creating

80+the commit. Fix any lint errors it reports, then commit.

12281

123−## Testing

82+# Git

12483

125−Tests live in `test/dynamo/`. Use `torch._dynamo.test_case.TestCase` as base

126−class — it calls `torch._dynamo.reset()` in setUp/tearDown and patches config

127−for strict error checking.

84+This refines the Bash tool's `# Git` guidance to "branch first" when on the

85+default branch:

12886

129−```bash

130−python test/dynamo/test_misc.py # whole file

131−python test/dynamo/test_misc.py MiscTests.test_foo # single test

132−python test/dynamo/test_misc.py -k test_foo # pattern match

133−```

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).

13495

135−### Common patterns

96+# Commit messages

13697

137−The default backend to `torch.compile()` is `backend="eager"`.

98+Don't commit unless the user explicitly asks you to.

13899

139−**CompileCounter** — count compilations and graph ops:

140−```python

141−cnt = torch._dynamo.testing.CompileCounter()

100+When writing a commit message, don't make a bullet list of the individual

101+changes. 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

103+entirely.

142104

143−@torch.compile(backend=cnt)

144−def fn(x):

145− return x + 1

105+The commit message should be clear, informative, and have a Test Plan section

106+that describes how you tested the change. If you are fixing a bug, the commit

107+message must explain the root cause of the bug and how the fix works.

108+If there were multiple potential paths you could have taken, please call them

109+out succinctly and justify the one you took.

146110

147−fn(torch.randn(10))

148−self.assertEqual(cnt.frame_count, 1)

149−self.assertEqual(cnt.op_count, 1)

150−```

111+When describing the testing strategy in a commit message, include the literal

112+commands that were run in fenced Markdown code blocks.

151113

152−**fullgraph=True** — assert no graph breaks:

153−```python

154−torch.compile(fn, backend="eager", fullgraph=True)(x)

155−```

114+Disclose that the PR was authored with an AI assistant.

156115

157−**EagerAndRecordGraphs** — inspect captured FX graphs:

158−```python

159−backend = torch._dynamo.testing.EagerAndRecordGraphs()

160−torch.compile(fn, backend=backend)(x)

161−graph = backend.graphs[0]

162−```

116+When the user asks you to amend a commit, check whether the commit message

117+still accurately describes the changes. If it doesn't and the commit is not a

118+ghstack commit, update the message. For ghstack commits, amending the message

119+is a no-op, so just remind the user to update the PR description if needed.

163120

164−**normalize_gm + assertExpectedInline** — snapshot test graph output:

165−```python

166−from torch._dynamo.testing import normalize_gm

167−self.assertExpectedInline(

168− normalize_gm(backend.graphs[0].print_readable(False)),

169− """\

170−expected output here

171−""",

172−)

173−```

121+If a commit message contains `ghstack-source-id` or `Pull-Request` trailers,

122+you MUST preserve them when rewriting or splitting commit messages. ghstack

123+will update the source id automatically when needed.

174124

175−Call `torch._dynamo.reset()` within a test when testing multiple compilation

176−scenarios in a single test method. The base class handles setUp/tearDown reset

177−automatically.

125+# ghstack Workflow

178126

179−## Debugging

127+ghstack commits follow a different workflow than the conventional GitHub branch

128+and PR workflow. First identify whether you're on a ghstack commit:

180129

181−### TORCH_LOGS

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).

182136

183−```bash

184−TORCH_LOGS="graph_breaks" python script.py # see graph breaks

185−TORCH_LOGS="guards,recompiles" python script.py # see guards and recompilation reasons

186−TORCH_LOGS="graph_code" python script.py # see captured FX graph code

187−TORCH_LOGS="+dynamo" python script.py # full debug logging

188−TORCH_LOGS="bytecode" python script.py # see bytecode transformations

189−```

137+Rules for working with ghstack:

190138

191−### Reproducing crashes

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.

192162

193−When writing a minimal repro for a Dynamo crash:

163+# Coding Style Guidelines

194164

195−1. **Capture `TORCH_LOGS="+dynamo"` for the known-failing case.** Look for

196− the frame ID (e.g. `[3/0_1]`), INLINING/FAILED/COMPILING/Restart events,

197− speculation behavior, and graph break reasons.

198−2. **Capture the same logs for your repro attempt** and diff against the

199− failing case. The divergence point tells you what condition you're missing.

200−3. **Match each condition from the crash traceback:**

201− - Does the code need a speculation checkpoint? Add tensor ops before the

202− failing code so `compile_subgraph` creates one.

203− - Does the function need to be inlined on retry (not skipped)? The graph

204− break type matters — `step_unsupported` keeps inlining on retry while

205− `unimplemented` may skip the function.

206− - Does the crash happen in a resume function? Add a `graph_break()` earlier

207− so PEP 523 compiles the resume as a fresh frame.

165+Follow these rules for all code changes in this repository:

208166

209−### Structured tracing (for production)

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.

210191

211−```bash

212−TORCH_TRACE=/path/to/dir python script.py # explicit trace directory

213−```

192+If uncertain, choose the simpler, more concise implementation.

214193

215−Analyze with `tlparse`.

194+# cuda.bindings Error Checking

216195

217−### Compile-time profiling

196+Use `torch.cuda._utils._check_cuda_bindings` to error-check `cuda.bindings`

197+runtime calls. Do not write inline error-checking helpers.

218198

219−`TORCH_COMPILE_DYNAMO_PROFILER=1` prints per-function cumtime/tottime

220−(cProfile-style) showing where Dynamo spends time during tracing. Set to a

221−file path instead to save a profile loadable by `snakeviz`.

199+# cuda.bindings Raw Handles

222200

223−### comptime.breakpoint() (`comptime.py`)

201+`cuda.bindings` runtime functions accept a raw handle passed as a Python `int`

202+directly as their handle argument. Whenever you already have an int handle --

203+from `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

205+construct a typed wrapper (`cudaGraph_t(init_value=...)`,

206+`cudaGraphExec_t(init_value=...)`, `cudaStream_t(init_value=...)`, etc.) just to

207+hand 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

210+typed object when you genuinely need it as a value in its own right.

224211

225−Drops into pdb during **compilation** to inspect Dynamo state. Call

226−`comptime.breakpoint()` in user code; in the pdb session use `ctx`

227−(`ComptimeContext`) to call `print_locals()`, `print_bt()`, `print_graph()`,

228−or `get_local("x").as_fake()`.

212+# Dynamo Config

229213

230−### Bytecode Debugger (`bytecode_debugger.py`)

214+Use `torch._dynamo.config.patch` for temporarily changing config. It can be used as a decorator on test methods or as a context manager:

231215

232−pdb-like debugger for stepping through Dynamo-generated bytecode. Useful for

233−debugging segfaults (no Python traceback) and codegen errors.

234−

235216 ```python

236−with torch._dynamo.bytecode_debugger.debug():

237− my_compiled_fn(x)

217+# Good - use patch as decorator on test method

218+@torch._dynamo.config.patch(force_compile_during_fx_trace=True)

219+def test_my_feature(self):

220+ # test code here

221+ pass

222+

223+# Good - use patch as context manager

224+with torch._dynamo.config.patch(force_compile_during_fx_trace=True):

225+ # test code here

226+ pass

227+

228+# Bad - manual save/restore

229+orig = torch._dynamo.config.force_compile_during_fx_trace

230+try:

231+ torch._dynamo.config.force_compile_during_fx_trace = True

232+ # test code here

233+finally:

234+ torch._dynamo.config.force_compile_during_fx_trace = orig

238235 ```

239236

240−**Programmatic breakpoints** (no graph break): call

241−`torch._dynamo.bytecode_debugger.breakpoint()` in user code, or

242−`codegen.extend_output(create_breakpoint())` in codegen. Auto-activates

243−without an explicit `debug()` wrapper.

237+# Fixing B950 line too long in multi-line string blocks

244238

245−**Segfault debugging**: `v` (verbose) then `c` (continue) — every instruction

246−is printed with `flush=True` before execution, so the last line before a crash

247−is the culprit. On exceptions, the debugger stops at the faulting instruction

248−automatically.

239+If B950 line too long triggers on a multi-line string block, you cannot fix it by

240+putting # noqa: B950 on that line directly, as that would change the meaning of the

241+string, nor can you fix it by line breaking the string (since you need the string

242+to stay the same). Instead, put # noqa: B950 on the same line as the terminating

243+triple quote.

249244

250−## C++ Runtime (`torch/csrc/dynamo/`)

245+Example:

251246

252−The C/C++ layer implements the PEP 523 frame evaluation hook, the cache, and

253−the guard evaluation tree. Performance-critical runtime on every Python frame.

247+```

248+ self.assertExpectedInline(

249+ foo(),

250+ """

251+this line is too long...

252+""", # noqa: B950

253+ )

254+```

254255

255−### Frame Evaluation

256+# Logging and Structured Tracing

256257

257−**`eval_frame.c`** — Installs a custom frame evaluation function via

258−`_PyInterpreterState_SetEvalFrameFunc`. A thread-local callback controls

259−behavior: `None` (disabled), `Py_False` (run-only / cache lookup), or a

260−callable (full Dynamo).

258+When adding debug logging for errors or diagnostic info, consider two user personas:

261259

262−**`eval_frame_cpp.cpp`** — `dynamo__custom_eval_frame` is called for every

263−frame: gets `ExtraState` from the code object, builds a `FrameLocalsMapping`

264−(O(1) access to locals without dict materialization), evaluates guards via

265−`run_root_guard_manager()` across all `CacheEntry`s (LRU ordered). On cache

266−hit, executes compiled code via a shadow frame (`dynamo_eval_custom_code_impl`

267−copies `localsplus` into a new frame with the compiled code object). On miss,

268−calls the Python callback to trigger compilation.

260+1. **Local development**: Users run locally and can access files on disk

261+2. **Production jobs**: Users can only access logs via `tlparse` from structured traces

269262

270−### Cache

263+For production debugging, use `trace_structured` to log artifacts:

271264

272−**`extra_state.cpp/.h`** — `ExtraState` is attached per code object via

273−`_PyCode_SetExtra`. Contains a `cache_entry_list` (LRU linked list),

274−`frame_state` (dynamic shapes detection), and `FrameExecStrategy`.

265+```python

266+from torch._logging import trace_structured

275267

276−**`cache_entry.cpp/.h`** — Each `CacheEntry` stores a `RootGuardManager*`

277−(raw C++ pointer for fast guard eval), the compiled code object, and the

278−backend.

268+# Log an artifact (graph, edge list, etc.)

269+trace_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+```

279278

280−### Guard Evaluation Tree (`guards.cpp`)

279+To check if structured tracing is enabled (for conditional messaging):

281280

282−Guards are organized as a C++ tree (~7800 lines) mirroring the data access

283−pattern. `RootGuardManager` is the root, receiving a `FrameLocalsMapping`.

284−Each `GuardManager` node has leaf guards and child accessors.

281+```python

282+from torch._logging._internal import trace_log

285283

286−**LeafGuard** subclasses: `TYPE_MATCH` (Py_TYPE pointer comparison),

287−`ID_MATCH`, `EQUALS_MATCH`, `TENSOR_MATCH` (dtype/device/shape/strides/dispatch

288−keys in C++), `DICT_VERSION`, `GLOBAL_STATE` (grad mode, autocast, etc.).

284+if trace_log.handlers:

285+ # Structured tracing is enabled, suggest tlparse in error messages

286+ msg += "[Use tlparse to extract debug artifacts]"

287+```

289288

290−**GuardAccessor** subclasses define tree edges: `FrameLocalsGuardAccessor`

291−(O(1) index), `GetAttrGuardAccessor`, `DictGetItemGuardAccessor`,

292−`GlobalsGuardAccessor`, etc.

289+**Best practices for error diagnostics:**

293290

294−Key optimizations: fail-fast accessor reordering, dict version tag matching to

295−skip subtrees, `FrameLocalsMapping` avoids dict construction,

296−`check_nopybind()` avoids pybind11 overhead.

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

297298

298−### Other C++ files

299+# cuda::ptx

299300

300−- `framelocals_mapping.cpp` — O(1) frame locals/cells/freevars access

301−- `cpython_defs.c` — copied CPython internals for frame manipulation

302−- `init.cpp` — `torch._C._dynamo` module and pybind11 bindings

303−- `debug_macros.h` — `DEBUG_TRACE`, `NULL_CHECK`, `INSPECT(...)` (drops into

304− pdb from C)

305−- `compiled_autograd.cpp/.h` — compiled autograd engine

301+When 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.

306320