open3d::ml::impl::sycl_voxelize_detail Namespace Reference

Functions
template<class T >
T	ReadScalar (sycl::queue &queue, const T *device_ptr, const std::vector< sycl::event > &deps={})
sycl::event	ComputeIndicesBatchesSYCL (sycl::queue &queue, int64_t *indices_batches, const int64_t *row_splits, int64_t batch_size)
template<class T , int NDIM>
sycl::event	ComputeHashSYCL (sycl::queue &queue, int64_t *hashes, int64_t num_points, const T *const points, const int64_t *const indices_batches, const MiniVec< T, NDIM > points_range_min_vec, const MiniVec< T, NDIM > points_range_max_vec, const MiniVec< T, NDIM > inv_voxel_size, const MiniVec< int64_t, NDIM > strides, int64_t batch_hash, int64_t invalid_hash, const std::vector< sycl::event > &deps={})
sycl::event	LimitCountsSYCL (sycl::queue &queue, int64_t *counts, int64_t num, int64_t limit, const std::vector< sycl::event > &deps={})
int64_t	RunLengthEncodeSYCL (sycl::queue &queue, const int64_t *const keys, int64_t num_keys, int64_t *unique_keys_out, int64_t *unique_counts_out, const std::vector< sycl::event > &deps={})
void	ComputeBatchIdSYCL (sycl::queue &queue, int64_t *hashes, int64_t num_voxels, int64_t batch_hash)
void	ComputeVoxelPerBatchSYCL (sycl::queue &queue, int64_t *num_voxels_per_batch, const int64_t *unique_batches_count, const int64_t *unique_batches, int64_t num_batches)
sycl::event	ComputeStartIdxSYCL (sycl::queue &queue, int64_t *start_idx, int64_t *points_count, const int64_t *num_voxels_prefix_sum, const int64_t *unique_hashes_count_prefix_sum, const int64_t *out_batch_splits, int64_t batch_size, int64_t max_points_per_voxel, const std::vector< sycl::event > &deps={})
template<class T , int NDIM>
void	ComputeVoxelCoordsSYCL (sycl::queue &queue, int32_t *voxel_coords, const T *const points, const int64_t *const point_indices, const int64_t *const prefix_sum, const MiniVec< T, NDIM > points_range_min_vec, const MiniVec< T, NDIM > inv_voxel_size, int64_t num_voxels)
void	CopyPointIndicesSYCL (sycl::queue &queue, int64_t *out, const int64_t *const point_indices, const int64_t *const prefix_sum_in, const int64_t *const prefix_sum_out, int64_t num_voxels)

Function Documentation

ComputeBatchIdSYCL()

void open3d::ml::impl::sycl_voxelize_detail::ComputeBatchIdSYCL ( sycl::queue &  queue, int64_t *  hashes, int64_t  num_voxels, int64_t  batch_hash  )

inline

hashes[i] /= batch_hash, i.e., converts a voxel hash to a batch id. Ports ComputeBatchIdKernel.

ComputeHashSYCL()

template<class T , int NDIM>

sycl::event open3d::ml::impl::sycl_voxelize_detail::ComputeHashSYCL ( sycl::queue &  queue, int64_t *  hashes, int64_t  num_points, const T *const  points, const int64_t *const  indices_batches, const MiniVec< T, NDIM >  points_range_min_vec, const MiniVec< T, NDIM >  points_range_max_vec, const MiniVec< T, NDIM >  inv_voxel_size, const MiniVec< int64_t, NDIM >  strides, int64_t  batch_hash, int64_t  invalid_hash, const std::vector< sycl::event > &  deps = {}  )

inline

Computes the hash (linear voxel index, offset by batch) for each point. Points outside [points_range_min, points_range_max] get invalid_hash. Ports ComputeHashKernel. Depends on deps and returns its completion event (non-blocking).

ComputeIndicesBatchesSYCL()

sycl::event open3d::ml::impl::sycl_voxelize_detail::ComputeIndicesBatchesSYCL ( sycl::queue &  queue, int64_t *  indices_batches, const int64_t *  row_splits, int64_t  batch_size  )

inline

Assigns each point its batch id (index into row_splits) by looping over batches on the device. Ports ComputeIndicesBatchesKernel.

ComputeStartIdxSYCL()

sycl::event open3d::ml::impl::sycl_voxelize_detail::ComputeStartIdxSYCL ( sycl::queue &  queue, int64_t *  start_idx, int64_t *  points_count, const int64_t *  num_voxels_prefix_sum, const int64_t *  unique_hashes_count_prefix_sum, const int64_t *  out_batch_splits, int64_t  batch_size, int64_t  max_points_per_voxel, const std::vector< sycl::event > &  deps = {}  )

inline

Computes the starting index and clamped point count for each valid voxel, used when the number of voxels exceeds max_voxels. Ports ComputeStartIdxKernel. Depends on deps and returns its completion event (non-blocking).

ComputeVoxelCoordsSYCL()

template<class T , int NDIM>

void open3d::ml::impl::sycl_voxelize_detail::ComputeVoxelCoordsSYCL ( sycl::queue &  queue, int32_t *  voxel_coords, const T *const  points, const int64_t *const  point_indices, const int64_t *const  prefix_sum, const MiniVec< T, NDIM >  points_range_min_vec, const MiniVec< T, NDIM >  inv_voxel_size, int64_t  num_voxels  )

inline

Computes integer voxel coordinates for each valid voxel from the position of its first (sorted) point. Ports ComputeVoxelCoordsKernel.

ComputeVoxelPerBatchSYCL()

void open3d::ml::impl::sycl_voxelize_detail::ComputeVoxelPerBatchSYCL ( sycl::queue &  queue, int64_t *  num_voxels_per_batch, const int64_t *  unique_batches_count, const int64_t *  unique_batches, int64_t  num_batches  )

inline

Scatters unique_batches_count into num_voxels_per_batch, indexed by unique_batches (distinct destination indices, so no atomics needed). Ports ComputeVoxelPerBatchKernel.

CopyPointIndicesSYCL()

void open3d::ml::impl::sycl_voxelize_detail::CopyPointIndicesSYCL ( sycl::queue &  queue, int64_t *  out, const int64_t *const  point_indices, const int64_t *const  prefix_sum_in, const int64_t *const  prefix_sum_out, int64_t  num_voxels  )

inline

Copies (limited-count) point indices for each valid voxel into the flat output array. Ports CopyPointIndicesKernel.

LimitCountsSYCL()

sycl::event open3d::ml::impl::sycl_voxelize_detail::LimitCountsSYCL ( sycl::queue &  queue, int64_t *  counts, int64_t  num, int64_t  limit, const std::vector< sycl::event > &  deps = {}  )

inline

Element-wise min(counts[i], limit). Ports LimitCountsKernel. Depends on deps and returns its completion event (non-blocking).

ReadScalar()

template<class T >

T open3d::ml::impl::sycl_voxelize_detail::ReadScalar ( sycl::queue &  queue, const T *  device_ptr, const std::vector< sycl::event > &  deps = {}  )

inline

Reads a single device scalar back to the host. Genuine host-sync point (the caller needs the value to make a branch/sizing decision), not a lazy default – so it blocks-waits after depending on deps.

RunLengthEncodeSYCL()

int64_t open3d::ml::impl::sycl_voxelize_detail::RunLengthEncodeSYCL ( sycl::queue &  queue, const int64_t *const  keys, int64_t  num_keys, int64_t *  unique_keys_out, int64_t *  unique_counts_out, const std::vector< sycl::event > &  deps = {}  )

inline

Run-length-encodes a sorted array of keys via oneapi::dpl::reduce_by_key (plan §6.4): reducing a constant-1 "value" sequence, grouped by equal consecutive keys, yields the unique keys plus their run lengths in a single device-parallel call. Replaces cub::DeviceRunLengthEncode::Encode (see file-level comment). reduce_by_key has no non-blocking oneDPL async equivalent, so this function is itself a genuine synchronization point (it blocks before returning); the returned run count is needed immediately by the caller anyway (to size subsequent allocations).

Parameters

unique_keys_out	Output buffer of size >= num_keys (upper bound on the number of runs).
unique_counts_out	Output buffer of size >= num_keys.

Returns

The number of runs (unique consecutive keys).