@triton.jit
def bit_split_kernel(x_ptr, output_hi_ptr, output_lo_ptr, n_elements,
    BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(x_ptr + offsets, mask=mask).to(tl.uint32, bitcast=True)
    output_hi = (x >> 16).to(tl.uint16).to(tl.bfloat16, bitcast=True)
    output_lo = x.to(tl.uint16).to(tl.bfloat16, bitcast=True)
    tl.store(output_hi_ptr + offsets, output_hi, mask=mask)
    tl.store(output_lo_ptr + offsets, output_lo, mask=mask)
