Repository: pytorch
Issue: misaligned address without persistent reductions from test_selecsls42b_misaligned_address
Labels: triaged,oncall: pt2,module: inductor,upstream triton

```
CUDA_LAUNCH_BLOCKING=1 TORCHINDUCTOR_PERSISTENT_REDUCTIONS=0 pytest test/inductor/test_cuda_repro.py -k test_selecsls42b_misaligned_address
```
Fails with:
```
RuntimeError: Triton Error [CUDA]: misaligned address
```

The failing kernel code is:
```py
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, DeviceProperties
triton_helpers.set_driver_to_gpu()

@triton_heuristics.reduction(
    size_hints=[1024, 128],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {'in_ptr0': '*i1', 'in_ptr1': '*fp32', 'in_ptr2': '*fp16', 'in_ptr3': '*fp32', 'in_ptr4': '*fp32', 'out_ptr0': '*fp32', 'out_ptr1': '*fp32', 'xnumel': 'i32', 'rnumel': 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=86, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=1536, multi_processor_count=82, warp_size=32), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5, 6, 7, 8), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_red_fused_convert_element_type_div_mul_sub_sum_where_0', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'num_load': 7, 'num_reduction': 2, 'backend_hash': '9EFB47F39A9F0EC578D4E029967452062C70DD8478024529C6C592E836AC91B7', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': False, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False}
)
@triton.jit
def triton_red_fused_convert_element_type_div_mul_sub_sum_where_0(in_ptr0, in_ptr1, in_ptr2, in_ptr3, in_ptr4, out_ptr0, out_ptr1, xnumel, rnumel, XBLOCK : tl.constexpr, RBLOCK : tl.constexpr):
    xnumel = 1024
    rnumel = 128
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    block_ptr0 = tl.make_block_ptr(in_ptr0, shape=[1024, 8, 16], strides=[16, 16384, 1], block_shape=[XBLOCK, ((15 + RBLOCK) // 16), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))], order=[2, 1, 0], offsets=[xoffset, 0, 0])
    tmp1 = tl.load(in_ptr1 + (0))
    tmp2 = tl.broadcast_to(tmp1, [XBLOCK, RBLOCK])
    block_ptr1 = tl.make_block_ptr(in_ptr2, shape=[1024, 8], strides=[1, 1024], block_shape=[XBLOCK, ((15 + RBLOCK) // 16)], order=[1, 0], offsets=[xoffset, 0])
    block_ptr2 = tl.make_block_ptr(in_ptr3, shape=[1024, 8, 16], strides=[16, 16384, 1], block_shape=[XBLOCK, ((15 + RBLOCK) // 16), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))], order=[2, 1, 0], offsets=[xoffset, 0, 0])
    tmp10 = tl.load(tl.make_block_ptr(in_ptr4, shape=[1024], strides=[1], block_shape=[XBLOCK], order=[0], offsets=[xoffset]), boundary_check=[0], eviction_policy='evict_last')[:, None]
    _tmp14 = tl.full([XBLOCK, RBLOCK], 0, tl.float32)
    block_ptr3 = tl.make_block_ptr(in_ptr0, shape=[1024, 8, 16], strides=[16, 16384, 1], block_shape=[XBLOCK, ((15 + RBLOCK) // 16), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))], order=[2, 1, 0], offsets=[xoffset, 0, 0])
    block_ptr4 = tl.make_block_ptr(in_ptr2, shape=[1024, 8], strides=[1, 1024], block_shape=[XBLOCK, ((15 + RBLOCK) // 16)], order=[1, 0], offsets=[xoffset, 0])
    _tmp22 = tl.full([XBLOCK, RBLOCK], 0, tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r1 = rindex % 16
        r2 = (rindex // 16)
        tmp0 = tl.reshape(tl.broadcast_to(tl.load(block_ptr0, boundary_check=[0, 1], padding_option='zero', eviction_policy='evict_last')[:, :, None, :], [XBLOCK, ((15 + RBLOCK) // 16), ((1) * ((1) <= (((15 + RBLOCK) // 16))) + (((15 + RBLOCK) // 16)) * ((((15 + RBLOCK) // 16)) < (1))), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))]), [XBLOCK, RBLOCK]).to(tl.int1)
        block_ptr0 = tl.advance(block_ptr0, [0, (RBLOCK // 16), RBLOCK % 16])
        tmp4 = tl.reshape(tl.broadcast_to(tl.load(block_ptr1, boundary_check=[0, 1], padding_option='zero', eviction_policy='evict_last')[:, :, None, None], [XBLOCK, ((15 + RBLOCK) // 16), ((1) * ((1) <= (((15 + RBLOCK) // 16))) + (((15 + RBLOCK) // 16)) * ((((15 + RBLOCK) // 16)) < (1))), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))]), [XBLOCK, RBLOCK]).to(tl.float32)
        block_ptr1 = tl.advance(block_ptr1, [0, (RBLOCK // 16)])
        tmp9 = tl.reshape(tl.broadcast_to(tl.load(block_ptr2, boundary_check=[0, 1], padding_option='zero', eviction_policy='evict_first')[:, :, None, :], [XBLOCK, ((15 + RBLOCK) // 16), ((1) * ((1) <= (((15 + RBLOCK) // 16))) + (((15 + RBLOCK) // 16)) * ((((15 + RBLOCK) // 16)) < (1))), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))]), [XBLOCK, RBLOCK])
        block_ptr2 = tl.advance(block_ptr2, [0, (RBLOCK // 16), RBLOCK % 16])
        tmp16 = tl.reshape(tl.broadcast_to(tl.load(block_ptr3, boundary_check=[0, 1], padding_option='zero', eviction_policy='evict_first')[:, :, None, :], [XBLOCK, ((15 + RBLOCK) // 16), ((1) * ((1) <= (((15 + RBLOCK) // 16))) + (((15 + RBLOCK) // 16)) * ((((15 + RBLOCK) // 16)) < (1))), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))]), [XBLOCK, RBLOCK]).to(tl.int1)
        block_ptr3 = tl.advance(block_ptr3, [0, (RBLOCK // 16), RBLOCK % 16])
        tmp17 = tl.reshape(tl.broadcast_to(tl.load(block_ptr4, boundary_check=[0, 1], padding_option='zero', eviction_policy='evict_last')[:, :, None, None], [XBLOCK, ((15 + RBLOCK) // 16), ((1) * ((1) <= (((15 + RBLOCK) // 16))) + (((15 + RBLOCK) // 16)) * ((((15 + RBLOCK) // 16)) < (1))), ((16) * ((16) <= (RBLOCK)) + (RBLOCK) * ((RBLOCK) < (16)))]), [XBLOCK, RBLOCK]).to(tl.float32)
        block_ptr4 = tl.advance(block_ptr4, [0, (RBLOCK // 16)])
        tmp3 = tmp2.to(tl.float32)
        tmp5 = 0.0625
        tmp6 = tmp4 * tmp5
        tmp7 = tl.where(tmp0, tmp3, tmp6)
        tmp8 = tmp7.to(tl.float32)
        tmp11 = tmp9 - tmp10
        tmp12 = tmp8 * tmp11
        tmp13 = tl.broadcast_to(tmp12, [XBLOCK, RBLOCK])
        tmp15 = _tmp14 + tmp13
        _tmp14 = tl.where(rmask & xmask, tmp15, _tmp14)
        tmp18 = tmp17 * tmp5
        tmp19 = tl.where(tmp16, tmp3, tmp18)
        tmp20 = tmp19.to(tl.float32)
        tmp21 = tl.broadcast_to(tmp20, [XBLOCK, RBLOCK])
        tmp23 = _tmp22 + tmp21
        _tmp22 = tl.where(rmask & xmask, tmp23, _tmp22)
    tmp14 = tl.sum(_tmp14, 1)[:, None]
    tmp22 = tl.sum(_tmp22, 1)[:, None]
    tl.store(tl.make_block_ptr(out_ptr0, shape=[1024], strides=[1], block_shape=[XBLOCK], order=[0], offsets=[xoffset]), tl.reshape(tmp14, [XBLOCK]).to(tl.float32), boundary_check=[0])
    tl.store(tl.make_block_ptr(out_ptr1, shape=[1024], strides=[1], block_shape=[XBLOCK], order=[0], offsets=[xoffset]), tl.reshape(tmp22, [XBLOCK]).to(tl.float32), boundary_check=[0])
```
However if you set `TORCHINDUCTOR_PERSISTENT_REDUCTIONS=1` it passes, which results in removing the `for` loop from the generated code.

I suspect this is due to a Triton bug and may be related to https://github.com/triton-lang/triton/issues/2836

@manman-ren @plotfi @bertmaher is this something your team could look at?



cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @yf225 @chenyang78 @kadeng @muchulee8 @ColinPeppler @amjames @desertfire @chauhang @aakhundov @bertmaher @int3 @davidberard98 @nmacchioni @embg @peterbell10