@triton.jit
def _sequential_rnn_scan_fwd_kernel(x_ptr, a_ptr, h0_ptr, out_ptr,
    stride_x_batch, stride_x_len, stride_x_dim, stride_a_batch,
    stride_a_len, stride_a_dim, stride_h0_batch, stride_h0_dim,
    stride_out_batch, stride_out_dim, seq_len: tl.constexpr, BLOCK_SIZE: tl
    .constexpr):
    pid_batch = tl.program_id(0)
    pid_dim = tl.program_id(1)
    x_ptr += pid_batch * stride_x_batch
    a_ptr += pid_batch * stride_a_batch
    ht_ptr = h0_ptr + pid_batch * stride_h0_batch
    out_ptr += pid_batch * stride_out_batch
    offsets = pid_dim * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    x_ptrs = x_ptr + offsets * stride_x_dim
    a_ptrs = a_ptr + offsets * stride_a_dim
    ht_ptrs = ht_ptr + offsets * stride_h0_dim
    out_ptrs = out_ptr + offsets * stride_out_dim
    h_t = tl.load(ht_ptrs).to(tl.float32)
    for t in range(seq_len):
        x_t = tl.load(x_ptrs).to(tl.float32)
        a_t = tl.load(a_ptrs).to(tl.float32)
        h_t = a_t * h_t + x_t
        if t < seq_len - 1:
            x_ptrs += stride_x_len
            a_ptrs += stride_a_len
    tl.store(out_ptrs, h_t.to(out_ptr.dtype.element_ty))
