Skip to content

Commit 1829ecf

Browse files
committed
Improve export_autograd_triton generated readability
1 parent 666c1b1 commit 1829ecf

7 files changed

Lines changed: 755 additions & 43 deletions

File tree

test/test_export_autograd_triton.py

Lines changed: 245 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,15 @@
1616
CleanTritonUnsupportedError,
1717
clean_triton_source,
1818
)
19+
from transformer_nuggets.export_autograd_triton.codegen import (
20+
generate_autograd_source,
21+
write_autograd_source,
22+
)
1923
from transformer_nuggets.export_autograd_triton.guards import tensor_guard_for
24+
from transformer_nuggets.export_autograd_triton.specs import (
25+
CapturedSpecialization,
26+
TensorGuardSpec,
27+
)
2028

2129
os.environ.setdefault("TORCHINDUCTOR_COMPILE_THREADS", "1")
2230

@@ -153,6 +161,37 @@ def _clone_args(args):
153161
return tuple(_clone_tensor(arg) if isinstance(arg, torch.Tensor) else arg for arg in args)
154162

155163

164+
def _fake_captured_specialization(
165+
forward_source="def call(args):\n return ()\n",
166+
name="spec_0",
167+
):
168+
return CapturedSpecialization(
169+
name=name,
170+
runtime_tensor_names=("x",),
171+
static_args=(),
172+
tensor_guards=(
173+
TensorGuardSpec(
174+
name="x",
175+
rank=2,
176+
shape=(4, 8),
177+
stride=(8, 1),
178+
dtype="torch.float32",
179+
device_type="cuda",
180+
device_index=0,
181+
),
182+
),
183+
forward_source=forward_source,
184+
backward_source=None,
185+
num_user_outputs=1,
186+
output_kind="single",
187+
needs_autograd=False,
188+
dynamic=False,
189+
differentiable_output_mask=(True,),
190+
forward_residual_names=(),
191+
backward_saved_input_names=(),
192+
)
193+
194+
156195
def _flatten_tensors(value):
157196
if isinstance(value, torch.Tensor):
158197
return [value]
@@ -228,6 +267,121 @@ def test_unbounded_dynamic_dim_guard_uses_none_for_symbolic_max():
228267
assert guard.shape[0] == {"symbol": "batch", "min": 0, "max": None}
229268

230269

270+
def test_generated_wrapper_source_is_labeled_and_splits_specs():
271+
source = generate_autograd_source(
272+
trig_pointwise,
273+
"trig_pointwise_compiled",
274+
[_fake_captured_specialization()],
275+
)
276+
277+
assert source.startswith('"""Generated by transformer_nuggets.export_autograd_triton.')
278+
assert "# Generated Triton artifact index. Paths are relative to this wrapper file." in source
279+
assert "# spec_0 forward: <this_module>_artifacts/spec_0_forward.py" in source
280+
assert "# spec_0 backward: <this_module>_artifacts/spec_0_backward.py" not in source
281+
assert (
282+
"# Calls the selected *_forward.py artifact listed in the artifact index above." in source
283+
)
284+
assert (
285+
"# Calls the selected *_backward.py artifact listed in the artifact index above." in source
286+
)
287+
assert "\n\n\n_SPEC_0" not in source
288+
assert "_SPEC_0 = {" in source
289+
assert "_SPECS = [\n _SPEC_0," in source
290+
assert "class _TrigPointwiseCompiledAutogradFunction(torch.autograd.Function):" in source
291+
assert "return _RUNTIME.autograd_forward(ctx, spec_id, runtime_tensors)" in source
292+
assert "return _RUNTIME.autograd_backward(ctx, *grad_outputs)" in source
293+
assert (
294+
"result = _TrigPointwiseCompiledAutogradFunction.apply(spec_id, *runtime_tensors)"
295+
in source
296+
)
297+
assert "return _RUNTIME.restore_output_container(spec_id, result)" in source
298+
assert "Run the exported autograd/Triton specialization" in source
299+
300+
301+
def test_write_autograd_source_labels_artifacts_and_removes_stale_files(tmp_path):
302+
generated_path = tmp_path / "generated_fake.py"
303+
artifact_dir = tmp_path / "generated_fake_artifacts"
304+
artifact_dir.mkdir()
305+
(artifact_dir / "stale.py").write_text("stale")
306+
307+
write_autograd_source(
308+
generated_path,
309+
"wrapper source",
310+
[_fake_captured_specialization()],
311+
source_backend="inductor",
312+
)
313+
314+
artifact_path = artifact_dir / "spec_0_forward.py"
315+
artifact_source = artifact_path.read_text()
316+
assert not (artifact_dir / "stale.py").exists()
317+
assert artifact_source.startswith('"""Generated Triton artifact.')
318+
assert "Runtime tensor order:\n - x" in artifact_source
319+
assert "Tensor guards:\n - x: shape=(4, 8), stride=(8, 1)" in artifact_source
320+
assert "The runtime imports this file and calls call(args)." in artifact_source
321+
322+
323+
def test_generated_wrapper_uses_output_path_in_artifact_index(tmp_path):
324+
generated_path = tmp_path / "generated_fake.py"
325+
326+
source = generate_autograd_source(
327+
trig_pointwise,
328+
"trig_pointwise_compiled",
329+
[_fake_captured_specialization()],
330+
artifact_dir_name=f"{generated_path.stem}_artifacts",
331+
)
332+
333+
assert "# spec_0 forward: generated_fake_artifacts/spec_0_forward.py" in source
334+
335+
336+
def test_write_clean_triton_artifact_strips_benchmark_tail_without_async_compile(tmp_path):
337+
generated_path = tmp_path / "generated_fake.py"
338+
write_autograd_source(
339+
generated_path,
340+
"wrapper source",
341+
[
342+
_fake_captured_specialization(
343+
forward_source="""def call(args):
344+
return args
345+
346+
347+
def get_args():
348+
return []
349+
350+
351+
def benchmark_compiled_module(args, times=10, repeat=10):
352+
return None
353+
354+
355+
if __name__ == "__main__":
356+
print(get_args())
357+
"""
358+
)
359+
],
360+
source_backend="clean_triton",
361+
)
362+
363+
artifact_source = (tmp_path / "generated_fake_artifacts" / "spec_0_forward.py").read_text()
364+
assert "def call(args):" in artifact_source
365+
assert "def get_args():" not in artifact_source
366+
assert "benchmark_compiled_module" not in artifact_source
367+
assert "__main__" not in artifact_source
368+
369+
370+
def test_artifact_header_escapes_docstring_metadata(tmp_path):
371+
generated_path = tmp_path / "generated_fake.py"
372+
write_autograd_source(
373+
generated_path,
374+
"wrapper source",
375+
[_fake_captured_specialization(name='bad"""name\nline')],
376+
source_backend="inductor",
377+
)
378+
379+
artifact_path = tmp_path / "generated_fake_artifacts" / "spec_0_bad_name_line_forward.py"
380+
artifact_source = artifact_path.read_text()
381+
compile(artifact_source, str(artifact_path), "exec")
382+
assert 'Specialization: bad\\"\\"\\"name\\nline' in artifact_source
383+
384+
231385
def test_export_static_specialization_forward_backward_signature_and_errors(tmp_path):
232386
_requires_export_runtime()
233387
x = torch.randn(4, 8, device="cuda", requires_grad=True)
@@ -659,11 +813,90 @@ def call(args):
659813
assert result.patched_launches == ["triton_poi_fused_sin_0", "triton_poi_fused_cos_1"]
660814
assert "triton.cdiv(triton_poi_fused_sin_0_xnumel, 32)" in result.source
661815
assert "triton.cdiv(triton_poi_fused_cos_1_xnumel, 64)" in result.source
662-
assert "buf0, triton_poi_fused_sin_0_xnumel, XBLOCK=32" in result.source
663-
assert "buf1, triton_poi_fused_cos_1_xnumel, XBLOCK=64" in result.source
816+
assert "buf0,\n triton_poi_fused_sin_0_xnumel,\n XBLOCK=32" in result.source
817+
assert "buf1,\n triton_poi_fused_cos_1_xnumel,\n XBLOCK=64" in result.source
664818
assert "[(1, 1, 1)]" not in result.source
665819

666820

821+
def test_clean_triton_source_removes_unused_imports_assignments_and_blank_runs():
822+
source = """
823+
import math
824+
import os
825+
import torch
826+
import torch
827+
from pathlib import Path
828+
from torch._C import unused_stream, _cuda_getCurrentRawStream as get_raw_stream
829+
830+
unused_alias = torch.ops.aten
831+
used_alias = torch.ops.aten
832+
833+
834+
835+
836+
def call(args):
837+
stream = get_raw_stream(0)
838+
return torch.empty(1), used_alias, stream
839+
"""
840+
841+
result = clean_triton_source(source, dynamic=False)
842+
843+
assert "import math" not in result.source
844+
assert "import os" not in result.source
845+
assert result.source.count("import torch") == 1
846+
assert "from pathlib import Path" not in result.source
847+
assert "unused_alias" not in result.source
848+
assert "used_alias = torch.ops.aten" in result.source
849+
assert "unused_stream" not in result.source
850+
assert "get_raw_stream" in result.source
851+
assert "\n\n\n" not in result.source
852+
853+
854+
def test_clean_triton_source_strips_standalone_benchmark_tail():
855+
source = """
856+
def call(args):
857+
return args
858+
859+
860+
def get_args():
861+
return []
862+
863+
864+
def benchmark_compiled_module(args, times=10, repeat=10):
865+
return None
866+
867+
868+
if __name__ == "__main__":
869+
print(get_args())
870+
"""
871+
872+
result = clean_triton_source(source, dynamic=False)
873+
874+
assert "def call(args):" in result.source
875+
assert "def get_args():" not in result.source
876+
assert "benchmark_compiled_module" not in result.source
877+
assert "__main__" not in result.source
878+
879+
880+
def test_clean_triton_source_does_not_strip_non_benchmark_get_args_helper():
881+
source = """
882+
def call(args):
883+
return get_args(args)
884+
885+
886+
def get_args(args):
887+
return args
888+
889+
890+
def still_needed(args):
891+
return args
892+
"""
893+
894+
result = clean_triton_source(source, dynamic=False)
895+
896+
assert "def get_args(args):" in result.source
897+
assert "def still_needed(args):" in result.source
898+
899+
667900
def test_dynamic_clean_triton_source_patches_repeated_kernel_launches_by_occurrence():
668901
original_source = """
669902
def call(args):
@@ -684,8 +917,14 @@ def call(args):
684917

685918
result = clean_triton_source(source, dynamic=True, original_source=original_source)
686919

687-
assert "triton.cdiv(s0, 4), 1, 1)](buf0, s0" in result.source
688-
assert "triton.cdiv(s1, 4), 1, 1)](buf1, s1" in result.source
920+
assert (
921+
"triton.cdiv(s0, 4),\n 1,\n 1,\n )](\n buf0,\n s0,"
922+
in result.source
923+
)
924+
assert (
925+
"triton.cdiv(s1, 4),\n 1,\n 1,\n )](\n buf1,\n s1,"
926+
in result.source
927+
)
689928

690929

691930
def test_dynamic_clean_triton_reduction(tmp_path):
@@ -756,7 +995,8 @@ def test_dynamic_clean_triton_matmul_dynamic_m(tmp_path):
756995
assert "@triton.jit" in artifact_source
757996
assert "async_compile.triton" not in artifact_source
758997
assert "triton_tem_" in artifact_source
759-
assert "[((" in artifact_source
998+
assert "[(\n" in artifact_source
999+
assert "((15 + s" in artifact_source
7601000
assert "[(4, 1, 1)]" not in artifact_source
7611001

7621002
for batch_size in (1, 7, 32):

transformer_nuggets/export_autograd_triton/README.md

Lines changed: 20 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -55,18 +55,29 @@ agent_space/generated_rms_norm_artifacts/
5555
spec_0_backward.py
5656
```
5757

58-
The top-level generated wrapper is intentionally small. It contains:
58+
The top-level generated wrapper is intentionally small and labeled, but keeps the autograd node
59+
visible. It contains:
5960

60-
- `_SPECS`: specialization metadata and tensor guards
61+
- a generated artifact index with direct relative paths to each forward/backward artifact
62+
- one `_SPEC_<n>` block per specialization, with metadata and tensor guards
63+
- `_SPECS`: the ordered specialization list used for dispatch
6164
- `_RUNTIME = ExportedAutogradRuntime(...)`
62-
- one public function such as `rms_norm_compiled(...)`
65+
- one explicit `torch.autograd.Function` subclass whose `forward` and `backward` call the
66+
generated artifact runners through the runtime helper
67+
- one documented public function such as `rms_norm_compiled(...)` that selects a specialization,
68+
gathers runtime tensors, and invokes the generated autograd function
69+
70+
Artifact files start with a short generated-artifact docstring that names the specialization,
71+
forward/backward direction, runtime tensor order, static arguments, tensor guards, output
72+
metadata, and saved-residual order. The artifact directory is recreated on each export so stale
73+
kernels from earlier runs do not hang around.
6374

6475
Shared runtime behavior lives in `runtime.py`, not in every generated file:
6576

6677
- artifact module loading
6778
- specialization dispatch
6879
- static/tensor guard checking
69-
- generated autograd function
80+
- autograd forward/backward helper plumbing
7081
- saved residual reordering for backward graphs
7182
- non-differentiable output marking
7283
- no-match diagnostics
@@ -89,7 +100,10 @@ Use this when:
89100

90101
Runs `torch.utils._get_clean_triton.get_clean_triton(...)` and then applies the dynamic-clean validation/patching layer in `clean_triton.py`.
91102

92-
The goal is readable inline `@triton.jit` artifacts without `async_compile.triton(...)` while preserving dynamic launch arguments.
103+
The goal is readable inline `@triton.jit` artifacts without `async_compile.triton(...)` while
104+
preserving dynamic launch arguments. The cleaner also removes unused imports/top-level aliases,
105+
standalone benchmark scaffolding (`get_args`, `benchmark_compiled_module`,
106+
`if __name__ == "__main__"`), and formats patched dynamic launches as multiline Triton calls.
93107

94108
Supported dynamic clean families today:
95109

@@ -229,6 +243,5 @@ python examples/export_autograd_triton/dynamic.py
229243
- Symbolic stride guards for broader dynamic layouts
230244
- Dynamic CUDA device ordinal support
231245
- More dynamic clean-Triton families and variants as Inductor output changes
232-
- Cleaner metadata formatting in generated wrappers if `_SPECS` grows large
246+
- Semantic names for common generated symbols such as `arg0_1`, `buf*`, and `s*`
233247
- Optional Nsight Compute-based bandwidth measurement for exact HBM traffic
234-
- Tests that compare generated wrapper readability/size directly

transformer_nuggets/export_autograd_triton/api.py

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,12 @@ def export_autograd_triton(
4646
]
4747
output_path = Path(out)
4848
public_name = exported_name or f"{fn.__name__}_compiled"
49-
source = generate_autograd_source(fn, public_name, captured_specializations)
49+
source = generate_autograd_source(
50+
fn,
51+
public_name,
52+
captured_specializations,
53+
artifact_dir_name=f"{output_path.stem}_artifacts",
54+
)
5055
write_autograd_source(output_path, source, captured_specializations, source_backend)
5156
return ExportedAutogradSource(
5257
output_path=output_path,

0 commit comments

Comments
 (0)