@triton.jit
def code_downscale_kernel(code_ptr, out_ptr, n_steps, N, BLK: tl.constexpr):
    pid = tl.program_id(0)
    offs_n = BLK * pid + tl.arange(0, BLK)
    mask_n = offs_n < N
    code = tl.load(code_ptr + offs_n, mask=mask_n)
    top16bit = code & 32767 << 48
    low16bit = code & (1 << 48) - 1
    low16bit >>= n_steps * 3
    new_code = top16bit | low16bit
    tl.store(out_ptr + offs_n, new_code, mask=mask_n)
