@triton.jit
def simulate_kernel(output, n_steps, seed, p, start, block_size: tl.constexpr):
    n_program = tl.num_programs(axis=0)
    pid = tl.program_id(axis=0)
    block_start = pid * block_size
    offsets = block_start + tl.arange(0, block_size)
    state = tl.full([block_size], start, dtype=tl.uint32)
    for _ in range(n_steps):
        this_seed = tl.randint(seed, pid)
        rand = tl.rand(this_seed, offsets)
        state = tl.where(state == 0, state, tl.where(rand < p, state - 1, 
            state + 1))
        pid += n_program
    fall_off = state == 0
    prob = tl.sum(fall_off.to(tl.int64))
    tl.store(output + tl.program_id(0), prob)
