Namespaces
namespace	impl

Data Structures
class	FixedRadiusIndex
	FixedRadiusIndex for nearest neighbor range search. More...

class	KnnDirectKernel
	Named kernel tag for KnnDirect (SYCL kernel naming). More...

class	KnnIndex

class	MemoryAllocation
	A class for managing memory segments within a memory allocation. More...

class	NanoFlannIndex

struct	NanoFlannIndexHolder
	NanoFlann Index Holder. More...

struct	NanoFlannIndexHolderBase
	Base struct for NanoFlann index holder. More...

class	NearestNeighborSearch
	Facade for batched KNN, fixed-radius, hybrid, and multi-radius search. More...

class	NeighborSearchAllocator

class	NNSIndex

Enumerations
enum	Metric { L1 , L2 , Linf }
	Supported metrics. More...

Functions
template<class T >
void	BuildSpatialHashTableCPU (const Tensor &points, double radius, const Tensor &points_row_splits, const Tensor &hash_table_splits, Tensor &hash_table_index, Tensor &hash_table_cell_splits)

template<class T , class TIndex >
void	FixedRadiusSearchCPU (const Tensor &points, const Tensor &queries, double radius, const Tensor &points_row_splits, const Tensor &queries_row_splits, const Tensor &hash_table_splits, const Tensor &hash_table_index, const Tensor &hash_table_cell_splits, const Metric metric, const bool ignore_query_point, const bool return_distances, const bool sort, Tensor &neighbors_index, Tensor &neighbors_row_splits, Tensor &neighbors_distance)

template<class T , class TIndex >
void	HybridSearchCPU (const Tensor &points, const Tensor &queries, double radius, int max_knn, const Tensor &points_row_splits, const Tensor &queries_row_splits, const Tensor &hash_table_splits, const Tensor &hash_table_index, const Tensor &hash_table_cell_splits, const Metric metric, Tensor &neighbors_index, Tensor &neighbors_count, Tensor &neighbors_distance)

template<class T >
sycl::event	BuildSpatialHashTableSYCLRaw (sycl::queue &queue, const T points_ptr, T inv_voxel_size, int batch_size, const int64_t host_points_row_splits, const uint32_t host_hash_table_splits, uint32_t cell_splits_ptr, size_t cell_splits_size, uint32_t *index_ptr)

template<class T >
void	BuildSpatialHashTableSYCL (const Tensor &points, double radius, const Tensor &points_row_splits, const Tensor &hash_table_splits, Tensor &hash_table_index, Tensor &hash_table_cell_splits)

template<class T >
void	CountNeighborsSYCL (sycl::queue &queue, uint32_t neighbors_count_ptr, const uint32_t const point_index_table, const uint32_t const hash_table_cell_splits, uint32_t hash_table_size, const T const query_points, int64_t num_queries, const T *const points, T inv_voxel_size, T radius, Metric metric, bool ignore_query_point, T threshold)

template<class T , class TIndex >
void	WriteNeighborsSYCL (sycl::queue &queue, TIndex indices, T distances, const int64_t const neighbors_row_splits, const uint32_t const point_index_table, const uint32_t const hash_table_cell_splits, uint32_t hash_table_size, const T const query_points, int64_t num_queries, const T *const points, T inv_voxel_size, T radius, Metric metric, bool ignore_query_point, T threshold, bool return_distances)

template<class T , class TIndex >
void	WriteNeighborsHybridSYCL (sycl::queue &queue, TIndex indices, T distances, TIndex counts, const uint32_t const point_index_table, const uint32_t const hash_table_cell_splits, uint32_t hash_table_size, const T const query_points, int64_t num_queries, const T *const points, T inv_voxel_size, T radius, Metric metric, T threshold, int max_knn)

template<class T , class TIndex >
sycl::event	SortNeighborsByDistanceSYCL (const Device &device, TIndex indices_ptr, T distances_ptr, const int64_t *row_splits_ptr, int64_t num_queries, int64_t num_indices)

void	ChooseTileSize (int64_t num_queries, int64_t num_points, int64_t element_size, int64_t tile_bytes, int64_t &tile_queries, int64_t &tile_points, int64_t max_tile_queries=128, int64_t tile_points_alignment=128)

template<typename T , typename TIndex , int K>
void	HeapifyDown (T d, TIndex idx, int root)

template<typename T , typename TIndex , int K>
void	HeapSort (T d, TIndex idx)
	Heap-sort a compile-time max-heap of size K into ascending order.

template<typename T , typename TIndex , int K>
void	UpdateTopKFromTile (sycl::queue &queue, const T neg2qp_ptr, int64_t distance_stride, const T point_norms_ptr, int64_t num_queries, int64_t num_points, TIndex point_offset, T best_dist_ptr, TIndex best_idx_ptr, bool use_threshold, T threshold)

template<typename T , typename TIndex , int K>
void	FinalizeTopK (sycl::queue &queue, int64_t num_queries, const T running_dist_ptr, const TIndex running_idx_ptr, T out_dist_ptr, TIndex out_idx_ptr, int64_t actual_k, const T *query_norms_ptr)

int64_t	KBucket (int64_t k)
	Return the smallest dispatch-bucket value ≥ k.

template<typename T , typename TIndex >
void	DispatchUpdateTopKFromTile (sycl::queue &queue, const T neg2qp_ptr, int64_t distance_stride, const T point_norms_ptr, int64_t num_queries, int64_t num_points, int64_t k_bucket, TIndex point_offset, T best_dist_ptr, TIndex best_idx_ptr, bool use_threshold, T threshold)
	Instantiate UpdateTopKFromTile for the given `k_bucket`.

template<typename T , typename TIndex >
void	DispatchFinalizeTopK (sycl::queue &queue, int64_t num_queries, const T running_dist_ptr, const TIndex running_idx_ptr, T out_dist_ptr, TIndex out_idx_ptr, int64_t actual_k, int64_t k_bucket, const T *query_norms_ptr)
	Instantiate FinalizeTopK for the given `k_bucket`.

template<typename T , typename TIndex , int NDIM, int K, int SG>
void	KnnDirect (sycl::queue &queue, const T points_ptr, const T queries_ptr, int64_t num_points, int64_t num_queries, int64_t actual_k, T out_dist_ptr, TIndex out_idx_ptr, int64_t subgroups_per_wg, int64_t tile_points)
	Launch direct-distance KNN for fixed compile-time NDIM, K, and SG.

template<typename T , typename TIndex , int NDIM, int SG>
void	DispatchKnnDirectKForSG (sycl::queue &queue, const T points_ptr, const T queries_ptr, int64_t num_points, int64_t num_queries, int64_t actual_k, T out_dist_ptr, TIndex out_idx_ptr, int64_t subgroups_per_wg, int64_t tile_points)

template<typename T , typename TIndex , int NDIM>
void	DispatchKnnDirectK (sycl::queue &queue, const T points_ptr, const T queries_ptr, int64_t num_points, int64_t num_queries, int64_t actual_k, T out_dist_ptr, TIndex out_idx_ptr, int64_t subgroups_per_wg, int64_t tile_points)

template<typename T , typename TIndex >
void	DispatchKnnDirect (sycl::queue &queue, const T points_ptr, const T queries_ptr, int64_t dim, int64_t num_points, int64_t num_queries, int64_t actual_k, T out_dist_ptr, TIndex out_idx_ptr, int64_t subgroups_per_wg=kKnnDirectSubgroupsPerWG, int64_t tile_points=kKnnDirectTilePoints)

bool	UseKnnDirect (int64_t dim, int64_t knn)
	True if (dim, knn) qualifies for the direct-distance SYCL KNN path.

template<typename T , typename TIndex >
void	DispatchSelectTopKQueries (sycl::queue &queue, const T distances_ptr, int64_t distance_query_stride, int64_t num_queries, int64_t num_points, int64_t knn, int64_t k_bucket, TIndex index_offset, TIndex out_indices_ptr, T out_distances_ptr, int64_t out_query_stride, bool use_threshold, const T query_norms_ptr, T radius_sq, T scalar_threshold)

template<typename T , typename TIndex >
void	SelectTopKQueries (const Device &device, const T distances_ptr, int64_t distance_query_stride, int64_t num_queries, int64_t num_points, int64_t knn, TIndex index_offset, TIndex scratch_indices_ptr, int64_t scratch_query_stride, TIndex out_indices_ptr, T out_distances_ptr, int64_t out_query_stride, bool use_threshold=false, const T *query_norms_ptr=nullptr, T radius_sq=T(0), T scalar_threshold=T(0))

template<typename T , typename TIndex >
void	MergeTopKQueries (const Device &device, const T curr_dist_ptr, const TIndex curr_idx_ptr, int64_t curr_stride, const T cand_dist_ptr, const TIndex cand_idx_ptr, int64_t cand_stride, int64_t num_queries, int64_t knn, TIndex scratch_ptr, int64_t scratch_stride, TIndex out_idx_ptr, T *out_dist_ptr, int64_t out_stride)

template<typename T , typename TIndex >
void	AddQueryNormsToDistances (const Device &device, int64_t num_queries, int64_t knn, const TIndex indices_ptr, T distances_ptr, const T *query_norms_ptr)

template<class T , class TIndex >
void	KnnSearchSYCL (const Tensor &points, const Tensor &points_row_splits, const Tensor &queries, const Tensor &queries_row_splits, int knn, Tensor &neighbors_index, Tensor &neighbors_row_splits, Tensor &neighbors_distance, int64_t tile_bytes, int64_t max_tile_queries, int64_t tile_points_alignment, bool force_addmm_path)

template<class T , class TIndex >
void	FixedRadiusSearchSYCL (const Tensor &points, const Tensor &queries, double radius, const Tensor &points_row_splits, const Tensor &queries_row_splits, const Tensor &hash_table_splits, const Tensor &hash_table_index, const Tensor &hash_table_cell_splits, const Metric metric, const bool ignore_query_point, const bool return_distances, const bool sort, Tensor &neighbors_index, Tensor &neighbors_row_splits, Tensor &neighbors_distance, int64_t)

template<class T , class TIndex >
void	HybridSearchSYCL (const Tensor &points, const Tensor &queries, double radius, int max_knn, const Tensor &points_row_splits, const Tensor &queries_row_splits, const Tensor &hash_table_splits, const Tensor &hash_table_index, const Tensor &hash_table_cell_splits, const Metric metric, Tensor &neighbors_index, Tensor &neighbors_count, Tensor &neighbors_distance, int64_t)

template void	BuildSpatialHashTableSYCL< float > (const Tensor &points, double radius, const Tensor &points_row_splits, const Tensor &hash_table_splits, Tensor &hash_table_index, Tensor &hash_table_cell_splits)

template void	BuildSpatialHashTableSYCL< double > (const Tensor &points, double radius, const Tensor &points_row_splits, const Tensor &hash_table_splits, Tensor &hash_table_index, Tensor &hash_table_cell_splits)

HOST_DEVICE size_t	SpatialHash (int x, int y, int z)
	Spatial hashing function for integer coordinates.

HOST_DEVICE size_t	SpatialHash (const utility::MiniVec< int, 3 > &xyz)

template<class TVecf >
HOST_DEVICE utility::MiniVec< int, 3 >	ComputeVoxelIndex (const TVecf &pos, const typename TVecf::Scalar_t &inv_voxel_size)

Variables
constexpr int64_t	kKnnDirectSubgroupSize = 16
	Default sub-group width for the direct KNN kernel (float path).

constexpr int64_t	kKnnDirectSubgroupsPerWG = 32
	Default sub-groups per work-group (512 work-items at SG=16).

constexpr int64_t	kKnnDirectTilePoints = 2048
	Default point tile size for SLM staging.

constexpr int64_t	kKnnDirectMaxDim = 8
	Maximum point dimension compiled for DispatchKnnDirect.

constexpr int64_t	kSYCLKnnDefaultTileBytes = 8LL * 1024 * 1024
	SYCL NNS defaults for KnnIndex and FixedRadiusIndex constructors.

constexpr int64_t	kSYCLKnnSmallKMax = 32
	Upper bound of k for the GRF-register heap path (eliminates scratch spill).

constexpr int64_t	kSYCLKnnMidKMax = 512

Enumeration Type Documentation

◆ Metric

enum open3d::core::nns::Metric

Supported metrics.

Enumerator
L1
L2
Linf

Function Documentation

◆ BuildSpatialHashTableCPU()

template<class T >

void open3d::core::nns::BuildSpatialHashTableCPU	(	const Tensor &	points,
		double	radius,
		const Tensor &	points_row_splits,
		const Tensor &	hash_table_splits,
		Tensor &	hash_table_index,
		Tensor &	hash_table_cell_splits
	)

Builds a spatial hash table for a fixed radius search of 3D points.

Template Parameters

T	Floating-point data type for the point positions.

Parameters

points	The tensor of 3D points. This tensor may be splitted into multiple batch items by defining points_row_splits_size accordingly.
radius	The radius that will be used for searching.
points_row_splits	Defines the start and end of the points in each batch item. The size of the tensor is batch_size+1. If there is only 1 batch item then this array is [0, num_points]
hash_table_splits	Tensor defining the start and end the hash table for each batch item. This is [0, number of cells] if there is only 1 batch item or [0, hash_table_cell_splits_size-1] which is the same.
hash_table_index	This is an output tensor storing the values of the hash table, which are the indices to the points. The size of the tensor must be equal to the number of points.
hash_table_cell_splits	This is an output tensor storing the start of each hash table entry. The size of this array defines the size of the hash table. The hash table size is hash_table_cell_splits_size - 1.

◆ BuildSpatialHashTableSYCL()

template<class T >

void open3d::core::nns::BuildSpatialHashTableSYCL	(	const Tensor &	points,
		double	radius,
		const Tensor &	points_row_splits,
		const Tensor &	hash_table_splits,
		Tensor &	hash_table_index,
		Tensor &	hash_table_cell_splits
	)

Builds the uniform-grid spatial hash table. points_row_splits and hash_table_splits are host (CPU) tensors; hash_table_index and hash_table_cell_splits are device output tensors already sized by FixedRadiusIndex::SetTensorData. Delegates to BuildSpatialHashTableSYCLRaw.

◆ BuildSpatialHashTableSYCL< double >()

template void open3d::core::nns::BuildSpatialHashTableSYCL< double >	(	const Tensor &	points,
		double	radius,
		const Tensor &	points_row_splits,
		const Tensor &	hash_table_splits,
		Tensor &	hash_table_index,
		Tensor &	hash_table_cell_splits
	)

◆ BuildSpatialHashTableSYCL< float >()

template void open3d::core::nns::BuildSpatialHashTableSYCL< float >	(	const Tensor &	points,
		double	radius,
		const Tensor &	points_row_splits,
		const Tensor &	hash_table_splits,
		Tensor &	hash_table_index,
		Tensor &	hash_table_cell_splits
	)

◆ BuildSpatialHashTableSYCLRaw()

template<class T >

sycl::event open3d::core::nns::BuildSpatialHashTableSYCLRaw	(	sycl::queue &	queue,
		const T *	points_ptr,
		T	inv_voxel_size,
		int	batch_size,
		const int64_t *	host_points_row_splits,
		const uint32_t *	host_hash_table_splits,
		uint32_t *	cell_splits_ptr,
		size_t	cell_splits_size,
		uint32_t *	index_ptr
	)

Builds a uniform spatial-hash grid ("cell list") for a fixed-radius search: count points per cell -> device inclusive scan -> scatter point indices into their cell's slot range. Mirrors BuildSpatialHashTableCUDA.

Raw-pointer variant: takes a SYCL queue directly plus host-accessible batch arrays, so both the Open3D Tensor API and the PyTorch XPU dispatch can share the same kernel implementation without tensor conversion.

host_points_row_splits and host_hash_table_splits are CPU arrays. cell_splits_ptr and index_ptr are device (USM or XPU) pointers.

Non-blocking: returns the event of the last enqueued command instead of waiting host-side. Note that Pass 3's sycl::buffer scratch (see the comment above it) still forces a host block at its own scope exit, so this function is not fully async yet; the event is still returned for API consistency and to chain Pass 1/2 without redundant host waits.

◆ ComputeVoxelIndex()

template<class TVecf >

HOST_DEVICE utility::MiniVec< int, 3 > open3d::core::nns::ComputeVoxelIndex	(	const TVecf &	pos,
		const typename TVecf::Scalar_t &	inv_voxel_size
	)

inline

Computes an integer voxel index for a 3D position.

Parameters

pos	A 3D position.
inv_voxel_size	The reciprocal of the voxel size

◆ CountNeighborsSYCL()

template<class T >

void open3d::core::nns::CountNeighborsSYCL	(	sycl::queue &	queue,
		uint32_t *	neighbors_count_ptr,
		const uint32_t *const	point_index_table,
		const uint32_t *const	hash_table_cell_splits,
		uint32_t	hash_table_size,
		const T *const	query_points,
		int64_t	num_queries,
		const T *const	points,
		T	inv_voxel_size,
		T	radius,
		Metric	metric,
		bool	ignore_query_point,
		T	threshold
	)

Counts, for every query, how many dataset points lie within radius, using the grid built by BuildSpatialHashTableSYCL. Mirrors CountNeighborsKernel (CUDA).

◆ FixedRadiusSearchCPU()

template<class T , class TIndex >

void open3d::core::nns::FixedRadiusSearchCPU	(	const Tensor &	points,
		const Tensor &	queries,
		double	radius,
		const Tensor &	points_row_splits,
		const Tensor &	queries_row_splits,
		const Tensor &	hash_table_splits,
		const Tensor &	hash_table_index,
		const Tensor &	hash_table_cell_splits,
		const Metric	metric,
		const bool	ignore_query_point,
		const bool	return_distances,
		const bool	sort,
		Tensor &	neighbors_index,
		Tensor &	neighbors_row_splits,
		Tensor &	neighbors_distance
	)

Fixed radius search. This function computes a list of neighbor indices for each query point. The lists are stored linearly and an exclusive prefix sum defines the start and end of list in the array. In addition the function optionally can return the distances for each neighbor in the same format as the indices to the neighbors.

Template Parameters

T	Floating-point data type for the point positions.

Parameters

points	Tensor with the 3D point positions. This must be the tensor that was used for building the spatial hash table.
queries	Tensor with the 3D query positions. This may be the same tensor as `points`.
radius	The search radius.
points_row_splits	Defines the start and end of the points in each batch item. The size of the tensor is batch_size+1. If there is only 1 batch item then this tensor is [0, num_points]
queries_row_splits	Defines the start and end of the queries in each batch item. The size of the tensor is batch_size+1. If there is only 1 batch item then this tensor is [0, num_queries]
hash_table_splits	Tensor defining the start and end the hash table for each batch item. This is [0, number of cells] if there is only 1 batch item or [0, hash_table_cell_splits_size-1] which is the same.
hash_table_index	This is an output of the function BuildSpatialHashTableCPU. This is tensor storing the values of the hash table, which are the indices to the points. The size of the tensor must be equal to the number of points.
hash_table_cell_splits	This is an output of the function BuildSpatialHashTableCPU. The row splits array describing the start and end of each cell.
metric	One of L1, L2, Linf. Defines the distance metric for the search.
ignore_query_point	If true then points with the same position as the query point will be ignored.
return_distances	If true then this function will return the distances for each neighbor to its query point in the same format as the indices. Note that for the L2 metric the squared distances will be returned!!
sort	If true then sort the results in ascending order of distance
neighbors_index	The output tensor that saves the resulting neighbor indices
neighbors_row_splits	Tensor defining the start and end the neighbor indices in each batch item. The size of the tensor is num_query_points + 1
neighbors_distance	The output tensor that saves the resulting neighbor distances.

◆ FixedRadiusSearchSYCL()

template<class T , class TIndex >

void open3d::core::nns::FixedRadiusSearchSYCL	(	const Tensor &	points,
		const Tensor &	queries,
		double	radius,
		const Tensor &	points_row_splits,
		const Tensor &	queries_row_splits,
		const Tensor &	hash_table_splits,
		const Tensor &	hash_table_index,
		const Tensor &	hash_table_cell_splits,
		const Metric	metric,
		const bool	ignore_query_point,
		const bool	return_distances,
		const bool	sort,
		Tensor &	neighbors_index,
		Tensor &	neighbors_row_splits,
		Tensor &	neighbors_distance,
		int64_t
	)

◆ HybridSearchCPU()

template<class T , class TIndex >

void open3d::core::nns::HybridSearchCPU	(	const Tensor &	points,
		const Tensor &	queries,
		double	radius,
		int	max_knn,
		const Tensor &	points_row_splits,
		const Tensor &	queries_row_splits,
		const Tensor &	hash_table_splits,
		const Tensor &	hash_table_index,
		const Tensor &	hash_table_cell_splits,
		const Metric	metric,
		Tensor &	neighbors_index,
		Tensor &	neighbors_count,
		Tensor &	neighbors_distance
	)

Hybrid search. This function computes a list of neighbor indices for each query point. The lists are stored linearly and if there is less neighbors than requested, the output tensor will be assigned with default values, -1 for indices and 0 for distances. In addition the function returns the number of neighbors for each query.

Template Parameters

T	Floating-point data type for the point positions.

Parameters

points	Tensor with the 3D point positions. This must be the tensor that was used for building the spatial hash table.
queries	Tensor with the 3D query positions. This may be the same tensor as `points`.
radius	The search radius.
max_knn	The maximum number of neighbor for each query
points_row_splits	Defines the start and end of the points in each batch item. The size of the tensor is batch_size+1. If there is only 1 batch item then this tensor is [0, num_points]
queries_row_splits	Defines the start and end of the queries in each batch item. The size of the tensor is batch_size+1. If there is only 1 batch item then this tensor is [0, num_queries]
hash_table_splits	Tensor defining the start and end the hash table for each batch item. This is [0, number of cells] if there is only 1 batch item or [0, hash_table_cell_splits_size-1] which is the same.
hash_table_index	This is an output of the function BuildSpatialHashTableCPU. This is tensor storing the values of the hash table, which are the indices to the points. The size of the tensor must be equal to the number of points.
hash_table_cell_splits	This is an output of the function BuildSpatialHashTableCPU. The row splits array describing the start and end of each cell.
metric	One of L1, L2, Linf. Defines the distance metric for the search.
neighbors_index	The output tensor that saves the resulting neighbor indices
neighbors_count	The output tensor that saves the number of neighbors for each query points
neighbors_distance	The output tensor that saves the resulting neighbor distances.

◆ HybridSearchSYCL()

template<class T , class TIndex >

void open3d::core::nns::HybridSearchSYCL	(	const Tensor &	points,
		const Tensor &	queries,
		double	radius,
		int	max_knn,
		const Tensor &	points_row_splits,
		const Tensor &	queries_row_splits,
		const Tensor &	hash_table_splits,
		const Tensor &	hash_table_index,
		const Tensor &	hash_table_cell_splits,
		const Metric	metric,
		Tensor &	neighbors_index,
		Tensor &	neighbors_count,
		Tensor &	neighbors_distance,
		int64_t
	)

◆ KnnSearchSYCL()

template<class T , class TIndex >

void open3d::core::nns::KnnSearchSYCL	(	const Tensor &	points,
		const Tensor &	points_row_splits,
		const Tensor &	queries,
		const Tensor &	queries_row_splits,
		int	knn,
		Tensor &	neighbors_index,
		Tensor &	neighbors_row_splits,
		Tensor &	neighbors_distance,
		int64_t	tile_bytes,
		int64_t	max_tile_queries,
		int64_t	tile_points_alignment,
		bool	force_addmm_path
	)

◆ SortNeighborsByDistanceSYCL()

template<class T , class TIndex >

sycl::event open3d::core::nns::SortNeighborsByDistanceSYCL	(	const Device &	device,
		TIndex *	indices_ptr,
		T *	distances_ptr,
		const int64_t *	row_splits_ptr,
		int64_t	num_queries,
		int64_t	num_indices
	)

Sorts each query's variable-length neighbor segment by ascending distance, entirely on device (no host round trip). Mirrors cub::DeviceSegmentedRadixSort::SortPairs (CUDA): like CUDA, ties are not secondarily ordered by index.

float uses a scalar uint64 radix key (query_id << 32) | bit_cast<uint32>(dist) so oneDPL's sort_by_key stays on the fast radix path (valid because distances are clamped >= 0, so their float32 bit patterns are monotonic as unsigned integers, and num_queries < 2^32 always holds here). double cannot use this trick: a monotonic transform of a double needs all 64 bits, leaving no room to also pack the segment id, so it falls back to a struct key + device comparator (oneDPL merge sort, still fully on device).

Non-blocking: returns the event of the last enqueued command. query_id is filled via oneapi::dpl::upper_bound (a single coalesced pass over row_splits_ptr, which is tiny and normally L1/L2-resident) instead of one work-item per query looping over its whole variable-length segment (worst-case divergence + fully uncoalesced writes). sort_by_key and upper_bound are synchronous oneDPL calls (no _async variant exists for either), so each blocks the host until its device work completes; they still need to run, in order, after the async Pass-1 kernels below, which is why an explicit barrier precedes each such call (item 20 rule, see SYCLUtils.h) rather than relying on submission order on this possibly out-of-order (PyTorch) queue. Only double's sycl::malloc_device scratch forces a genuine host wait (Phase 2 rule: raw USM feeding an external free must be provably complete first).

◆ SpatialHash() [1/2]

HOST_DEVICE size_t open3d::core::nns::SpatialHash ( const utility::MiniVec< int, 3 > & xyz )

inline

◆ SpatialHash() [2/2]

HOST_DEVICE size_t open3d::core::nns::SpatialHash	(	int	x,
		int	y,
		int	z
	)

inline

Spatial hashing function for integer coordinates.

◆ WriteNeighborsHybridSYCL()

template<class T , class TIndex >

void open3d::core::nns::WriteNeighborsHybridSYCL	(	sycl::queue &	queue,
		TIndex *	indices,
		T *	distances,
		TIndex *	counts,
		const uint32_t *const	point_index_table,
		const uint32_t *const	hash_table_cell_splits,
		uint32_t	hash_table_size,
		const T *const	query_points,
		int64_t	num_queries,
		const T *const	points,
		T	inv_voxel_size,
		T	radius,
		Metric	metric,
		T	threshold,
		int	max_knn
	)

Single-pass hybrid search: simultaneously counts all points within radius and keeps a running top-max_knn (by ascending distance) per query in fixed-size output slots. Mirrors WriteNeighborsHybridKernel (CUDA), including its per-query bubble sort of the (small, bounded by max_knn) result slice – no device-wide sort is needed here since the output size is already capped. Supports L1, L2, and Linf (via metric and IsNeighbor), matching CUDA's NeighborTest<METRIC>. As with fixed- radius search: for L2, threshold and the returned/compared distances are SQUARED; for L1/Linf they are the metric distance directly (see FixedRadiusThreshold in KnnSearchOpsSYCL.cpp, which the caller uses to compute threshold consistently with this).

◆ WriteNeighborsSYCL()

template<class T , class TIndex >

void open3d::core::nns::WriteNeighborsSYCL	(	sycl::queue &	queue,
		TIndex *	indices,
		T *	distances,
		const int64_t *const	neighbors_row_splits,
		const uint32_t *const	point_index_table,
		const uint32_t *const	hash_table_cell_splits,
		uint32_t	hash_table_size,
		const T *const	query_points,
		int64_t	num_queries,
		const T *const	points,
		T	inv_voxel_size,
		T	radius,
		Metric	metric,
		bool	ignore_query_point,
		T	threshold,
		bool	return_distances
	)

Writes neighbor indices (and optionally distances) for every query into the offsets given by neighbors_row_splits (an exclusive prefix sum over per-query counts). Mirrors WriteNeighborsIndicesAndDistancesKernel (CUDA). Output is unsorted within each query's segment; use SortNeighborsByDistanceSYCL afterward if sort=true was requested.

Variable Documentation

◆ kSYCLKnnDefaultTileBytes

constexpr int64_t open3d::core::nns::kSYCLKnnDefaultTileBytes = 8LL * 1024 * 1024

constexpr

SYCL NNS defaults for KnnIndex and FixedRadiusIndex constructors.

Default distance-tile budget in bytes (8 MiB, tuned for iGPU last-level cache). Discrete GPUs cache is larger and we can increase to 16-32 MiB.

◆ kSYCLKnnMidKMax

constexpr int64_t open3d::core::nns::kSYCLKnnMidKMax = 512

constexpr

Upper bound of k for the proportional scratch-resident heap path. Larger k uses sequential oneDPL partial_sort.

◆ kSYCLKnnSmallKMax

constexpr int64_t open3d::core::nns::kSYCLKnnSmallKMax = 32

constexpr

Upper bound of k for the GRF-register heap path (eliminates scratch spill).

Namespaces

Data Structures

Enumerations

Functions

Variables

Enumeration Type Documentation

◆ Metric

Function Documentation

◆ BuildSpatialHashTableCPU()

◆ BuildSpatialHashTableSYCL()

◆ BuildSpatialHashTableSYCL< double >()

◆ BuildSpatialHashTableSYCL< float >()

◆ BuildSpatialHashTableSYCLRaw()

◆ ComputeVoxelIndex()

◆ CountNeighborsSYCL()

◆ FixedRadiusSearchCPU()

◆ FixedRadiusSearchSYCL()

◆ HybridSearchCPU()

◆ HybridSearchSYCL()

◆ KnnSearchSYCL()

◆ SortNeighborsByDistanceSYCL()

◆ SpatialHash() [1/2]

◆ SpatialHash() [2/2]

◆ WriteNeighborsHybridSYCL()

◆ WriteNeighborsSYCL()

Variable Documentation

◆ kSYCLKnnDefaultTileBytes

◆ kSYCLKnnMidKMax

◆ kSYCLKnnSmallKMax