pytorch - ✅(Solved) Fix [Inductor] User kernel pointer arg with no `tl.load`/`tl.store` incorrectly eliminated when epilogue fusion enabled [1 pull requests, 1 participants]

pytorch2026-04-29 13:10:48

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.

GitHub issue URL

Helpful · Quick feedback

GitHub stats

pytorch/pytorch#181864•Fetched 2026-04-30 06:18:07

View on GitHub

Comments

Participants

Timeline

Reactions

Author

jjvraw

Participants

jjvraw

Timeline (top)

cross-referenced ×2mentioned ×2subscribed ×2labeled ×1

Error Message

W0429 14:34:16.841000 648526 torch/_inductor/utils.py:1390] on error, temporary cache dir kept at /tmp/tmppc6o3c0z Traceback (most recent call last): File "/home/.../pytorch/min_repro.py", line 31, in <module> print(compiled(x)) ^^^^^^^^^^^ File "/home/.../pytorch/torch/_dynamo/eval_frame.py", line 1095, in compile_wrapper result = fn(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^ File "/home/.../pytorch/min_repro.py", line 20, in fn def fn(x): File "/home/.../pytorch/torch/_dynamo/eval_frame.py", line 1353, in _fn return fn(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^ File "/home/.../pytorch/torch/_functorch/aot_autograd.py", line 1273, in forward return compiled_fn(full_args) ^^^^^^^^^^^^^^^^^^^^^^ File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 939, in runtime_wrapper all_outs = compiled_invoker.run(args, on_before_call=exit_prologue) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 537, in run return call_func_at_runtime_with_args( ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/.../pytorch/torch/_functorch/_aot_autograd/utils.py", line 126, in call_func_at_runtime_with_args out = normalize_as_list(f(args)) ^^^^^^^ File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 2604, in call return self.compiled_fn(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/home/.../pytorch/torch/_inductor/output_code.py", line 725, in call return self.current_callable(inputs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/tmp/tmppc6o3c0z/fa/cfarxhfy46l3rqfx6lhbjdfoo7b345ff5d776q6zccqkyodubqok.py", line 91, in call kernel_0.run(arg0_1, buf3, buf1, 4, 1, 1, stream=stream0) ^^^^ NameError: name 'buf1' is not defined. Did you mean: 'buf3'?

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: 69% 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 #181868: [Inductor] Remove `UserDefinedTritonKernel.get_read_writes` override

Repository: pytorch/pytorch
Author: jjvraw
State: open | merged: False
Link: https://github.com/pytorch/pytorch/pull/181868

Description (problem / solution / changelog)

It is unclear why we need this override. InputsKernel.get_read_writes already produces correct reads/writes from self.inputs and self.get_outputs(), and fusion logic does not rely upon it.

Fixes: #181864

Test case is from the related issue above.

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, +27/-0)
torch/_inductor/ir.py (modified, +0/-38)

Code Example

import torch
import triton
import triton.language as tl

torch._inductor.config.epilogue_fusion_user_defined_triton_kernel = True # causes crash

@triton.jit
def kernel(
    X,
    Y,
    unused, 
    n: tl.constexpr,
    BLOCK: tl.constexpr,
):
    idx = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
    mask = idx < n
    x = tl.load(X + idx, mask=mask)
    tl.store(Y + idx, x * 2, mask=mask)

def fn(x):
    y = torch.empty_like(x)
    unused = torch.empty(x.numel(), device="cuda", dtype=torch.float32)
    kernel[(triton.cdiv(x.numel(), 1024),)](
        x, y, unused, x.numel(), BLOCK=1024
    )
    return y.relu()

x = torch.randn(4096, device="cuda")
with torch._inductor.utils.fresh_inductor_cache():
    compiled = torch.compile(fn)
    print(compiled(x))

---

W0429 14:34:16.841000 648526 torch/_inductor/utils.py:1390] on error, temporary cache dir kept at /tmp/tmppc6o3c0z
Traceback (most recent call last):
  File "/home/.../pytorch/min_repro.py", line 31, in <module>
    print(compiled(x))
          ^^^^^^^^^^^
  File "/home/.../pytorch/torch/_dynamo/eval_frame.py", line 1095, in compile_wrapper
    result = fn(*args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/min_repro.py", line 20, in fn
    def fn(x):
  File "/home/.../pytorch/torch/_dynamo/eval_frame.py", line 1353, in _fn
    return fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/aot_autograd.py", line 1273, in forward
    return compiled_fn(full_args)
           ^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 939, in runtime_wrapper
    all_outs = compiled_invoker.run(args, on_before_call=exit_prologue)
               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 537, in run
    return call_func_at_runtime_with_args(
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/utils.py", line 126, in call_func_at_runtime_with_args
    out = normalize_as_list(f(args))
                            ^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 2604, in __call__
    return self.compiled_fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_inductor/output_code.py", line 725, in __call__
    return self.current_callable(inputs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/tmp/tmppc6o3c0z/fa/cfarxhfy46l3rqfx6lhbjdfoo7b345ff5d776q6zccqkyodubqok.py", line 91, in call
    kernel_0.run(arg0_1, buf3, buf1, 4, 1, 1, stream=stream0)
                               ^^^^
NameError: name 'buf1' is not defined. Did you mean: 'buf3'?

---

PyTorch version: 2.13.0a0+git70f40a4
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:                      69%
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.13.0a0+git70f40a4
[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.13.0a0+git70f40a4  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

As title.

Issue lies in UserDefinedTritonKernel.get_read_writes override (introduced in #173662). Read/writes are now derived from tl.store/tl.load ops directly. Tensors without a corresponding store/load are omitted from reads, eliminated during dead-node elimination, and not allocated. I believe use of TMA descriptors will produce the same error as the repro.

https://github.com/pytorch/pytorch/pull/176832

Min Repro

import torch
import triton
import triton.language as tl

torch._inductor.config.epilogue_fusion_user_defined_triton_kernel = True # causes crash

@triton.jit
def kernel(
    X,
    Y,
    unused, 
    n: tl.constexpr,
    BLOCK: tl.constexpr,
):
    idx = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
    mask = idx < n
    x = tl.load(X + idx, mask=mask)
    tl.store(Y + idx, x * 2, mask=mask)

def fn(x):
    y = torch.empty_like(x)
    unused = torch.empty(x.numel(), device="cuda", dtype=torch.float32)
    kernel[(triton.cdiv(x.numel(), 1024),)](
        x, y, unused, x.numel(), BLOCK=1024
    )
    return y.relu()

x = torch.randn(4096, device="cuda")
with torch._inductor.utils.fresh_inductor_cache():
    compiled = torch.compile(fn)
    print(compiled(x))

Error logs

W0429 14:34:16.841000 648526 torch/_inductor/utils.py:1390] on error, temporary cache dir kept at /tmp/tmppc6o3c0z
Traceback (most recent call last):
  File "/home/.../pytorch/min_repro.py", line 31, in <module>
    print(compiled(x))
          ^^^^^^^^^^^
  File "/home/.../pytorch/torch/_dynamo/eval_frame.py", line 1095, in compile_wrapper
    result = fn(*args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/min_repro.py", line 20, in fn
    def fn(x):
  File "/home/.../pytorch/torch/_dynamo/eval_frame.py", line 1353, in _fn
    return fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/aot_autograd.py", line 1273, in forward
    return compiled_fn(full_args)
           ^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 939, in runtime_wrapper
    all_outs = compiled_invoker.run(args, on_before_call=exit_prologue)
               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 537, in run
    return call_func_at_runtime_with_args(
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/utils.py", line 126, in call_func_at_runtime_with_args
    out = normalize_as_list(f(args))
                            ^^^^^^^
  File "/home/.../pytorch/torch/_functorch/_aot_autograd/runtime_wrappers.py", line 2604, in __call__
    return self.compiled_fn(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/.../pytorch/torch/_inductor/output_code.py", line 725, in __call__
    return self.current_callable(inputs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/tmp/tmppc6o3c0z/fa/cfarxhfy46l3rqfx6lhbjdfoo7b345ff5d776q6zccqkyodubqok.py", line 91, in call
    kernel_0.run(arg0_1, buf3, buf1, 4, 1, 1, stream=stream0)
                               ^^^^
NameError: name 'buf1' is not defined. Did you mean: 'buf3'?

Versions

PyTorch version: 2.13.0a0+git70f40a4
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:                      69%
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.13.0a0+git70f40a4
[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.13.0a0+git70f40a4  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 introduction of UserDefinedTritonKernel.get_read_writes override in #173662, which omits tensors without a corresponding store/load, and can be fixed by adjusting the get_read_writes method to include these tensors.

Guidance

Review the UserDefinedTritonKernel.get_read_writes method to ensure it correctly handles tensors without a corresponding store/load.
Verify that the unused tensor is properly allocated and passed to the kernel function.
Check the PyTorch and Triton versions for any known issues or updates that may address this problem.
Consider reverting the changes introduced in #173662 to see if it resolves the issue.

Example

No specific code example is provided as the issue seems to be related to the internal implementation of UserDefinedTritonKernel.get_read_writes. However, reviewing the get_read_writes method and ensuring it correctly handles all tensors, including those without a corresponding store/load, may help resolve the issue.

Notes

The issue may be specific to the PyTorch and Triton versions used, and updating or downgrading these libraries may help resolve the problem. Additionally, the introduction of UserDefinedTritonKernel.get_read_writes override in #173662 may have unintended consequences, and reverting these changes may be necessary to resolve the issue.

Recommendation

Apply a workaround by adjusting the get_read_writes method to include tensors without a corresponding store/load, as this is likely the root cause of the issue. If this is not possible, consider reverting the changes introduced in #173662 to see if it resolves the issue.

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?

#api #tokenizer error #prompt formatting #chain error #conversation history

Still need to ship something?

×6

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

Data

Security

Network

Code

UI/UX

Text

System

Multimedia

Protocol

API

Engineering

pytorch - ✅(Solved) Fix [Inductor] User kernel pointer arg with no `tl.load`/`tl.store` incorrectly eliminated when epilogue fusion enabled [1 pull requests, 1 participants]

Recommended Tools

GitHub issue graph ai analysis

Error Message

Fix Action

Fix / Workaround

PR fix notes

PR #181868: [Inductor] Remove `UserDefinedTritonKernel.get_read_writes` override

Description (problem / solution / changelog)

Changed files

Code Example

🐛 Describe the bug

Min Repro

Error logs

Versions

extent analysis

TL;DR

Guidance

Example

Notes

Recommendation

Still need to ship something?

TRENDING

pytorch - ✅(Solved) Fix [Inductor] User kernel pointer arg with no `tl.load`/`tl.store` incorrectly eliminated when epilogue fusion enabled [1 pull requests, 1 participants]

Recommended Tools

GitHub issue graph ai analysis

Error Message

Fix Action

Fix / Workaround

PR fix notes

PR #181868: [Inductor] Remove UserDefinedTritonKernel.get_read_writes override

Description (problem / solution / changelog)

Changed files

Code Example

🐛 Describe the bug

Min Repro

Error logs

Versions

extent analysis

TL;DR

Guidance

Example

Notes

Recommendation

Still need to ship something?

RELATED_DISCOVERY

TRENDING

PR #181868: [Inductor] Remove `UserDefinedTritonKernel.get_read_writes` override