Issue Station
Sign In
Tools hereToolbox

pytorch - ✅(Solved) Fix [Inductor] User-defined kernel epilogue fusion not guarded against non-unary epilogues [6 pull requests, 1 comments, 1 participants]

Official PRs (…)
ON THIS PAGE

Recommended Tools

×6

Utilities matched from this issue’s tags and category — try them while you read without losing context.

GitHub issue graph ai analysis

Paste a GitHub issue URL. We fetch that issue, discover linked issues from bodies/comments/timeline, collect linked pull requests, and produce a structured English report.

The report is written in English Markdown for sharing and archival.

Helpful · Quick feedback

Loading…
GitHub stats
pytorch/pytorch#179233•Fetched 2026-04-08 02:32:45
View on GitHub
Comments
1
Participants
1
Timeline
45
Reactions
0
Author
jjvrawjjvraw
Participants
jjvrawjjvraw
Timeline (top)
mentioned ×19subscribed ×19cross-referenced ×3labeled ×3

Error Message

W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last): W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ttir_module, ordered_arg_names = generate_ttir( W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ^^^^^^^^^^^^^^ W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] a = kwargs[name] W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ~~~~~~^^^^^^ W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE' W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last): W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ttir_module, ordered_arg_names = generate_ttir( W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ^^^^^^^^^^^^^^ W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] a = kwargs[name] W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ~~~~~~^^^^^^ W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE' W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last): W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ttir_module, ordered_arg_names = generate_ttir( W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ^^^^^^^^^^^^^^ W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] a = kwargs[name] W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ~~~~~~^^^^^^ W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE' W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last): W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ttir_module, ordered_arg_names = generate_ttir( W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ^^^^^^^^^^^^^^ W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] a = kwargs[name] W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] ~~~~~~^^^^^^ W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE' E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] Triton compilation failed: add_kernel_0 E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr): E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] pid = tl.program_id(0) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] mask = offs < n_elements E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] x = tl.load(in_ptr0 + offs, mask=mask) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] y = tl.load(in_ptr1 + offs, mask=mask) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] tl.store(out_ptr + offs, x + y, mask=mask) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] metadata: {'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'out_ptr': '*fp32', 'n_elements': 'i32', 'BLOCK_SIZE': 'constexpr'}, 'device': 0, 'constants': {}, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]]}], 'device_type': 'cuda', 'num_warps': 4, 'num_stages': 3, 'debug': True, 'cc': 89} E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] Traceback (most recent call last): E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 891, in _precompile_config E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] binary = triton.compile(*compile_args, **compile_kwargs) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 307, in compile E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] module = src.make_ir(target, options, codegen_fns, module_map, context) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 80, in make_ir E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns, E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] triton.compiler.errors.CompilationError: at 3:11: E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr): E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] pid = tl.program_id(0) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] ^ E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] AttributeError("'NoneType' object has no attribute 'type'") Traceback (most recent call last): File "/home/jjvraw/.../pytorch/repro.py", line 25, in <module> torch.compile(fn)(a, b, c) File "/home/jjvraw/.../pytorch/torch/_dynamo/eval_frame.py", line 1050, in compile_wrapper raise e.remove_dynamo_frames() from None # see TORCHDYNAMO_VERBOSE=1 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1053, in _compile_fx_inner raise InductorError(e, currentframe()).with_traceback( File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1037, in _compile_fx_inner mb_compiled_graph = fx_codegen_and_compile( ^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1802, in fx_codegen_and_compile return scheme.codegen_and_compile(gm, example_inputs, inputs_to_check, graph_kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1574, in codegen_and_compile compiled_module = graph.compile_to_module() ^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2606, in compile_to_module return self._compile_to_module() ^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2616, in _compile_to_module mod = self._compile_to_module_lines(wrapper_code) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2691, in _compile_to_module_lines mod = PyCodeCache.load_by_key_path( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/codecache.py", line 3917, in load_by_key_path mod = _reload_python_module(key, path, set_sys_modules=in_toplevel) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/compile_tasks.py", line 35, in _reload_python_module exec(code, mod.dict, mod.dict) File "/tmp/torchinductor_jjvraw/ba/cbax55aypfukxnzzwfn4vodmyommxg6owube5knhl5zfb4thn63g.py", line 38, in <module> add_kernel_0 = async_compile.triton('add_kernel', ''' ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/async_compile.py", line 491, in triton kernel.precompile( File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 512, in precompile self._precompile_worker() File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 534, in _precompile_worker compile_results.append(self._precompile_config(c)) ^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 891, in _precompile_config binary = triton.compile(*compile_args, **compile_kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 307, in compile module = src.make_ir(target, options, codegen_fns, module_map, context) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 80, in make_ir return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns, ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ torch._inductor.exc.InductorError: CompilationError: at 3:11: def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(0) offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) ^ AttributeError("'NoneType' object has no attribute 'type'")

Set TORCHDYNAMO_VERBOSE=1 for the internal stack trace (please do this especially if you're reporting a bug to PyTorch). For even more developer context, set TORCH_LOGS="+dynamo"

Fix Action

Fix / Workaround

CPU: Architecture: x86_64 CPU op-mode(s): 32-bit, 64-bit Address sizes: 48 bits physical, 48 bits virtual Byte Order: Little Endian CPU(s): 32 On-line CPU(s) list: 0-31 Vendor ID: AuthenticAMD Model name: AMD Ryzen 9 7950X3D 16-Core Processor CPU family: 25 Model: 97 Thread(s) per core: 2 Core(s) per socket: 16 Socket(s): 1 Stepping: 2 Frequency boost: enabled CPU(s) scaling MHz: 63% CPU max MHz: 5763.0000 CPU min MHz: 545.0000 BogoMIPS: 8383.69 Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ht syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm constant_tsc rep_good amd_lbr_v2 nopl xtopology nonstop_tsc cpuid extd_apicid aperfmperf rapl pni pclmulqdq monitor ssse3 fma cx16 sse4_1 sse4_2 movbe popcnt aes xsave avx f16c rdrand lahf_lm cmp_legacy svm extapic cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw ibs skinit wdt tce topoext perfctr_core perfctr_nb bpext perfctr_llc mwaitx cpb cat_l3 cdp_l3 hw_pstate ssbd mba perfmon_v2 ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase bmi1 avx2 smep bmi2 erms invpcid cqm rdt_a avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves cqm_llc cqm_occup_llc cqm_mbm_total cqm_mbm_local user_shstk avx512_bf16 clzero irperf xsaveerptr rdpru wbnoinvd cppc arat npt lbrv svm_lock nrip_save tsc_scale vmcb_clean flushbyasid decodeassists pausefilter pfthreshold avic vgif x2avic v_spec_ctrl vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq rdpid overflow_recov succor smca fsrm flush_l1d amd_lbr_pmc_freeze Virtualization: AMD-V L1d cache: 512 KiB (16 instances) L1i cache: 512 KiB (16 instances) L2 cache: 16 MiB (16 instances) L3 cache: 128 MiB (2 instances) NUMA node(s): 1 NUMA node0 CPU(s): 0-31 Vulnerability Gather data sampling: Not affected Vulnerability Ghostwrite: Not affected Vulnerability Indirect target selection: Not affected Vulnerability Itlb multihit: Not affected Vulnerability L1tf: Not affected Vulnerability Mds: Not affected Vulnerability Meltdown: Not affected Vulnerability Mmio stale data: Not affected Vulnerability Reg file data sampling: Not affected Vulnerability Retbleed: Not affected Vulnerability Spec rstack overflow: Mitigation; Safe RET Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl Vulnerability Spectre v1: Mitigation; usercopy/swapgs barriers and __user pointer sanitization Vulnerability Spectre v2: Mitigation; Enhanced / Automatic IBRS; IBPB conditional; STIBP always-on; PBRSB-eIBRS Not affected; BHI Not affected Vulnerability Srbds: Not affected Vulnerability Tsa: Mitigation; Clear CPU buffers Vulnerability Tsx async abort: Not affected Vulnerability Vmscape: Mitigation; IBPB before exit to userspace

PR fix notes

PR #179370: [Inductor] Guard user-defined Triton kernel fusion against non-unary epilogues

Description (problem / solution / changelog)

Summary

This PR fixes a bug where torch.compile would crash when attempting to fuse an epilogue onto a user-defined Triton kernel if that epilogue required additional tensor dependencies (e.g., out = user_kernel(a, b); return out + c).

Currently, the Scheduler only checks if the epilogue's indexing logic is representable (purely derived from loads). However, for user-defined kernels, the function signature is fixed. If an epilogue requires a read from a buffer other than the kernel's mutated output, the fusion fails during codegen because the required tensor pointer is not available in the Triton kernel arguments.

Changes

  • torch/_inductor/scheduler.py: Added a dependency guard in the can_fuse logic for user-defined kernels. It now explicitly rejects fusion if the epilogue node (node2) reads from any buffer other than the producer's (node1) output.
  • test/inductor/test_triton_kernels.py: Added a regression test case involving a binary epilogue (out + c) to ensure it correctly skips fusion and executes safely.
  • agent_space/repro_user_kernel_epilogue_fusion.py: Created a standalone reproduction script to validate the fix.

Validation

  • Verified that the patched scheduler.py correctly identifies non-unary reads as unsafe for fusion.
  • Confirmed syntax and basic integrity of the logic via a standalone mock-scheduler test.
  • Ensured the fix prevents the KeyError: 'BLOCK_SIZE' and CompilationError reported in issue #179233.

Fixes

Fixes #179233

cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo

Changed files

  • test/inductor/test_triton_kernels.py (modified, +29/-0)
  • torch/_inductor/scheduler.py (modified, +10/-3)

PR #179149: [Inductor][RFC] Symbolic Analysis of User-Defined Triton Kernels

Description (problem / solution / changelog)

cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo @jansel

Motivation

Epilogue fusion for user-defined kernels (#173662) is currently scoped to UB tensors, with the rationale epilogue(UB) == UB, assuming no in-bound masks. This covers a wide range of common cases but, of course, excludes in-place kernels or any operating on non-empty output tensors.

Extracting index expressions from user-defined kernels opens a few directions:

  1. Loosen the UB restriction by reasoning formally about read/write access patterns.
  2. Improve fusion-scoring/cost modeling.
  3. Prologue fusion.

This PR implements symbolic index expression extraction from user-defined Triton kernels via TTIR traversal as a foundation for the above, with the intention of gathering feedback on these directions.

Dependencies

To expose iteration bounds, we rely on MLIR bindings (https://github.com/triton-lang/triton/pull/8892, https://github.com/triton-lang/triton/pull/9866) that are called during TTIR parsing. Some of which are not cherry-picked on release/3.7.x / the current Triton pin. See https://github.com/triton-lang/triton/pull/9901 for diff against the current release branch.

TTIR Traversal

While flattening the TTIR module into a linear def-use chain, we store attributes for certain MLIR ops in the Op dataclass that are relevant to expression building. For example, tt.make_range stores start and end, which are used to represent the bounds of an iteration symbol.

To handle loop constructs, specifically scf.for, we mirror the iter_arg-yield relationship (see scf.for docs) to express any pointer loop-carried arithmetic in closed-form. This is done through two new Op dataclasses, InductionVar and IterArg, which are handled per-case during expression building.

Expression Building

The general strategy is to categorise MLIR ops: Perform upward traversal from read/write sinks, as before, recursively build expression while traversing through arithmetic ops. We always mint a new symbol when encountering a leaf (other than arith.constant), while maintaining a recursion cache keyed on both the SSA-index and the current shape. This prevents both redundant symbols from being minted and, more importantly, handles cases where the same SSA value can produce different expressions depending on how its output has been broadcast or reshaped by shape-context ops.

Certain expressions may accumulate nested divisions and modulos, that can be simplified away during traversal depending on kernel configuration:

Compute an upper bound B, such that expr < B for range of expr.

  • arith.remsi(expr, k): if B <= k, expr % k == expr
  • arith.divsi(expr, k): if B <= k, expr // k == 0

This is particularly useful for PID grouping.

As a fallback of symbolic analysis, we revert to conservative analysis/mutation tracking. This behaves exactly the same as what's currently upstream. This may be due to unsupported ops (intentionally or unintentionally), in which case SymbolicFailure is raised, or for any other reason which raises an Exception.

UserTritonDep

Regardless of whether symbolic or conservative analysis is used, identify_accessed_tensors always returns a new dependency type, UserTritonDep. This is more akin to a MemoryDep containing fields for both index and mask.

Tension arises during downstream usage of UserTritonDep when comparing it against a MemoryDep in an epilogue/prologue. We have UserTritonDep expressed in terms of kernel configuration, i.e., grid config, program ID, and other runtime scalars captured transitively (BLOCK_SIZE), and MemoryDep in terms of tensor shapes and strides. The scheduler's current index and bound equality checks will not suffice here.

There are various ways of approaching this:

  • ISL-like handling to compare iteration-spaces.
  • Extending Inductor's sizevars to support greater coverage of expression normalisation and bound checking.

Both of which I've explored locally.

I believe the use of UserTritonDep sets the scheduler up to have specific branching for user-defined scheduler nodes as well.

Unrelated Changes

Changes to Scheduler.can_fuse are patches for bugs related to upstream user-defined kernel fusion:

  • #179233
  • #179232

Will find min-repros and open corresponding issues.

Test Coverage

See the newly added tests in test_triton_kernel for examples of index expressions. There is full test coverage on test_triton_kernel as well, tested locally.

Benchmarks

Microbenchmarking identify_accessed_tensors over a few kernels from Liger, and Triton Tutorials. Indication of the introduced overhead of symbolic analysis. Each result is the minimum over 5 x 200 calls. With the current operator handler coverage, the following kernels result in a fallback to "conservative analysis":

  • cross_entropy - data-dependent pointer (tl.load(X_ptr + y) where y is a runtime value)
  • rope_forward - unhandled arith.select (tl.where). Requires predicate handlers.
  • triton_softmax - unhandled tt.get_num_programs. Trivial to add, similar to existing tl.program_id handling.
<img width="1500" height="750" alt="bench_liger_comparison" src="https://github.com/user-attachments/assets/c8e0db4e-5fb3-40d8-9dfa-06b989676b33" /> <img width="1500" height="750" alt="bench_triton_comparison" src="https://github.com/user-attachments/assets/0fe88127-5428-45e4-9f74-426eb1a073c9" />

cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo

Changed files

  • test/inductor/test_triton_kernels.py (modified, +698/-6)
  • torch/_higher_order_ops/triton_kernel_wrap.py (modified, +807/-297)
  • torch/_inductor/dependencies.py (modified, +48/-0)
  • torch/_inductor/ir.py (modified, +1/-0)
  • torch/_inductor/scheduler.py (modified, +38/-22)

PR #179735: [Inductor] Prevent user-kernel fusion with non-unary epilogues

Description (problem / solution / changelog)

Fixes: #179233

cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo

Changed files

  • test/inductor/test_triton_kernels.py (modified, +26/-0)
  • torch/_inductor/codegen/triton.py (modified, +5/-1)
  • torch/_inductor/scheduler.py (modified, +9/-3)

Code Example

# the epilogue depends on expressions which may not available in the user triton kernel
# (e.g. indexing exprs used not in a load)
node2_inner_fn_free_symbols = node2.node.data.inner_fn_free_symbols()
for symbol in node2_inner_fn_free_symbols:
    usages = node2.node.data.collect_inner_fn_symbol_usage(symbol)
    if any(usage != "load" for usage in usages):
        return False

---

import torch
import triton
import triton.language as tl

torch._inductor.config.epilogue_fusion_user_defined_triton_kernel = True

@triton.jit
def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offs < n_elements
    x = tl.load(in_ptr0 + offs, mask=mask)
    y = tl.load(in_ptr1 + offs, mask=mask)
    tl.store(out_ptr + offs, x + y, mask=mask)

def fn(a, b, c):
    out = torch.empty_like(a)
    add_kernel[(a.numel(),)](a, b, out, a.numel(), block_size=1)
    return out + c


a = torch.randn(10, device="cuda")
b = torch.randn(10, device="cuda")
c = torch.randn(10, device="cuda")
torch.compile(fn)(a, b, c)

---

W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] Triton compilation failed: add_kernel_0
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     pid = tl.program_id(0)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     mask = offs < n_elements
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     x = tl.load(in_ptr0 + offs, mask=mask)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     y = tl.load(in_ptr1 + offs, mask=mask)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     tl.store(out_ptr + offs, x + y, mask=mask)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] metadata: {'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'out_ptr': '*fp32', 'n_elements': 'i32', 'BLOCK_SIZE': 'constexpr'}, 'device': 0, 'constants': {}, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]]}], 'device_type': 'cuda', 'num_warps': 4, 'num_stages': 3, 'debug': True, 'cc': 89}
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] Traceback (most recent call last):
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]   File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 891, in _precompile_config
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     binary = triton.compile(*compile_args, **compile_kwargs)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]   File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 307, in compile
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     module = src.make_ir(target, options, codegen_fns, module_map, context)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]   File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 80, in make_ir
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns,
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] triton.compiler.errors.CompilationError: at 3:11:
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     pid = tl.program_id(0)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]            ^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] AttributeError("'NoneType' object has no attribute 'type'")
Traceback (most recent call last):
  File "/home/jjvraw/.../pytorch/repro.py", line 25, in <module>
    torch.compile(fn)(a, b, c)
  File "/home/jjvraw/.../pytorch/torch/_dynamo/eval_frame.py", line 1050, in compile_wrapper
    raise e.remove_dynamo_frames() from None  # see TORCHDYNAMO_VERBOSE=1
    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1053, in _compile_fx_inner
    raise InductorError(e, currentframe()).with_traceback(
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1037, in _compile_fx_inner
    mb_compiled_graph = fx_codegen_and_compile(
                        ^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1802, in fx_codegen_and_compile
    return scheme.codegen_and_compile(gm, example_inputs, inputs_to_check, graph_kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1574, in codegen_and_compile
    compiled_module = graph.compile_to_module()
                      ^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2606, in compile_to_module
    return self._compile_to_module()
           ^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2616, in _compile_to_module
    mod = self._compile_to_module_lines(wrapper_code)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2691, in _compile_to_module_lines
    mod = PyCodeCache.load_by_key_path(
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/codecache.py", line 3917, in load_by_key_path
    mod = _reload_python_module(key, path, set_sys_modules=in_toplevel)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/compile_tasks.py", line 35, in _reload_python_module
    exec(code, mod.__dict__, mod.__dict__)
  File "/tmp/torchinductor_jjvraw/ba/cbax55aypfukxnzzwfn4vodmyommxg6owube5knhl5zfb4thn63g.py", line 38, in <module>
    add_kernel_0 = async_compile.triton('add_kernel', '''
                   ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/async_compile.py", line 491, in triton
    kernel.precompile(
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 512, in precompile
    self._precompile_worker()
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 534, in _precompile_worker
    compile_results.append(self._precompile_config(c))
                           ^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 891, in _precompile_config
    binary = triton.compile(*compile_args, **compile_kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 307, in compile
    module = src.make_ir(target, options, codegen_fns, module_map, context)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 80, in make_ir
    return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns,
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
torch._inductor.exc.InductorError: CompilationError: at 3:11:
def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
           ^
AttributeError("'NoneType' object has no attribute 'type'")

Set TORCHDYNAMO_VERBOSE=1 for the internal stack trace (please do this especially if you're reporting a bug to PyTorch). For even more developer context, set TORCH_LOGS="+dynamo"

---

Collecting environment information...
PyTorch version: 2.12.0a0+gitd386e0b
Is debug build: False
CUDA used to build PyTorch: 12.8
ROCM used to build PyTorch: N/A

OS: Ubuntu 25.04 (x86_64)
GCC version: (Ubuntu 14.2.0-19ubuntu2) 14.2.0
Clang version: 20.1.2 (0ubuntu1)
CMake version: version 4.3.1
Libc version: glibc-2.41

Python version: 3.11.15 (main, Mar 11 2026, 17:20:07) [GCC 14.3.0] (64-bit runtime)
Python platform: Linux-6.14.0-37-generic-x86_64-with-glibc2.41
Is CUDA available: True
CUDA runtime version: 12.8.61
CUDA_MODULE_LOADING set to:
GPU models and configuration: GPU 0: NVIDIA GeForce RTX 4090
Nvidia driver version: 580.95.05
cuDNN version: Could not collect
Is XPU available: False
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True
Caching allocator config: N/A

CPU:
Architecture:                            x86_64
CPU op-mode(s):                          32-bit, 64-bit
Address sizes:                           48 bits physical, 48 bits virtual
Byte Order:                              Little Endian
CPU(s):                                  32
On-line CPU(s) list:                     0-31
Vendor ID:                               AuthenticAMD
Model name:                              AMD Ryzen 9 7950X3D 16-Core Processor
CPU family:                              25
Model:                                   97
Thread(s) per core:                      2
Core(s) per socket:                      16
Socket(s):                               1
Stepping:                                2
Frequency boost:                         enabled
CPU(s) scaling MHz:                      63%
CPU max MHz:                             5763.0000
CPU min MHz:                             545.0000
BogoMIPS:                                8383.69
Flags:                                   fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ht syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm constant_tsc rep_good amd_lbr_v2 nopl xtopology nonstop_tsc cpuid extd_apicid aperfmperf rapl pni pclmulqdq monitor ssse3 fma cx16 sse4_1 sse4_2 movbe popcnt aes xsave avx f16c rdrand lahf_lm cmp_legacy svm extapic cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw ibs skinit wdt tce topoext perfctr_core perfctr_nb bpext perfctr_llc mwaitx cpb cat_l3 cdp_l3 hw_pstate ssbd mba perfmon_v2 ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase bmi1 avx2 smep bmi2 erms invpcid cqm rdt_a avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves cqm_llc cqm_occup_llc cqm_mbm_total cqm_mbm_local user_shstk avx512_bf16 clzero irperf xsaveerptr rdpru wbnoinvd cppc arat npt lbrv svm_lock nrip_save tsc_scale vmcb_clean flushbyasid decodeassists pausefilter pfthreshold avic vgif x2avic v_spec_ctrl vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq rdpid overflow_recov succor smca fsrm flush_l1d amd_lbr_pmc_freeze
Virtualization:                          AMD-V
L1d cache:                               512 KiB (16 instances)
L1i cache:                               512 KiB (16 instances)
L2 cache:                                16 MiB (16 instances)
L3 cache:                                128 MiB (2 instances)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-31
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Mitigation; Safe RET
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; Enhanced / Automatic IBRS; IBPB conditional; STIBP always-on; PBRSB-eIBRS Not affected; BHI Not affected
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Mitigation; Clear CPU buffers
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Mitigation; IBPB before exit to userspace

Versions of relevant libraries:
[pip3] bert_pytorch==0.0.1a4
[pip3] functorch==1.14.0a0+b71aa0b
[pip3] intel-cmplr-lib-ur==2025.3.3
[pip3] intel-openmp==2025.3.3
[pip3] mkl==2025.3.1
[pip3] mkl-include==2025.3.1
[pip3] mypy==1.20.0
[pip3] mypy_extensions==1.1.0
[pip3] numpy==2.4.3
[pip3] onemkl-license==2025.3.1
[pip3] onnx==1.21.0
[pip3] pytorch-labs-segment-anything-fast==0.2
[pip3] tbb==2022.3.1
[pip3] tcmlib==1.4.1
[pip3] torch==2.12.0a0+gitd386e0b
[pip3] torch_geometric==2.4.0
[pip3] torchao==0.17.0
[pip3] torchaudio==2.11.0a0+c0cbdb9
[pip3] torchdata==0.12.0a0+93b65f7
[pip3] torchmetrics==1.0.3
[pip3] torchmultimodal==0.1.0b0
[pip3] torchrec==1.7.0a0+2d6eb73
[pip3] torchvision==0.27.0a0+4e58149
[pip3] triton==3.7.0
[pip3] umf==1.0.3
[conda] bert-pytorch                          0.0.1a4              pypi_0           pypi
[conda] functorch                             1.14.0a0+b71aa0b     pypi_0           pypi
[conda] intel-cmplr-lib-ur                    2025.3.3             pypi_0           pypi
[conda] intel-openmp                          2025.3.3             pypi_0           pypi
[conda] mkl                                   2025.3.1             pypi_0           pypi
[conda] mkl-include                           2025.3.1             pypi_0           pypi
[conda] numpy                                 2.4.3                pypi_0           pypi
[conda] onemkl-license                        2025.3.1             pypi_0           pypi
[conda] pytorch-labs-segment-anything-fast    0.2                  pypi_0           pypi
[conda] tbb                                   2022.3.1             pypi_0           pypi
[conda] tcmlib                                1.4.1                pypi_0           pypi
[conda] torch                                 2.12.0a0+gitd386e0b  pypi_0           pypi
[conda] torch-geometric                       2.4.0                pypi_0           pypi
[conda] torchao                               0.17.0               pypi_0           pypi
[conda] torchaudio                            2.11.0a0+c0cbdb9     pypi_0           pypi
[conda] torchdata                             0.12.0a0+93b65f7     pypi_0           pypi
[conda] torchmetrics                          1.0.3                pypi_0           pypi
[conda] torchmultimodal                       0.1.0b0              pypi_0           pypi
[conda] torchrec                              1.7.0a0+2d6eb73      pypi_0           pypi
[conda] torchvision                           0.27.0a0+4e58149     pypi_0           pypi
[conda] triton                                3.7.0                pypi_0           pypi
[conda] umf                                   1.0.3                pypi_0           pypi
RAW_BUFFERClick to expand / collapse

🐛 Describe the bug

The saftey check for user kernel epilogue fusion in Scheduler.can_fuse contains the following handling for index symbols, ensuring index symbols are purely derived from loads.

# the epilogue depends on expressions which may not available in the user triton kernel
# (e.g. indexing exprs used not in a load)
node2_inner_fn_free_symbols = node2.node.data.inner_fn_free_symbols()
for symbol in node2_inner_fn_free_symbols:
    usages = node2.node.data.collect_inner_fn_symbol_usage(symbol)
    if any(usage != "load" for usage in usages):
        return False

However, it does not account for the case where the epilogue reads from additional tensor buffers beyond the mutated output.

Min Repro

import torch
import triton
import triton.language as tl

torch._inductor.config.epilogue_fusion_user_defined_triton_kernel = True

@triton.jit
def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offs < n_elements
    x = tl.load(in_ptr0 + offs, mask=mask)
    y = tl.load(in_ptr1 + offs, mask=mask)
    tl.store(out_ptr + offs, x + y, mask=mask)

def fn(a, b, c):
    out = torch.empty_like(a)
    add_kernel[(a.numel(),)](a, b, out, a.numel(), block_size=1)
    return out + c


a = torch.randn(10, device="cuda")
b = torch.randn(10, device="cuda")
c = torch.randn(10, device="cuda")
torch.compile(fn)(a, b, c)

Error logs

W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.592000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.600000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Encountered an exception in identify_accessed_tensors, assuming every input is mutated
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] Traceback (most recent call last):
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 1111, in identify_accessed_tensors
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     ttir_module, ordered_arg_names = generate_ttir(
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]                                      ^^^^^^^^^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]   File "/home/jjvraw/.../pytorch/torch/_higher_order_ops/triton_kernel_wrap.py", line 303, in generate_ttir
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]     a = kwargs[name]
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0]         ~~~~~~^^^^^^
W0403 09:55:15.866000 4168291 torch/_higher_order_ops/triton_kernel_wrap.py:1151] [0/0] KeyError: 'BLOCK_SIZE'
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] Triton compilation failed: add_kernel_0
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     pid = tl.program_id(0)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     mask = offs < n_elements
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     x = tl.load(in_ptr0 + offs, mask=mask)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     y = tl.load(in_ptr1 + offs, mask=mask)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     tl.store(out_ptr + offs, x + y, mask=mask)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] metadata: {'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'out_ptr': '*fp32', 'n_elements': 'i32', 'BLOCK_SIZE': 'constexpr'}, 'device': 0, 'constants': {}, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]]}], 'device_type': 'cuda', 'num_warps': 4, 'num_stages': 3, 'debug': True, 'cc': 89}
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] Traceback (most recent call last):
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]   File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 891, in _precompile_config
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     binary = triton.compile(*compile_args, **compile_kwargs)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]   File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 307, in compile
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     module = src.make_ir(target, options, codegen_fns, module_map, context)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]   File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 80, in make_ir
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns,
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] triton.compiler.errors.CompilationError: at 3:11:
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     pid = tl.program_id(0)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]     offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0]            ^
E0403 09:55:15.930000 4168291 torch/_inductor/runtime/triton_heuristics.py:893] [0/0] AttributeError("'NoneType' object has no attribute 'type'")
Traceback (most recent call last):
  File "/home/jjvraw/.../pytorch/repro.py", line 25, in <module>
    torch.compile(fn)(a, b, c)
  File "/home/jjvraw/.../pytorch/torch/_dynamo/eval_frame.py", line 1050, in compile_wrapper
    raise e.remove_dynamo_frames() from None  # see TORCHDYNAMO_VERBOSE=1
    ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1053, in _compile_fx_inner
    raise InductorError(e, currentframe()).with_traceback(
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1037, in _compile_fx_inner
    mb_compiled_graph = fx_codegen_and_compile(
                        ^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1802, in fx_codegen_and_compile
    return scheme.codegen_and_compile(gm, example_inputs, inputs_to_check, graph_kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/compile_fx.py", line 1574, in codegen_and_compile
    compiled_module = graph.compile_to_module()
                      ^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2606, in compile_to_module
    return self._compile_to_module()
           ^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2616, in _compile_to_module
    mod = self._compile_to_module_lines(wrapper_code)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2691, in _compile_to_module_lines
    mod = PyCodeCache.load_by_key_path(
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/codecache.py", line 3917, in load_by_key_path
    mod = _reload_python_module(key, path, set_sys_modules=in_toplevel)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/compile_tasks.py", line 35, in _reload_python_module
    exec(code, mod.__dict__, mod.__dict__)
  File "/tmp/torchinductor_jjvraw/ba/cbax55aypfukxnzzwfn4vodmyommxg6owube5knhl5zfb4thn63g.py", line 38, in <module>
    add_kernel_0 = async_compile.triton('add_kernel', '''
                   ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/async_compile.py", line 491, in triton
    kernel.precompile(
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 512, in precompile
    self._precompile_worker()
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 534, in _precompile_worker
    compile_results.append(self._precompile_config(c))
                           ^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/runtime/triton_heuristics.py", line 891, in _precompile_config
    binary = triton.compile(*compile_args, **compile_kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 307, in compile
    module = src.make_ir(target, options, codegen_fns, module_map, context)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/Git/triton/python/triton/compiler/compiler.py", line 80, in make_ir
    return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns,
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
torch._inductor.exc.InductorError: CompilationError: at 3:11:
def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
           ^
AttributeError("'NoneType' object has no attribute 'type'")

Set TORCHDYNAMO_VERBOSE=1 for the internal stack trace (please do this especially if you're reporting a bug to PyTorch). For even more developer context, set TORCH_LOGS="+dynamo"

Versions

Collecting environment information...
PyTorch version: 2.12.0a0+gitd386e0b
Is debug build: False
CUDA used to build PyTorch: 12.8
ROCM used to build PyTorch: N/A

OS: Ubuntu 25.04 (x86_64)
GCC version: (Ubuntu 14.2.0-19ubuntu2) 14.2.0
Clang version: 20.1.2 (0ubuntu1)
CMake version: version 4.3.1
Libc version: glibc-2.41

Python version: 3.11.15 (main, Mar 11 2026, 17:20:07) [GCC 14.3.0] (64-bit runtime)
Python platform: Linux-6.14.0-37-generic-x86_64-with-glibc2.41
Is CUDA available: True
CUDA runtime version: 12.8.61
CUDA_MODULE_LOADING set to:
GPU models and configuration: GPU 0: NVIDIA GeForce RTX 4090
Nvidia driver version: 580.95.05
cuDNN version: Could not collect
Is XPU available: False
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True
Caching allocator config: N/A

CPU:
Architecture:                            x86_64
CPU op-mode(s):                          32-bit, 64-bit
Address sizes:                           48 bits physical, 48 bits virtual
Byte Order:                              Little Endian
CPU(s):                                  32
On-line CPU(s) list:                     0-31
Vendor ID:                               AuthenticAMD
Model name:                              AMD Ryzen 9 7950X3D 16-Core Processor
CPU family:                              25
Model:                                   97
Thread(s) per core:                      2
Core(s) per socket:                      16
Socket(s):                               1
Stepping:                                2
Frequency boost:                         enabled
CPU(s) scaling MHz:                      63%
CPU max MHz:                             5763.0000
CPU min MHz:                             545.0000
BogoMIPS:                                8383.69
Flags:                                   fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ht syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm constant_tsc rep_good amd_lbr_v2 nopl xtopology nonstop_tsc cpuid extd_apicid aperfmperf rapl pni pclmulqdq monitor ssse3 fma cx16 sse4_1 sse4_2 movbe popcnt aes xsave avx f16c rdrand lahf_lm cmp_legacy svm extapic cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw ibs skinit wdt tce topoext perfctr_core perfctr_nb bpext perfctr_llc mwaitx cpb cat_l3 cdp_l3 hw_pstate ssbd mba perfmon_v2 ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase bmi1 avx2 smep bmi2 erms invpcid cqm rdt_a avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves cqm_llc cqm_occup_llc cqm_mbm_total cqm_mbm_local user_shstk avx512_bf16 clzero irperf xsaveerptr rdpru wbnoinvd cppc arat npt lbrv svm_lock nrip_save tsc_scale vmcb_clean flushbyasid decodeassists pausefilter pfthreshold avic vgif x2avic v_spec_ctrl vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq rdpid overflow_recov succor smca fsrm flush_l1d amd_lbr_pmc_freeze
Virtualization:                          AMD-V
L1d cache:                               512 KiB (16 instances)
L1i cache:                               512 KiB (16 instances)
L2 cache:                                16 MiB (16 instances)
L3 cache:                                128 MiB (2 instances)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-31
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Mitigation; Safe RET
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; Enhanced / Automatic IBRS; IBPB conditional; STIBP always-on; PBRSB-eIBRS Not affected; BHI Not affected
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Mitigation; Clear CPU buffers
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Mitigation; IBPB before exit to userspace

Versions of relevant libraries:
[pip3] bert_pytorch==0.0.1a4
[pip3] functorch==1.14.0a0+b71aa0b
[pip3] intel-cmplr-lib-ur==2025.3.3
[pip3] intel-openmp==2025.3.3
[pip3] mkl==2025.3.1
[pip3] mkl-include==2025.3.1
[pip3] mypy==1.20.0
[pip3] mypy_extensions==1.1.0
[pip3] numpy==2.4.3
[pip3] onemkl-license==2025.3.1
[pip3] onnx==1.21.0
[pip3] pytorch-labs-segment-anything-fast==0.2
[pip3] tbb==2022.3.1
[pip3] tcmlib==1.4.1
[pip3] torch==2.12.0a0+gitd386e0b
[pip3] torch_geometric==2.4.0
[pip3] torchao==0.17.0
[pip3] torchaudio==2.11.0a0+c0cbdb9
[pip3] torchdata==0.12.0a0+93b65f7
[pip3] torchmetrics==1.0.3
[pip3] torchmultimodal==0.1.0b0
[pip3] torchrec==1.7.0a0+2d6eb73
[pip3] torchvision==0.27.0a0+4e58149
[pip3] triton==3.7.0
[pip3] umf==1.0.3
[conda] bert-pytorch                          0.0.1a4              pypi_0           pypi
[conda] functorch                             1.14.0a0+b71aa0b     pypi_0           pypi
[conda] intel-cmplr-lib-ur                    2025.3.3             pypi_0           pypi
[conda] intel-openmp                          2025.3.3             pypi_0           pypi
[conda] mkl                                   2025.3.1             pypi_0           pypi
[conda] mkl-include                           2025.3.1             pypi_0           pypi
[conda] numpy                                 2.4.3                pypi_0           pypi
[conda] onemkl-license                        2025.3.1             pypi_0           pypi
[conda] pytorch-labs-segment-anything-fast    0.2                  pypi_0           pypi
[conda] tbb                                   2022.3.1             pypi_0           pypi
[conda] tcmlib                                1.4.1                pypi_0           pypi
[conda] torch                                 2.12.0a0+gitd386e0b  pypi_0           pypi
[conda] torch-geometric                       2.4.0                pypi_0           pypi
[conda] torchao                               0.17.0               pypi_0           pypi
[conda] torchaudio                            2.11.0a0+c0cbdb9     pypi_0           pypi
[conda] torchdata                             0.12.0a0+93b65f7     pypi_0           pypi
[conda] torchmetrics                          1.0.3                pypi_0           pypi
[conda] torchmultimodal                       0.1.0b0              pypi_0           pypi
[conda] torchrec                              1.7.0a0+2d6eb73      pypi_0           pypi
[conda] torchvision                           0.27.0a0+4e58149     pypi_0           pypi
[conda] triton                                3.7.0                pypi_0           pypi
[conda] umf                                   1.0.3                pypi_0           pypi

cc @chauhang @penguinwu

extent analysis

TL;DR

The issue is likely caused by the missing handling of additional tensor buffers in the epilogue fusion safety check, and a KeyError exception is raised due to the missing 'BLOCK_SIZE' key.

Guidance

  1. Check the BLOCK_SIZE parameter: Verify that the BLOCK_SIZE parameter is correctly passed to the add_kernel function.
  2. Update the safety check: Modify the safety check in Scheduler.can_fuse to account for additional tensor buffers beyond the mutated output.
  3. Verify the torch.compile configuration: Ensure that the torch.compile configuration is correct and the epilogue_fusion_user_defined_triton_kernel option is properly set.

Example

# Update the safety check to handle additional tensor buffers
def Scheduler.can_fuse(node1, node2):
    # ...
    node2_inner_fn_free_symbols = node2.node.data.inner_fn_free_symbols()
    for symbol in node2_inner_fn_free_symbols:
        usages = node2.node.data.collect_inner_fn_symbol_usage(symbol)
        if any(usage != "load" for usage in usages):
            return False
    # Add a check for additional tensor buffers
    if node2.node.data.has_additional_buffers():
        # Handle the additional buffers
        pass
    return True

Notes

The provided code snippet and error logs suggest that the issue is related to the epilogue fusion safety check and the missing handling of additional tensor buffers. However, without more information about the Scheduler.can_fuse function and the torch.compile configuration, it is difficult to provide a more specific solution.

Recommendation

Apply a workaround by modifying the Scheduler.can_fuse function to handle additional tensor buffers, and verify that the torch.compile configuration is correct. If the issue persists, consider upgrading to a newer version of PyTorch or seeking further assistance from the PyTorch community.

Vote matrix · Quick signals

Works
Did the solution work? Tap to confirm.
Easy Fix
Was it a quick fix?
Time Saver
Did it save you time?
Blocking
Was it severely blocking?
Common Issue
Are others likely hitting this too?
Flaky / Intermittent
Is it intermittent?
Verified / Reproducible
Can you reproduce it reliably?
Loading…
#api#prompt template#agent execution#callback error#memory management

Still need to ship something?

×6

Another batch ranked right after the header list — different links, same matching logic.

Back to top recommendations

TRENDING

pytorch - ✅(Solved) Fix [Inductor] User-defined kernel epilogue fusion not guarded against non-unary epilogues [6 pull requests, 1 comments, 1 participants]