@triton.jit
def wkv_triton_vanilla_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_ab, 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_ab,
    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_ab, gstate_s_c, tsz, chans, 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_ptr = state_ptr + b_idx * state_s_b
    beta_ptr = state_ptr + b_idx * state_s_b + state_s_ab
    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_out_ptr = gstate_out_ptr + b_idx * gstate_out_s_b
    gbeta_out_ptr = gstate_out_ptr + b_idx * gstate_out_s_b + gstate_out_s_ab
    galpha = tl.load(galpha_out_ptr + gstate_out_s_c * cs, mask=cmask).to(tl
        .float32)
    gbeta = 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)
    ew = tl.exp(w)
    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
            )
        alpha_prev = tl.load(alpha_ptr + tc * state_s_t + state_s_c * cs,
            mask=cmask).to(tl.float32)
        beta_prev = tl.load(beta_ptr + tc * state_s_t + state_s_c * cs,
            mask=cmask).to(tl.float32)
        euk = tl.exp(u + kt)
        ek = tl.exp(kt)
        denom = beta_prev + euk
        denom_sq = denom * denom
        gwkvt = tl.load(gwkv_ptr + tc * gwkv_s_t + gwkv_s_c * cs, mask=cmask
            ).to(tl.float32)
        guk = gwkvt * euk * (beta_prev * vt - alpha_prev) / denom_sq
        gu += guk
        gk = guk
        gv = gwkvt * euk / denom
        galpha_wkv = gwkvt / denom
        gbeta_wkv = -gwkvt * (euk * vt + alpha_prev) / denom_sq
        gw += galpha * ew * alpha_prev
        gk += galpha * ek * vt
        gv += galpha * ek
        gw += gbeta * ew * beta_prev
        gk += gbeta * ek
        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)
        galpha = galpha * ew + galpha_wkv
        gbeta = gbeta * ew + gbeta_wkv
    galpha_ptr = gstate_ptr + b_idx * gstate_s_b
    gbeta_ptr = gstate_ptr + b_idx * gstate_s_b + gstate_s_ab
    tl.store(galpha_ptr + gstate_s_c * cs, galpha, mask=cmask)
    tl.store(gbeta_ptr + gstate_s_c * cs, gbeta, 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)
