@triton.jit
def wkv_triton_log_space_backward_kernel(w_ptr, w_s_c, u_ptr, u_s_c, k_ptr,
    k_s_b, k_s_t, k_s_c, v_ptr, v_s_b, v_s_t, v_s_c, state_ptr, state_s_b,
    state_s_abe, state_s_t, state_s_c, gwkv_ptr, gwkv_s_b, gwkv_s_t,
    gwkv_s_c, gstate_out_ptr, gstate_out_s_b, gstate_out_s_abe,
    gstate_out_s_c, gw_ptr, gw_s_c, gu_ptr, gu_s_c, gk_ptr, gk_s_b, gk_s_t,
    gk_s_c, gv_ptr, gv_s_b, gv_s_t, gv_s_c, gstate_ptr, gstate_s_b,
    gstate_s_abe, gstate_s_c, tsz, chans, eps: tl.constexpr, BLOCK_SIZE_C:
    tl.constexpr):
    b_idx = tl.program_id(0)
    c_idx = tl.program_id(1)
    cs = c_idx * BLOCK_SIZE_C + tl.arange(0, BLOCK_SIZE_C)
    cmask = cs < chans
    k_ptr = k_ptr + b_idx * k_s_b
    v_ptr = v_ptr + b_idx * v_s_b
    alpha_p_ptr = state_ptr + b_idx * state_s_b
    alpha_m_ptr = state_ptr + b_idx * state_s_b + state_s_abe
    beta_ptr = state_ptr + b_idx * state_s_b + 2 * state_s_abe
    gk_ptr = gk_ptr + b_idx * gk_s_b
    gv_ptr = gv_ptr + b_idx * gv_s_b
    gwkv_ptr = gwkv_ptr + b_idx * gwkv_s_b
    galpha_p_out_ptr = gstate_out_ptr + b_idx * gstate_out_s_b
    galpha_m_out_ptr = (gstate_out_ptr + b_idx * gstate_out_s_b +
        gstate_out_s_abe)
    gbeta_out_ptr = (gstate_out_ptr + b_idx * gstate_out_s_b + 2 *
        gstate_out_s_abe)
    gln_alpha_p = tl.load(galpha_p_out_ptr + gstate_out_s_c * cs, mask=cmask
        ).to(tl.float32)
    gln_alpha_m = tl.load(galpha_m_out_ptr + gstate_out_s_c * cs, mask=cmask
        ).to(tl.float32)
    gln_beta = tl.load(gbeta_out_ptr + gstate_out_s_c * cs, mask=cmask).to(tl
        .float32)
    w = tl.load(w_ptr + w_s_c * cs, mask=cmask).to(tl.float32)
    u = tl.load(u_ptr + u_s_c * cs, mask=cmask).to(tl.float32)
    gw = tl.zeros_like(w)
    gu = tl.zeros_like(u)
    for t in range(tsz):
        tc = tsz - t - 1
        kt = tl.load(k_ptr + tc * k_s_t + k_s_c * cs, mask=cmask).to(tl.float32
            )
        vt = tl.load(v_ptr + tc * v_s_t + v_s_c * cs, mask=cmask).to(tl.float32
            )
        vt_p = tl.maximum(vt, 0) + eps
        vt_m = tl.maximum(-vt, 0) + eps
        ln_v_p = tl.log(vt_p)
        ln_v_m = tl.log(vt_m)
        ln_alpha_p_prev = tl.load(alpha_p_ptr + tc * state_s_t + state_s_c *
            cs, mask=cmask).to(tl.float32)
        ln_alpha_m_prev = tl.load(alpha_m_ptr + tc * state_s_t + state_s_c *
            cs, mask=cmask).to(tl.float32)
        ln_beta_prev = tl.load(beta_ptr + tc * state_s_t + state_s_c * cs,
            mask=cmask).to(tl.float32)
        uk = u + kt
        ukv_p = uk + ln_v_p
        ukv_m = uk + ln_v_m
        ukb = logaddexp(uk, ln_beta_prev)
        wkv_p = tl.exp(logaddexp(ukv_p, ln_alpha_p_prev) - ukb)
        wkv_m = tl.exp(logaddexp(ukv_m, ln_alpha_m_prev) - ukb)
        gwkvt = tl.load(gwkv_ptr + tc * gwkv_s_t + gwkv_s_c * cs, mask=cmask
            ).to(tl.float32)
        gln_wkv_p = gwkvt * wkv_p
        gln_wkv_m = gwkvt * -wkv_m
        e_num_p = tl.exp(ln_alpha_p_prev - ukv_p)
        e_num_m = tl.exp(ln_alpha_m_prev - ukv_m)
        e_den = tl.exp(ln_beta_prev - uk)
        grad_wkv_den_p = gln_wkv_p / (1 + e_den)
        grad_wkv_den_m = gln_wkv_m / (1 + e_den)
        gkv_p = gln_wkv_p / (1 + e_num_p)
        gkv_m = gln_wkv_m / (1 + e_num_m)
        grad_uk = gkv_p + gkv_m - grad_wkv_den_p - grad_wkv_den_m
        gu += grad_uk
        gk = grad_uk
        gv = tl.where(vt > 0, gkv_p / vt_p, gkv_m / -vt_m)
        gln_alpha_wkv_p = gln_wkv_p / (1 + 1 / e_num_p)
        gln_alpha_wkv_m = gln_wkv_m / (1 + 1 / e_num_m)
        gln_beta_wkv = -gln_wkv_p / (1 + 1 / e_den) - gln_wkv_m / (1 + 1 /
            e_den)
        e_alpha_p = tl.exp(kt + ln_v_p - (w + ln_alpha_p_prev))
        e_alpha_m = tl.exp(kt + ln_v_m - (w + ln_alpha_m_prev))
        gwa_p = gln_alpha_p / (1 + e_alpha_p)
        gwa_m = gln_alpha_m / (1 + e_alpha_m)
        gkv_p = gln_alpha_p / (1 + 1 / e_alpha_p)
        gkv_m = gln_alpha_m / (1 + 1 / e_alpha_m)
        gw += gwa_p + gwa_m
        gk += gkv_p + gkv_m
        gv += tl.where(vt > 0, gkv_p / vt_p, -gkv_m / vt_m)
        e_beta = tl.exp(kt - (w + ln_beta_prev))
        gwb = gln_beta / (1 + e_beta)
        gw += gwb
        gk += gln_beta / (1 + 1 / e_beta)
        tl.store(gk_ptr + tc * gk_s_t + gk_s_c * cs, gk, mask=cmask)
        tl.store(gv_ptr + tc * gv_s_t + gv_s_c * cs, gv, mask=cmask)
        gln_alpha_p = gwa_p + gln_alpha_wkv_p
        gln_alpha_m = gwa_m + gln_alpha_wkv_m
        gln_beta = gwb + gln_beta_wkv
    gln_alpha_p_ptr = gstate_ptr + b_idx * gstate_s_b
    gln_alpha_m_ptr = gstate_ptr + b_idx * gstate_s_b + gstate_s_abe
    gln_beta_ptr = gstate_ptr + b_idx * gstate_s_b + 2 * gstate_s_abe
    tl.store(gln_alpha_p_ptr + gstate_s_c * cs, gln_alpha_p, mask=cmask)
    tl.store(gln_alpha_m_ptr + gstate_s_c * cs, gln_alpha_m, mask=cmask)
    tl.store(gln_beta_ptr + gstate_s_c * cs, gln_beta, mask=cmask)
    tl.atomic_add(gw_ptr + gw_s_c * cs, gw, mask=cmask)
    tl.atomic_add(gu_ptr + gu_s_c * cs, gu, mask=cmask)
