Skip to content

[graph_trainer] Add CUDA graph kernel annotation pass#2926

Merged
yushangdi merged 1 commit into
mainfrom
sy_cudagraph_annotation
Apr 21, 2026
Merged

[graph_trainer] Add CUDA graph kernel annotation pass#2926
yushangdi merged 1 commit into
mainfrom
sy_cudagraph_annotation

Conversation

@yushangdi

@yushangdi yushangdi commented Apr 10, 2026

Copy link
Copy Markdown
Contributor

Need pytorch/pytorch#179867 and cuda 13.1 driver/compat

[graph_trainer] Add CUDA graph kernel annotation pass

Summary

Adds a compiler pass that automatically labels CUDA graph kernels with their originating nn.Module path in profiler traces. When enabled, Perfetto/chrome traces show which component each kernel belongs to (e.g., layers.0.attention.wq, layers.0.feed_forward.w1, norm). Both forward and backward kernels are annotated.

No model code changes required — module paths are derived from the existing nn.Module hierarchy.

How it works

  1. annotate_module_fqns(model) walks the nn.Module tree and wraps each submodule's forward with torch.fx.traceback.annotate_fn({"module_fqn": fqn}). This sets node.meta["custom"]["module_fqn"] on all FX nodes during tracing — surviving dynamo, make_fx, and export tracers. Backward nodes receive the metadata via _copy_fwd_metadata_to_bw_nodes.

  2. insert_kernel_annotations_pass reads the module_fqn metadata from FX nodes and inserts mark_kernels() enter/exit calls at module boundaries in the graph. Bundled with cudagraph_pass — runs automatically whenever cudagraph is used. No-op when cuda-python or cuda-compat >= 13.1 is unavailable.

  3. CUDAGraphWrapper captures the annotations during CUDA graph recording when the pass is active, and registers a trace post-processor with the core profiler.

  4. Profiler integrationprofiling.py exposes register_trace_post_processor() so experiment code can hook into trace export without core importing from experiments. The annotation post-processor merges kernel annotations into the trace automatically.

Usage

The annotation pass runs automatically in aot_fx_trace mode when cudagraph is compatible. Profiler traces come out pre-annotated:

NGPU=1 MODULE=graph_trainer.llama3 CONFIG=graph_trainer_llama3_debugmodel ./run_train.sh \
    --compile.mode aot_fx_trace \
    --profiling.enable_profiling \
    --profiling.profile_freq 4

Open the exported trace at outputs/.../rank0_trace.json in https://ui.perfetto.dev to see module labels on each kernel.

Example output (Llama3 debug model, 6 layers)

tok_embeddings                      → embedding gather
layers.0.attention_norm             → RMSNorm
layers.0.attention.wq/.wk/.wv      → Q/K/V projections
layers.0.attention                  → RoPE, reshapes
layers.0.attention.inner_attention  → flash_fwd_kernel
layers.0.attention.wo               → output projection
layers.0.ffn_norm                   → RMSNorm
layers.0.feed_forward.w1/.w3       → gate/up projections
layers.0.feed_forward               → silu + mul
layers.0.feed_forward.w2            → down projection
norm                                → final RMSNorm
output                              → output linear

Requirements

Requires PyTorch with torch.cuda._graph_annotations module and enable_annotations kwarg on torch.cuda.graph(). Falls back gracefully when unavailable (pass is a no-op).

Files changed

File Change
common_utils.py Add annotate_module_fqns() and _MODULE_FQN constant
passes.py Add insert_kernel_annotations_pass (bundled with cudagraph pass), with alternative approaches documented
cudagraph.py Add enable_annotations flag, enable_cudagraph_annotations(), get_cudagraph_annotations(), trace post-processor registration
llama3/parallelize.py Call annotate_module_fqns(model) in annotate_llama()
deepseek_v3/parallelize.py Call annotate_module_fqns(model) in annotate_deepseekv3()
qwen3/parallelize.py Call annotate_module_fqns(model) in annotate_qwen3()
tools/profiling.py Add register_trace_post_processor() callback API (no experiment imports)
models/common/attention.py Remove unused import os + DISABLE_LLVM_OPT
tests/test_passes.py 5 unit tests: transformer-like model with trace_train_step, pass insertion, no-op when unavailable, no-op without metadata, xfail for same-class instances
tests/profiler_tests.py E2E test: trace_train_step fwd+bwd → pass → cudagraph → profile → verify module_fqn in trace for both fwd and bwd kernels
.github/workflows/integration_test_8gpu_graph_trainer.yaml Add annotation tests to CI

Test plan

# Unit tests (single GPU):
python -m torchtitan.experiments.graph_trainer.tests.test_passes TestAnnotateModuleFqns -v

# E2E profiler test (single GPU, requires cuda-compat >= 13.1):
LD_LIBRARY_PATH=$CONDA_PREFIX/cuda-compat:$LD_LIBRARY_PATH \
    python -m torchtitan.experiments.graph_trainer.tests.profiler_tests -v

# Full training run (single GPU):
NGPU=1 LD_LIBRARY_PATH=$CONDA_PREFIX/cuda-compat:$LD_LIBRARY_PATH \
    torchrun --nproc_per_node=1 --rdzv_backend c10d --rdzv_endpoint="localhost:0" \
    -m torchtitan.train --module graph_trainer.llama3 --config graph_trainer_llama3_debugmodel \
    --compile.mode aot_fx_trace \
    --training.steps 4 --profiler.enable_profiling --profiler.profile_freq 4

Authored with Claude.

Example:

NGPU=4 LD_LIBRARY_PATH=/home/shangdiy/.conda/envs/pytorch-3.12/cuda-compat:$LD_LIBRARY_PATH \
  torchrun --nproc_per_node=1 --rdzv_backend c10d --rdzv_endpoint="localhost:0" \
      -m torchtitan.train \
      --module graph_trainer.llama3 \
      --config graph_trainer_llama3_debugmodel \
      --compile.mode aot_fx_trace \
      --training.steps 4 \
      --profiler.enable_profiling \
      --profiler.save_traces_folder agent_space/traces \
      --profiler.profile_freq 4 

Result:

https://www.internalfb.com/intern/perfetto/open_trace/?manifold_path=perfetto_internal_traces%2Ftree%2Fshared_trace%2Fshangdiy_22d012fd-21e5-412f-aa06-d17a32426030_rank0_trace.json

Screenshot 2026-04-14 at 1 48 16 PM

@meta-cla meta-cla Bot added the CLA Signed This label is managed by the Meta Open Source bot. label Apr 10, 2026
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch 3 times, most recently from 5ed6f0d to 7cdf3e2 Compare April 14, 2026 20:49
@yushangdi yushangdi changed the title add cudagraph annotation Add cudagraph kernel annotation Apr 14, 2026
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch from 7cdf3e2 to 8faf83e Compare April 14, 2026 20:57
@yushangdi yushangdi changed the title Add cudagraph kernel annotation [graph_trainer] Add CUDA graph kernel annotation pass Apr 14, 2026
@yushangdi yushangdi marked this pull request as ready for review April 14, 2026 20:59
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch from 8faf83e to b28d5ca Compare April 14, 2026 21:07
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch from b28d5ca to 92dc19d Compare April 14, 2026 21:18

@yiming0416 yiming0416 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What happens when cuda<13.1? Would this be a no-op?

FWIW currently torchtitan's CI runs with pytorch cuda-13.0

@yushangdi

Copy link
Copy Markdown
Contributor Author

What happens when cuda<13.1? Would this be a no-op?

FWIW currently torchtitan's CI runs with pytorch cuda-13.0

yeah, this will be a no-op

@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch 2 times, most recently from 3b5c4f5 to 040f22d Compare April 14, 2026 22:46
Comment thread .gitignore Outdated
Comment thread torchtitan/experiments/graph_trainer/passes.py Outdated
]

if pass_names and "insert_kernel_annotations" in pass_names:
passes.append(insert_kernel_annotations_pass)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

add comment on it's ordering requirement, does this need to run before cudagraph?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

added. it needs to run before cudagraph pass

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

✅ commit 9ddd440bd — CRITICAL 진짜 버그 확정. store.upsertCrawlPayload로 source_url/source_type/source_contents 3개 컬럼만 명시적으로 INSERT INTO ... ON DUPLICATE KEY UPDATE 처리. 기존 row의 extracted_result_json/inspection_status 보존됨.

Comment thread torchtitan/experiments/graph_trainer/passes.py
Comment thread torchtitan/experiments/graph_trainer/passes.py Outdated
Comment on lines +292 to +293
Requires ``cuda-python`` package and CUDA toolkit/driver >= 13.1
(or cuda-compat >= 13.1). Returns the graph unchanged when unavailable.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

what happens if the pass is run on env doesn't meet the requirement ?
can it be a no-op pass in such case?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yeah, it will be a no-op pass

Comment thread torchtitan/experiments/graph_trainer/tests/test_profiler.py
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch from 77eeaf6 to 10d8db5 Compare April 15, 2026 20:46
Comment thread torchtitan/tools/profiling.py Outdated
Comment thread torchtitan/experiments/graph_trainer/tests/test_passes.py Outdated
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch 5 times, most recently from af19dbb to 49777a1 Compare April 20, 2026 16:46
@yushangdi

yushangdi commented Apr 20, 2026

Copy link
Copy Markdown
Contributor Author

Seeing some test failures like below on 8 GPU Model test. Likely OOM-style CUBLAS allocation failure, not caused by this PR. They're also present on a PR that only changes .md files: #3026

2026-04-20T02:59:10.1706429Z [rank6]:  triton_flex_attention_backward_28 5.1978 ms 85.4% BLOCKS_ARE_CONTIGUOUS=False, BLOCK_M1=32, BLOCK_M2=32, BLOCK_N1=32, BLOCK_N2=32, FLOAT32_PRECISION="'ieee'", GQA_SHARED_HEADS=2, HAS_FULL_BLOCKS=True, IS_DIVISIBLE=False, OUTPUT_LOGSUMEXP=True, OUTPUT_MAX=False, PRESCALE_QK=False, QK_HEAD_DIM=64, QK_HEAD_DIM_ROUNDED=64, ROWS_GUARANTEED_SAFE=False, SAFE_HEAD_DIM=True, SM_SCALE=0.125, SPARSE_KV_BLOCK_SIZE=128, SPARSE_Q_BLOCK_SIZE=128, USE_TMA=False, V_HEAD_DIM=64, V_HEAD_DIM_ROUNDED=64, WRITE_DQ=True, num_stages=5, num_warps=4
2026-04-20T02:59:10.1708181Z [rank6]:  triton_flex_attention_backward_27 5.2163 ms 85.1% BLOCKS_ARE_CONTIGUOUS=False, BLOCK_M1=32, BLOCK_M2=32, BLOCK_N1=32, BLOCK_N2=32, FLOAT32_PRECISION="'ieee'", GQA_SHARED_HEADS=2, HAS_FULL_BLOCKS=True, IS_DIVISIBLE=False, OUTPUT_LOGSUMEXP=True, OUTPUT_MAX=False, PRESCALE_QK=False, QK_HEAD_DIM=64, QK_HEAD_DIM_ROUNDED=64, ROWS_GUARANTEED_SAFE=False, SAFE_HEAD_DIM=True, SM_SCALE=0.125, SPARSE_KV_BLOCK_SIZE=128, SPARSE_Q_BLOCK_SIZE=128, USE_TMA=False, V_HEAD_DIM=64, V_HEAD_DIM_ROUNDED=64, WRITE_DQ=True, num_stages=4, num_warps=4
2026-04-20T02:59:10.1708638Z [rank6]:SingleProcess AUTOTUNE benchmarking takes 34.3047 seconds and 4.7301 seconds precompiling for 13 choices
2026-04-20T02:59:10.1708775Z [rank7]:[rank7]: Traceback (most recent call last):
2026-04-20T02:59:10.1709022Z [rank7]:[rank7]:   File "<frozen runpy>", line 198, in _run_module_as_main
2026-04-20T02:59:10.1709204Z [rank7]:[rank7]:   File "<frozen runpy>", line 88, in _run_code
2026-04-20T02:59:10.1709446Z [rank7]:[rank7]:   File "/pytorch/torchtitan/torchtitan/train.py", line 68, in <module>
2026-04-20T02:59:10.1709546Z [rank7]:[rank7]:     main()
2026-04-20T02:59:10.1709783Z [rank7]:[rank7]:   File "/pytorch/torchtitan/torchtitan/train.py", line 55, in main
2026-04-20T02:59:10.1709895Z [rank7]:[rank7]:     trainer.train()
2026-04-20T02:59:10.1710452Z [rank7]:[rank7]:   File "/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/distributed/elastic/multiprocessing/errors/__init__.py", line 367, in wrapper
2026-04-20T02:59:10.1710578Z [rank7]:[rank7]:     return f(*args, **kwargs)
2026-04-20T02:59:10.1710690Z [rank7]:[rank7]:            ^^^^^^^^^^^^^^^^^^
2026-04-20T02:59:10.1710927Z [rank7]:[rank7]:   File "/pytorch/torchtitan/torchtitan/trainer.py", line 834, in train
2026-04-20T02:59:10.1711055Z [rank7]:[rank7]:     self.train_step(data_iterator)
2026-04-20T02:59:10.1711321Z [rank7]:[rank7]:   File "/pytorch/torchtitan/torchtitan/trainer.py", line 752, in train_step
2026-04-20T02:59:10.1711470Z [rank7]:[rank7]:     loss = self.forward_backward_step(
2026-04-20T02:59:10.1711592Z [rank7]:[rank7]:            ^^^^^^^^^^^^^^^^^^^^^^^^^^^
2026-04-20T02:59:10.1711951Z [rank7]:[rank7]:   File "/pytorch/torchtitan/torchtitan/trainer.py", line 710, in forward_backward_step
2026-04-20T02:59:10.1712057Z [rank7]:[rank7]:     loss.backward()
2026-04-20T02:59:10.1712605Z [rank7]:[rank7]:   File "/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/_tensor.py", line 631, in backward
2026-04-20T02:59:10.1712754Z [rank7]:[rank7]:     torch.autograd.backward(
2026-04-20T02:59:10.1713146Z [rank7]:[rank7]:   File "/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/autograd/__init__.py", line 379, in backward
2026-04-20T02:59:10.1713253Z [rank7]:[rank7]:     _engine_run_backward(
2026-04-20T02:59:10.1713690Z [rank7]:[rank7]:   File "/opt/conda/envs/py_3.12/lib/python3.12/site-packages/torch/autograd/graph.py", line 882, in _engine_run_backward
2026-04-20T02:59:10.1714058Z [rank7]:[rank7]:     return Variable._execution_engine.run_backward(  # Calls into the C++ engine to run the backward pass
2026-04-20T02:59:10.1714269Z [rank7]:[rank7]:            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
2026-04-20T02:59:10.1714597Z [rank7]:[rank7]: RuntimeError: CUDA error: CUBLAS_STATUS_ALLOC_FAILED when calling `cublasCreate(handle)

@tianyu-l tianyu-l left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

sgtm

@felipemello1 related to your observability work

@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch 3 times, most recently from dac49c5 to 7fb2787 Compare April 20, 2026 21:17
if is_cudagraph_compatible(traced_result.gm):
if not precompiled:
pre_passes = compile_time_passes(traced_result)
if cudagraph_compatible:

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

hmm... does this have to be a compile time pass?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

technically no, but it needs to happen before the custom_codegen_pass, and custom_codegen_pass is in compile_time_passes

Comment thread torchtitan/experiments/graph_trainer/deepseek_v3/parallelize.py Outdated
Comment thread torchtitan/tools/profiler.py Outdated
save_memory_snapshot_folder: str = "memory_snapshot"
"""Memory snapshot files location."""

trace_post_processors: list[Function.Config] = field(default_factory=list)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

do you expect this to actually grow into a list of post processors?

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

also, I am wondering if we actually need to introduce this callback?

I think ppl would always want processed trace file, as long as enable_cudagraph_annotations() is true

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

not sure, but could be possible?

for example, if we want custom_cdegen_pass to attach more info to kernels other than just file name and line number, we could tag each region with an id, and then post-process the trace and attach full stack traces/node names/etc.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ok, but i don't think it need to be a list of processor, hook with single entry seems better.
if you need to do multiple thing in the callback, just do it inside the hook

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

makes sense. updated to a single entry now.

@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch from 7fb2787 to 306411e Compare April 20, 2026 23:28
Comment thread torchtitan/experiments/graph_trainer/passes.py Outdated
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch 3 times, most recently from dbea916 to 4d2b30e Compare April 20, 2026 23:56
Adds `insert_kernel_annotations_pass` that labels CUDA graph kernels with
their originating nn.Module path in profiler traces.

How it works:
1. `annotate_module_fqns(model)` wraps each submodule's forward with
   `torch.fx.traceback.annotate_fn({"module_fqn": fqn})`, setting
   `node.meta["custom"]["module_fqn"]` on FX nodes during tracing.
2. `insert_kernel_annotations_pass` reads the `module_fqn` metadata and
   inserts `mark_kernels` enter/exit calls at module boundaries.  Bundled
   with the cudagraph pass.
3. `CUDAGraphWrapper` captures annotations during graph recording when the
   pass is active.
4. `Profiler.Config.trace_post_processors` (new field) runs
   `cudagraph_annotate_trace_post_processor()` on each exported trace so
   profiler traces automatically carry `module_fqn` fields on graphed
   kernel events.

Requires PyTorch with `torch.cuda._graph_annotations` and the
`enable_annotations` kwarg on `torch.cuda.graph()`.  Falls back gracefully
when unavailable (pass is a no-op).

Authored with Claude.
@yushangdi yushangdi force-pushed the sy_cudagraph_annotation branch from 4d2b30e to 56ceea7 Compare April 21, 2026 00:09
@yushangdi yushangdi merged commit fc54b89 into main Apr 21, 2026
18 of 20 checks passed
tianyu-l pushed a commit that referenced this pull request Apr 21, 2026
## Summary

Fixes a crash introduced by #2926 where `trace_post_processor:
Function.Config | None` in `Profiler.Config` causes `tyro.cli()` to fail
because tyro cannot parse `Callable[..., Any]` types from CLI arguments.

**Error:**
```
Invalid input to tyro.cli()
Unsupported type annotation for field profiler.trace-post-processor.fn
with type collections.abc.Callable[..., typing.Any]
```

**Fix:** Annotate `trace_post_processor` with `tyro.conf.Suppress`,
matching the existing pattern used for:
- `model_spec` in `trainer.py:62`
- `sample_processor` in `text_datasets.py:502`

The field defaults to `None` and is only set programmatically by
`CUDAGraphWrapper`, never via CLI.

## Test plan

- [ ] Verify `python -m torchtitan.train ...` no longer crashes at
config parsing
- [ ] Verify graph_trainer still works (trace_post_processor set
programmatically)
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ciflow/8gpu CLA Signed This label is managed by the Meta Open Source bot.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants