Issue Station
Sign In
Tools hereToolbox

pytorch - 💡(How to fix) Fix [Inductor] Performance regression of phlippe_densenet model due to non-fused Cat kernels from new fusion mode

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…

Error Message

Error logs

Fix Action

Fix / Workaround

CPU: Architecture: x86_64 CPU op-mode(s): 32-bit, 64-bit Address sizes: 46 bits physical, 48 bits virtual Byte Order: Little Endian CPU(s): 24 On-line CPU(s) list: 0-23 Vendor ID: GenuineIntel Model name: Intel(R) Core(TM) Ultra 9 285K CPU family: 6 Model: 198 Thread(s) per core: 1 Core(s) per socket: 1 Socket(s): 24 Stepping: 2 CPU max MHz: 5100.0000 CPU min MHz: 800.0000 BogoMIPS: 7372.80 Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf tsc_known_freq pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault intel_ppin ssbd ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid rdt_a rdseed adx smap clflushopt clwb intel_pt sha_ni xsaveopt xsavec xgetbv1 xsaves split_lock_detect user_shstk avx_vnni lam wbnoinvd dtherm ida arat pln pts hwp hwp_notify hwp_act_window hwp_epp hwp_pkg_req hfi vnmi umip pku ospke waitpkg gfni vaes vpclmulqdq tme rdpid bus_lock_detect movdiri movdir64b fsrm md_clear serialize arch_lbr ibt flush_l1d arch_capabilities Virtualization: VT-x L1d cache: 768 KiB (20 instances) L1i cache: 1.3 MiB (20 instances) L2 cache: 40 MiB (12 instances) L3 cache: 36 MiB (1 instance) NUMA node(s): 1 NUMA node0 CPU(s): 0-23 Vulnerability Gather data sampling: 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: Not affected 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; RSB filling; PBRSB-eIBRS Not affected; BHI Not affected Vulnerability Srbds: Not affected Vulnerability Tsx async abort: Not affected

Code Example

python \
benchmarks/dynamo/torchbench.py \
--performance \
--bfloat16 \
--inference \
-d cuda \
-n 10 \
--only phlippe_densenet \
--backend=inductor \
--cold-start-latency \
--timeout 10800 \
--disable-cudagraphs \
--output test.csv

---

buf7 = extern_kernels.convolution(buf5, buf6, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
            assert_size_stride(buf7, (128, 16, 32, 32), (16384, 1, 512, 16), 'torch.ops.aten.convolution.default')
            del buf5
            assert_size_stride(arg13_1, (48, ), (1, ))
            assert_size_stride(arg14_1, (48, ), (1, ))
            assert_size_stride(arg15_1, (48, ), (1, ))
            assert_size_stride(arg16_1, (48, ), (1, ))
            buf9 = empty_strided_cuda((128, 48, 32, 32), (49152, 1, 1536, 48), torch.bfloat16)
            # Topologically Sorted Source Nodes: [input_1, out, input_8, input_9], Original ATen: [aten.convolution, aten.cat, aten._native_batch_norm_legit_no_training, aten.relu]
            # [Provenance debug handles] triton_poi_fused__native_batch_norm_legit_no_training_cat_convolution_relu_5:9
            raw_stream0 = get_raw_stream(0)
            triton_poi_fused__native_batch_norm_legit_no_training_cat_convolution_relu_5.run(buf7, buf2, arg1_1, arg13_1, arg14_1, arg15_1, arg16_1, buf9, 6291456, stream=raw_stream0)
            del arg13_1
            del arg14_1
            del arg15_1
            del arg16_1
            assert_size_stride(arg17_1, (32, 48, 1, 1), (48, 1, 1, 1))
            # Topologically Sorted Source Nodes: [input_8, input_9, input_10], Original ATen: [aten._native_batch_norm_legit_no_training, aten.relu, aten.convolution]
            # [Provenance debug handles] extern_kernels.convolution:10
            buf10 = extern_kernels.convolution(buf9, arg17_1, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)

---

@triton.jit
def triton_poi_fused__native_batch_norm_legit_no_training_cat_convolution_relu_5(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, in_ptr5, in_ptr6, out_ptr1, xnumel, XBLOCK : tl.constexpr):
    xnumel = 6291456
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)[:]
    x0 = (xindex % 48)
    x1 = xindex // 48
    x2 = xindex
    tmp16 = tl.load(in_ptr3 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp19 = tl.load(in_ptr4 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp28 = tl.load(in_ptr5 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp31 = tl.load(in_ptr6 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp0 = x0
    tmp1 = tl.full([1], 0, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = tl.full([1], 16, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (16*x1 + (x0)), tmp4, eviction_policy='evict_last', other=0.0).to(tl.float32)
    tmp6 = tmp0 >= tmp3
    tmp7 = tl.full([1], 48, tl.int64)
    tmp8 = tmp0 < tmp7
    tmp9 = tl.load(in_ptr1 + (32*x1 + ((-16) + x0)), tmp6, eviction_policy='evict_last', other=0.0).to(tl.float32)
    tmp10 = tl.load(in_ptr2 + ((-16) + x0), tmp6, eviction_policy='evict_last', other=0.0).to(tl.float32)
    tmp11 = tmp9 + tmp10
    tmp12 = tl.full(tmp11.shape, 0.0, tmp11.dtype)
    tmp13 = tl.where(tmp6, tmp11, tmp12)
    tmp14 = tl.where(tmp4, tmp5, tmp13)
    tmp15 = tmp14.to(tl.float32)
    tmp17 = tmp16.to(tl.float32)
    tmp18 = tmp15 - tmp17
    tmp20 = tmp19.to(tl.float32)
    tmp21 = tl.full([1], 1e-05, tl.float32)
    tmp22 = tmp20 + tmp21
    tmp23 = tl.sqrt_rn(tmp22)
    tmp24 = tl.full([1], 1.0, tl.float32)
    tmp25 = (tmp24 / tmp23)
    tmp26 = tmp25 * tmp24
    tmp27 = tmp18 * tmp26
    tmp29 = tmp28.to(tl.float32)
    tmp30 = tmp27 * tmp29
    tmp32 = tmp31.to(tl.float32)
    tmp33 = tmp30 + tmp32
    tmp34 = tmp33.to(tl.float32)
    tmp35 = tl.full([1], 0, tl.int32)
    tmp36 = triton_helpers.maximum(tmp35, tmp34)
    tl.store(out_ptr1 + (x2), tmp36, None)

---

buf7 = extern_kernels.convolution(buf5, buf6, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
            assert_size_stride(buf7, (128, 16, 32, 32), (16384, 1, 512, 16), 'torch.ops.aten.convolution.default')
            del buf5
            buf8 = reinterpret_tensor(buf10, (128, 16, 32, 32), (49152, 1, 1536, 48), 0)  # alias
            # Topologically Sorted Source Nodes: [input_1, out], Original ATen: [aten.convolution, aten.cat]
            # [Provenance debug handles] triton_poi_fused_cat_convolution_5:9
            raw_stream0 = get_raw_stream(0)
            triton_poi_fused_cat_convolution_5.run(buf7, buf8, 2097152, stream=raw_stream0)
            del buf7
            assert_size_stride(arg13_1, (48, ), (1, ))
            assert_size_stride(arg14_1, (48, ), (1, ))
            assert_size_stride(arg15_1, (48, ), (1, ))
            assert_size_stride(arg16_1, (48, ), (1, ))
            buf11 = empty_strided_cuda((128, 48, 32, 32), (49152, 1, 1536, 48), torch.bfloat16)
            buf18 = empty_strided_cuda((128, 64, 32, 32), (65536, 1, 2048, 64), torch.bfloat16)
            buf17 = reinterpret_tensor(buf18, (128, 48, 32, 32), (65536, 1, 2048, 64), 16)  # alias
            # Topologically Sorted Source Nodes: [input_8, input_9, out_1], Original ATen: [aten._native_batch_norm_legit_no_training, aten.relu, aten.cat]
            # [Provenance debug handles] triton_poi_fused__native_batch_norm_legit_no_training_cat_relu_6:10
            raw_stream0 = get_raw_stream(0)
            triton_poi_fused__native_batch_norm_legit_no_training_cat_relu_6.run(buf10, arg13_1, arg14_1, arg15_1, arg16_1, buf11, buf17, 6291456, stream=raw_stream0)
            del arg13_1
            del arg14_1
            del arg15_1
            del arg16_1
            del buf10
            del buf8
            del buf9
            assert_size_stride(arg17_1, (32, 48, 1, 1), (48, 1, 1, 1))
            # Topologically Sorted Source Nodes: [input_8, input_9, input_10], Original ATen: [aten._native_batch_norm_legit_no_training, aten.relu, aten.convolution]
            # [Provenance debug handles] extern_kernels.convolution:11
            buf12 = extern_kernels.convolution(buf11, arg17_1, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)

---

@triton.jit
def triton_poi_fused_cat_convolution_5(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 2097152
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)[:]
    x2 = xindex
    x0 = (xindex % 16)
    x1 = xindex // 16
    tmp0 = tl.load(in_ptr0 + (x2), None).to(tl.float32)
    tl.store(out_ptr0 + (x0 + 48*x1), tmp0, None)

---

@triton.jit
def triton_poi_fused__native_batch_norm_legit_no_training_cat_relu_6(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, out_ptr1, xnumel, XBLOCK : tl.constexpr):
    xnumel = 6291456
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)[:]
    x2 = xindex
    x0 = (xindex % 48)
    x1 = xindex // 48
    tmp0 = tl.load(in_ptr0 + (x2), None).to(tl.float32)
    tmp2 = tl.load(in_ptr1 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp5 = tl.load(in_ptr2 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp14 = tl.load(in_ptr3 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp17 = tl.load(in_ptr4 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp1 = tmp0.to(tl.float32)
    tmp3 = tmp2.to(tl.float32)
    tmp4 = tmp1 - tmp3
    tmp6 = tmp5.to(tl.float32)
    tmp7 = tl.full([1], 1e-05, tl.float32)
    tmp8 = tmp6 + tmp7
    tmp9 = tl.sqrt_rn(tmp8)
    tmp10 = tl.full([1], 1.0, tl.float32)
    tmp11 = (tmp10 / tmp9)
    tmp12 = tmp11 * tmp10
    tmp13 = tmp4 * tmp12
    tmp15 = tmp14.to(tl.float32)
    tmp16 = tmp13 * tmp15
    tmp18 = tmp17.to(tl.float32)
    tmp19 = tmp16 + tmp18
    tmp20 = tmp19.to(tl.float32)
    tmp21 = tl.full([1], 0, tl.int32)
    tmp22 = triton_helpers.maximum(tmp21, tmp20)
    tl.store(out_ptr0 + (x2), tmp22, None)
    tl.store(out_ptr1 + (x0 + 64*x1), tmp0, None)

---

dev,name,batch_size,speedup,abs_latency,compilation_latency,compression_ratio,eager_peak_mem,dynamo_peak_mem,calls_captured,unique_graphs,graph_breaks,unique_graph_breaks,autograd_captures,autograd_compiles,cudagraph_skips
cuda,phlippe_densenet,128,1.712117,1.531227,0.927166,1.396972,0.120151,0.086008,186,1,0,0,0,0,0

---

dev,name,batch_size,speedup,abs_latency,compilation_latency,compression_ratio,eager_peak_mem,dynamo_peak_mem,calls_captured,unique_graphs,graph_breaks,unique_graph_breaks,autograd_captures,autograd_compiles,cudagraph_skips
cuda,phlippe_densenet,128,1.419691,1.784112,6.071082,0.707248,0.120151,0.169885,186,1,0,0,0,0,0
RAW_BUFFERClick to expand / collapse

🐛 Describe the bug

On 4080super GPUs, the phlippe_densenet model exhibited a performance regression, with absolute latency rising by 16.5%.

The regression started appearing after a recent inductor change that introduced a special Cat kernel fusion mode (#175729).

Reproduce

Use torchbench

python \
benchmarks/dynamo/torchbench.py \
--performance \
--bfloat16 \
--inference \
-d cuda \
-n 10 \
--only phlippe_densenet \
--backend=inductor \
--cold-start-latency \
--timeout 10800 \
--disable-cudagraphs \
--output test.csv

Before (fused triton_poi_fused__native_batch_norm_legit_no_training_cat_convolution_relu)

buf7 = extern_kernels.convolution(buf5, buf6, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
            assert_size_stride(buf7, (128, 16, 32, 32), (16384, 1, 512, 16), 'torch.ops.aten.convolution.default')
            del buf5
            assert_size_stride(arg13_1, (48, ), (1, ))
            assert_size_stride(arg14_1, (48, ), (1, ))
            assert_size_stride(arg15_1, (48, ), (1, ))
            assert_size_stride(arg16_1, (48, ), (1, ))
            buf9 = empty_strided_cuda((128, 48, 32, 32), (49152, 1, 1536, 48), torch.bfloat16)
            # Topologically Sorted Source Nodes: [input_1, out, input_8, input_9], Original ATen: [aten.convolution, aten.cat, aten._native_batch_norm_legit_no_training, aten.relu]
            # [Provenance debug handles] triton_poi_fused__native_batch_norm_legit_no_training_cat_convolution_relu_5:9
            raw_stream0 = get_raw_stream(0)
            triton_poi_fused__native_batch_norm_legit_no_training_cat_convolution_relu_5.run(buf7, buf2, arg1_1, arg13_1, arg14_1, arg15_1, arg16_1, buf9, 6291456, stream=raw_stream0)
            del arg13_1
            del arg14_1
            del arg15_1
            del arg16_1
            assert_size_stride(arg17_1, (32, 48, 1, 1), (48, 1, 1, 1))
            # Topologically Sorted Source Nodes: [input_8, input_9, input_10], Original ATen: [aten._native_batch_norm_legit_no_training, aten.relu, aten.convolution]
            # [Provenance debug handles] extern_kernels.convolution:10
            buf10 = extern_kernels.convolution(buf9, arg17_1, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
@triton.jit
def triton_poi_fused__native_batch_norm_legit_no_training_cat_convolution_relu_5(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, in_ptr5, in_ptr6, out_ptr1, xnumel, XBLOCK : tl.constexpr):
    xnumel = 6291456
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)[:]
    x0 = (xindex % 48)
    x1 = xindex // 48
    x2 = xindex
    tmp16 = tl.load(in_ptr3 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp19 = tl.load(in_ptr4 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp28 = tl.load(in_ptr5 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp31 = tl.load(in_ptr6 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp0 = x0
    tmp1 = tl.full([1], 0, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = tl.full([1], 16, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (16*x1 + (x0)), tmp4, eviction_policy='evict_last', other=0.0).to(tl.float32)
    tmp6 = tmp0 >= tmp3
    tmp7 = tl.full([1], 48, tl.int64)
    tmp8 = tmp0 < tmp7
    tmp9 = tl.load(in_ptr1 + (32*x1 + ((-16) + x0)), tmp6, eviction_policy='evict_last', other=0.0).to(tl.float32)
    tmp10 = tl.load(in_ptr2 + ((-16) + x0), tmp6, eviction_policy='evict_last', other=0.0).to(tl.float32)
    tmp11 = tmp9 + tmp10
    tmp12 = tl.full(tmp11.shape, 0.0, tmp11.dtype)
    tmp13 = tl.where(tmp6, tmp11, tmp12)
    tmp14 = tl.where(tmp4, tmp5, tmp13)
    tmp15 = tmp14.to(tl.float32)
    tmp17 = tmp16.to(tl.float32)
    tmp18 = tmp15 - tmp17
    tmp20 = tmp19.to(tl.float32)
    tmp21 = tl.full([1], 1e-05, tl.float32)
    tmp22 = tmp20 + tmp21
    tmp23 = tl.sqrt_rn(tmp22)
    tmp24 = tl.full([1], 1.0, tl.float32)
    tmp25 = (tmp24 / tmp23)
    tmp26 = tmp25 * tmp24
    tmp27 = tmp18 * tmp26
    tmp29 = tmp28.to(tl.float32)
    tmp30 = tmp27 * tmp29
    tmp32 = tmp31.to(tl.float32)
    tmp33 = tmp30 + tmp32
    tmp34 = tmp33.to(tl.float32)
    tmp35 = tl.full([1], 0, tl.int32)
    tmp36 = triton_helpers.maximum(tmp35, tmp34)
    tl.store(out_ptr1 + (x2), tmp36, None)

After (triton_poi_fused_cat_convolution + triton_poi_fused__native_batch_norm_legit_no_training_cat_relu)

buf7 = extern_kernels.convolution(buf5, buf6, stride=(1, 1), padding=(1, 1), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
            assert_size_stride(buf7, (128, 16, 32, 32), (16384, 1, 512, 16), 'torch.ops.aten.convolution.default')
            del buf5
            buf8 = reinterpret_tensor(buf10, (128, 16, 32, 32), (49152, 1, 1536, 48), 0)  # alias
            # Topologically Sorted Source Nodes: [input_1, out], Original ATen: [aten.convolution, aten.cat]
            # [Provenance debug handles] triton_poi_fused_cat_convolution_5:9
            raw_stream0 = get_raw_stream(0)
            triton_poi_fused_cat_convolution_5.run(buf7, buf8, 2097152, stream=raw_stream0)
            del buf7
            assert_size_stride(arg13_1, (48, ), (1, ))
            assert_size_stride(arg14_1, (48, ), (1, ))
            assert_size_stride(arg15_1, (48, ), (1, ))
            assert_size_stride(arg16_1, (48, ), (1, ))
            buf11 = empty_strided_cuda((128, 48, 32, 32), (49152, 1, 1536, 48), torch.bfloat16)
            buf18 = empty_strided_cuda((128, 64, 32, 32), (65536, 1, 2048, 64), torch.bfloat16)
            buf17 = reinterpret_tensor(buf18, (128, 48, 32, 32), (65536, 1, 2048, 64), 16)  # alias
            # Topologically Sorted Source Nodes: [input_8, input_9, out_1], Original ATen: [aten._native_batch_norm_legit_no_training, aten.relu, aten.cat]
            # [Provenance debug handles] triton_poi_fused__native_batch_norm_legit_no_training_cat_relu_6:10
            raw_stream0 = get_raw_stream(0)
            triton_poi_fused__native_batch_norm_legit_no_training_cat_relu_6.run(buf10, arg13_1, arg14_1, arg15_1, arg16_1, buf11, buf17, 6291456, stream=raw_stream0)
            del arg13_1
            del arg14_1
            del arg15_1
            del arg16_1
            del buf10
            del buf8
            del buf9
            assert_size_stride(arg17_1, (32, 48, 1, 1), (48, 1, 1, 1))
            # Topologically Sorted Source Nodes: [input_8, input_9, input_10], Original ATen: [aten._native_batch_norm_legit_no_training, aten.relu, aten.convolution]
            # [Provenance debug handles] extern_kernels.convolution:11
            buf12 = extern_kernels.convolution(buf11, arg17_1, stride=(1, 1), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1, bias=None)
@triton.jit
def triton_poi_fused_cat_convolution_5(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 2097152
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)[:]
    x2 = xindex
    x0 = (xindex % 16)
    x1 = xindex // 16
    tmp0 = tl.load(in_ptr0 + (x2), None).to(tl.float32)
    tl.store(out_ptr0 + (x0 + 48*x1), tmp0, None)
@triton.jit
def triton_poi_fused__native_batch_norm_legit_no_training_cat_relu_6(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, out_ptr1, xnumel, XBLOCK : tl.constexpr):
    xnumel = 6291456
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)[:]
    x2 = xindex
    x0 = (xindex % 48)
    x1 = xindex // 48
    tmp0 = tl.load(in_ptr0 + (x2), None).to(tl.float32)
    tmp2 = tl.load(in_ptr1 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp5 = tl.load(in_ptr2 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp14 = tl.load(in_ptr3 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp17 = tl.load(in_ptr4 + (x0), None, eviction_policy='evict_last').to(tl.float32)
    tmp1 = tmp0.to(tl.float32)
    tmp3 = tmp2.to(tl.float32)
    tmp4 = tmp1 - tmp3
    tmp6 = tmp5.to(tl.float32)
    tmp7 = tl.full([1], 1e-05, tl.float32)
    tmp8 = tmp6 + tmp7
    tmp9 = tl.sqrt_rn(tmp8)
    tmp10 = tl.full([1], 1.0, tl.float32)
    tmp11 = (tmp10 / tmp9)
    tmp12 = tmp11 * tmp10
    tmp13 = tmp4 * tmp12
    tmp15 = tmp14.to(tl.float32)
    tmp16 = tmp13 * tmp15
    tmp18 = tmp17.to(tl.float32)
    tmp19 = tmp16 + tmp18
    tmp20 = tmp19.to(tl.float32)
    tmp21 = tl.full([1], 0, tl.int32)
    tmp22 = triton_helpers.maximum(tmp21, tmp20)
    tl.store(out_ptr0 + (x2), tmp22, None)
    tl.store(out_ptr1 + (x0 + 64*x1), tmp0, None)

Error logs

The torchbench output result is

  1. Before
dev,name,batch_size,speedup,abs_latency,compilation_latency,compression_ratio,eager_peak_mem,dynamo_peak_mem,calls_captured,unique_graphs,graph_breaks,unique_graph_breaks,autograd_captures,autograd_compiles,cudagraph_skips
cuda,phlippe_densenet,128,1.712117,1.531227,0.927166,1.396972,0.120151,0.086008,186,1,0,0,0,0,0
  1. After
dev,name,batch_size,speedup,abs_latency,compilation_latency,compression_ratio,eager_peak_mem,dynamo_peak_mem,calls_captured,unique_graphs,graph_breaks,unique_graph_breaks,autograd_captures,autograd_compiles,cudagraph_skips
cuda,phlippe_densenet,128,1.419691,1.784112,6.071082,0.707248,0.120151,0.169885,186,1,0,0,0,0,0

Versions

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

OS: Ubuntu 22.04.2 LTS (x86_64) GCC version: (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0 Clang version: Could not collect CMake version: version 4.3.2 Libc version: glibc-2.35

Python version: 3.12.13 | packaged by conda-forge | (main, Mar 5 2026, 16:50:00) [GCC 14.3.0] (64-bit runtime) Python platform: Linux-6.8.0-52-generic-x86_64-with-glibc2.35 Is CUDA available: True CUDA runtime version: 13.0.88 CUDA_MODULE_LOADING set to: GPU models and configuration: GPU 0: NVIDIA GeForce RTX 4080 SUPER GPU 1: NVIDIA GeForce RTX 4070 Ti SUPER

Nvidia driver version: 580.126.20 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: 46 bits physical, 48 bits virtual Byte Order: Little Endian CPU(s): 24 On-line CPU(s) list: 0-23 Vendor ID: GenuineIntel Model name: Intel(R) Core(TM) Ultra 9 285K CPU family: 6 Model: 198 Thread(s) per core: 1 Core(s) per socket: 1 Socket(s): 24 Stepping: 2 CPU max MHz: 5100.0000 CPU min MHz: 800.0000 BogoMIPS: 7372.80 Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf tsc_known_freq pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault intel_ppin ssbd ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid rdt_a rdseed adx smap clflushopt clwb intel_pt sha_ni xsaveopt xsavec xgetbv1 xsaves split_lock_detect user_shstk avx_vnni lam wbnoinvd dtherm ida arat pln pts hwp hwp_notify hwp_act_window hwp_epp hwp_pkg_req hfi vnmi umip pku ospke waitpkg gfni vaes vpclmulqdq tme rdpid bus_lock_detect movdiri movdir64b fsrm md_clear serialize arch_lbr ibt flush_l1d arch_capabilities Virtualization: VT-x L1d cache: 768 KiB (20 instances) L1i cache: 1.3 MiB (20 instances) L2 cache: 40 MiB (12 instances) L3 cache: 36 MiB (1 instance) NUMA node(s): 1 NUMA node0 CPU(s): 0-23 Vulnerability Gather data sampling: 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: Not affected 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; RSB filling; PBRSB-eIBRS Not affected; BHI Not affected Vulnerability Srbds: Not affected Vulnerability Tsx async abort: Not affected

Versions of relevant libraries: [pip3] bert_pytorch==0.0.1a4 [pip3] functorch==1.14.0a0+b71aa0b [pip3] mypy==2.1.0 [pip3] mypy_extensions==1.1.0 [pip3] numpy==1.26.4 [pip3] nvidia-cublas==13.1.1.3 [pip3] nvidia-cuda-cupti==13.0.85 [pip3] nvidia-cuda-nvrtc==13.0.88 [pip3] nvidia-cuda-runtime==13.0.96 [pip3] nvidia-cudnn-cu13==9.20.0.48 [pip3] nvidia-cufft==12.0.0.61 [pip3] nvidia-curand==10.4.0.35 [pip3] nvidia-cusolver==12.0.4.66 [pip3] nvidia-cusparse==12.6.3.3 [pip3] nvidia-cusparselt-cu13==0.8.1 [pip3] nvidia-nccl-cu13==2.29.7 [pip3] nvidia-nvjitlink==13.0.88 [pip3] nvidia-nvtx==13.0.85 [pip3] onnx==1.21.0 [pip3] optree==0.19.1 [pip3] pytorch-labs-segment-anything-fast==0.2 [pip3] torch==2.12.0+cu130 [pip3] torch_geometric==2.4.0 [pip3] torchao==0.17.0 [pip3] torchaudio==2.11.0+cu130 [pip3] torchbench==0.1 [pip3] torchmultimodal==0.1.0b0 [pip3] torchvision==0.27.0+cu130 [pip3] triton==3.7.0 [conda] bert-pytorch 0.0.1a4 pypi_0 pypi [conda] functorch 1.14.0a0+b71aa0b pypi_0 pypi [conda] numpy 1.26.4 pypi_0 pypi [conda] nvidia-cublas 13.1.1.3 pypi_0 pypi [conda] nvidia-cuda-cupti 13.0.85 pypi_0 pypi [conda] nvidia-cuda-nvrtc 13.0.88 pypi_0 pypi [conda] nvidia-cuda-runtime 13.0.96 pypi_0 pypi [conda] nvidia-cudnn-cu13 9.20.0.48 pypi_0 pypi [conda] nvidia-cufft 12.0.0.61 pypi_0 pypi [conda] nvidia-curand 10.4.0.35 pypi_0 pypi [conda] nvidia-cusolver 12.0.4.66 pypi_0 pypi [conda] nvidia-cusparse 12.6.3.3 pypi_0 pypi [conda] nvidia-cusparselt-cu13 0.8.1 pypi_0 pypi [conda] nvidia-nccl-cu13 2.29.7 pypi_0 pypi [conda] nvidia-nvjitlink 13.0.88 pypi_0 pypi [conda] nvidia-nvtx 13.0.85 pypi_0 pypi [conda] optree 0.19.1 pypi_0 pypi [conda] pytorch-labs-segment-anything-fast 0.2 pypi_0 pypi [conda] torch 2.12.0+cu130 pypi_0 pypi [conda] torch-geometric 2.4.0 pypi_0 pypi [conda] torchao 0.17.0 pypi_0 pypi [conda] torchaudio 2.11.0+cu130 pypi_0 pypi [conda] torchbench 0.1 pypi_0 pypi [conda] torchmultimodal 0.1.0b0 pypi_0 pypi [conda] torchvision 0.27.0+cu130 pypi_0 pypi [conda] triton 3.7.0 pypi_0 pypi

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

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…

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 - 💡(How to fix) Fix [Inductor] Performance regression of phlippe_densenet model due to non-fused Cat kernels from new fusion mode