[BUG][mthreads]kernel crashes on device illegal memory access + core dump

## Environment
- Triton 3.6.0 (FlagTree mthreads3.6), MUSA backend, MTT S5000
- torch_musa 2.7.1, MUSA 4.3.5, Python 3.10
## Symptom
A valid Triton kernel is compiled successfully by the MTGPU backend, but the compiled kernel **crashes on device**: the launch +`torch.musa.synchronize()` raises `RuntimeError: MUSA error: an illegal memory access was encountered` and the driver drops a `core_*.mudmp`. All injected code is provably semantics no-op — the branch predicates are compile-time uniform(`opaque_true = 0<1 = True`, `opaque_false = 0>1 = False`), the three template blocks only write to a dedicated scratch("switch") buffer, and the untouched main path is line-for-line the original kernel (its store is merely wrapped in three always-true nested `if opaque_true:` levels). Reproduced **6/6** in the campaign (3 async + 3 blocking) and deterministically in a standalone reproducer.Unlike the sibling issue-000009 (same kernel family, silent wrong output), this variant is **not reducible**: 16 ablation variants (removing any single block / statement / nesting level / duplicate assignment) all stop the crash. The trigger is a brittle combination that shapes the LLIR/scf layout — removing any piece makes the backend take a different codegen path and the crash disappears.

**FlagTree mthreads3.6 output:**
original = SOLVE_OK
original == CPU max_pool2d ref: True
transformed = DEVICE_ERROR: RuntimeError: MUSA error: an illegal memory access was encountered

**CoreDump**:

```
import glob
import sys

import torch
import triton
import triton.language as tl

# functional[3] params of 42_2d_max_pooling
N, C, H, W = 1, 1, 10, 10
KS, STRIDE, PADDING = 5, 2, 2
BLOCK_SIZE_KERNEL = triton.next_power_of_2(KS) # 8
BLOCK_SIZE_OUTPUT = 1024
OUT_SIZE_H = (H + 2 * PADDING - KS) // STRIDE + 1 # 5
OUT_SIZE_W = (W + 2 * PADDING - KS) // STRIDE + 1 # 5
NUM_OUT = OUT_SIZE_H * OUT_SIZE_W # 25

# ============================ original kernel (correct) ============================
@triton.jit
def max_pooling_kernel(
input: tl.tensor, output: tl.tensor,
input_size_h: int, input_size_w: int,
output_size_h: int, output_size_w: int,
kernel_size: int, stride: int, padding: int,
BLOCK_SIZE_KERNEL: tl.constexpr, BLOCK_SIZE_OUTPUT: tl.constexpr,
):
pid = tl.program_id(0)
slice = tl.program_id(1)
output_offset = pid * BLOCK_SIZE_OUTPUT + tl.arange(0, BLOCK_SIZE_OUTPUT)
output_idx_h = output_offset // output_size_w
output_idx_w = output_offset % output_size_w
input_idx_h = output_idx_h[:, None] * stride + (tl.arange(0, BLOCK_SIZE_KERNEL) - padding)[None, :]
input_idx_w = output_idx_w[:, None] * stride + (tl.arange(0, BLOCK_SIZE_KERNEL) - padding)[None, :]
input_ptrs = input_idx_h[:, :, None] * input_size_w + input_idx_w[:, None, :]
input_mask_h = (input_idx_h >= 0) & (input_idx_h < input_size_h)
input_mask_w = (input_idx_w >= 0) & (input_idx_w < input_size_w)
input_mask = input_mask_h[:, :, None] & input_mask_w[:, None, :]
kernel_mask = tl.arange(0, BLOCK_SIZE_KERNEL) < kernel_size
kernel_mask = kernel_mask[:, None] & kernel_mask[None, :]
input_ptrs = input_ptrs + slice * input_size_h * input_size_w
input_data = tl.load(input + input_ptrs, mask=input_mask & kernel_mask, other=-float("inf"))
output_data = tl.max(
tl.reshape(input_data, (BLOCK_SIZE_OUTPUT, BLOCK_SIZE_KERNEL * BLOCK_SIZE_KERNEL)),
axis=1, keep_dims=False,
)
output_mask = output_offset < output_size_h * output_size_w
output_ptr = output_offset + slice * output_size_h * output_size_w
tl.store(output + output_ptr, output_data.to(dtype=output.dtype.element_ty), mask=output_mask)

# ============ variant: original main path + semantics-no-op injection ==============
# opaque prologue; 3 template blocks writing only to the scratch buffer (a16
# histogram block with dead `return`, a17 uint32/uniform-float block, a8
# bit-ops block); duplicate assignments; main store wrapped in 3 nested
# always-true `if opaque_true:` levels. Nothing here executes at runtime
# (opaque_true=0<1=True, opaque_false=0>1=False).
@triton.jit
def _fz_outlined_max_pooling_kernel_9(output_offset, output_size_h, output_size_w):
return output_offset < output_size_h * output_size_w

@triton.jit
def _fz_outlined_max_pooling_kernel_19(opaque_one):
return tl.arange(0, 8) + opaque_one

@triton.jit
def max_pooling_kernel_tx(
_fz_switch_ptr, _fz_switch_output_ptr,
input: tl.tensor, output: tl.tensor,
input_size_h: int, input_size_w: int,
output_size_h: int, output_size_w: int,
kernel_size: int, stride: int, padding: int,
BLOCK_SIZE_KERNEL: tl.constexpr, BLOCK_SIZE_OUTPUT: tl.constexpr,
):
opaque_values = tl.load(_fz_switch_ptr + tl.arange(0, 2))
(opaque_zero, opaque_one) = tl.split(opaque_values)
opaque_true = opaque_zero < opaque_one
opaque_false = opaque_zero > opaque_one
pid = tl.program_id(0)
slice = tl.program_id(1)
output_offset = pid * BLOCK_SIZE_OUTPUT + tl.arange(0, BLOCK_SIZE_OUTPUT)
output_idx_h = output_offset // output_size_w
output_idx_w = output_offset % output_size_w
input_idx_h = output_idx_h[:, None] * stride + (tl.arange(0, BLOCK_SIZE_KERNEL) - padding)[None, :]
input_idx_w = output_idx_w[:, None] * stride + (tl.arange(0, BLOCK_SIZE_KERNEL) - padding)[None, :]
input_ptrs = input_idx_h[:, :, None] * input_size_w + input_idx_w[:, None, :]
input_ptrs = input_idx_h[:, :, None] * input_size_w + input_idx_w[:, None, :]
input_mask_h = (input_idx_h >= 0) & (input_idx_h < input_size_h)
input_ptrs = input_idx_h[:, :, None] * input_size_w + input_idx_w[:, None, :]
input_mask_w = (input_idx_w >= 0) & (input_idx_w < input_size_w)
input_mask = input_mask_h[:, :, None] & input_mask_w[:, None, :]
kernel_mask = tl.arange(0, BLOCK_SIZE_KERNEL) < kernel_size
kernel_mask = kernel_mask[:, None] & kernel_mask[None, :]
input_ptrs = input_ptrs + slice * input_size_h * input_size_w
input_data = tl.load(input + input_ptrs, mask=input_mask & kernel_mask, other=-float('inf'))
output_data = tl.max(tl.reshape(input_data, (BLOCK_SIZE_OUTPUT, BLOCK_SIZE_KERNEL * BLOCK_SIZE_KERNEL)), axis=1,
keep_dims=False)
output_mask = _fz_outlined_max_pooling_kernel_9(output_offset, output_size_h, output_size_w)
_fz_a15_split_23 = output_offset + slice * output_size_h * output_size_w
if opaque_false:
_fz_a16_hs_v = tl.arange(0, 8) & 3
_fz_a16_hs_m = (_fz_a16_hs_v + opaque_one) % 2 == 0
_fz_a16_hs_m = (_fz_a16_hs_v + opaque_one) % 2 == 0
_fz_a16_hs_h = tl.histogram((_fz_a16_hs_v + opaque_one) % 4, 4, mask=_fz_a16_hs_m)
_fz_a16_hs_c = tl.cumsum(_fz_a16_hs_h, axis=0)
_fz_a16_hs_s = tl.sum(_fz_a16_hs_c)
_fz_a16_tpl_result = (_fz_a16_hs_s + opaque_values).to(tl.float32)
tl.store(_fz_switch_output_ptr + 30 + tl.arange(0, 2), _fz_a16_tpl_result, mask=(tl.program_id(axis=0) == 0) &
(tl.program_id(axis=1) == 0) & (tl.program_id(axis=2) == 0))
return
output_ptr = _fz_a15_split_23
output_ptr = output_offset + slice * output_size_h * output_size_w
if opaque_true:
_fz_a17_uc_x = tl.arange(0, 8) + opaque_one
_fz_a17_uc_z = tl.zeros_like(_fz_a17_uc_x.to(tl.float32))
_fz_a17_uc_u = (tl.arange(0, 8) * 100003).to(tl.uint32) + opaque_one.to(tl.uint32)
_fz_a17_uc_f = tl.uint_to_uniform_float(_fz_a17_uc_u)
_fz_a17_uc_w = tl.where(_fz_a17_uc_f < 0.5, _fz_a17_uc_z, _fz_a17_uc_f)
_fz_a17_uc_s = tl.sum(_fz_a17_uc_w)
_fz_a17_tpl_result = _fz_a17_uc_s + opaque_values.to(tl.float32)
tl.store(_fz_switch_output_ptr + 32 + tl.arange(0, 2), _fz_a17_tpl_result, mask=(tl.program_id(axis=0) == 0) &
(tl.program_id(axis=1) == 0) & (tl.program_id(axis=2) == 0))
if opaque_false:
_fz_a8_bits_x = _fz_outlined_max_pooling_kernel_19(opaque_one)
_fz_a8_bits_y = _fz_a8_bits_x << 2 | _fz_a8_bits_x >> 1
_fz_a8_bits_z = _fz_a8_bits_y ^ _fz_a8_bits_x & 3
_fz_a8_bits_w = tl.umulhi(_fz_a8_bits_z, _fz_a8_bits_x)
_fz_a8_bits_s = tl.sum(tl.cdiv(_fz_a8_bits_w + opaque_one, 2))
_fz_a8_tpl_result = (_fz_a8_bits_s + opaque_values).to(tl.float32)
tl.store(_fz_switch_output_ptr + 14 + tl.arange(0, 2), _fz_a8_tpl_result, mask=(tl.program_id(axis=0) == 0) &
(tl.program_id(axis=1) == 0) & (tl.program_id(axis=2) == 0))
_fz_a2_loop_once = 0
if opaque_true:
_fz_a4_loop_once = 0
if opaque_true:
if opaque_true:
_fz_a18_split_0 = 0
_fz_a5_loop_once = _fz_a18_split_0
_fz_a5_loop_once = 0
else:
_fz_a5_loop_once = 0
if opaque_true:
tl.store(output + output_ptr, output_data.to(dtype=output.dtype.element_ty), mask=output_mask)
else:
pass
else:
pass
else:
pass

def main() -> int:
torch.manual_seed(0)
input_t = torch.randn(N * C * H * W, dtype=torch.float32, device="musa")
orig = torch.zeros(NUM_OUT, dtype=torch.float32, device="musa")
grid = (triton.cdiv(NUM_OUT, BLOCK_SIZE_OUTPUT), N * C) # (1, 1)
max_pooling_kernel[grid](input_t, orig, H, W, OUT_SIZE_H, OUT_SIZE_W, KS, STRIDE, PADDING,
BLOCK_SIZE_KERNEL=BLOCK_SIZE_KERNEL, BLOCK_SIZE_OUTPUT=BLOCK_SIZE_OUTPUT)
torch.musa.synchronize()
print("original = SOLVE_OK")

switch_ptr = torch.tensor([0, 1], dtype=torch.int32, device="musa")
switch_out = torch.empty(400, dtype=torch.float32, device="musa")
before = set(glob.glob("core_*.mudmp"))
dev_err = ""
try:
max_pooling_kernel_tx[grid](switch_ptr, switch_out, input_t,
torch.zeros(NUM_OUT, dtype=torch.float32, device="musa"),
H, W, OUT_SIZE_H, OUT_SIZE_W, KS, STRIDE, PADDING,
BLOCK_SIZE_KERNEL=BLOCK_SIZE_KERNEL, BLOCK_SIZE_OUTPUT=BLOCK_SIZE_OUTPUT)
torch.musa.synchronize()
print("transformed = RUN_OK (not expected)")
return 1
except Exception as exc:
dev_err = str(exc)
print("transformed = DEVICE_ERROR:", type(exc).__name__ + ":", dev_err.splitlines()[0])
after = set(glob.glob("core_*.mudmp"))
if after - before:
print("new core mudmp:", sorted(after - before)[0])
return 0 if "illegal memory access" in dev_err.lower() else 1

if __name__ == "__main__":
sys.exit(main())
```

[core_2026-08-19_00_03_41.159_2ac895ab3f25_11900.txt](https://github.com/user-attachments/files/31188044/core_2026-08-19_00_03_41.159_2ac895ab3f25_11900.txt)

Start by running the standalone Python reproducer in the issue with Triton 3.6.0, the MUSA backend, and the stated torch_musa/MUSA versions. Compare max_pooling_kernel with max_pooling_kernel_tx, inspect the generated backend code and the linked core dump, and use the reported 16 ablations to isolate the codegen trigger. Done means the transformed kernel no longer causes illegal device memory access while the original behavior remains correct.

Written by the indexing model from the issue text.