FixedRadiusSearchSYCLImpl.h File Reference

(bba9247 (Tue Jul 28 21:42:19 2026 -0700))

SYCL device kernels: uniform-grid fixed-radius and hybrid neighbor search. More...

#include <oneapi/dpl/algorithm>
#include <oneapi/dpl/execution>
#include <oneapi/dpl/numeric>
#include <sycl/sycl.hpp>
#include <type_traits>
#include "open3d/core/SYCLContext.h"
#include "open3d/core/Tensor.h"
#include "open3d/core/nns/NeighborSearchCommon.h"
#include "open3d/utility/MiniVec.h"

Go to the source code of this file.

Namespaces
namespace	open3d
namespace	open3d::core
namespace	open3d::core::nns

Functions
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)
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)
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)
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, T threshold, int max_knn)
template<class T , class TIndex >
void	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)

Detailed Description

SYCL device kernels: uniform-grid fixed-radius and hybrid neighbor search.

Included only from KnnSearchOpsSYCL.cpp (not public API). Algorithm matches CUDA FixedRadiusSearchImpl.cuh; shared geometry in NeighborSearchCommon.h (SpatialHash, ComputeVoxelIndex). See nns/SYCL_DESIGN.md for overview.

Grid build (\ref BuildSpatialHashTableSYCL)

1. Bucket dataset into a uniform spatial-hash grid with cell size 2 * radius (any neighbor within radius lies in the query cell or one of seven corner-adjacent cells — 8 bins visited per query, deduplicated).
2. Count points per cell (per batch).
3. Inclusive scan of counts → CSR offsets (hash_table_cell_splits; oneDPL; CUDA uses CUB DeviceScan::InclusiveSum).
4. Scatter point indices into cell ranges (hash_table_index).

Runs on SYCL CPU and GPU. Host driver uses one in-order queue with minimal sync between batch loops.

Query kernels

Mode	Kernels	Passes
Fixed-radius	CountNeighborsSYCL, WriteNeighborsSYCL	Count →

scan on host → allocate → gather | | Hybrid | WriteNeighborsHybridSYCL | Single pass: running top-max_knn + in-radius count, then bubble sort | | Optional sort | SortNeighborsByDistanceSYCL (sort=true) | Segmented sort per query segment |

Metrics: L1, L2, Linf (same as CUDA). L2 compares squared distance to radius²; L1/Linf compare metric distance to radius.

**Sort (sort=true):** oneDPL sort_by_key — float uses packed radix key; double uses struct key + comparator (needs full 64-bit distance). Ties are not secondarily ordered by neighbor index (CUDA parity).

Primitives (CUDA → SYCL)

Role	CUDA	SYCL (this file)
Prefix sum	CUB inclusive scan	oneDPL inclusive scan
Segmented sort	CUB segmented radix sort	oneDPL sort_by_key
Query parallelism	1 thread / query	1 work-item / query

(parallel_for) |

Variable Documentation

dist

T dist

query_id

int64_t query_id