[SGLANG] Triton Attention UT failed

[airMeng]

Describe the bug

When testing the latest sglang on PVC, the triton attention kernel UT failed

git clone https://github.com/sgl-project/sglang.git
cd sglang
pip install -e "python[all_xpu]"
cd test/srt
pytest test_triton_attention_backend.py

The results will be

FAILED test_triton_attention_kernels.py::TestTritonAttention::test_extend_attention - AssertionError: False is not true
FAILED test_triton_attention_kernels.py::TestTritonAttention::test_grouped_decode_attention - AssertionError: False is not true
======================================================================================================================= 2 failed, 2 passed in 195.69s (0:03:15) ========================================================================================================================

Environment details

Triton: 3.2.0+git7bd27ed6
Torch: 2.7.0a0+git8326d27-cp310-cp310
GPU: 1550 with drivers 1032.21

@EikanWang @etaf @liangan1 for awareness

[zhuyuhua-v]

DeepSeek-R1 with sglang also failed in the Triton attention kernel, which uses RadixAttention - extend_attention_fwd https://github.com/sgl-project/sglang/blob/main/python/sglang/srt/layers/attention/triton_backend.py#L401.
log:
RadixAttention <sglang.srt.layers.attention.triton_backend.TritonAttnBackend object at 0x7f62a0dcdf60> Segmentation fault (core dumped)

[airMeng]

@dchigarev here is the ttgir for the 2 kernels:

extend_cache.tar.gz
group_decode_cache.tar.gz

[liangan1]

@dchigarev, we are working on the evaluation of the fp8 kernel of deepseek-v3 and it is a block issue now, is there any updates about it?

[dchigarev]

@dchigarev, we are working on the evaluation of the fp8 kernel of deepseek-v3 and it is a block issue now, is there any updates about it?

No updates yet. I'm starting to work on this ticket today and will post updates as soon as I find something

[dchigarev]

1. test_grouped_decode_attention

I was able to reproduce the failure for the TestTritonAttention::test_grouped_decode_attention on PVC a1100. The test fail when the number of heads (H_KV) is either 1 or 2, but it passes for any H_KV value that is greater than 2 (e.g 4, 8, 16, ...). However, a low H_KV value doesn't seem to be an issue when the sequence length is large (500). Below is the complete list of passing and failing cases for test_grouped_decode_attention:

Passed cases

seq_len=5, (B=2, H_Q=16, H_KV=16, D=64, D_V=64)
seq_len=100, (B=2, H_Q=16, H_KV=16, D=64, D_V=64)
seq_len=128, (B=2, H_Q=16, H_KV=16, D=64, D_V=64)
seq_len=500, (B=2, H_Q=16, H_KV=16, D=64, D_V=64)
seq_len=500, (B=2, H_Q=16, H_KV=1, D=64, D_V=64)
seq_len=500, (B=2, H_Q=64, H_KV=1, D=13, D_V=13)
seq_len=500, (B=2, H_Q=128, H_KV=1, D=80, D_V=80)
seq_len=500, (B=2, H_Q=128, H_KV=2, D=512, D_V=512)
seq_len=500, (B=2, H_Q=128, H_KV=1, D=576, D_V=512)

Failed cases

seq_len=5, (B=2, H_Q=16, H_KV=1, D=64, D_V=64)
seq_len=5, (B=2, H_Q=64, H_KV=1, D=13, D_V=13)
seq_len=5, (B=2, H_Q=128, H_KV=1, D=80, D_V=80)
seq_len=5, (B=2, H_Q=128, H_KV=2, D=512, D_V=512)
seq_len=5, (B=2, H_Q=128, H_KV=1, D=576, D_V=512)
seq_len=100, (B=2, H_Q=16, H_KV=1, D=64, D_V=64)
seq_len=100, (B=2, H_Q=64, H_KV=1, D=13, D_V=13)
seq_len=100, (B=2, H_Q=128, H_KV=1, D=80, D_V=80)
seq_len=100, (B=2, H_Q=128, H_KV=2, D=512, D_V=512)
seq_len=100, (B=2, H_Q=128, H_KV=1, D=576, D_V=512)
seq_len=128, (B=2, H_Q=16, H_KV=1, D=64, D_V=64)
seq_len=128, (B=2, H_Q=64, H_KV=1, D=13, D_V=13)
seq_len=128, (B=2, H_Q=128, H_KV=1, D=80, D_V=80)
seq_len=128, (B=2, H_Q=128, H_KV=2, D=512, D_V=512)
seq_len=128, (B=2, H_Q=128, H_KV=1, D=576, D_V=512)

The problematic kernel appears to be this one. The kernel itself seems to be written correctly, since it works fine on Nvidia even with H_KV = 1, 2. I will try to narrow down the problem and come up with a simpler reproducer.

2. test_extend_attention

I was not able to reproduce the failure on PVC A1100. I will try to get access to an A1550 and run the test on it later.

[liangan1]

@dchigarev Thanks for your updates. for Deepseek-v3, the H_KV=1 is a critical path and it is used when the w_k_b_proj is absorbed.

[dchigarev]

test_grouped_decode_attention
So this tl.where statement appears to be causing the problem. The input matrix qk has a shape of 16x32 and the where statement basically means "take the first value in each row of qk and replace the rest with float("-inf")". But for some reason the mask is applied to each block of 16 columns independently, leading to an incorrect result of the tl.where statement:

# Now we take the first value in each row from 'qk' and replace the rest with '-inf'
# We except that each row of 'qk_masked' now looks like this:
# [128, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf]
# But in reality it looks like this:
# [128, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, 128, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf]
#                                                                                                  ^
#                                                                                           idx=16 |
qk_masked = tl.where(mask_qk_x[:, None] & mask_qk_y[None, :], qk, float("-inf"))

The problem is probably related to the dpas shapes being 8x16, 16x16, and 8x16, as the issue occurs only when the input and output of tl.where come from/to a tl.dot statement.

That said, the minimal reproducer is the following:

(the repro works fine on NVIDIA but fails on Intel XPU)

reproducer

import torch
import os
import triton
import triton.language as tl

@triton.jit
def where_repro(
    c0_masked_out,
    c0_masked_stride_x,
    c1_out,
    c1_stride_x,
    BY : tl.constexpr,
    BZ : tl.constexpr,
):
    bx = tl.program_id(0)  # X grid index
    by = tl.program_id(1)  # Y grid index
    bz = tl.program_id(2)  # Z grid index

    A0_NROWS : tl.constexpr = 16; A0_NCOLS : tl.constexpr = 64
    B0_NROWS : tl.constexpr = 64; B0_NCOLS : tl.constexpr = 32
    C0_NROWS : tl.constexpr = A0_NROWS; C0_NCOLS : tl.constexpr = B0_NCOLS

    B1_NROWS : tl.constexpr = 32; B1_NCOLS : tl.constexpr = 64
    C1_NROWS : tl.constexpr = C0_NROWS; C1_NCOLS : tl.constexpr = B1_NCOLS

    a0 = tl.full(shape=(A0_NROWS, A0_NCOLS), value=2.0, dtype=tl.float32)
    b0 = tl.full(shape=(B0_NROWS, B0_NCOLS), value=1.0, dtype=tl.float32)
    c0 = tl.dot(a0, b0) # c0 is a matrix [16x32] with all its values equal '128.0'

    mask_c0_x = tl.arange(0, C0_NROWS) >= 0 # full 'True' mask
    mask_c0_y = tl.arange(0, C0_NCOLS) < 1 # only take the first element in each row

    # Now we take the first value in each row from 'c0' and replace the rest with '-inf'
    # We except that each row of 'c0_masked' now looks like this:
    # [128, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf]
    # But in reality it looks like this:
    # [128, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, 128, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf, -inf]
    #                                                                                                 ^
    #                                                                                          idx=16 |
    c0_masked = tl.where(mask_c0_x[:, None] & mask_c0_y[None, :], c0, float("-inf"))

    # ======== STORE c0_masked =========
    grid_idx = bx * (BY) * (BZ) + by * (BZ) + bz
    c0_blck_row_idx = grid_idx * C0_NROWS * c0_masked_stride_x + tl.arange(0, C0_NROWS)[:, None] * c0_masked_stride_x
    c0_blck_col_idx = tl.arange(0, C0_NCOLS)[None, :]

    tl.store(c0_masked_out + c0_blck_row_idx + c0_blck_col_idx, c0_masked)

    # There has to be a `tl.dot(c0_masked, ...)` statement for the bug to be reproducible.
    # ======== COMPUTE C1 = C0_MASKED @ B1 =========

    b1 = tl.full(shape=(B1_NROWS, B1_NCOLS), value=1.0, dtype=tl.float32)
    c1 = tl.dot(c0_masked, b1)

    # ======== STORE C1 =========
    c1_blck_row_idx = grid_idx * C1_NROWS * c1_stride_x + tl.arange(0, C1_NROWS)[:, None] * c1_stride_x
    c1_blck_col_idx = tl.arange(0, C1_NCOLS)[None, :]

    tl.store(c1_out + c1_blck_row_idx + c1_blck_col_idx, c1)


def compare_tensors(tensor1: torch.Tensor, tensor2: torch.Tensor):    
    diff_indices = (tensor1 != tensor2).nonzero(as_tuple=True)
    
    if len(diff_indices[0]) == 0:
        print("Tensors are equal")
        return
    
    for index in zip(*diff_indices):
        print(f"Index {index}: tensor1={tensor1[index].item()}, tensor2={tensor2[index].item()}")


grid = (2, 1, 8)
DEVICE = os.environ.get("SRT_DEVICE", "xpu")

# ===== DEFINE c0_masked_out =========
C0_NROWS = 16
C0_NCOLS = 32
c0_masked_out = torch.full((grid[0] * grid[1] * grid[2] * C0_NROWS, C0_NCOLS), -1, dtype=torch.float32, device=DEVICE)

# ===== DEFINE c1_out =========
C1_NROWS = 16
C1_NCOLS = 64
c1_out = torch.full((grid[0] * grid[1] * grid[2] * C1_NROWS, C1_NCOLS), -1, dtype=torch.float32, device=DEVICE)

where_repro[grid](
    c0_masked_out,
    c0_masked_out.stride(0),
    c1_out,
    c1_out.stride(0),
    BY=grid[1],
    BZ=grid[2],
)

torch_ref = torch.full((C0_NROWS * grid[0] * grid[1] * grid[2], C0_NCOLS), 128.0, dtype=torch.float32)
torch_ref = torch.where(torch.arange(C0_NCOLS) < 1, torch_ref, float("-inf"))

# This comparison passes on Nvidia but fails on Intel XPU
compare_tensors(c0_masked_out.cpu(), torch_ref)
breakpoint()
print("Compare tensors on your own...")

P.S. I would appreciate if someone more experienced with Triton could give me a hint on what I should look at to investigate the problem (for example, which stage/pass of the pipeline is responsible for dpas/arith.select lowering) @anmyachev @alexbaden @etiotto

[dchigarev]

Attaching IR dumps as well:

where_repro.ttir

#loc = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0)
module {
  tt.func public @where_repro(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0), %arg1: i32 {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0), %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0), %arg3: i32 {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0)) attributes {noinline = false} {
    %cst = arith.constant dense<0.000000e+00> : tensor<16x64xf32> loc(#loc1)
    %cst_0 = arith.constant dense<1.000000e+00> : tensor<32x64xf32> loc(#loc1)
    %c16_i32 = arith.constant 16 : i32 loc(#loc1)
    %c8_i32 = arith.constant 8 : i32 loc(#loc1)
    %cst_1 = arith.constant dense<0xFF800000> : tensor<16x32xf32> loc(#loc1)
    %cst_2 = arith.constant dense<1> : tensor<32xi32> loc(#loc1)
    %cst_3 = arith.constant dense<0> : tensor<16xi32> loc(#loc1)
    %cst_4 = arith.constant dense<0.000000e+00> : tensor<16x32xf32> loc(#loc1)
    %cst_5 = arith.constant dense<1.000000e+00> : tensor<64x32xf32> loc(#loc1)
    %cst_6 = arith.constant dense<2.000000e+00> : tensor<16x64xf32> loc(#loc1)
    %0 = tt.get_program_id x : i32 loc(#loc2)
    %1 = tt.get_program_id y : i32 loc(#loc3)
    %2 = tt.get_program_id z : i32 loc(#loc4)
    %3 = tt.dot %cst_6, %cst_5, %cst_4, inputPrecision = tf32 : tensor<16x64xf32> * tensor<64x32xf32> -> tensor<16x32xf32> loc(#loc5)
    %4 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc6)
    %5 = arith.cmpi sge, %4, %cst_3 : tensor<16xi32> loc(#loc7)
    %6 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc8)
    %7 = arith.cmpi slt, %6, %cst_2 : tensor<32xi32> loc(#loc9)
    %8 = tt.expand_dims %5 {axis = 1 : i32} : tensor<16xi1> -> tensor<16x1xi1> loc(#loc10)
    %9 = tt.expand_dims %7 {axis = 0 : i32} : tensor<32xi1> -> tensor<1x32xi1> loc(#loc11)
    %10 = tt.broadcast %8 : tensor<16x1xi1> -> tensor<16x32xi1> loc(#loc12)
    %11 = tt.broadcast %9 : tensor<1x32xi1> -> tensor<16x32xi1> loc(#loc12)
    %12 = arith.andi %10, %11 : tensor<16x32xi1> loc(#loc12)
    %13 = arith.select %12, %3, %cst_1 : tensor<16x32xi1>, tensor<16x32xf32> loc(#loc13)
    %14 = arith.muli %0, %c8_i32 : i32 loc(#loc14)
    %15 = arith.muli %1, %c8_i32 : i32 loc(#loc15)
    %16 = arith.addi %14, %15 : i32 loc(#loc16)
    %17 = arith.addi %16, %2 : i32 loc(#loc17)
    %18 = arith.muli %17, %c16_i32 : i32 loc(#loc18)
    %19 = arith.muli %18, %arg1 : i32 loc(#loc19)
    %20 = tt.expand_dims %4 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc20)
    %21 = tt.splat %arg1 : i32 -> tensor<16x1xi32> loc(#loc21)
    %22 = arith.muli %20, %21 : tensor<16x1xi32> loc(#loc21)
    %23 = tt.splat %19 : i32 -> tensor<16x1xi32> loc(#loc22)
    %24 = arith.addi %23, %22 : tensor<16x1xi32> loc(#loc22)
    %25 = tt.expand_dims %6 {axis = 0 : i32} : tensor<32xi32> -> tensor<1x32xi32> loc(#loc23)
    %26 = tt.splat %arg0 : !tt.ptr<f32> -> tensor<16x1x!tt.ptr<f32>> loc(#loc24)
    %27 = tt.addptr %26, %24 : tensor<16x1x!tt.ptr<f32>>, tensor<16x1xi32> loc(#loc24)
    %28 = tt.broadcast %27 : tensor<16x1x!tt.ptr<f32>> -> tensor<16x32x!tt.ptr<f32>> loc(#loc25)
    %29 = tt.broadcast %25 : tensor<1x32xi32> -> tensor<16x32xi32> loc(#loc25)
    %30 = tt.addptr %28, %29 : tensor<16x32x!tt.ptr<f32>>, tensor<16x32xi32> loc(#loc25)
    tt.store %30, %13 : tensor<16x32x!tt.ptr<f32>> loc(#loc26)
    %31 = tt.dot %13, %cst_0, %cst, inputPrecision = tf32 : tensor<16x32xf32> * tensor<32x64xf32> -> tensor<16x64xf32> loc(#loc27)
    %32 = arith.muli %18, %arg3 : i32 loc(#loc28)
    %33 = tt.splat %arg3 : i32 -> tensor<16x1xi32> loc(#loc29)
    %34 = arith.muli %20, %33 : tensor<16x1xi32> loc(#loc29)
    %35 = tt.splat %32 : i32 -> tensor<16x1xi32> loc(#loc30)
    %36 = arith.addi %35, %34 : tensor<16x1xi32> loc(#loc30)
    %37 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32> loc(#loc31)
    %38 = tt.expand_dims %37 {axis = 0 : i32} : tensor<64xi32> -> tensor<1x64xi32> loc(#loc32)
    %39 = tt.splat %arg2 : !tt.ptr<f32> -> tensor<16x1x!tt.ptr<f32>> loc(#loc33)
    %40 = tt.addptr %39, %36 : tensor<16x1x!tt.ptr<f32>>, tensor<16x1xi32> loc(#loc33)
    %41 = tt.broadcast %40 : tensor<16x1x!tt.ptr<f32>> -> tensor<16x64x!tt.ptr<f32>> loc(#loc34)
    %42 = tt.broadcast %38 : tensor<1x64xi32> -> tensor<16x64xi32> loc(#loc34)
    %43 = tt.addptr %41, %42 : tensor<16x64x!tt.ptr<f32>>, tensor<16x64xi32> loc(#loc34)
    tt.store %43, %31 : tensor<16x64x!tt.ptr<f32>> loc(#loc35)
    tt.return loc(#loc36)
  } loc(#loc)
} loc(#loc)
#loc1 = loc(unknown)
#loc2 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":15:23)
#loc3 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":16:23)
#loc4 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":17:23)
#loc5 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":28:20)
#loc6 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":30:29)
#loc7 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":30:42)
#loc8 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":31:29)
#loc9 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":31:41)
#loc10 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:35)
#loc11 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:56)
#loc12 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:46)
#loc13 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:70)
#loc14 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:28)
#loc15 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:40)
#loc16 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:34)
#loc17 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:46)
#loc18 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:33)
#loc19 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:44)
#loc20 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:88)
#loc21 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:99)
#loc22 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:65)
#loc23 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":45:45)
#loc24 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":47:29)
#loc25 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":47:47)
#loc26 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":47:64)
#loc27 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":53:27)
#loc28 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":56:44)
#loc29 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":56:92)
#loc30 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":56:58)
#loc31 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":57:35)
#loc32 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":57:45)
#loc33 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:22)
#loc34 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:40)
#loc35 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:57)
#loc36 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:4)

where_repro.ttgir

#blocked = #ttg.blocked<{sizePerThread = [1, 4], threadsPerWarp = [2, 8], warpsPerCTA = [4, 1], order = [1, 0]}>
#blocked1 = #ttg.blocked<{sizePerThread = [1, 4], threadsPerWarp = [1, 16], warpsPerCTA = [4, 1], order = [1, 0]}>
#loc = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0)
#mma = #triton_intel_gpu.dpas<{repeatCount = 8, systolicDepth = 8, executionSize = 16, opsPerChan = 1, threadsPerWarp = 16, warpsPerCTA = [4, 1], repCluster = [1, 1], A = [8, 8], B = [8, 16], C = [8, 16]}>
module attributes {triton_intel_gpu.min_sg_size = 16 : i32, triton_intel_gpu.support_bf16_conversion, triton_intel_gpu.support_dpas, triton_intel_gpu.support_sg_2d_block, triton_intel_gpu.target_arch = "spir64", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "xpu", "ttg.threads-per-warp" = 16 : i32} {
  tt.func public @where_repro(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0), %arg1: i32 {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0), %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0), %arg3: i32 {tt.divisibility = 16 : i32} loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":7:0)) attributes {noinline = false} {
    %cst = arith.constant dense<0xFF800000> : tensor<16x32xf32, #mma> loc(#loc1)
    %c8_i32 = arith.constant 8 : i32 loc(#loc1)
    %c16_i32 = arith.constant 16 : i32 loc(#loc1)
    %cst_0 = arith.constant dense<1.000000e+00> : tensor<32x64xf32, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 1}>> loc(#loc1)
    %cst_1 = arith.constant dense<0.000000e+00> : tensor<16x64xf32, #mma> loc(#loc1)
    %cst_2 = arith.constant dense<1.000000e+00> : tensor<64x32xf32, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 1}>> loc(#loc1)
    %cst_3 = arith.constant dense<2.000000e+00> : tensor<16x64xf32, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 1}>> loc(#loc1)
    %cst_4 = arith.constant dense<0.000000e+00> : tensor<16x32xf32, #mma> loc(#loc1)
    %cst_5 = arith.constant dense<1> : tensor<32xi32, #ttg.slice<{dim = 0, parent = #mma}>> loc(#loc1)
    %cst_6 = arith.constant dense<0> : tensor<16xi32, #ttg.slice<{dim = 1, parent = #mma}>> loc(#loc1)
    %0 = tt.get_program_id x : i32 loc(#loc2)
    %1 = tt.get_program_id y : i32 loc(#loc3)
    %2 = tt.get_program_id z : i32 loc(#loc4)
    %3 = tt.dot %cst_3, %cst_2, %cst_4, inputPrecision = tf32 : tensor<16x64xf32, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 1}>> * tensor<64x32xf32, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 1}>> -> tensor<16x32xf32, #mma> loc(#loc5)
    %4 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #mma}>> loc(#loc6)
    %5 = arith.cmpi sge, %4, %cst_6 : tensor<16xi32, #ttg.slice<{dim = 1, parent = #mma}>> loc(#loc7)
    %6 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 0, parent = #mma}>> loc(#loc8)
    %7 = arith.cmpi slt, %6, %cst_5 : tensor<32xi32, #ttg.slice<{dim = 0, parent = #mma}>> loc(#loc9)
    %8 = tt.expand_dims %5 {axis = 1 : i32} : tensor<16xi1, #ttg.slice<{dim = 1, parent = #mma}>> -> tensor<16x1xi1, #mma> loc(#loc10)
    %9 = tt.expand_dims %7 {axis = 0 : i32} : tensor<32xi1, #ttg.slice<{dim = 0, parent = #mma}>> -> tensor<1x32xi1, #mma> loc(#loc11)
    %10 = tt.broadcast %8 : tensor<16x1xi1, #mma> -> tensor<16x32xi1, #mma> loc(#loc12)
    %11 = tt.broadcast %9 : tensor<1x32xi1, #mma> -> tensor<16x32xi1, #mma> loc(#loc12)
    %12 = arith.andi %10, %11 : tensor<16x32xi1, #mma> loc(#loc12)
    %13 = arith.select %12, %3, %cst : tensor<16x32xi1, #mma>, tensor<16x32xf32, #mma> loc(#loc13)
    %14 = arith.muli %0, %c8_i32 : i32 loc(#loc14)
    %15 = arith.muli %1, %c8_i32 : i32 loc(#loc15)
    %16 = arith.addi %14, %15 : i32 loc(#loc16)
    %17 = arith.addi %16, %2 : i32 loc(#loc17)
    %18 = arith.muli %17, %c16_i32 : i32 loc(#loc18)
    %19 = arith.muli %18, %arg1 : i32 loc(#loc19)
    %20 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc20)
    %21 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> loc(#loc20)
    %22 = tt.expand_dims %20 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc20)
    %23 = tt.expand_dims %21 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked1}>> -> tensor<16x1xi32, #blocked1> loc(#loc20)
    %24 = tt.splat %arg1 : i32 -> tensor<16x1xi32, #blocked> loc(#loc21)
    %25 = arith.muli %22, %24 : tensor<16x1xi32, #blocked> loc(#loc21)
    %26 = tt.splat %19 : i32 -> tensor<16x1xi32, #blocked> loc(#loc22)
    %27 = arith.addi %26, %25 : tensor<16x1xi32, #blocked> loc(#loc22)
    %28 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc23)
    %29 = tt.expand_dims %28 {axis = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x32xi32, #blocked> loc(#loc23)
    %30 = tt.splat %arg0 : !tt.ptr<f32> -> tensor<16x1x!tt.ptr<f32>, #blocked> loc(#loc24)
    %31 = tt.addptr %30, %27 : tensor<16x1x!tt.ptr<f32>, #blocked>, tensor<16x1xi32, #blocked> loc(#loc24)
    %32 = tt.broadcast %31 : tensor<16x1x!tt.ptr<f32>, #blocked> -> tensor<16x32x!tt.ptr<f32>, #blocked> loc(#loc25)
    %33 = tt.broadcast %29 : tensor<1x32xi32, #blocked> -> tensor<16x32xi32, #blocked> loc(#loc25)
    %34 = tt.addptr %32, %33 : tensor<16x32x!tt.ptr<f32>, #blocked>, tensor<16x32xi32, #blocked> loc(#loc25)
    %35 = ttg.convert_layout %13 : tensor<16x32xf32, #mma> -> tensor<16x32xf32, #blocked> loc(#loc26)
    tt.store %34, %35 : tensor<16x32x!tt.ptr<f32>, #blocked> loc(#loc26)
    %36 = ttg.convert_layout %13 : tensor<16x32xf32, #mma> -> tensor<16x32xf32, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 1}>> loc(#loc13)
    %37 = tt.dot %36, %cst_0, %cst_1, inputPrecision = tf32 : tensor<16x32xf32, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 1}>> * tensor<32x64xf32, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 1}>> -> tensor<16x64xf32, #mma> loc(#loc27)
    %38 = arith.muli %18, %arg3 : i32 loc(#loc28)
    %39 = tt.splat %arg3 : i32 -> tensor<16x1xi32, #blocked1> loc(#loc29)
    %40 = arith.muli %23, %39 : tensor<16x1xi32, #blocked1> loc(#loc29)
    %41 = tt.splat %38 : i32 -> tensor<16x1xi32, #blocked1> loc(#loc30)
    %42 = arith.addi %41, %40 : tensor<16x1xi32, #blocked1> loc(#loc30)
    %43 = tt.make_range {end = 64 : i32, start = 0 : i32} : tensor<64xi32, #ttg.slice<{dim = 0, parent = #blocked1}>> loc(#loc31)
    %44 = tt.expand_dims %43 {axis = 0 : i32} : tensor<64xi32, #ttg.slice<{dim = 0, parent = #blocked1}>> -> tensor<1x64xi32, #blocked1> loc(#loc31)
    %45 = tt.splat %arg2 : !tt.ptr<f32> -> tensor<16x1x!tt.ptr<f32>, #blocked1> loc(#loc32)
    %46 = tt.addptr %45, %42 : tensor<16x1x!tt.ptr<f32>, #blocked1>, tensor<16x1xi32, #blocked1> loc(#loc32)
    %47 = tt.broadcast %46 : tensor<16x1x!tt.ptr<f32>, #blocked1> -> tensor<16x64x!tt.ptr<f32>, #blocked1> loc(#loc33)
    %48 = tt.broadcast %44 : tensor<1x64xi32, #blocked1> -> tensor<16x64xi32, #blocked1> loc(#loc33)
    %49 = tt.addptr %47, %48 : tensor<16x64x!tt.ptr<f32>, #blocked1>, tensor<16x64xi32, #blocked1> loc(#loc33)
    %50 = ttg.convert_layout %37 : tensor<16x64xf32, #mma> -> tensor<16x64xf32, #blocked1> loc(#loc34)
    tt.store %49, %50 : tensor<16x64x!tt.ptr<f32>, #blocked1> loc(#loc34)
    tt.return loc(#loc35)
  } loc(#loc)
} loc(#loc)
#loc1 = loc(unknown)
#loc2 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":15:23)
#loc3 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":16:23)
#loc4 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":17:23)
#loc5 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":28:20)
#loc6 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":30:29)
#loc7 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":30:42)
#loc8 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":31:29)
#loc9 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":31:41)
#loc10 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:35)
#loc11 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:56)
#loc12 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:46)
#loc13 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":40:70)
#loc14 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:28)
#loc15 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:40)
#loc16 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:34)
#loc17 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":43:46)
#loc18 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:33)
#loc19 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:44)
#loc20 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:88)
#loc21 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:99)
#loc22 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":44:65)
#loc23 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":45:45)
#loc24 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":47:29)
#loc25 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":47:47)
#loc26 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":47:64)
#loc27 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":53:27)
#loc28 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":56:44)
#loc29 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":56:92)
#loc30 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":56:58)
#loc31 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":57:45)
#loc32 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:22)
#loc33 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:40)
#loc34 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:57)
#loc35 = loc("/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton/where_repro.py":59:4)

where_repro.llir

; ModuleID = 'LLVMDialectModule'
source_filename = "LLVMDialectModule"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
target triple = "spir64-unknown-unknown"

; Function Attrs: convergent mustprogress nounwind willreturn
declare !dbg !9 spir_func void @_Z7barrierj(i32) local_unnamed_addr #0

; Function Attrs: mustprogress nofree nosync nounwind willreturn memory(none)
declare !dbg !12 spir_func i64 @_Z12get_local_idj(i32) local_unnamed_addr #1

; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
declare !dbg !13 spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float>, <8 x float>, <8 x float>) local_unnamed_addr #2

; Function Attrs: mustprogress nofree nosync nounwind willreturn memory(none)
declare !dbg !14 spir_func i64 @_Z12get_group_idj(i32) local_unnamed_addr #1

; Function Attrs: mustprogress nounwind willreturn
define spir_kernel void @where_repro(ptr addrspace(1) writeonly captures(none) %0, i32 %1, ptr addrspace(1) writeonly captures(none) %2, i32 %3, ptr addrspace(3) %4) local_unnamed_addr #3 !dbg !15 !intel_reqd_sub_group_size !16 !max_work_group_size !17 {
  %6 = tail call spir_func i64 @_Z12get_group_idj(i32 0) #4, !dbg !18
  %7 = trunc i64 %6 to i32, !dbg !18
  %8 = tail call spir_func i64 @_Z12get_group_idj(i32 1) #4, !dbg !19
  %9 = trunc i64 %8 to i32, !dbg !19
  %10 = tail call spir_func i64 @_Z12get_group_idj(i32 2) #4, !dbg !20
  %11 = trunc i64 %10 to i32, !dbg !20
  %12 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> zeroinitializer) #5
  %13 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> %12) #5
  %14 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> %13) #5
  %15 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> %14) #5
  %16 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> %15) #5
  %17 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> %16) #5
  %18 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> %17) #5
  %19 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> splat (float 2.000000e+00), <8 x float> splat (float 1.000000e+00), <8 x float> %18) #5
  %20 = extractelement <8 x float> %19, i64 0, !dbg !21
  %21 = extractelement <8 x float> %19, i64 1, !dbg !21
  %22 = extractelement <8 x float> %19, i64 2, !dbg !21
  %23 = extractelement <8 x float> %19, i64 3, !dbg !21
  %24 = extractelement <8 x float> %19, i64 4, !dbg !21
  %25 = extractelement <8 x float> %19, i64 5, !dbg !21
  %26 = extractelement <8 x float> %19, i64 6, !dbg !21
  %27 = extractelement <8 x float> %19, i64 7, !dbg !21
  %28 = tail call spir_func i64 @_Z12get_local_idj(i32 0) #4, !dbg !22
  %29 = trunc i64 %28 to i32, !dbg !22
  %30 = lshr i32 %29, 4, !dbg !22
  %31 = lshr i32 %29, 1, !dbg !22
  %32 = and i32 %31, 8, !dbg !22
  %33 = and i32 %29, 15, !dbg !23
  %34 = and i32 %29, 7, !dbg !23
  %35 = icmp eq i32 %33, 0, !dbg !24
  %36 = select i1 %35, float %20, float 0xFFF0000000000000, !dbg !25
  %37 = select i1 %35, float %21, float 0xFFF0000000000000, !dbg !25
  %38 = select i1 %35, float %22, float 0xFFF0000000000000, !dbg !25
  %39 = select i1 %35, float %23, float 0xFFF0000000000000, !dbg !25
  %40 = select i1 %35, float %24, float 0xFFF0000000000000, !dbg !25
  %41 = select i1 %35, float %25, float 0xFFF0000000000000, !dbg !25
  %42 = select i1 %35, float %26, float 0xFFF0000000000000, !dbg !25
  %43 = select i1 %35, float %27, float 0xFFF0000000000000, !dbg !25
  %44 = add i32 %9, %7, !dbg !26
  %45 = shl i32 %44, 7, !dbg !27
  %46 = shl i32 %11, 4, !dbg !27
  %47 = add i32 %45, %46, !dbg !27
  %48 = lshr i32 %29, 3, !dbg !28
  %.lobit = and i32 %48, 1, !dbg !28
  %49 = and i32 %48, 7, !dbg !28
  %50 = and i32 %30, 3, !dbg !28
  %51 = or disjoint i32 %49, %47, !dbg !29
  %52 = mul i32 %51, %1, !dbg !29
  %53 = or disjoint i32 %51, 8, !dbg !29
  %54 = mul i32 %53, %1, !dbg !29
  %55 = shl i32 %29, 2, !dbg !30
  %56 = and i32 %55, 28, !dbg !30
  %57 = sext i32 %52 to i64, !dbg !31
  %58 = getelementptr float, ptr addrspace(1) %0, i64 %57, !dbg !31
  %59 = sext i32 %54 to i64, !dbg !31
  %60 = getelementptr float, ptr addrspace(1) %0, i64 %59, !dbg !31
  %61 = zext nneg i32 %56 to i64, !dbg !32
  %62 = getelementptr float, ptr addrspace(1) %58, i64 %61, !dbg !32
  %63 = getelementptr float, ptr addrspace(1) %60, i64 %61, !dbg !32
  %64 = shl nuw nsw i32 %30, 3, !dbg !33
  %65 = and i32 %64, 8, !dbg !33
  %66 = mul nuw nsw i32 %65, 36, !dbg !33
  %67 = or disjoint i32 %66, %33, !dbg !33
  %68 = zext nneg i32 %67 to i64, !dbg !33
  %69 = getelementptr float, ptr addrspace(3) %4, i64 %68, !dbg !33
  %70 = insertelement <1 x float> poison, float %36, i64 0, !dbg !33
  store <1 x float> %70, ptr addrspace(3) %69, align 4, !dbg !33
  %71 = or disjoint i32 %65, 1, !dbg !33
  %72 = mul nuw nsw i32 %71, 36, !dbg !33
  %73 = add nuw nsw i32 %72, %33, !dbg !33
  %74 = zext nneg i32 %73 to i64, !dbg !33
  %75 = getelementptr float, ptr addrspace(3) %4, i64 %74, !dbg !33
  %76 = insertelement <1 x float> poison, float %37, i64 0, !dbg !33
  store <1 x float> %76, ptr addrspace(3) %75, align 4, !dbg !33
  %77 = or disjoint i32 %65, 2, !dbg !33
  %78 = mul nuw nsw i32 %77, 36, !dbg !33
  %79 = add nuw nsw i32 %78, %33, !dbg !33
  %80 = zext nneg i32 %79 to i64, !dbg !33
  %81 = getelementptr float, ptr addrspace(3) %4, i64 %80, !dbg !33
  %82 = insertelement <1 x float> poison, float %38, i64 0, !dbg !33
  store <1 x float> %82, ptr addrspace(3) %81, align 4, !dbg !33
  %83 = or disjoint i32 %65, 3, !dbg !33
  %84 = mul nuw nsw i32 %83, 36, !dbg !33
  %85 = add nuw nsw i32 %84, %33, !dbg !33
  %86 = zext nneg i32 %85 to i64, !dbg !33
  %87 = getelementptr float, ptr addrspace(3) %4, i64 %86, !dbg !33
  %88 = insertelement <1 x float> poison, float %39, i64 0, !dbg !33
  store <1 x float> %88, ptr addrspace(3) %87, align 4, !dbg !33
  %89 = or disjoint i32 %65, 4, !dbg !33
  %90 = mul nuw nsw i32 %89, 36, !dbg !33
  %91 = or disjoint i32 %90, %33, !dbg !33
  %92 = zext nneg i32 %91 to i64, !dbg !33
  %93 = getelementptr float, ptr addrspace(3) %4, i64 %92, !dbg !33
  %94 = insertelement <1 x float> poison, float %40, i64 0, !dbg !33
  store <1 x float> %94, ptr addrspace(3) %93, align 4, !dbg !33
  %95 = or disjoint i32 %65, 5, !dbg !33
  %96 = mul nuw nsw i32 %95, 36, !dbg !33
  %97 = add nuw nsw i32 %96, %33, !dbg !33
  %98 = zext nneg i32 %97 to i64, !dbg !33
  %99 = getelementptr float, ptr addrspace(3) %4, i64 %98, !dbg !33
  %100 = insertelement <1 x float> poison, float %41, i64 0, !dbg !33
  store <1 x float> %100, ptr addrspace(3) %99, align 4, !dbg !33
  %101 = or disjoint i32 %65, 6, !dbg !33
  %102 = mul nuw nsw i32 %101, 36, !dbg !33
  %103 = add nuw nsw i32 %102, %33, !dbg !33
  %104 = zext nneg i32 %103 to i64, !dbg !33
  %105 = getelementptr float, ptr addrspace(3) %4, i64 %104, !dbg !33
  %106 = insertelement <1 x float> poison, float %42, i64 0, !dbg !33
  store <1 x float> %106, ptr addrspace(3) %105, align 4, !dbg !33
  %107 = or disjoint i32 %65, 7, !dbg !33
  %108 = mul nuw nsw i32 %107, 36, !dbg !33
  %109 = add nuw nsw i32 %108, %33, !dbg !33
  %110 = zext nneg i32 %109 to i64, !dbg !33
  %111 = getelementptr float, ptr addrspace(3) %4, i64 %110, !dbg !33
  %112 = insertelement <1 x float> poison, float %43, i64 0, !dbg !33
  store <1 x float> %112, ptr addrspace(3) %111, align 4, !dbg !33
  %113 = or disjoint i32 %33, 16, !dbg !33
  %114 = or disjoint i32 %66, %113, !dbg !33
  %115 = zext nneg i32 %114 to i64, !dbg !33
  %116 = getelementptr float, ptr addrspace(3) %4, i64 %115, !dbg !33
  store <1 x float> %70, ptr addrspace(3) %116, align 4, !dbg !33
  %117 = add nuw nsw i32 %72, %113, !dbg !33
  %118 = zext nneg i32 %117 to i64, !dbg !33
  %119 = getelementptr float, ptr addrspace(3) %4, i64 %118, !dbg !33
  store <1 x float> %76, ptr addrspace(3) %119, align 4, !dbg !33
  %120 = add nuw nsw i32 %78, %113, !dbg !33
  %121 = zext nneg i32 %120 to i64, !dbg !33
  %122 = getelementptr float, ptr addrspace(3) %4, i64 %121, !dbg !33
  store <1 x float> %82, ptr addrspace(3) %122, align 4, !dbg !33
  %123 = add nuw nsw i32 %84, %113, !dbg !33
  %124 = zext nneg i32 %123 to i64, !dbg !33
  %125 = getelementptr float, ptr addrspace(3) %4, i64 %124, !dbg !33
  store <1 x float> %88, ptr addrspace(3) %125, align 4, !dbg !33
  %126 = add nuw nsw i32 %90, %113, !dbg !33
  %127 = zext nneg i32 %126 to i64, !dbg !33
  %128 = getelementptr float, ptr addrspace(3) %4, i64 %127, !dbg !33
  store <1 x float> %94, ptr addrspace(3) %128, align 4, !dbg !33
  %129 = add nuw nsw i32 %96, %113, !dbg !33
  %130 = zext nneg i32 %129 to i64, !dbg !33
  %131 = getelementptr float, ptr addrspace(3) %4, i64 %130, !dbg !33
  store <1 x float> %100, ptr addrspace(3) %131, align 4, !dbg !33
  %132 = add nuw nsw i32 %102, %113, !dbg !33
  %133 = zext nneg i32 %132 to i64, !dbg !33
  %134 = getelementptr float, ptr addrspace(3) %4, i64 %133, !dbg !33
  store <1 x float> %106, ptr addrspace(3) %134, align 4, !dbg !33
  %135 = add nuw nsw i32 %108, %113, !dbg !33
  %136 = zext nneg i32 %135 to i64, !dbg !33
  %137 = getelementptr float, ptr addrspace(3) %4, i64 %136, !dbg !33
  store <1 x float> %112, ptr addrspace(3) %137, align 4, !dbg !33
  tail call spir_func void @_Z7barrierj(i32 1) #6, !dbg !33
  %138 = lshr i32 %33, 3, !dbg !33
  %139 = shl nuw nsw i32 %50, 1, !dbg !33
  %140 = or disjoint i32 %139, %138, !dbg !33
  %141 = mul nuw nsw i32 %140, 36, !dbg !33
  %142 = add nuw nsw i32 %141, %56, !dbg !33
  %143 = zext nneg i32 %142 to i64, !dbg !33
  %144 = getelementptr float, ptr addrspace(3) %4, i64 %143, !dbg !33
  %145 = load <4 x i32>, ptr addrspace(3) %144, align 16, !dbg !33
  %146 = getelementptr i8, ptr addrspace(3) %144, i64 1152, !dbg !33
  %147 = load <4 x i32>, ptr addrspace(3) %146, align 16, !dbg !33
  store <4 x i32> %145, ptr addrspace(1) %62, align 16, !dbg !33
  store <4 x i32> %147, ptr addrspace(1) %63, align 16, !dbg !33
  tail call spir_func void @_Z7barrierj(i32 1) #6, !dbg !25
  %148 = mul nuw nsw i32 %65, 33, !dbg !25
  %149 = add nuw nsw i32 %148, %33, !dbg !25
  %150 = zext nneg i32 %149 to i64, !dbg !25
  %151 = getelementptr float, ptr addrspace(3) %4, i64 %150, !dbg !25
  store <1 x float> %70, ptr addrspace(3) %151, align 4, !dbg !25
  %152 = mul nuw nsw i32 %71, 33, !dbg !25
  %153 = add nuw nsw i32 %152, %33, !dbg !25
  %154 = zext nneg i32 %153 to i64, !dbg !25
  %155 = getelementptr float, ptr addrspace(3) %4, i64 %154, !dbg !25
  store <1 x float> %76, ptr addrspace(3) %155, align 4, !dbg !25
  %156 = mul nuw nsw i32 %77, 33, !dbg !25
  %157 = add nuw nsw i32 %156, %33, !dbg !25
  %158 = zext nneg i32 %157 to i64, !dbg !25
  %159 = getelementptr float, ptr addrspace(3) %4, i64 %158, !dbg !25
  store <1 x float> %82, ptr addrspace(3) %159, align 4, !dbg !25
  %160 = mul nuw nsw i32 %83, 33, !dbg !25
  %161 = add nuw nsw i32 %160, %33, !dbg !25
  %162 = zext nneg i32 %161 to i64, !dbg !25
  %163 = getelementptr float, ptr addrspace(3) %4, i64 %162, !dbg !25
  store <1 x float> %88, ptr addrspace(3) %163, align 4, !dbg !25
  %164 = mul nuw nsw i32 %89, 33, !dbg !25
  %165 = add nuw nsw i32 %164, %33, !dbg !25
  %166 = zext nneg i32 %165 to i64, !dbg !25
  %167 = getelementptr float, ptr addrspace(3) %4, i64 %166, !dbg !25
  store <1 x float> %94, ptr addrspace(3) %167, align 4, !dbg !25
  %168 = mul nuw nsw i32 %95, 33, !dbg !25
  %169 = add nuw nsw i32 %168, %33, !dbg !25
  %170 = zext nneg i32 %169 to i64, !dbg !25
  %171 = getelementptr float, ptr addrspace(3) %4, i64 %170, !dbg !25
  store <1 x float> %100, ptr addrspace(3) %171, align 4, !dbg !25
  %172 = mul nuw nsw i32 %101, 33, !dbg !25
  %173 = add nuw nsw i32 %172, %33, !dbg !25
  %174 = zext nneg i32 %173 to i64, !dbg !25
  %175 = getelementptr float, ptr addrspace(3) %4, i64 %174, !dbg !25
  store <1 x float> %106, ptr addrspace(3) %175, align 4, !dbg !25
  %176 = mul nuw nsw i32 %107, 33, !dbg !25
  %177 = add nuw nsw i32 %176, %33, !dbg !25
  %178 = zext nneg i32 %177 to i64, !dbg !25
  %179 = getelementptr float, ptr addrspace(3) %4, i64 %178, !dbg !25
  store <1 x float> %112, ptr addrspace(3) %179, align 4, !dbg !25
  %180 = add nuw nsw i32 %148, %113, !dbg !25
  %181 = zext nneg i32 %180 to i64, !dbg !25
  %182 = getelementptr float, ptr addrspace(3) %4, i64 %181, !dbg !25
  store <1 x float> %70, ptr addrspace(3) %182, align 4, !dbg !25
  %183 = add nuw nsw i32 %152, %113, !dbg !25
  %184 = zext nneg i32 %183 to i64, !dbg !25
  %185 = getelementptr float, ptr addrspace(3) %4, i64 %184, !dbg !25
  store <1 x float> %76, ptr addrspace(3) %185, align 4, !dbg !25
  %186 = add nuw nsw i32 %156, %113, !dbg !25
  %187 = zext nneg i32 %186 to i64, !dbg !25
  %188 = getelementptr float, ptr addrspace(3) %4, i64 %187, !dbg !25
  store <1 x float> %82, ptr addrspace(3) %188, align 4, !dbg !25
  %189 = add nuw nsw i32 %160, %113, !dbg !25
  %190 = zext nneg i32 %189 to i64, !dbg !25
  %191 = getelementptr float, ptr addrspace(3) %4, i64 %190, !dbg !25
  store <1 x float> %88, ptr addrspace(3) %191, align 4, !dbg !25
  %192 = add nuw nsw i32 %164, %113, !dbg !25
  %193 = zext nneg i32 %192 to i64, !dbg !25
  %194 = getelementptr float, ptr addrspace(3) %4, i64 %193, !dbg !25
  store <1 x float> %94, ptr addrspace(3) %194, align 4, !dbg !25
  %195 = add nuw nsw i32 %168, %113, !dbg !25
  %196 = zext nneg i32 %195 to i64, !dbg !25
  %197 = getelementptr float, ptr addrspace(3) %4, i64 %196, !dbg !25
  store <1 x float> %100, ptr addrspace(3) %197, align 4, !dbg !25
  %198 = add nuw nsw i32 %172, %113, !dbg !25
  %199 = zext nneg i32 %198 to i64, !dbg !25
  %200 = getelementptr float, ptr addrspace(3) %4, i64 %199, !dbg !25
  store <1 x float> %106, ptr addrspace(3) %200, align 4, !dbg !25
  %201 = add nuw nsw i32 %176, %113, !dbg !25
  %202 = zext nneg i32 %201 to i64, !dbg !25
  %203 = getelementptr float, ptr addrspace(3) %4, i64 %202, !dbg !25
  store <1 x float> %112, ptr addrspace(3) %203, align 4, !dbg !25
  tail call spir_func void @_Z7barrierj(i32 1) #6, !dbg !25
  %204 = lshr i32 %29, 2, !dbg !25
  %205 = and i32 %204, 8, !dbg !25
  %206 = and i32 %29, 31, !dbg !25
  %207 = mul nuw nsw i32 %205, 33, !dbg !25
  %208 = add nuw nsw i32 %207, %206, !dbg !25
  %209 = zext nneg i32 %208 to i64, !dbg !25
  %210 = getelementptr float, ptr addrspace(3) %4, i64 %209, !dbg !25
  %211 = load <1 x float>, ptr addrspace(3) %210, align 4, !dbg !25
  %212 = getelementptr i8, ptr addrspace(3) %210, i64 132, !dbg !25
  %213 = load <1 x float>, ptr addrspace(3) %212, align 4, !dbg !25
  %214 = getelementptr i8, ptr addrspace(3) %210, i64 264, !dbg !25
  %215 = load <1 x float>, ptr addrspace(3) %214, align 4, !dbg !25
  %216 = getelementptr i8, ptr addrspace(3) %210, i64 396, !dbg !25
  %217 = load <1 x float>, ptr addrspace(3) %216, align 4, !dbg !25
  %218 = getelementptr i8, ptr addrspace(3) %210, i64 528, !dbg !25
  %219 = load <1 x float>, ptr addrspace(3) %218, align 4, !dbg !25
  %220 = getelementptr i8, ptr addrspace(3) %210, i64 660, !dbg !25
  %221 = load <1 x float>, ptr addrspace(3) %220, align 4, !dbg !25
  %222 = getelementptr i8, ptr addrspace(3) %210, i64 792, !dbg !25
  %223 = load <1 x float>, ptr addrspace(3) %222, align 4, !dbg !25
  %224 = getelementptr i8, ptr addrspace(3) %210, i64 924, !dbg !25
  %225 = load <1 x float>, ptr addrspace(3) %224, align 4, !dbg !25
  tail call spir_func void @_Z7barrierj(i32 1) #6, !dbg !25
  %226 = shl nuw nsw i32 %205, 5, !dbg !25
  %227 = or disjoint i32 %226, %206, !dbg !25
  %228 = zext nneg i32 %227 to i64, !dbg !25
  %229 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %228, !dbg !25
  store <1 x float> %211, ptr addrspace(3) %229, align 4, !dbg !25
  %230 = or disjoint i32 %227, 32, !dbg !25
  %231 = zext nneg i32 %230 to i64, !dbg !25
  %232 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %231, !dbg !25
  store <1 x float> %213, ptr addrspace(3) %232, align 4, !dbg !25
  %233 = or disjoint i32 %227, 64, !dbg !25
  %234 = zext nneg i32 %233 to i64, !dbg !25
  %235 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %234, !dbg !25
  store <1 x float> %215, ptr addrspace(3) %235, align 4, !dbg !25
  %236 = or disjoint i32 %227, 96, !dbg !25
  %237 = zext nneg i32 %236 to i64, !dbg !25
  %238 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %237, !dbg !25
  store <1 x float> %217, ptr addrspace(3) %238, align 4, !dbg !25
  %239 = or disjoint i32 %227, 128, !dbg !25
  %240 = zext nneg i32 %239 to i64, !dbg !25
  %241 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %240, !dbg !25
  store <1 x float> %219, ptr addrspace(3) %241, align 4, !dbg !25
  %242 = or disjoint i32 %227, 160, !dbg !25
  %243 = zext nneg i32 %242 to i64, !dbg !25
  %244 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %243, !dbg !25
  store <1 x float> %221, ptr addrspace(3) %244, align 4, !dbg !25
  %245 = or disjoint i32 %227, 192, !dbg !25
  %246 = zext nneg i32 %245 to i64, !dbg !25
  %247 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %246, !dbg !25
  store <1 x float> %223, ptr addrspace(3) %247, align 4, !dbg !25
  %248 = or disjoint i32 %227, 224, !dbg !25
  %249 = zext nneg i32 %248 to i64, !dbg !25
  %250 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %249, !dbg !25
  store <1 x float> %225, ptr addrspace(3) %250, align 4, !dbg !25
  tail call spir_func void @_Z7barrierj(i32 1) #6, !dbg !25
  %251 = or disjoint i32 %.lobit, %32, !dbg !25
  %252 = shl nuw nsw i32 %251, 5, !dbg !25
  %253 = or disjoint i32 %252, %34, !dbg !25
  %254 = zext nneg i32 %253 to i64, !dbg !25
  %255 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %254, !dbg !25
  %256 = load <1 x float>, ptr addrspace(3) %255, align 4, !dbg !25
  %257 = or disjoint i32 %252, 64, !dbg !25
  %258 = or disjoint i32 %257, %34, !dbg !25
  %259 = zext nneg i32 %258 to i64, !dbg !25
  %260 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %259, !dbg !25
  %261 = load <1 x float>, ptr addrspace(3) %260, align 4, !dbg !25
  %262 = or disjoint i32 %252, 128, !dbg !25
  %263 = or disjoint i32 %262, %34, !dbg !25
  %264 = zext nneg i32 %263 to i64, !dbg !25
  %265 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %264, !dbg !25
  %266 = load <1 x float>, ptr addrspace(3) %265, align 4, !dbg !25
  %267 = shufflevector <1 x float> %266, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %268 = or disjoint i32 %252, 192, !dbg !25
  %269 = or disjoint i32 %268, %34, !dbg !25
  %270 = zext nneg i32 %269 to i64, !dbg !25
  %271 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %270, !dbg !25
  %272 = load <1 x float>, ptr addrspace(3) %271, align 4, !dbg !25
  %273 = shufflevector <1 x float> %272, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %274 = or disjoint i32 %34, 8, !dbg !25
  %275 = or disjoint i32 %252, %274, !dbg !25
  %276 = zext nneg i32 %275 to i64, !dbg !25
  %277 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %276, !dbg !25
  %278 = load <1 x float>, ptr addrspace(3) %277, align 4, !dbg !25
  %279 = or disjoint i32 %257, %274, !dbg !25
  %280 = zext nneg i32 %279 to i64, !dbg !25
  %281 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %280, !dbg !25
  %282 = load <1 x float>, ptr addrspace(3) %281, align 4, !dbg !25
  %283 = or disjoint i32 %262, %274, !dbg !25
  %284 = zext nneg i32 %283 to i64, !dbg !25
  %285 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %284, !dbg !25
  %286 = load <1 x float>, ptr addrspace(3) %285, align 4, !dbg !25
  %287 = shufflevector <1 x float> %286, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %288 = or disjoint i32 %268, %274, !dbg !25
  %289 = zext nneg i32 %288 to i64, !dbg !25
  %290 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %289, !dbg !25
  %291 = load <1 x float>, ptr addrspace(3) %290, align 4, !dbg !25
  %292 = shufflevector <1 x float> %291, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %293 = or disjoint i32 %34, 16, !dbg !25
  %294 = or disjoint i32 %252, %293, !dbg !25
  %295 = zext nneg i32 %294 to i64, !dbg !25
  %296 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %295, !dbg !25
  %297 = load <1 x float>, ptr addrspace(3) %296, align 4, !dbg !25
  %298 = or disjoint i32 %257, %293, !dbg !25
  %299 = zext nneg i32 %298 to i64, !dbg !25
  %300 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %299, !dbg !25
  %301 = load <1 x float>, ptr addrspace(3) %300, align 4, !dbg !25
  %302 = or disjoint i32 %262, %293, !dbg !25
  %303 = zext nneg i32 %302 to i64, !dbg !25
  %304 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %303, !dbg !25
  %305 = load <1 x float>, ptr addrspace(3) %304, align 4, !dbg !25
  %306 = shufflevector <1 x float> %305, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %307 = or disjoint i32 %268, %293, !dbg !25
  %308 = zext nneg i32 %307 to i64, !dbg !25
  %309 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %308, !dbg !25
  %310 = load <1 x float>, ptr addrspace(3) %309, align 4, !dbg !25
  %311 = shufflevector <1 x float> %310, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %312 = or disjoint i32 %34, 24, !dbg !25
  %313 = or disjoint i32 %252, %312, !dbg !25
  %314 = zext nneg i32 %313 to i64, !dbg !25
  %315 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %314, !dbg !25
  %316 = load <1 x float>, ptr addrspace(3) %315, align 4, !dbg !25
  %317 = or disjoint i32 %257, %312, !dbg !25
  %318 = zext nneg i32 %317 to i64, !dbg !25
  %319 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %318, !dbg !25
  %320 = load <1 x float>, ptr addrspace(3) %319, align 4, !dbg !25
  %321 = or disjoint i32 %262, %312, !dbg !25
  %322 = zext nneg i32 %321 to i64, !dbg !25
  %323 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %322, !dbg !25
  %324 = load <1 x float>, ptr addrspace(3) %323, align 4, !dbg !25
  %325 = shufflevector <1 x float> %324, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %326 = or disjoint i32 %268, %312, !dbg !25
  %327 = zext nneg i32 %326 to i64, !dbg !25
  %328 = getelementptr inbounds nuw float, ptr addrspace(3) %4, i64 %327, !dbg !25
  %329 = load <1 x float>, ptr addrspace(3) %328, align 4, !dbg !25
  %330 = shufflevector <1 x float> %329, <1 x float> poison, <4 x i32> <i32 0, i32 poison, i32 poison, i32 poison>
  %331 = shufflevector <1 x float> %256, <1 x float> %261, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>, !dbg !34
  %332 = shufflevector <4 x float> %331, <4 x float> %267, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>, !dbg !34
  %333 = shufflevector <4 x float> %332, <4 x float> %273, <4 x i32> <i32 0, i32 1, i32 2, i32 4>, !dbg !34
  %334 = shufflevector <1 x float> %278, <1 x float> %282, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>, !dbg !34
  %335 = shufflevector <4 x float> %334, <4 x float> %287, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>, !dbg !34
  %336 = shufflevector <4 x float> %335, <4 x float> %292, <4 x i32> <i32 0, i32 1, i32 2, i32 4>, !dbg !34
  %337 = shufflevector <1 x float> %297, <1 x float> %301, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>, !dbg !34
  %338 = shufflevector <4 x float> %337, <4 x float> %306, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>, !dbg !34
  %339 = shufflevector <4 x float> %338, <4 x float> %311, <4 x i32> <i32 0, i32 1, i32 2, i32 4>, !dbg !34
  %340 = shufflevector <1 x float> %316, <1 x float> %320, <4 x i32> <i32 0, i32 1, i32 poison, i32 poison>, !dbg !34
  %341 = shufflevector <4 x float> %340, <4 x float> %325, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>, !dbg !34
  %342 = shufflevector <4 x float> %341, <4 x float> %330, <4 x i32> <i32 0, i32 1, i32 2, i32 4>, !dbg !34
  %343 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> %333, <8 x float> splat (float 1.000000e+00), <8 x float> zeroinitializer) #5
  %344 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> %336, <8 x float> splat (float 1.000000e+00), <8 x float> %343) #5
  %345 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> %339, <8 x float> splat (float 1.000000e+00), <8 x float> %344) #5
  %346 = tail call spir_func <8 x float> @_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_(<4 x float> %342, <8 x float> splat (float 1.000000e+00), <8 x float> %345) #5
  %347 = or disjoint i32 %50, %47, !dbg !35
  %348 = mul i32 %347, %3, !dbg !35
  %349 = or disjoint i32 %347, 4, !dbg !35
  %350 = mul i32 %349, %3, !dbg !35
  %351 = or disjoint i32 %347, 8, !dbg !35
  %352 = mul i32 %351, %3, !dbg !35
  %353 = or disjoint i32 %347, 12, !dbg !35
  %354 = mul i32 %353, %3, !dbg !35
  %355 = and i32 %55, 60, !dbg !36
  %356 = sext i32 %348 to i64, !dbg !37
  %357 = getelementptr float, ptr addrspace(1) %2, i64 %356, !dbg !37
  %358 = sext i32 %350 to i64, !dbg !37
  %359 = getelementptr float, ptr addrspace(1) %2, i64 %358, !dbg !37
  %360 = sext i32 %352 to i64, !dbg !37
  %361 = getelementptr float, ptr addrspace(1) %2, i64 %360, !dbg !37
  %362 = sext i32 %354 to i64, !dbg !37
  %363 = getelementptr float, ptr addrspace(1) %2, i64 %362, !dbg !37
  %364 = zext nneg i32 %355 to i64, !dbg !38
  %365 = getelementptr float, ptr addrspace(1) %357, i64 %364, !dbg !38
  %366 = getelementptr float, ptr addrspace(1) %359, i64 %364, !dbg !38
  %367 = getelementptr float, ptr addrspace(1) %361, i64 %364, !dbg !38
  %368 = getelementptr float, ptr addrspace(1) %363, i64 %364, !dbg !38
  tail call spir_func void @_Z7barrierj(i32 1) #6, !dbg !39
  %369 = mul nuw nsw i32 %65, 68, !dbg !39
  %370 = or disjoint i32 %369, %33, !dbg !39
  %371 = zext nneg i32 %370 to i64, !dbg !39
  %372 = getelementptr float, ptr addrspace(3) %4, i64 %371, !dbg !39
  %373 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> zeroinitializer, !dbg !39
  store <1 x float> %373, ptr addrspace(3) %372, align 4, !dbg !39
  %374 = mul nuw nsw i32 %71, 68, !dbg !39
  %375 = add nuw nsw i32 %374, %33, !dbg !39
  %376 = zext nneg i32 %375 to i64, !dbg !39
  %377 = getelementptr float, ptr addrspace(3) %4, i64 %376, !dbg !39
  %378 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> <i32 1>, !dbg !39
  store <1 x float> %378, ptr addrspace(3) %377, align 4, !dbg !39
  %379 = mul nuw nsw i32 %77, 68, !dbg !39
  %380 = add nuw nsw i32 %379, %33, !dbg !39
  %381 = zext nneg i32 %380 to i64, !dbg !39
  %382 = getelementptr float, ptr addrspace(3) %4, i64 %381, !dbg !39
  %383 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> <i32 2>, !dbg !39
  store <1 x float> %383, ptr addrspace(3) %382, align 4, !dbg !39
  %384 = mul nuw nsw i32 %83, 68, !dbg !39
  %385 = add nuw nsw i32 %384, %33, !dbg !39
  %386 = zext nneg i32 %385 to i64, !dbg !39
  %387 = getelementptr float, ptr addrspace(3) %4, i64 %386, !dbg !39
  %388 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> <i32 3>, !dbg !39
  store <1 x float> %388, ptr addrspace(3) %387, align 4, !dbg !39
  %389 = mul nuw nsw i32 %89, 68, !dbg !39
  %390 = or disjoint i32 %389, %33, !dbg !39
  %391 = zext nneg i32 %390 to i64, !dbg !39
  %392 = getelementptr float, ptr addrspace(3) %4, i64 %391, !dbg !39
  %393 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> <i32 4>, !dbg !39
  store <1 x float> %393, ptr addrspace(3) %392, align 4, !dbg !39
  %394 = mul nuw nsw i32 %95, 68, !dbg !39
  %395 = add nuw nsw i32 %394, %33, !dbg !39
  %396 = zext nneg i32 %395 to i64, !dbg !39
  %397 = getelementptr float, ptr addrspace(3) %4, i64 %396, !dbg !39
  %398 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> <i32 5>, !dbg !39
  store <1 x float> %398, ptr addrspace(3) %397, align 4, !dbg !39
  %399 = mul nuw nsw i32 %101, 68, !dbg !39
  %400 = add nuw nsw i32 %399, %33, !dbg !39
  %401 = zext nneg i32 %400 to i64, !dbg !39
  %402 = getelementptr float, ptr addrspace(3) %4, i64 %401, !dbg !39
  %403 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> <i32 6>, !dbg !39
  store <1 x float> %403, ptr addrspace(3) %402, align 4, !dbg !39
  %404 = mul nuw nsw i32 %107, 68, !dbg !39
  %405 = add nuw nsw i32 %404, %33, !dbg !39
  %406 = zext nneg i32 %405 to i64, !dbg !39
  %407 = getelementptr float, ptr addrspace(3) %4, i64 %406, !dbg !39
  %408 = shufflevector <8 x float> %346, <8 x float> poison, <1 x i32> <i32 7>, !dbg !39
  store <1 x float> %408, ptr addrspace(3) %407, align 4, !dbg !39
  %409 = or disjoint i32 %369, %113, !dbg !39
  %410 = zext nneg i32 %409 to i64, !dbg !39
  %411 = getelementptr float, ptr addrspace(3) %4, i64 %410, !dbg !39
  store <1 x float> %373, ptr addrspace(3) %411, align 4, !dbg !39
  %412 = add nuw nsw i32 %374, %113, !dbg !39
  %413 = zext nneg i32 %412 to i64, !dbg !39
  %414 = getelementptr float, ptr addrspace(3) %4, i64 %413, !dbg !39
  store <1 x float> %378, ptr addrspace(3) %414, align 4, !dbg !39
  %415 = add nuw nsw i32 %379, %113, !dbg !39
  %416 = zext nneg i32 %415 to i64, !dbg !39
  %417 = getelementptr float, ptr addrspace(3) %4, i64 %416, !dbg !39
  store <1 x float> %383, ptr addrspace(3) %417, align 4, !dbg !39
  %418 = add nuw nsw i32 %384, %113, !dbg !39
  %419 = zext nneg i32 %418 to i64, !dbg !39
  %420 = getelementptr float, ptr addrspace(3) %4, i64 %419, !dbg !39
  store <1 x float> %388, ptr addrspace(3) %420, align 4, !dbg !39
  %421 = add nuw nsw i32 %389, %113, !dbg !39
  %422 = zext nneg i32 %421 to i64, !dbg !39
  %423 = getelementptr float, ptr addrspace(3) %4, i64 %422, !dbg !39
  store <1 x float> %393, ptr addrspace(3) %423, align 4, !dbg !39
  %424 = add nuw nsw i32 %394, %113, !dbg !39
  %425 = zext nneg i32 %424 to i64, !dbg !39
  %426 = getelementptr float, ptr addrspace(3) %4, i64 %425, !dbg !39
  store <1 x float> %398, ptr addrspace(3) %426, align 4, !dbg !39
  %427 = add nuw nsw i32 %399, %113, !dbg !39
  %428 = zext nneg i32 %427 to i64, !dbg !39
  %429 = getelementptr float, ptr addrspace(3) %4, i64 %428, !dbg !39
  store <1 x float> %403, ptr addrspace(3) %429, align 4, !dbg !39
  %430 = add nuw nsw i32 %404, %113, !dbg !39
  %431 = zext nneg i32 %430 to i64, !dbg !39
  %432 = getelementptr float, ptr addrspace(3) %4, i64 %431, !dbg !39
  store <1 x float> %408, ptr addrspace(3) %432, align 4, !dbg !39
  %433 = or disjoint i32 %33, 32, !dbg !39
  %434 = add nuw nsw i32 %369, %433, !dbg !39
  %435 = zext nneg i32 %434 to i64, !dbg !39
  %436 = getelementptr float, ptr addrspace(3) %4, i64 %435, !dbg !39
  store <1 x float> %373, ptr addrspace(3) %436, align 4, !dbg !39
  %437 = add nuw nsw i32 %374, %433, !dbg !39
  %438 = zext nneg i32 %437 to i64, !dbg !39
  %439 = getelementptr float, ptr addrspace(3) %4, i64 %438, !dbg !39
  store <1 x float> %378, ptr addrspace(3) %439, align 4, !dbg !39
  %440 = add nuw nsw i32 %379, %433, !dbg !39
  %441 = zext nneg i32 %440 to i64, !dbg !39
  %442 = getelementptr float, ptr addrspace(3) %4, i64 %441, !dbg !39
  store <1 x float> %383, ptr addrspace(3) %442, align 4, !dbg !39
  %443 = add nuw nsw i32 %384, %433, !dbg !39
  %444 = zext nneg i32 %443 to i64, !dbg !39
  %445 = getelementptr float, ptr addrspace(3) %4, i64 %444, !dbg !39
  store <1 x float> %388, ptr addrspace(3) %445, align 4, !dbg !39
  %446 = add nuw nsw i32 %389, %433, !dbg !39
  %447 = zext nneg i32 %446 to i64, !dbg !39
  %448 = getelementptr float, ptr addrspace(3) %4, i64 %447, !dbg !39
  store <1 x float> %393, ptr addrspace(3) %448, align 4, !dbg !39
  %449 = add nuw nsw i32 %394, %433, !dbg !39
  %450 = zext nneg i32 %449 to i64, !dbg !39
  %451 = getelementptr float, ptr addrspace(3) %4, i64 %450, !dbg !39
  store <1 x float> %398, ptr addrspace(3) %451, align 4, !dbg !39
  %452 = add nuw nsw i32 %399, %433, !dbg !39
  %453 = zext nneg i32 %452 to i64, !dbg !39
  %454 = getelementptr float, ptr addrspace(3) %4, i64 %453, !dbg !39
  store <1 x float> %403, ptr addrspace(3) %454, align 4, !dbg !39
  %455 = add nuw nsw i32 %404, %433, !dbg !39
  %456 = zext nneg i32 %455 to i64, !dbg !39
  %457 = getelementptr float, ptr addrspace(3) %4, i64 %456, !dbg !39
  store <1 x float> %408, ptr addrspace(3) %457, align 4, !dbg !39
  %458 = or disjoint i32 %33, 48, !dbg !39
  %459 = add nuw nsw i32 %369, %458, !dbg !39
  %460 = zext nneg i32 %459 to i64, !dbg !39
  %461 = getelementptr float, ptr addrspace(3) %4, i64 %460, !dbg !39
  store <1 x float> %373, ptr addrspace(3) %461, align 4, !dbg !39
  %462 = add nuw nsw i32 %374, %458, !dbg !39
  %463 = zext nneg i32 %462 to i64, !dbg !39
  %464 = getelementptr float, ptr addrspace(3) %4, i64 %463, !dbg !39
  store <1 x float> %378, ptr addrspace(3) %464, align 4, !dbg !39
  %465 = add nuw nsw i32 %379, %458, !dbg !39
  %466 = zext nneg i32 %465 to i64, !dbg !39
  %467 = getelementptr float, ptr addrspace(3) %4, i64 %466, !dbg !39
  store <1 x float> %383, ptr addrspace(3) %467, align 4, !dbg !39
  %468 = add nuw nsw i32 %384, %458, !dbg !39
  %469 = zext nneg i32 %468 to i64, !dbg !39
  %470 = getelementptr float, ptr addrspace(3) %4, i64 %469, !dbg !39
  store <1 x float> %388, ptr addrspace(3) %470, align 4, !dbg !39
  %471 = add nuw nsw i32 %389, %458, !dbg !39
  %472 = zext nneg i32 %471 to i64, !dbg !39
  %473 = getelementptr float, ptr addrspace(3) %4, i64 %472, !dbg !39
  store <1 x float> %393, ptr addrspace(3) %473, align 4, !dbg !39
  %474 = add nuw nsw i32 %394, %458, !dbg !39
  %475 = zext nneg i32 %474 to i64, !dbg !39
  %476 = getelementptr float, ptr addrspace(3) %4, i64 %475, !dbg !39
  store <1 x float> %398, ptr addrspace(3) %476, align 4, !dbg !39
  %477 = add nuw nsw i32 %399, %458, !dbg !39
  %478 = zext nneg i32 %477 to i64, !dbg !39
  %479 = getelementptr float, ptr addrspace(3) %4, i64 %478, !dbg !39
  store <1 x float> %403, ptr addrspace(3) %479, align 4, !dbg !39
  %480 = add nuw nsw i32 %404, %458, !dbg !39
  %481 = zext nneg i32 %480 to i64, !dbg !39
  %482 = getelementptr float, ptr addrspace(3) %4, i64 %481, !dbg !39
  store <1 x float> %408, ptr addrspace(3) %482, align 4, !dbg !39
  tail call spir_func void @_Z7barrierj(i32 1) #6, !dbg !39
  %483 = shl nuw nsw i32 %33, 2, !dbg !39
  %484 = mul nuw nsw i32 %50, 68, !dbg !39
  %485 = add nuw nsw i32 %484, %483, !dbg !39
  %486 = zext nneg i32 %485 to i64, !dbg !39
  %487 = getelementptr float, ptr addrspace(3) %4, i64 %486, !dbg !39
  %488 = load <4 x i32>, ptr addrspace(3) %487, align 16, !dbg !39
  %489 = getelementptr i8, ptr addrspace(3) %487, i64 1088, !dbg !39
  %490 = load <4 x i32>, ptr addrspace(3) %489, align 16, !dbg !39
  %491 = getelementptr i8, ptr addrspace(3) %487, i64 2176, !dbg !39
  %492 = load <4 x i32>, ptr addrspace(3) %491, align 16, !dbg !39
  %493 = getelementptr i8, ptr addrspace(3) %487, i64 3264, !dbg !39
  %494 = load <4 x i32>, ptr addrspace(3) %493, align 16, !dbg !39
  store <4 x i32> %488, ptr addrspace(1) %365, align 16, !dbg !39
  store <4 x i32> %490, ptr addrspace(1) %366, align 16, !dbg !39
  store <4 x i32> %492, ptr addrspace(1) %367, align 16, !dbg !39
  store <4 x i32> %494, ptr addrspace(1) %368, align 16, !dbg !39
  ret void, !dbg !40
}

attributes #0 = { convergent mustprogress nounwind willreturn }
attributes #1 = { mustprogress nofree nosync nounwind willreturn memory(none) }
attributes #2 = { convergent mustprogress nofree nounwind willreturn memory(none) }
attributes #3 = { mustprogress nounwind willreturn }
attributes #4 = { nounwind willreturn memory(none) }
attributes #5 = { convergent nounwind willreturn memory(none) }
attributes #6 = { convergent nounwind willreturn }

!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.dbg.cu = !{!4}
!opencl.spir.version = !{!6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6}
!spirv.Source = !{!7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7}
!llvm.ident = !{!8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8}

!0 = !{i32 2, !"Debug Info Version", i32 3}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 1, !"sycl-device", i32 1}
!3 = !{i32 7, !"frame-pointer", i32 2}
!4 = distinct !DICompileUnit(language: DW_LANG_C, file: !5, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
!5 = !DIFile(filename: "where_repro.py", directory: "/home/jovyan/sglang/test/srt/../../../intel-xpu-backend-for-triton")
!6 = !{i32 1, i32 2}
!7 = !{i32 3, i32 100000}
!8 = !{!"Intel(R) oneAPI DPC++/C++ Compiler 2025.0.0 (2025.0.0.20241008)"}
!9 = !DISubprogram(name: "_Z7barrierj", linkageName: "_Z7barrierj", scope: !5, file: !5, line: 7, type: !10, scopeLine: 7, spFlags: DISPFlagOptimized)
!10 = !DISubroutineType(cc: DW_CC_normal, types: !11)
!11 = !{}
!12 = !DISubprogram(name: "_Z12get_local_idj", linkageName: "_Z12get_local_idj", scope: !5, file: !5, line: 7, type: !10, scopeLine: 7, spFlags: DISPFlagOptimized)
!13 = !DISubprogram(name: "_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_", linkageName: "_Z39intel_sub_group_tf32_tf32_matrix_mad_k8Dv4_fDv8_fS0_", scope: !5, file: !5, line: 7, type: !10, scopeLine: 7, spFlags: DISPFlagOptimized)
!14 = !DISubprogram(name: "_Z12get_group_idj", linkageName: "_Z12get_group_idj", scope: !5, file: !5, line: 7, type: !10, scopeLine: 7, spFlags: DISPFlagOptimized)
!15 = distinct !DISubprogram(name: "where_repro", linkageName: "where_repro", scope: !5, file: !5, line: 7, type: !10, scopeLine: 7, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !4)
!16 = !{i32 16}
!17 = !{i64 64, i64 1, i64 1}
!18 = !DILocation(line: 15, column: 23, scope: !15)
!19 = !DILocation(line: 16, column: 23, scope: !15)
!20 = !DILocation(line: 17, column: 23, scope: !15)
!21 = !DILocation(line: 28, column: 20, scope: !15)
!22 = !DILocation(line: 30, column: 29, scope: !15)
!23 = !DILocation(line: 31, column: 29, scope: !15)
!24 = !DILocation(line: 31, column: 41, scope: !15)
!25 = !DILocation(line: 40, column: 70, scope: !15)
!26 = !DILocation(line: 43, column: 34, scope: !15)
!27 = !DILocation(line: 44, column: 33, scope: !15)
!28 = !DILocation(line: 44, column: 88, scope: !15)
!29 = !DILocation(line: 44, column: 65, scope: !15)
!30 = !DILocation(line: 45, column: 45, scope: !15)
!31 = !DILocation(line: 47, column: 29, scope: !15)
!32 = !DILocation(line: 47, column: 47, scope: !15)
!33 = !DILocation(line: 47, column: 64, scope: !15)
!34 = !DILocation(line: 53, column: 27, scope: !15)
!35 = !DILocation(line: 56, column: 58, scope: !15)
!36 = !DILocation(line: 57, column: 45, scope: !15)
!37 = !DILocation(line: 59, column: 22, scope: !15)
!38 = !DILocation(line: 59, column: 40, scope: !15)
!39 = !DILocation(line: 59, column: 57, scope: !15)
!40 = !DILocation(line: 59, column: 4, scope: !15)

[alexbaden]

Great reproducer - I think the issue is occurring because we are chaining the two dot ops through the select:

    %3 = tt.dot %cst_3, %cst_2, %cst_4, inputPrecision = tf32 : tensor<16x64xf32, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 1}>> * tensor<64x32xf32, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 1}>> -> tensor<16x32xf32, #mma> loc(#loc5)

    %13 = arith.select %12, %3, %cst : tensor<16x32xi1, #mma>, tensor<16x32xf32, #mma> loc(#loc13)

%35 = ttg.convert_layout %13 {allocation.offset = 0 : i32} : tensor<16x32xf32, #mma> -> tensor<16x32xf32, #blocked> loc(#loc26)
    tt.store %34, %35 : tensor<16x32x!tt.ptr<f32>, #blocked> loc(#loc26)
    %36 = ttg.convert_layout %13 : tensor<16x32xf32, #mma> -> tensor<16x32xf32, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 1}>> loc(#loc13)
    %37 = tt.dot %36, %cst_0, %cst_1, inputPrecision = tf32 : tensor<16x32xf32, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 1}>> * tensor<32x64xf32, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 1}>> -> tensor<16x64xf32, #mma> loc(#loc27)

if I remove the first dot op, then the test passes.

I think the issue could be in either the conversion from the tensor<16x32xf32, #mma> to the dot operand layout (%36) or in the output of the first dot op passing into the select. I looked at the layout conversions for mma to dot operand layout and they seem reasonable in the Intel dialect, but the tf32 DPAS layouts are tricky. I also tried falling back to the legacy path instead of using the new linear layout path, but the legacy path has the same issue - which probably points to the src or dst layouts not accurately reflecting the output.

[dchigarev]

After some debugging, I found that the mask values for the second 16-column block in arith.select were computed correctly but were simply ignored and replaced with the mask values for the first block.

It appears that this maybeDeduplicate function does something silly, as it assumes the mask values of the second block are duplicates of those in the first block and drops them.

Simply removing this line, which calls the maybeDeduplicate function, makes both the reproducer and the corresponding test from sglang (test_grouped_decode_attention) work as expected.

---

@airMeng I'll look into how to properly fix the issue, and will merge the fix into the main branch as soon as it’s ready. For now, you can try my triton branch with a temporary fix on your side: https://github.com/intel/intel-xpu-backend-for-triton/tree/dchigarev/3380_preview

In order to install it:

git clone https://github.com/intel/intel-xpu-backend-for-triton.git
cd intel-xpu-backend-for-triton
git checkout dchigarev/3380_preview
./scripts/compile-triton.sh --triton