@triton.jit
def _encode_hilbert_unpadded_kernel(xyz_ptr, batch_idx_ptr, code_ptr,
    space_size, x_offset, y_offset, z_offset, str_xyz_n, str_xyz_c, N, BLK:
    tl.constexpr, ASSIGN_BATCH_INDEX: tl.constexpr):
    pid = tl.program_id(0)
    offs_n = pid * BLK + tl.arange(0, BLK)
    mask = offs_n < N
    xyz_ptrs = xyz_ptr + offs_n * str_xyz_n
    fx = tl.load(xyz_ptrs + x_offset * str_xyz_c, mask=mask)
    fy = tl.load(xyz_ptrs + y_offset * str_xyz_c, mask=mask)
    fz = tl.load(xyz_ptrs + z_offset * str_xyz_c, mask=mask)
    ret = _calculate_hilbert_distance(fx, fy, fz, space_size)
    if ASSIGN_BATCH_INDEX:
        batch_idx_ptrs = batch_idx_ptr + offs_n
        batch_idx = tl.load(batch_idx_ptrs, mask=mask).to(tl.int64)
        ret |= batch_idx << 48
    code_ptrs = code_ptr + offs_n
    tl.store(code_ptrs, ret, mask=mask)
