@triton.jit
def bwd_sequential_scan(grad_output, v, f, h, B, L, C, BLOCK_M: tl.constexpr):
    offset_b = tl.program_id(0)
    if offset_b >= B:
        return
    offset_n = tl.program_id(1)
    ptr = tl.arange(0, BLOCK_M) + offset_b * L * C + (L - 1
        ) * C + offset_n * BLOCK_M
    grad_h = tl.zeros([BLOCK_M], dtype=tl.float32)
    for time_step in range(L - 1, -1, -1):
        grad = tl.load(grad_output + ptr).to(tl.float32)
        grad_h += grad
        decay = tl.load(f + ptr).to(tl.float32)
        input = tl.load(v + ptr).to(tl.float32)
        grad_v = (1 - decay) * grad_h
        tl.store(v + ptr, grad_v.to(v.dtype.element_ty))
        hidden_state = tl.load(h + ptr - C, mask=ptr >= offset_b * L * C +
            C, other=0.0).to(tl.float32)
        grad_f = grad_h * (hidden_state - input)
        tl.store(f + ptr, grad_f.to(f.dtype.element_ty))
        grad_h *= decay
        ptr -= C
