@triton.jit
def _backward_sequence_kernel(delta_states_all_outside,
    delta_states_last_outside, R, states_all, gates_all,
    delta_states_initial, delta_Wx, delta_R, delta_b, T: tl.constexpr, NS:
    tl.constexpr, B: tl.constexpr, NH: tl.constexpr, DH: tl.constexpr, NGI:
    tl.constexpr, NGR: tl.constexpr, siz_B: tl.constexpr, DTYPE: tl.
    constexpr=tl.float32, backward_recurrent_clip_val: tl.constexpr=-1.0):
    idx_b_NH, idx_b_B = tl.program_id(0), tl.program_id(1)
    str_matR_B = NH * NGR * DH * DH
    str_matR_NH = NGR * DH * DH
    str_matR_NGR = DH * DH
    str_matStatesAll_NH = (T + 1) * NS * B * DH
    str_matStatesAll_T = NS * B * DH
    str_matGatesAll_NH = T * NGI * B * DH
    str_matGatesAll_T = NGI * B * DH
    str_delta_states_all_outside_NH = T * NS * B * DH
    str_delta_states_all_outside_T = NS * B * DH
    str_matDeltaWx_NH = T * NGI * B * DH
    str_matDeltaWx_T = NGI * B * DH
    matDeltaHtrans_last_ptr = tl.make_block_ptr(base=
        delta_states_last_outside + idx_b_NH * NS * B * DH + 0 * B * DH,
        shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
        block_shape=(siz_B, DH), order=(0, 1))
    matDeltaH_tplus1 = tl.load(matDeltaHtrans_last_ptr).to(tl.float32)
    matDeltaCtrans_last_ptr = tl.make_block_ptr(base=
        delta_states_last_outside + idx_b_NH * NS * B * DH + 1 * B * DH,
        shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
        block_shape=(siz_B, DH), order=(0, 1))
    matDeltaC_tplus1 = tl.load(matDeltaCtrans_last_ptr).to(tl.float32)
    matDeltaNtrans_last_ptr = tl.make_block_ptr(base=
        delta_states_last_outside + idx_b_NH * NS * B * DH + 2 * B * DH,
        shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
        block_shape=(siz_B, DH), order=(0, 1))
    matDeltaN_tplus1 = tl.load(matDeltaNtrans_last_ptr).to(tl.float32)
    matR_i_ptr = tl.make_block_ptr(base=R + idx_b_NH * str_matR_NH + 0 *
        str_matR_NGR, shape=(DH, DH), strides=(DH, 1), offsets=(0, 0),
        block_shape=(DH, DH), order=(0, 1))
    matR_i = tl.load(matR_i_ptr)
    matR_f_ptr = tl.make_block_ptr(base=R + idx_b_NH * str_matR_NH + 1 *
        str_matR_NGR, shape=(DH, DH), strides=(DH, 1), offsets=(0, 0),
        block_shape=(DH, DH), order=(0, 1))
    matR_f = tl.load(matR_f_ptr)
    matR_z_ptr = tl.make_block_ptr(base=R + idx_b_NH * str_matR_NH + 2 *
        str_matR_NGR, shape=(DH, DH), strides=(DH, 1), offsets=(0, 0),
        block_shape=(DH, DH), order=(0, 1))
    matR_z = tl.load(matR_z_ptr)
    matR_o_ptr = tl.make_block_ptr(base=R + idx_b_NH * str_matR_NH + 3 *
        str_matR_NGR, shape=(DH, DH), strides=(DH, 1), offsets=(0, 0),
        block_shape=(DH, DH), order=(0, 1))
    matR_o = tl.load(matR_o_ptr)
    matDeltaR_i = tl.zeros((DH, DH), dtype=tl.float32)
    matDeltaR_f = tl.zeros((DH, DH), dtype=tl.float32)
    matDeltaR_z = tl.zeros((DH, DH), dtype=tl.float32)
    matDeltaR_o = tl.zeros((DH, DH), dtype=tl.float32)
    vecDeltaB_i = tl.zeros((DH,), dtype=tl.float32)
    vecDeltaB_f = tl.zeros((DH,), dtype=tl.float32)
    vecDeltaB_z = tl.zeros((DH,), dtype=tl.float32)
    vecDeltaB_o = tl.zeros((DH,), dtype=tl.float32)
    for idx_t in range(T - 1, -1, -1):
        matG_i_ptr = tl.make_block_ptr(base=gates_all + idx_b_NH *
            str_matGatesAll_NH + idx_t * str_matGatesAll_T + 0 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matG_ibar = tl.load(matG_i_ptr).to(tl.float32)
        matG_f_ptr = tl.make_block_ptr(base=gates_all + idx_b_NH *
            str_matGatesAll_NH + idx_t * str_matGatesAll_T + 1 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matG_fbar = tl.load(matG_f_ptr).to(tl.float32)
        matG_z_ptr = tl.make_block_ptr(base=gates_all + idx_b_NH *
            str_matGatesAll_NH + idx_t * str_matGatesAll_T + 2 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matG_z = tl.load(matG_z_ptr)
        matG_o_ptr = tl.make_block_ptr(base=gates_all + idx_b_NH *
            str_matGatesAll_NH + idx_t * str_matGatesAll_T + 3 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matG_o = tl.load(matG_o_ptr)
        matC_t_ptr = tl.make_block_ptr(base=states_all + idx_b_NH *
            str_matStatesAll_NH + (idx_t + 1) * str_matStatesAll_T + 1 * B *
            DH, shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0
            ), block_shape=(siz_B, DH), order=(0, 1))
        matC_t = tl.load(matC_t_ptr)
        matN_t_ptr = tl.make_block_ptr(base=states_all + idx_b_NH *
            str_matStatesAll_NH + (idx_t + 1) * str_matStatesAll_T + 2 * B *
            DH, shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0
            ), block_shape=(siz_B, DH), order=(0, 1))
        matN_t = tl.load(matN_t_ptr)
        matM_t_ptr = tl.make_block_ptr(base=states_all + idx_b_NH *
            str_matStatesAll_NH + (idx_t + 1) * str_matStatesAll_T + 3 * B *
            DH, shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0
            ), block_shape=(siz_B, DH), order=(0, 1))
        matM_t = tl.load(matM_t_ptr).to(tl.float32)
        matH_tminus1_ptr = tl.make_block_ptr(base=states_all + idx_b_NH *
            str_matStatesAll_NH + idx_t * str_matStatesAll_T + 0 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matH_tminus1 = tl.load(matH_tminus1_ptr)
        matC_tminus1_ptr = tl.make_block_ptr(base=states_all + idx_b_NH *
            str_matStatesAll_NH + idx_t * str_matStatesAll_T + 1 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matC_tminus1 = tl.load(matC_tminus1_ptr)
        matN_tminus1_ptr = tl.make_block_ptr(base=states_all + idx_b_NH *
            str_matStatesAll_NH + idx_t * str_matStatesAll_T + 2 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matN_tminus1 = tl.load(matN_tminus1_ptr)
        matM_tminus1_ptr = tl.make_block_ptr(base=states_all + idx_b_NH *
            str_matStatesAll_NH + idx_t * str_matStatesAll_T + 3 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        matM_tminus1 = tl.load(matM_tminus1_ptr).to(tl.float32)
        matDeltaHtrans_out_t_ptr = tl.make_block_ptr(base=
            delta_states_all_outside + idx_b_NH *
            str_delta_states_all_outside_NH + idx_t *
            str_delta_states_all_outside_T + 0 * B * DH, shape=(B, DH),
            strides=(DH, 1), offsets=(idx_b_B * siz_B, 0), block_shape=(
            siz_B, DH), order=(0, 1))
        matDeltaHtrans_out_t = tl.load(matDeltaHtrans_out_t_ptr)
        matDeltaCtrans_out_t_ptr = tl.make_block_ptr(base=
            delta_states_all_outside + idx_b_NH *
            str_delta_states_all_outside_NH + idx_t *
            str_delta_states_all_outside_T + 1 * B * DH, shape=(B, DH),
            strides=(DH, 1), offsets=(idx_b_B * siz_B, 0), block_shape=(
            siz_B, DH), order=(0, 1))
        matDeltaCtrans_out_t = tl.load(matDeltaCtrans_out_t_ptr)
        matDeltaNtrans_out_t_ptr = tl.make_block_ptr(base=
            delta_states_all_outside + idx_b_NH *
            str_delta_states_all_outside_NH + idx_t *
            str_delta_states_all_outside_T + 1 * B * DH, shape=(B, DH),
            strides=(DH, 1), offsets=(idx_b_B * siz_B, 0), block_shape=(
            siz_B, DH), order=(0, 1))
        matDeltaNtrans_out_t = tl.load(matDeltaNtrans_out_t_ptr)
        matDeltaH_t = matDeltaHtrans_out_t + matDeltaH_tplus1
        matDeltaC_t = matDeltaCtrans_out_t + matDeltaC_tplus1
        matDeltaN_t = matDeltaNtrans_out_t + matDeltaN_tplus1
        matDeltaC_t = matDeltaC_t + matDeltaH_t * (matG_o / matN_t)
        matDeltaN_t = matDeltaN_t - matDeltaH_t * (matG_o * matC_t / (
            matN_t * matN_t))
        matG_i = tl.exp(matG_ibar - matM_t)
        matG_logfplusm = matM_tminus1 + tl.log(tl.sigmoid(matG_fbar))
        matG_f = tl.exp(matG_logfplusm - matM_t)
        matDeltaGI = (matDeltaC_t * matG_z + matDeltaN_t) * matG_i
        matDeltaGF = (matDeltaC_t * matC_tminus1 + matDeltaN_t * matN_tminus1
            ) * matG_f * tl.sigmoid(-matG_fbar)
        matDeltaGZ = matDeltaC_t * matG_i * (1 - matG_z * matG_z)
        matDeltaGO = matDeltaH_t * (matC_t / matN_t) * (1 - matG_o) * matG_o
        matDeltaC_tminus1 = matDeltaC_t * matG_f
        matDeltaN_tminus1 = matDeltaN_t * matG_f
        matDeltaH_tminus1 = tl.dot(matDeltaGI.to(DTYPE), matR_i)
        matDeltaH_tminus1 += tl.dot(matDeltaGF.to(DTYPE), matR_f)
        matDeltaH_tminus1 += tl.dot(matDeltaGZ.to(DTYPE), matR_z)
        matDeltaH_tminus1 += tl.dot(matDeltaGO.to(DTYPE), matR_o)
        matDeltaR_i += tl.dot(tl.trans(matDeltaGI.to(DTYPE)), matH_tminus1)
        matDeltaR_f += tl.dot(tl.trans(matDeltaGF.to(DTYPE)), matH_tminus1)
        matDeltaR_z += tl.dot(tl.trans(matDeltaGZ.to(DTYPE)), matH_tminus1)
        matDeltaR_o += tl.dot(tl.trans(matDeltaGO.to(DTYPE)), matH_tminus1)
        vecDeltaB_i += tl.sum(matDeltaGI, axis=0)
        vecDeltaB_f += tl.sum(matDeltaGF, axis=0)
        vecDeltaB_z += tl.sum(matDeltaGZ, axis=0)
        vecDeltaB_o += tl.sum(matDeltaGO, axis=0)
        matDeltaGI_ptr = tl.make_block_ptr(base=delta_Wx + idx_b_NH *
            str_matDeltaWx_NH + idx_t * str_matDeltaWx_T + 0 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        tl.store(matDeltaGI_ptr, matDeltaGI.to(DTYPE))
        matDeltaGF_ptr = tl.make_block_ptr(base=delta_Wx + idx_b_NH *
            str_matDeltaWx_NH + idx_t * str_matDeltaWx_T + 1 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        tl.store(matDeltaGF_ptr, matDeltaGF.to(DTYPE))
        matDeltaGZ_ptr = tl.make_block_ptr(base=delta_Wx + idx_b_NH *
            str_matDeltaWx_NH + idx_t * str_matDeltaWx_T + 2 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        tl.store(matDeltaGZ_ptr, matDeltaGZ.to(DTYPE))
        matDeltaGO_ptr = tl.make_block_ptr(base=delta_Wx + idx_b_NH *
            str_matDeltaWx_NH + idx_t * str_matDeltaWx_T + 3 * B * DH,
            shape=(B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0),
            block_shape=(siz_B, DH), order=(0, 1))
        tl.store(matDeltaGO_ptr, matDeltaGO.to(DTYPE))
        matDeltaH_tplus1 = matDeltaH_tminus1
        matDeltaC_tplus1 = matDeltaC_tminus1
        matDeltaN_tplus1 = matDeltaN_tminus1
    matDeltaHtrans_initial_ptr = tl.make_block_ptr(base=
        delta_states_initial + idx_b_NH * NS * B * DH + 0 * B * DH, shape=(
        B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0), block_shape=
        (siz_B, DH), order=(0, 1))
    tl.store(matDeltaHtrans_initial_ptr, matDeltaH_tplus1.to(DTYPE))
    matDeltaCtrans_initial_ptr = tl.make_block_ptr(base=
        delta_states_initial + idx_b_NH * NS * B * DH + 1 * B * DH, shape=(
        B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0), block_shape=
        (siz_B, DH), order=(0, 1))
    tl.store(matDeltaCtrans_initial_ptr, matDeltaC_tplus1.to(DTYPE))
    matDeltaNtrans_initial_ptr = tl.make_block_ptr(base=
        delta_states_initial + idx_b_NH * NS * B * DH + 2 * B * DH, shape=(
        B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0), block_shape=
        (siz_B, DH), order=(0, 1))
    tl.store(matDeltaNtrans_initial_ptr, matDeltaN_tplus1.to(DTYPE))
    matDeltaMtrans_initial_ptr = tl.make_block_ptr(base=
        delta_states_initial + idx_b_NH * NS * B * DH + 3 * B * DH, shape=(
        B, DH), strides=(DH, 1), offsets=(idx_b_B * siz_B, 0), block_shape=
        (siz_B, DH), order=(0, 1))
    tl.store(matDeltaMtrans_initial_ptr, tl.zeros((siz_B, DH), dtype=DTYPE))
    matDeltaR_i_ptr = tl.make_block_ptr(base=delta_R + idx_b_B * str_matR_B +
        idx_b_NH * str_matR_NH + 0 * str_matR_NGR, shape=(DH, DH), strides=
        (DH, 1), offsets=(0, 0), block_shape=(DH, DH), order=(0, 1))
    tl.store(matDeltaR_i_ptr, matDeltaR_i.to(DTYPE))
    matDeltaR_f_ptr = tl.make_block_ptr(base=delta_R + idx_b_B * str_matR_B +
        idx_b_NH * str_matR_NH + 1 * str_matR_NGR, shape=(DH, DH), strides=
        (DH, 1), offsets=(0, 0), block_shape=(DH, DH), order=(0, 1))
    tl.store(matDeltaR_f_ptr, matDeltaR_f.to(DTYPE))
    matDeltaR_z_ptr = tl.make_block_ptr(base=delta_R + idx_b_B * str_matR_B +
        idx_b_NH * str_matR_NH + 2 * str_matR_NGR, shape=(DH, DH), strides=
        (DH, 1), offsets=(0, 0), block_shape=(DH, DH), order=(0, 1))
    tl.store(matDeltaR_z_ptr, matDeltaR_z.to(DTYPE))
    matDeltaR_o_ptr = tl.make_block_ptr(base=delta_R + idx_b_B * str_matR_B +
        idx_b_NH * str_matR_NH + 3 * str_matR_NGR, shape=(DH, DH), strides=
        (DH, 1), offsets=(0, 0), block_shape=(DH, DH), order=(0, 1))
    tl.store(matDeltaR_o_ptr, matDeltaR_o.to(DTYPE))
    vecDeltaB_i_ptr = (delta_b + idx_b_B * NH * NGI * DH + idx_b_NH * NGI *
        DH + 0 * DH + tl.arange(0, DH))
    tl.store(vecDeltaB_i_ptr, vecDeltaB_i.to(DTYPE))
    vecDeltaB_f_ptr = (delta_b + idx_b_B * NH * NGI * DH + idx_b_NH * NGI *
        DH + 1 * DH + tl.arange(0, DH))
    tl.store(vecDeltaB_f_ptr, vecDeltaB_f.to(DTYPE))
    vecDeltaB_z_ptr = (delta_b + idx_b_B * NH * NGI * DH + idx_b_NH * NGI *
        DH + 2 * DH + tl.arange(0, DH))
    tl.store(vecDeltaB_z_ptr, vecDeltaB_z.to(DTYPE))
    vecDeltaB_o_ptr = (delta_b + idx_b_B * NH * NGI * DH + idx_b_NH * NGI *
        DH + 3 * DH + tl.arange(0, DH))
    tl.store(vecDeltaB_o_ptr, vecDeltaB_o.to(DTYPE))
