@triton.jit
def _sparse_attention_prefill_fwd_kernel(Q, K, V, sm_scale, Out, lut,
    stride_qz, stride_qh, stride_qm, stride_qk, stride_kz, stride_kh,
    stride_kn, stride_kk, stride_vz, stride_vh, stride_vk, stride_vn,
    stride_oz, stride_oh, stride_om, stride_on, stride_lz, stride_lh,
    stride_lx, Z, H, N_CTX, LT, NNZ: tl.constexpr, BLOCK_M: tl.constexpr,
    BLOCK_DMODEL: tl.constexpr, BLOCK_N: tl.constexpr):
    start_m = tl.program_id(0)
    off_hz = tl.program_id(1)
    lut_indicator = tl.program_id(1) % H
    qvk_offset = off_hz * stride_qh
    lut_offset = lut_indicator * stride_lz
    Q_block_ptr = tl.make_block_ptr(base=Q + qvk_offset, shape=(N_CTX,
        BLOCK_DMODEL), strides=(stride_qm, stride_qk), offsets=(start_m *
        BLOCK_M, 0), block_shape=(BLOCK_M, BLOCK_DMODEL), order=(1, 0))
    K_block_ptr = tl.make_block_ptr(base=K + qvk_offset, shape=(
        BLOCK_DMODEL, N_CTX), strides=(stride_kk, stride_kn), offsets=(0, 0
        ), block_shape=(BLOCK_DMODEL, BLOCK_N), order=(0, 1))
    V_block_ptr = tl.make_block_ptr(base=V + qvk_offset, shape=(N_CTX,
        BLOCK_DMODEL), strides=(stride_vk, stride_vn), offsets=(0, 0),
        block_shape=(BLOCK_N, BLOCK_DMODEL), order=(1, 0))
    O_block_ptr = tl.make_block_ptr(base=Out + qvk_offset, shape=(N_CTX,
        BLOCK_DMODEL), strides=(stride_om, stride_on), offsets=(start_m *
        BLOCK_M, 0), block_shape=(BLOCK_M, BLOCK_DMODEL), order=(1, 0))
    offs_m = start_m * BLOCK_M + tl.arange(0, BLOCK_M)
    offs_n = tl.arange(0, BLOCK_N)
    m_i = tl.zeros([BLOCK_M], dtype=tl.float32) - float('inf')
    l_i = tl.zeros([BLOCK_M], dtype=tl.float32)
    acc = tl.zeros([BLOCK_M, BLOCK_DMODEL], dtype=tl.float32)
    qk_scale = sm_scale * 1.44269504
    q = tl.load(Q_block_ptr, boundary_check=(0, 1), padding_option='zero')
    q = (q * qk_scale).to(tl.float16)
    last_nnz_id = -1
    for nnz_id in range(NNZ):
        present_nnz_id = tl.load(lut + lut_offset + start_m * stride_lh + 
            nnz_id * stride_lx)
        start_n = present_nnz_id * BLOCK_N
        start_n = tl.multiple_of(start_n, BLOCK_N)
        present_nnz_id = present_nnz_id.to(tl.int32)
        k = tl.load(tl.advance(K_block_ptr, (0, start_n)), boundary_check=(
            0, 1), padding_option='zero')
        qk = tl.zeros([BLOCK_M, BLOCK_N], dtype=tl.float32)
        qk += tl.dot(q, k)
        if LT:
            qk = tl.where(offs_m[:, None] >= start_n + offs_n[None, :], qk,
                float('-inf'))
        qk = tl.where((offs_m[:, None] < N_CTX) & ((start_n + offs_n)[None,
            :] < N_CTX), qk, float('-inf'))
        m_ij = tl.max(qk, 1)
        p = tl.math.exp2(qk - m_ij[:, None])
        p = tl.where(m_ij[:, None] == tl.full((BLOCK_M, BLOCK_N), float(
            '-inf'), tl.float32), 0.0, tl.math.exp2(qk - m_ij[:, None]))
        p = p * (last_nnz_id != present_nnz_id)
        l_ij = tl.sum(p, 1)
        m_i_new = tl.maximum(m_i, m_ij)
        alpha = tl.math.exp2(m_i - m_i_new)
        beta = tl.math.exp2(m_ij - m_i_new)
        l_i *= alpha
        l_i_new = l_i + beta * l_ij
        p_scale = beta / l_i_new
        p = p * p_scale[:, None]
        acc_scale = l_i / l_i_new
        acc = acc * acc_scale[:, None]
        v = tl.load(tl.advance(V_block_ptr, (start_n, 0)), boundary_check=(
            0, 1), padding_option='zero')
        p = p.to(tl.float16)
        acc += tl.dot(p, v)
        l_i = l_i_new
        m_i = m_i_new
        last_nnz_id = present_nnz_id
    tl.store(O_block_ptr, acc.to(tl.float16), boundary_check=(0, 1))
