@triton.jit
def chunk_mlstm_bwd_kernel_dqkvif(q, k, v, C, m, m_total, norm, i, f, dh,
    dC, dq, dk, dv, s_qk_h, s_qk_t, s_qk_d, s_vh_h, s_vh_t, s_vh_d, s_C_h,
    s_C_t, scale, B: tl.constexpr, H: tl.constexpr, T: tl.constexpr, K: tl.
    constexpr, V: tl.constexpr, BT: tl.constexpr, BK: tl.constexpr, BV: tl.
    constexpr, NT: tl.constexpr):
    i_k, i_t, i_bC = tl.program_id(0), tl.program_id(1), tl.program_id(2)
    n_bh = tl.num_programs(2)
    o_i = tl.arange(0, BT)
    p_q = tl.make_block_ptr(q + i_bC * s_qk_h, (K, T), (s_qk_d, s_qk_t), (
        i_k * BK, i_t * BT), (BK, BT), (0, 1))
    p_k = tl.make_block_ptr(k + i_bC * s_qk_h, (T, K), (s_qk_t, s_qk_d), (
        i_t * BT, i_k * BK), (BT, BK), (1, 0))
    b_q = tl.load(p_q, boundary_check=(0, 1))
    b_k = tl.load(p_k, boundary_check=(0, 1))
    b_s = tl.dot(b_k, b_q, allow_tf32=False)
    p_f = f + i_bC * T + i_t * BT + tl.arange(0, BT)
    p_i = i + i_bC * T + i_t * BT + tl.arange(0, BT)
    b_f = tl.load(p_f)
    b_f_last = tl.load(f + i_bC * T + i_t * BT + BT - 1)
    b_m = tl.load(m + i_bC * (NT + 1) + i_t)
    b_m_total = tl.load(m_total + i_bC * T + i_t * BT + tl.arange(0, BT))
    b_norm = tl.load(norm + i_bC * T + i_t * BT + tl.arange(0, BT))
    b_i = tl.load(p_i)
    mask = tl.math.exp2(b_i[:, None] + b_f[None, :] - b_f[:, None] -
        b_m_total[None, :])
    mask = tl.where(o_i[:, None] <= o_i[None, :], mask * scale, 0)
    b_s = b_s * mask
    b_m_next = tl.load(m + i_bC * (NT + 1) + i_t + 1)
    b_dq = tl.zeros([BT, BK], dtype=tl.float32)
    b_dk = tl.zeros([BT, BK], dtype=tl.float32)
    b_ds = tl.zeros([BT, BT], dtype=tl.float32)
    for i_v in range(tl.cdiv(V, BV)):
        p_v = tl.make_block_ptr(v + i_bC * s_vh_h, (T, V), (s_vh_t, s_vh_d),
            (i_t * BT, i_v * BV), (BT, BV), (1, 0))
        p_C = tl.make_block_ptr(C + i_bC * s_C_h, (V, NT * K), (1, s_C_t),
            (i_v * BV, i_t * K + i_k * BK), (BV, BK), (0, 1))
        p_dh = tl.make_block_ptr(dh + i_bC * s_vh_h, (T, V), (s_vh_t,
            s_vh_d), (i_t * BT, i_v * BV), (BT, BV), (1, 0))
        p_dC = tl.make_block_ptr(dC + i_bC * s_C_h, (NT * K, V), (s_C_t, 1),
            (i_t * K + i_k * BK, i_v * BV), (BK, BV), (1, 0))
        p_dv = tl.make_block_ptr(dv + (i_k * n_bh + i_bC) * s_vh_h, (T, V),
            (s_vh_t, s_vh_d), (i_t * BT, i_v * BV), (BT, BV), (1, 0))
        b_v = tl.load(p_v, boundary_check=(0, 1))
        b_dh = tl.load(p_dh, boundary_check=(0, 1))
        b_C = tl.load(p_C, boundary_check=(0, 1))
        b_dC = tl.load(p_dC, boundary_check=(0, 1))
        b_ds += tl.dot(b_dh, tl.trans(b_v), allow_tf32=False)
        b_dq += tl.dot(b_dh, b_C, allow_tf32=False) * scale
        b_dk += tl.dot(b_v, tl.trans(b_dC), allow_tf32=False)
        b_dv = tl.dot(b_k, b_dC, allow_tf32=False) * tl.math.exp2(b_i - b_f +
            b_f_last - b_m_next)[:, None] + tl.dot(b_s.to(b_q.dtype) /
            b_norm[None, :], b_dh, allow_tf32=False)
        tl.store(p_dv, b_dv.to(p_dv.dtype.element_ty), boundary_check=(0, 1))
    b_dq = b_dq * tl.math.exp2(b_f + b_m - b_m_total)[:, None] / b_norm[:, None
        ]
    b_dk = b_dk * tl.math.exp2(b_i - b_f + b_f_last - b_m_next)[:, None]
    b_ds = b_ds * tl.trans(mask)
    b_ds = b_ds.to(b_k.dtype)
    b_dq += tl.dot(b_ds, b_k, allow_tf32=False) / b_norm[:, None]
    b_dk += tl.trans(tl.dot(b_q / b_norm[None, :], b_ds, allow_tf32=False))
    p_dq = tl.make_block_ptr(dq + i_bC * s_qk_h, (T, K), (s_qk_t, s_qk_d),
        (i_t * BT, i_k * BK), (BT, BK), (1, 0))
    p_dk = tl.make_block_ptr(dk + i_bC * s_qk_h, (T, K), (s_qk_t, s_qk_d),
        (i_t * BT, i_k * BK), (BT, BK), (1, 0))
    tl.store(p_dq, b_dq.to(p_dq.dtype.element_ty), boundary_check=(0, 1))
    tl.store(p_dk, b_dk.to(p_dk.dtype.element_ty), boundary_check=(0, 1))
