@triton.jit
def chunk_mlstm_bwd_kernel_dC(q, f, m, m_total, norm, dh, dC, final_dC,
    final_m, initial_dC, initial_m, 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, 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_v, i_bC = tl.program_id(0), tl.program_id(1), tl.program_id(2)
    p_dC = tl.make_block_ptr(final_dC + i_bC * K * V, (K, V), (s_C_t, 1), (
        i_k * BK, i_v * BV), (BK, BV), (1, 0))
    b_dC = tl.load(p_dC, boundary_check=(0, 1))
    b_m = tl.load(final_m + i_bC)
    for i_t in range(NT - 1, -1, -1):
        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_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 + i_t * K * V, (K, V), (
            s_C_t, 1), (i_k * BK, i_v * BV), (BK, BV), (1, 0))
        tl.store(p_dC, b_dC.to(p_dC.dtype.element_ty), boundary_check=(0, 1))
        b_f_last = tl.load(f + i_bC * T + i_t * BT + BT - 1)
        b_f = tl.load(f + i_bC * T + i_t * BT + tl.arange(0, BT))
        b_m_p = 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_q = tl.load(p_q, boundary_check=(0, 1))
        b_q = (b_q * scale * tl.math.exp2(b_f + b_m_p - b_m_total)[None, :]
            ).to(b_q.dtype)
        b_dh = tl.load(p_dh, boundary_check=(0, 1))
        b_dh /= b_norm[:, None]
        b_dC *= tl.math.exp2(b_f_last + b_m_p - b_m)
        b_dC += tl.dot(b_q, b_dh.to(b_q.dtype), allow_tf32=False)
        b_m = b_m_p
    p_initial_dC = tl.make_block_ptr(initial_dC + i_bC * K * V, (K, V), (V,
        1), (i_k * BK, i_v * BV), (BK, BV), (1, 0))
    tl.store(p_initial_dC, b_dC.to(p_initial_dC.dtype.element_ty),
        boundary_check=(0, 1))
    tl.store(initial_m + i_bC, b_m)
