@triton.jit
def triton_sum_dim0_in_fp32_kernel(xs_ptr, out_ptr, stride_a, a, b,
    BLOCK_SIZE_A: tl.constexpr, BLOCK_SIZE_B: tl.constexpr):
    pid = tl.program_id(0)
    offsets_b = tl.arange(0, BLOCK_SIZE_B) + pid * BLOCK_SIZE_B
    all_out = tl.zeros((BLOCK_SIZE_B,), dtype=tl.float32)
    for i in range(0, a, BLOCK_SIZE_A):
        offsets_a = tl.arange(0, BLOCK_SIZE_A) + i
        xs = tl.load(xs_ptr + offsets_a[:, None] * stride_a + offsets_b[
            None, :], mask=(offsets_a < a)[:, None] & (offsets_b < b)[None,
            :], other=0)
        xs = xs.to(tl.float32)
        out = tl.sum(xs, axis=0)
        all_out += out
    tl.store(out_ptr + offsets_b, all_out, mask=offsets_b < b)
