Issue Station
Sign In
Tools hereToolbox

pytorch - ✅(Solved) Fix [Inductor] User-defined kernel fusion incorrectly attempted after buffer reuse decisions [5 pull requests, 1 comments, 2 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#179232•Fetched 2026-04-08 02:32:46
View on GitHub
Comments
1
Participants
2
Timeline
51
Reactions
0
Author
jjvrawjjvraw
Participants
anijain2305anijain2305jjvrawjjvraw
Timeline (top)
mentioned ×20subscribed ×20labeled ×5cross-referenced ×3

Error Message

Traceback (most recent call last): File "/home/jjvraw/.../pytorch/repro.py", line 26, in <module> compiled(a, b) ^^^^^^^^^^^^^^ 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 2612, in _compile_to_module self.codegen_with_cpp_wrapper() if self.cpp_wrapper else self.codegen() ^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2544, in codegen self._update_scheduler() File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2538, in _update_scheduler self.scheduler = Scheduler(self.operations) ^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 3086, in init self._init(nodes) File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 3189, in _init self.nodes = self.fuse_nodes(self.nodes) ^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4037, in fuse_nodes nodes = self.fuse_nodes_once(nodes, is_reorder_round=False) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 5000, in fuse_nodes_once self._try_fusion_pairs( File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4911, in _try_fusion_pairs self.fuse_two_nodes(node1, node2, fused_nodes) File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4731, in fuse_two_nodes node3 = self.get_backend(device).fuse(node1, node2) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 7886, in fuse return FusedExternTritonKernelSchedulerNode.epilogue_fuse(node1, node2) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 2295, in epilogue_fuse assert len(node1.unmet_dependencies) == 1 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ torch._inductor.exc.InductorError: AssertionError:

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: 74% 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 #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 #179803: [Inductor] Handle ordering constraints for user-kernel fusion

Description (problem / solution / changelog)

Fixes: #179232

Use mutation_outputs to retrieve the intermediate buffer. Previously, unmet_dependencies was used, with the assumption that only the mutated buffer is stored. unmet_dependencies may, however, contain more than just the mutated buffer due to ordering constraints/mutation tracking.

Additionally, we add the epilogue source code to the cache key to prevent collisions between different user-kernel variants.

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, +56/-0)
  • torch/_inductor/codegen/wrapper.py (modified, +4/-0)
  • torch/_inductor/scheduler.py (modified, +11/-9)

Code Example

import triton
import triton.language as tl
import torch

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):
    out = torch.empty_like(a)
    add_kernel[(a.numel(),)](a, b, out, a.numel(), BLOCK_SIZE=1)
    out2 = torch.empty_like(a)
    add_kernel[(a.numel(),)](out, b, out2, a.numel(), BLOCK_SIZE=1)
    return out2.relu()

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

---

Traceback (most recent call last):
  File "/home/jjvraw/.../pytorch/repro.py", line 26, in <module>
    compiled(a, b)
    ^^^^^^^^^^^^^^
  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 2612, in _compile_to_module
    self.codegen_with_cpp_wrapper() if self.cpp_wrapper else self.codegen()
                                                             ^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2544, in codegen
    self._update_scheduler()
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2538, in _update_scheduler
    self.scheduler = Scheduler(self.operations)
                     ^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 3086, in __init__
    self._init(nodes)
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 3189, in _init
    self.nodes = self.fuse_nodes(self.nodes)
                 ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4037, in fuse_nodes
    nodes = self.fuse_nodes_once(nodes, is_reorder_round=False)
            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 5000, in fuse_nodes_once
    self._try_fusion_pairs(
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4911, in _try_fusion_pairs
    self.fuse_two_nodes(node1, node2, fused_nodes)
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4731, in fuse_two_nodes
    node3 = self.get_backend(device).fuse(node1, node2)
            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 7886, in fuse
    return FusedExternTritonKernelSchedulerNode.epilogue_fuse(node1, node2)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 2295, in epilogue_fuse
    assert len(node1.unmet_dependencies) == 1
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
torch._inductor.exc.InductorError: AssertionError:

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"

---

python3 collect_env.py
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:                      74%
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 fusion check for user kernel epilogue fusion in Scheduler.can_fuse does not account for buffer reuse decisions made prior to fusion.

import triton
import triton.language as tl
import torch

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):
    out = torch.empty_like(a)
    add_kernel[(a.numel(),)](a, b, out, a.numel(), BLOCK_SIZE=1)
    out2 = torch.empty_like(a)
    add_kernel[(a.numel(),)](out, b, out2, a.numel(), BLOCK_SIZE=1)
    return out2.relu()

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

Error logs

Traceback (most recent call last):
  File "/home/jjvraw/.../pytorch/repro.py", line 26, in <module>
    compiled(a, b)
    ^^^^^^^^^^^^^^
  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 2612, in _compile_to_module
    self.codegen_with_cpp_wrapper() if self.cpp_wrapper else self.codegen()
                                                             ^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2544, in codegen
    self._update_scheduler()
  File "/home/jjvraw/.../pytorch/torch/_inductor/graph.py", line 2538, in _update_scheduler
    self.scheduler = Scheduler(self.operations)
                     ^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 3086, in __init__
    self._init(nodes)
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 3189, in _init
    self.nodes = self.fuse_nodes(self.nodes)
                 ^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4037, in fuse_nodes
    nodes = self.fuse_nodes_once(nodes, is_reorder_round=False)
            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 5000, in fuse_nodes_once
    self._try_fusion_pairs(
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4911, in _try_fusion_pairs
    self.fuse_two_nodes(node1, node2, fused_nodes)
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 4731, in fuse_two_nodes
    node3 = self.get_backend(device).fuse(node1, node2)
            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 7886, in fuse
    return FusedExternTritonKernelSchedulerNode.epilogue_fuse(node1, node2)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/jjvraw/.../pytorch/torch/_inductor/scheduler.py", line 2295, in epilogue_fuse
    assert len(node1.unmet_dependencies) == 1
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
torch._inductor.exc.InductorError: AssertionError:

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

python3 collect_env.py
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:                      74%
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 @voznesenskym @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @aakhundov @coconutruben @jataylo

extent analysis

TL;DR

The issue is likely due to the fusion check in Scheduler.can_fuse not accounting for buffer reuse decisions, causing an InductorError when compiling a PyTorch model.

Guidance

  • The error occurs when the Scheduler attempts to fuse two nodes, but the can_fuse method does not correctly handle buffer reuse.
  • To mitigate this issue, you can try disabling epilogue fusion by setting torch._inductor.config.epilogue_fusion_user_defined_triton_kernel to False.
  • Alternatively, you can attempt to modify the Scheduler.can_fuse method to correctly account for buffer reuse decisions.
  • Verify that the issue is resolved by re-running the PyTorch model compilation with the suggested changes.

Example

torch._inductor.config.epilogue_fusion_user_defined_triton_kernel = False

Notes

  • The provided code snippet and error logs suggest that the issue is related to the fusion of nodes in the PyTorch compiler.
  • The InductorError exception is raised when the Scheduler attempts to fuse two nodes, indicating a potential issue with the fusion logic.
  • Disabling epilogue fusion may impact performance, so it is recommended to investigate and resolve the underlying issue with buffer reuse decisions.

Recommendation

Apply workaround: Disable epilogue fusion by setting torch._inductor.config.epilogue_fusion_user_defined_triton_kernel to False, as this may allow the PyTorch model to compile successfully.

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#agent execution#callback error#memory management#API rate limit

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 fusion incorrectly attempted after buffer reuse decisions [5 pull requests, 1 comments, 2 participants]