open3d::ml::impl Namespace Reference

Namespaces
namespace	sycl_gemm_detail
namespace	sycl_voxelize_detail

Data Structures
class	Accumulator
class	AccumulatorBackprop
class	Adaptor
	Adaptor for nanoflann. More...
struct	InterpolationVec
	Class for computing interpolation weights. More...
struct	InterpolationVec< T, VECSIZE, InterpolationMode::LINEAR >
	Implementation for LINEAR (uses coordinate clamping) More...
struct	InterpolationVec< T, VECSIZE, InterpolationMode::LINEAR_BORDER >
	Implementation for LINEAR_BORDER (uses zero border instead of clamping) More...
struct	InterpolationVec< T, VECSIZE, InterpolationMode::NEAREST_NEIGHBOR >
	Implementation for NEAREST_NEIGHBOR. More...
class	MemoryAllocation
	A class for managing memory segments within a memory allocation. More...
struct	SelectNanoflannAdaptor
struct	SelectNanoflannAdaptor< L1, T >
struct	SelectNanoflannAdaptor< L2, T >

Enumerations
enum class	InterpolationMode { LINEAR , LINEAR_BORDER , NEAREST_NEIGHBOR }
enum class	CoordinateMapping { BALL_TO_CUBE_RADIAL , BALL_TO_CUBE_VOLUME_PRESERVING , IDENTITY }
enum	Metric { L1 , L2 , Linf }
	Supported metrics. More...
enum	AccumulationFn { AVERAGE = 0 , NEAREST_NEIGHBOR , MAX , CENTER }

Functions
template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool POINT_IMPORTANCE>
void	_CConvComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, const TReal *out_positions, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, bool normalize)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, const TReal *out_positions, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize)
template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool POINT_IMPORTANCE>
void	_CConvBackropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TReal *out_positions, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, const TFeat *out_features_gradient, bool normalize)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvBackpropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TReal *out_positions, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, const TFeat *out_features_gradient, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvBackpropFilterSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *filter_backprop, const std::vector< int > &filter_dims, TIndex num_out, const TReal *out_positions, TIndex num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, const TFeat *out_features_gradient, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, bool allow_tf32)
template<class TFeat , class TReal , class TIndex >
void	FillColumn (const cudaStream_t &stream, TFeat *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, const TReal *const __restrict__ out_positions, TIndex num_inp, const TReal *const __restrict__ inp_positions, const TFeat *const __restrict__ inp_features, const TFeat *const __restrict__ inp_importance, size_t neighbors_index_size, const TIndex *const __restrict__ neighbors_index, const TFeat *const __restrict__ neighbors_importance, const int64_t *const __restrict__ neighbors_row_splits, const TReal *const __restrict__ extents, const TReal *const __restrict__ offsets, const std::vector< int > &filter_dims, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize)
template<class TFeat , class TReal , class TIndex >
void	FillColumnTranspose (const cudaStream_t &stream, TFeat *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, const TReal *const __restrict__ out_positions, TIndex num_inp, const TReal *const __restrict__ inp_positions, const TFeat *const __restrict__ inp_features, const TFeat *const __restrict__ inp_neighbors_importance_sum, const int64_t *const __restrict__ inp_neighbors_prefix_sum, size_t neighbors_index_size, const TIndex *const __restrict__ neighbors_index, const TFeat *const __restrict__ neighbors_importance, const int64_t *const __restrict__ neighbors_row_splits, const TReal *const __restrict__ extents, const TReal *const __restrict__ offsets, const std::vector< int > &filter_dims, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize)
template<class T >
void	MultiplyColumns (const cudaStream_t &stream, size_t rows, size_t cols, T *__restrict__ col_major_matrix, const T *const __restrict__ vector)
template<class T >
void	MultiplyAndCopyColumns (const cudaStream_t &stream, size_t rows, size_t cols, T *__restrict__ out_ptr, const T *const __restrict__ col_major_matrix, const T *const __restrict__ vector)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvComputeFeaturesSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, TIndex num_out, const TReal *out_positions, TIndex num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, bool allow_tf32)
template<class T >
sycl::event	MultiplyColumnsSYCL (sycl::queue &queue, size_t rows, size_t cols, T *col_major_matrix, const T *const vector, const std::vector< sycl::event > &deps)
template<class T >
sycl::event	MultiplyAndCopyColumnsSYCL (sycl::queue &queue, size_t rows, size_t cols, T *out_ptr, const T *const col_major_matrix, const T *const vector, const std::vector< sycl::event > &deps)
template sycl::event	MultiplyColumnsSYCL< float > (sycl::queue &queue, size_t rows, size_t cols, float *col_major_matrix, const float *const vector, const std::vector< sycl::event > &deps)
template sycl::event	MultiplyAndCopyColumnsSYCL< float > (sycl::queue &queue, size_t rows, size_t cols, float *out_ptr, const float *const col_major_matrix, const float *const vector, const std::vector< sycl::event > &deps)
template<class TFeat , class TReal , class TIndex >
sycl::event	FillColumnSYCL (sycl::queue &queue, TFeat *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, const TReal *const out_positions, TIndex num_inp, const TReal *const inp_positions, const TFeat *const inp_features, const TFeat *const inp_importance, size_t neighbors_index_size, const TIndex *const neighbors_index, const TFeat *const neighbors_importance, const int64_t *const neighbors_row_splits, const TReal *const extents, const TReal *const offsets, const std::vector< int > &filter_dims, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, const std::vector< sycl::event > &deps)
template<class TFeat , class TReal , class TIndex >
sycl::event	FillColumnTransposeSYCL (sycl::queue &queue, TFeat *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, const TReal *const out_positions, TIndex num_inp, const TReal *const inp_positions, const TFeat *const inp_features, const TFeat *const inp_neighbors_importance_sum, const int64_t *const inp_neighbors_prefix_sum, size_t neighbors_index_size, const TIndex *const neighbors_index, const TFeat *const neighbors_importance, const int64_t *const neighbors_row_splits, const TReal *const extents, const TReal *const offsets, const std::vector< int > &filter_dims, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, const std::vector< sycl::event > &deps)
template sycl::event	FillColumnSYCL< float, float, int32_t > (sycl::queue &queue, float *columns, int in_channels, int32_t begin_idx, int32_t end_idx, int32_t num_out, const float *const out_positions, int32_t num_inp, const float *const inp_positions, const float *const inp_features, const float *const inp_importance, size_t neighbors_index_size, const int32_t *const neighbors_index, const float *const neighbors_importance, const int64_t *const neighbors_row_splits, const float *const extents, const float *const offsets, const std::vector< int > &filter_dims, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, const std::vector< sycl::event > &deps)
template sycl::event	FillColumnTransposeSYCL< float, float, int32_t > (sycl::queue &queue, float *columns, int in_channels, int32_t begin_idx, int32_t end_idx, int32_t num_out, const float *const out_positions, int32_t num_inp, const float *const inp_positions, const float *const inp_features, const float *const inp_neighbors_importance_sum, const int64_t *const inp_neighbors_prefix_sum, size_t neighbors_index_size, const int32_t *const neighbors_index, const float *const neighbors_importance, const int64_t *const neighbors_row_splits, const float *const extents, const float *const offsets, const std::vector< int > &filter_dims, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, const std::vector< sycl::event > &deps)
template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool NORMALIZE>
void	_CConvTransposeComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, const TReal *out_positions, const TFeat *out_importance, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvTransposeComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, const TReal *out_positions, const TFeat *out_importance, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize)
template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool NORMALIZE>
void	_CConvTransposeBackpropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TReal *out_positions, const TFeat *out_importance, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, const TFeat *out_features_gradient)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvTransposeBackpropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TReal *out_positions, const TFeat *out_importance, size_t num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, const TFeat *out_features_gradient, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvTransposeBackpropFilterSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *filter_backprop, const std::vector< int > &filter_dims, TIndex num_out, const TReal *out_positions, const TFeat *out_importance, TIndex num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, const TFeat *out_features_gradient, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, bool allow_tf32)
template<class TFeat , class TOut , class TReal , class TIndex >
void	CConvTransposeComputeFeaturesSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, TIndex num_out, const TReal *out_positions, const TFeat *out_importance, TIndex num_inp, const TReal *inp_positions, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_prefix_sum, size_t neighbors_index_size, const TIndex *neighbors_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TReal *extents, const TReal *offsets, InterpolationMode interpolation, CoordinateMapping coordinate_mapping, bool align_corners, bool individual_extent, bool isotropic_extent, bool normalize, bool allow_tf32)
template<class T , int VECSIZE>
void	MapSphereToCylinder (Eigen::Array< T, VECSIZE, 1 > &x, Eigen::Array< T, VECSIZE, 1 > &y, Eigen::Array< T, VECSIZE, 1 > &z)
template<class T , int VECSIZE>
void	MapCylinderToCube (Eigen::Array< T, VECSIZE, 1 > &x, Eigen::Array< T, VECSIZE, 1 > &y, Eigen::Array< T, VECSIZE, 1 > &z)
template<bool ALIGN_CORNERS, CoordinateMapping MAPPING, class T , int VECSIZE>
void	ComputeFilterCoordinates (Eigen::Array< T, VECSIZE, 1 > &x, Eigen::Array< T, VECSIZE, 1 > &y, Eigen::Array< T, VECSIZE, 1 > &z, const Eigen::Array< int, 3, 1 > &filter_size, const Eigen::Array< T, VECSIZE, 3 > &inv_extents, const Eigen::Array< T, 3, 1 > &offset)
template<class T >
void	MapSphereToCylinderSYCL (T &x, T &y, T &z)
template<class T >
void	MapCylinderToCubeSYCL (T &x, T &y, T &z)
template<bool ALIGN_CORNERS, CoordinateMapping MAPPING, class T >
void	ComputeFilterCoordinatesSYCL (T &x, T &y, T &z, const int &filter_size_x, const int &filter_size_y, const int &filter_size_z, const T &inv_extent_x, const T &inv_extent_y, const T &inv_extent_z, const T &offset_x, const T &offset_y, const T &offset_z)
template<InterpolationMode INTERPOLATION, class T >
void	InterpolateSYCL (T *w, int *idx, const T &x, const T &y, const T &z, const int &filter_size_x, const int &filter_size_y, const int &filter_size_z, int num_channels=1)
template<class LayoutA = cutlass::layout::ColumnMajor, class LayoutB = cutlass::layout::ColumnMajor>
void	GemmColumnMajorCUDA (const cudaStream_t &stream, int m, int n, int k, float alpha, const float *A, int lda, const float *B, int ldb, float beta, float *C, int ldc)
template<class LayoutA , class LayoutB >
sycl::event	GemmColumnMajorSYCL (sycl::queue &queue, int m, int n, int k, float alpha, const float *A, int64_t lda, const float *B, int64_t ldb, float beta, float *C, int64_t ldc, bool allow_tf32, const std::vector< sycl::event > &deps)
template sycl::event	GemmColumnMajorSYCL< cutlass::layout::ColumnMajor, cutlass::layout::ColumnMajor > (sycl::queue &, int, int, int, float, const float *, int64_t, const float *, int64_t, float, float *, int64_t, bool, const std::vector< sycl::event > &)
template sycl::event	GemmColumnMajorSYCL< cutlass::layout::ColumnMajor, cutlass::layout::RowMajor > (sycl::queue &, int, int, int, float, const float *, int64_t, const float *, int64_t, float, float *, int64_t, bool, const std::vector< sycl::event > &)
template<class TIndex , class TAttr >
void	InvertNeighborsListCPU (const TIndex *const inp_neighbors_index, const TAttr *const inp_neighbors_attributes, const int num_attributes_per_neighbor, const int64_t *const inp_neighbors_row_splits, const size_t inp_num_queries, TIndex *out_neighbors_index, TAttr *out_neighbors_attributes, const size_t index_size, int64_t *out_neighbors_row_splits, const size_t out_num_queries)
template<class T >
sycl::event	CountIndexOccurrencesSYCL (sycl::queue &queue, uint32_t *count, size_t count_size, const T *indices, size_t indices_size, const std::vector< sycl::event > &deps={})
template<class TIndex , class TAttr >
sycl::event	FillNeighborsIndexAndAttributesSYCL (sycl::queue &queue, uint32_t *count, size_t count_size, TIndex *out_neighbors_index, TAttr *out_neighbors_attributes, const TIndex *inp_neighbors_index, const TAttr *inp_neighbors_attributes, int num_attributes_per_neighbor, size_t index_size, const int64_t *inp_neighbors_row_splits, size_t inp_num_queries, const int64_t *out_neighbors_row_splits, size_t out_num_queries, const std::vector< sycl::event > &deps={})
template<class TIndex , class TAttr >
sycl::event	InvertNeighborsListSYCL (sycl::queue &queue, uint32_t *count_buf, const TIndex *inp_neighbors_index, const TAttr *inp_neighbors_attributes, int num_attributes_per_neighbor, const int64_t *inp_neighbors_row_splits, size_t inp_num_queries, TIndex *out_neighbors_index, TAttr *out_neighbors_attributes, size_t index_size, int64_t *out_neighbors_row_splits, size_t out_num_queries)
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)
template<class T >
void	RaggedToDenseCPU (const T *const values, const int64_t *const row_splits, const size_t row_splits_size, const size_t out_col_size, const T *const default_value, const size_t default_value_size, T *out_values)
template<class T >
void	RaggedToDenseSYCL (sycl::queue &queue, const T *const values, const int64_t *const row_splits, const size_t row_splits_size, const size_t out_col_size, const T *const default_value, const size_t default_value_size, T *out_values)
template<class T >
void	ReduceSubarraysSumCPU (const T *const values, const size_t values_size, const int64_t *const row_splits, const size_t num_arrays, T *out_sums)
template<class T >
void	ReduceSubarraysSumSYCL (sycl::queue &queue, const T *const values, const size_t values_size, const int64_t *const row_splits, const size_t num_arrays, T *out_sums)
template<class T , int NDIM, class OUTPUT_ALLOCATOR >
void	VoxelizeCPU (const size_t num_points, const T *const points, const size_t batch_size, const int64_t *const row_splits, const T *const voxel_size, const T *const points_range_min, const T *const points_range_max, const int64_t max_points_per_voxel, const int64_t max_voxels, OUTPUT_ALLOCATOR &output_allocator)
template<class T , int NDIM, class OUTPUT_ALLOCATOR >
void	VoxelizeSYCL (sycl::queue &queue, size_t num_points, const T *const points, const size_t batch_size, const int64_t *const row_splits, const T *const voxel_size, const T *const points_range_min, const T *const points_range_max, const int64_t max_points_per_voxel, const int64_t max_voxels, OUTPUT_ALLOCATOR &output_allocator)
template<class T >
bool	CheckVoxelSize (std::string &err, const size_t num_positions, const T *const positions, const T voxel_size)
	Function for debugging. Checks if the voxel size is too small.
template<class TDerived >
Eigen::Vector3i	ComputeVoxelIndex (const Eigen::ArrayBase< TDerived > &pos, const typename TDerived::Scalar &inv_voxel_size)
template<class TReal , class TFeat , class ACCUMULATOR , class OUTPUT_ALLOCATOR >
void	_VoxelPooling (size_t num_inp, const TReal *const inp_positions, int in_channels, const TFeat *inp_features, TReal voxel_size, OUTPUT_ALLOCATOR &output_allocator)
template<class TReal , class TFeat , class ACCUMULATOR , AccumulationFn FEAT_FN>
void	_VoxelPoolingBackprop (TFeat *features_backprop, size_t num_inp, const TReal *const inp_positions, int in_channels, const TFeat *const inp_features, size_t num_pooled, const TReal *const pooled_positions, const TFeat *const pooled_features_gradient, TReal voxel_size)
template<class TReal , class TFeat , class OUTPUT_ALLOCATOR >
void	VoxelPooling (size_t num_inp, const TReal *const inp_positions, int in_channels, const TFeat *inp_features, TReal voxel_size, OUTPUT_ALLOCATOR &output_allocator, AccumulationFn position_fn, AccumulationFn feature_fn)
template<class TReal , class TFeat >
void	VoxelPoolingBackprop (TFeat *features_backprop, size_t num_inp, const TReal *const inp_positions, int in_channels, const TFeat *const inp_features, size_t num_pooled, const TReal *const pooled_positions, const TFeat *const pooled_features_gradient, TReal voxel_size, AccumulationFn position_fn, AccumulationFn feature_fn)
template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool POINT_IMPORTANCE>
void	_SparseConvComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, size_t num_inp, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, bool normalize)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, size_t num_inp, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, bool normalize)
template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool POINT_IMPORTANCE>
void	_SparseConvBackropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, size_t num_inp, const TFeat *inp_features, const TFeat *inp_importance, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient, bool normalize)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvBackpropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, size_t num_inp, const TFeat *inp_features, const TFeat *inp_importance, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient, bool normalize)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvBackpropFilterSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *filter_backprop, const std::vector< int > &filter_dims, TIndex num_out, TIndex num_inp, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient, bool normalize, bool allow_tf32)
template<class TReal , class TIndex , class TKernelIndex >
void	FillColumn (const cudaStream_t &stream, TReal *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, TIndex num_inp, const TReal *const __restrict__ inp_features, const TReal *const __restrict__ inp_importance, size_t neighbors_index_size, const TIndex *const __restrict__ neighbors_index, const TKernelIndex *const __restrict__ neighbors_kernel_index, const TReal *const __restrict__ neighbors_importance, const int64_t *const __restrict__ neighbors_row_splits, const int num_kernel_elements, bool normalize)
template<class TReal , class TIndex , class TKernelIndex >
void	FillColumnTranspose (const cudaStream_t &stream, TReal *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, TIndex num_inp, const TReal *const __restrict__ inp_features, const TReal *const __restrict__ inp_neighbors_importance_sum, const int64_t *const __restrict__ inp_neighbors_prefix_sum, size_t neighbors_index_size, const TIndex *const __restrict__ neighbors_index, const TKernelIndex *const __restrict__ neighbors_kernel_index, const TReal *const __restrict__ neighbors_importance, const int64_t *const __restrict__ neighbors_row_splits, const int num_kernel_elements, bool normalize)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvComputeFeaturesSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, TIndex num_out, TIndex num_inp, const TFeat *inp_features, const TFeat *inp_importance, size_t neighbors_index_size, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, bool normalize, bool allow_tf32)
template<class TReal , class TIndex , class TKernelIndex >
sycl::event	FillColumnSYCL (sycl::queue &queue, TReal *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, TIndex num_inp, const TReal *const inp_features, const TReal *const inp_importance, size_t neighbors_index_size, const TIndex *const neighbors_index, const TKernelIndex *const neighbors_kernel_index, const TReal *const neighbors_importance, const int64_t *const neighbors_row_splits, const int num_kernel_elements, bool normalize, const std::vector< sycl::event > &deps)
template<class TReal , class TIndex , class TKernelIndex >
sycl::event	FillColumnTransposeSYCL (sycl::queue &queue, TReal *columns, int in_channels, TIndex begin_idx, TIndex end_idx, TIndex num_out, TIndex num_inp, const TReal *const inp_features, const TReal *const inp_neighbors_importance_sum, const int64_t *const inp_neighbors_prefix_sum, size_t neighbors_index_size, const TIndex *const neighbors_index, const TKernelIndex *const neighbors_kernel_index, const TReal *const neighbors_importance, const int64_t *const neighbors_row_splits, const int num_kernel_elements, bool normalize, const std::vector< sycl::event > &deps)
template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool NORMALIZE>
void	_SparseConvTransposeComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, const TFeat *out_importance, size_t num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, const TIndex *neighbor_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbor_importance, const int64_t *neighbors_row_splits)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvTransposeComputeFeaturesCPU (TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, size_t num_out, const TFeat *out_importance, size_t num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, const TIndex *neighbor_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbor_importance, const int64_t *neighbors_row_splits, bool normalize)
template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool NORMALIZE>
void	_SparseConvTransposeBackpropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TFeat *out_importance, size_t num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvTransposeBackpropFilterCPU (TOut *filter_backprop, const std::vector< int > &filter_dims, size_t num_out, const TFeat *out_importance, size_t num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient, bool normalize)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvTransposeBackpropFilterSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *filter_backprop, const std::vector< int > &filter_dims, TIndex num_out, const TFeat *out_importance, TIndex num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_row_splits, size_t neighbors_index_size, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, const TFeat *out_features_gradient, bool normalize, bool allow_tf32)
template<class TFeat , class TOut , class TIndex , class TKernelIndex >
void	SparseConvTransposeComputeFeaturesSYCL (sycl::queue &queue, void *temp, size_t &temp_size, size_t &max_temp_size, int texture_alignment, TOut *out_features, const std::vector< int > &filter_dims, const TFeat *filter, TIndex num_out, const TFeat *out_importance, TIndex num_inp, const TFeat *inp_features, const TFeat *inp_neighbors_importance_sum, const int64_t *inp_neighbors_prefix_sum, size_t neighbors_index_size, const TIndex *neighbors_index, const TKernelIndex *neighbors_kernel_index, const TFeat *neighbors_importance, const int64_t *neighbors_row_splits, bool normalize, bool allow_tf32)

Enumeration Type Documentation

AccumulationFn

enum open3d::ml::impl::AccumulationFn

Enumerator
AVERAGE
NEAREST_NEIGHBOR
MAX
CENTER

CoordinateMapping

enum class open3d::ml::impl::CoordinateMapping

strong

Coordinate Mapping functions

• BALL_TO_CUBE_RADIAL uses radial stretching to map a sphere to a cube.
• BALL_TO_CUBE_VOLUME_PRESERVING is using a more expensive volume preserving mapping to map a sphere to a cube.
• IDENTITY no mapping is applied to the coordinates.

Enumerator
BALL_TO_CUBE_RADIAL
BALL_TO_CUBE_VOLUME_PRESERVING
IDENTITY

InterpolationMode

enum class open3d::ml::impl::InterpolationMode

strong

Interpolation modes LINEAR is a standard trilinear interpolation with coordinate clamping LINEAR_BORDER uses a zero border instead of clamping NEAREST_NEIGHBOR no interpolation, use nearest neighbor

Enumerator
LINEAR
LINEAR_BORDER
NEAREST_NEIGHBOR

Metric

enum open3d::ml::impl::Metric

Supported metrics.

Enumerator
L1
L2
Linf

Function Documentation

_CConvBackropFilterCPU()

template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool POINT_IMPORTANCE>

void open3d::ml::impl::_CConvBackropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		const TReal *	out_positions,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		const TFeat *	out_features_gradient,
		bool	normalize
	)

_CConvComputeFeaturesCPU()

template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool POINT_IMPORTANCE>

void open3d::ml::impl::_CConvComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		const TReal *	out_positions,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		bool	normalize
	)

Implementation of CConvComputeFeatures with template parameters for configuration.

_CConvTransposeBackpropFilterCPU()

template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool NORMALIZE>

void open3d::ml::impl::_CConvTransposeBackpropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		const TReal *	out_positions,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		const TFeat *	out_features_gradient
	)

Implementation of CConvTransposeBackpropFilterCPU with template parameters for configuration.

_CConvTransposeComputeFeaturesCPU()

template<class TFeat , class TOut , class TReal , class TIndex , InterpolationMode INTERPOLATION, CoordinateMapping MAPPING, bool ALIGN_CORNERS, bool INDIVIDUAL_EXTENT, bool ISOTROPIC_EXTENT, bool NORMALIZE>

void open3d::ml::impl::_CConvTransposeComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		const TReal *	out_positions,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets
	)

Implementation of CConvComputeFeatures with template parameters for configuration.

_SparseConvBackropFilterCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool POINT_IMPORTANCE>

void open3d::ml::impl::_SparseConvBackropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TFeat *	out_features_gradient,
		bool	normalize
	)

_SparseConvComputeFeaturesCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool POINT_IMPORTANCE>

void open3d::ml::impl::_SparseConvComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		bool	normalize
	)

Implementation of SparseConvComputeFeatures with template parameters for configuration.

_SparseConvTransposeBackpropFilterCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool NORMALIZE>

void open3d::ml::impl::_SparseConvTransposeBackpropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TFeat *	out_features_gradient
	)

Implementation of SparseConvTransposeBackpropFilterCPU with template parameters for configuration.

_SparseConvTransposeComputeFeaturesCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex , bool NORMALIZE>

void open3d::ml::impl::_SparseConvTransposeComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		const TIndex *	neighbor_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbor_importance,
		const int64_t *	neighbors_row_splits
	)

Implementation of SparseConvTransposeComputeFeatures with template parameters for configuration.

_VoxelPooling()

template<class TReal , class TFeat , class ACCUMULATOR , class OUTPUT_ALLOCATOR >

void open3d::ml::impl::_VoxelPooling	(	size_t	num_inp,
		const TReal *const	inp_positions,
		int	in_channels,
		const TFeat *	inp_features,
		TReal	voxel_size,
		OUTPUT_ALLOCATOR &	output_allocator
	)

_VoxelPoolingBackprop()

template<class TReal , class TFeat , class ACCUMULATOR , AccumulationFn FEAT_FN>

void open3d::ml::impl::_VoxelPoolingBackprop	(	TFeat *	features_backprop,
		size_t	num_inp,
		const TReal *const	inp_positions,
		int	in_channels,
		const TFeat *const	inp_features,
		size_t	num_pooled,
		const TReal *const	pooled_positions,
		const TFeat *const	pooled_features_gradient,
		TReal	voxel_size
	)

CConvBackpropFilterCPU()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvBackpropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		const TReal *	out_positions,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		const TFeat *	out_features_gradient,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize
	)

Computes the backprop for the filter of a continuous convolution.

Template Parameters

TFeat	Type for the features and weights
TOut	Type for the output features
TReal	Type for point positions and extents
TIndex	Type for neighbor indexing

Parameters

filter_backprop	Output array for the computed filter gradient with shape [depth,height,width, inp channels, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be 5. The order is [depth, height, width, inp channels, out channels].
num_out	The number of output points.
out_positions	The positions of the output points. The shape is [num_out, 3].
num_inp	The number of input points.
inp_positions	The positions of the input points. The shape is [num_inp, 3].
inp_features	The input features with shape [num_inp, in_channels].
inp_importance	Optional importance for each input point with shape [num_inp]. Set to null to disable.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
extents	The spatial extents of the filter in coordinate units. extents can be a scalar or a 1D array of shape [num_out] or a 2D array of shape [num_out,3]. The shape depends on individual_extent and isotropic_extent.
offsets	A single 3D vector used in the filter coordinate computation. The shape is [3].
interpolation	The interpolation mode. Either LINEAR or NEAREST_NEIGHBOR.
coordinate_mapping	The coordinate mapping function. One of IDENTITY, BALL_TO_CUBE_RADIAL, BALL_TO_CUBE_VOLUME_PRESERVING.
align_corners	If true then the voxel centers of the outer voxels of the filter array are mapped to the boundary of the filter shape. If false then the boundary of the filter array is mapped to the boundary of the filter shape.
individual_extent	If true each output point has an individual extent.
isotropic_extent	If true each then the extent is isotropic for each output point.
normalize	If true then the output features are normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

CConvBackpropFilterSYCL()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvBackpropFilterSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		TIndex	num_out,
		const TReal *	out_positions,
		TIndex	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		const TFeat *	out_features_gradient,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of CConvBackpropFilterCUDA. See ContinuousConvBackpropFilter .cuh for the full parameter documentation; semantics are identical.

CConvComputeFeaturesCPU()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		const TReal *	out_positions,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize
	)

Computes the output features of a continuous convolution.

Template Parameters

TFeat	Type for the features and weights
TOut	Type for the output features
TReal	Type for point positions and extents
TIndex	Type for neighbor indexing

Parameters

out_features	Output array for the computed features with shape [num_out, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be 5. The order is [depth, height, width, inp channels, out channels].
filter	Pointer to the filter values.
num_out	The number of output points.
out_positions	The positions of the output points. The shape is [num_out, 3].
num_inp	The number of input points.
inp_positions	The positions of the input points. The shape is [num_inp, 3].
inp_features	The input features with shape [num_inp, in_channels].
inp_importance	Optional importance for each input point with shape [num_inp]. Set to null to disable.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
extents	The spatial extents of the filter in coordinate units. extents can be a scalar or a 1D array of shape [num_out] or a 2D array of shape [num_out,3]. The shape depends on individual_extent and isotropic_extent.
offsets	A single 3D vector used in the filter coordinate computation. The shape is [3].
interpolation	The interpolation mode. Either LINEAR or NEAREST_NEIGHBOR.
coordinate_mapping	The coordinate mapping function. One of IDENTITY, BALL_TO_CUBE_RADIAL, BALL_TO_CUBE_VOLUME_PRESERVING.
align_corners	If true then the voxel centers of the outer voxels of the filter array are mapped to the boundary of the filter shape. If false then the boundary of the filter array is mapped to the boundary of the filter shape.
individual_extent	If true each output point has an individual extent.
isotropic_extent	If true each then the extent is isotropic for each output point.
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

CConvComputeFeaturesSYCL()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvComputeFeaturesSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		TIndex	num_out,
		const TReal *	out_positions,
		TIndex	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of CConvComputeFeaturesCUDA. See ContinuousConv.cuh for the full parameter documentation; semantics are identical.

CConvTransposeBackpropFilterCPU()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvTransposeBackpropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		const TReal *	out_positions,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		const TFeat *	out_features_gradient,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize
	)

Computes the backprop for the filter of a transpose continuous convolution.

Template Parameters

TFeat	Type for the features and weights
TOut	Type for the output features
TReal	Type for point positions and extents
TIndex	Type for neighbor indexing

Parameters

filter_backprop	Output array for the computed filter gradient with shape [depth,height,width, inp channels, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be 5. The order is [depth, height, width, inp channels, out channels].
num_out	The number of output points.
out_positions	The positions of the output points. The shape is [num_out, 3].
out_importance	Optional importance for each output point with shape [num_out]. Set to null to disable.
num_inp	The number of input points.
inp_positions	The positions of the input points. The shape is [num_inp, 3].
inp_features	The input features with shape [num_inp, in_channels].
inp_neighbors_importance_sum	The sum of the neighbors_importance values for each input with shape [num_inp].
inp_neighbors_row_splits	The prefix sum which defines the start and end of the sublists in inp_neighbors_index. The size of the array is num_inp + 1.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
extents	The spatial extents of the filter in coordinate units. extents can be a scalar or a 1D array of shape [num_out] or a 2D array of shape [num_out,3]. The shape depends on individual_extent and isotropic_extent.
offsets	A single 3D vector used in the filter coordinate computation. The shape is [3].
out_features_gradient	The gradient from the features with shape [num_out, out_channels]
interpolation	The interpolation mode. Either LINEAR or NEAREST_NEIGHBOR.
coordinate_mapping	The coordinate mapping function. One of IDENTITY, BALL_TO_CUBE_RADIAL, BALL_TO_CUBE_VOLUME_PRESERVING.
align_corners	If true then the voxel centers of the outer voxels of the filter array are mapped to the boundary of the filter shape. If false then the boundary of the filter array is mapped to the boundary of the filter shape.
individual_extent	If true each output point has an individual extent.
isotropic_extent	If true each then the extent is isotropic for each output point.
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

CConvTransposeBackpropFilterSYCL()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvTransposeBackpropFilterSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		TIndex	num_out,
		const TReal *	out_positions,
		const TFeat *	out_importance,
		TIndex	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		const TFeat *	out_features_gradient,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of CConvTransposeBackpropFilterCUDA. See ContinuousConvTransposeBackpropFilter.cuh for the full parameter documentation; semantics are identical.

CConvTransposeComputeFeaturesCPU()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvTransposeComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		const TReal *	out_positions,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize
	)

Computes the output features of a transpose continuous convolution.

Template Parameters

TFeat	Type for the features and weights
TOut	Type for the output features
TReal	Type for point positions and extents
TIndex	Type for neighbor indexing

Parameters

out_features	Output array for the computed features with shape [num_out, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be 5. The order is [depth, height, width, inp channels, out channels].
filter	Pointer to the filter values.
num_out	The number of output points.
out_positions	The positions of the output points. The shape is [num_out, 3].
out_importance	Optional importance for each output point with shape [num_out]. Set to null to disable.
num_inp	The number of input points.
inp_positions	The positions of the input points. The shape is [num_inp, 3].
inp_features	The input features with shape [num_inp, in_channels].
inp_neighbors_importance_sum	The sum of the neighbors_importance values for each input with shape [num_inp].
inp_neighbors_row_splits	The prefix sum which defines the start and end of the sublists in inp_neighbors_index. The size of the array is num_inp + 1.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
extents	The spatial extents of the filter in coordinate units. extents can be a scalar or a 1D array of shape [num_out] or a 2D array of shape [num_out,3]. The shape depends on individual_extent and isotropic_extent.
offsets	A single 3D vector used in the filter coordinate computation. The shape is [3].
interpolation	The interpolation mode. Either LINEAR or NEAREST_NEIGHBOR.
coordinate_mapping	The coordinate mapping function. One of IDENTITY, BALL_TO_CUBE_RADIAL, BALL_TO_CUBE_VOLUME_PRESERVING.
align_corners	If true then the voxel centers of the outer voxels of the filter array are mapped to the boundary of the filter shape. If false then the boundary of the filter array is mapped to the boundary of the filter shape.
individual_extent	If true each output point has an individual extent.
isotropic_extent	If true each then the extent is isotropic for each output point.
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

CConvTransposeComputeFeaturesSYCL()

template<class TFeat , class TOut , class TReal , class TIndex >

void open3d::ml::impl::CConvTransposeComputeFeaturesSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		TIndex	num_out,
		const TReal *	out_positions,
		const TFeat *	out_importance,
		TIndex	num_inp,
		const TReal *	inp_positions,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_prefix_sum,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TReal *	extents,
		const TReal *	offsets,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of CConvTransposeComputeFeaturesCUDA. See ContinuousConvTranspose.cuh for the full parameter documentation; semantics are identical.

CheckVoxelSize()

template<class T >

bool open3d::ml::impl::CheckVoxelSize	(	std::string &	err,
		const size_t	num_positions,
		const T *const	positions,
		const T	voxel_size
	)

Function for debugging. Checks if the voxel size is too small.

ComputeFilterCoordinates()

template<bool ALIGN_CORNERS, CoordinateMapping MAPPING, class T , int VECSIZE>

void open3d::ml::impl::ComputeFilterCoordinates ( Eigen::Array< T, VECSIZE, 1 > &  x, Eigen::Array< T, VECSIZE, 1 > &  y, Eigen::Array< T, VECSIZE, 1 > &  z, const Eigen::Array< int, 3, 1 > &  filter_size, const Eigen::Array< T, VECSIZE, 3 > &  inv_extents, const Eigen::Array< T, 3, 1 > &  offset  )

inline

Computes the filter coordinates. The input to this function are coordinates relative to the point where the convolution is evaluated. Coordinates are usually in the range [-extent/2,extent/2] with extent as the edge length of the bounding box of the filter shape. The output is a coordinate within the filter array, i.e. the range is [0, filter_size.xyz], if the point was inside the filter shape.

The simplest filter shape is a cuboid (MAPPING=IDENTITY) and the transformation is simply [-extent/2,extent/2] -> [0, filter_size.xyz]. The other type of shape that is implemented is a sphere with MAPPING=BALL_TO_CUBE_RADIAL or MAPPING=BALL_TO_CUBE_VOLUME_PRESERVING.

Template Parameters

ALIGN_CORNERS	If true then the voxel centers of the outer voxels of the filter array are mapped to the boundary of the filter shape. If false then the boundary of the filter array is mapped to the boundary of the filter shape.
MAPPING	The mapping that is applied to the input coordinates. BALL_TO_CUBE_RADIAL uses radial stretching to map a sphere to a cube. BALL_TO_CUBE_VOLUME_PRESERVING is using a more expensive volume preserving mapping to map a sphere to a cube. IDENTITY no mapping is applied to the coordinates.

Parameters

x	x coordinates. Input and output variable.
y	y coordinates. Input and output variable.
z	z coordinates. Input and output variable.
filter_size	The spatial size of the filter array in voxels.
inv_extents	The reciproval of the spatial extent of the filter in coordinate units.
offset	An offset for shifting the center. Can be used to implement discrete filters with even filter size.

ComputeFilterCoordinatesSYCL()

template<bool ALIGN_CORNERS, CoordinateMapping MAPPING, class T >

void open3d::ml::impl::ComputeFilterCoordinatesSYCL ( T &  x, T &  y, T &  z, const int &  filter_size_x, const int &  filter_size_y, const int &  filter_size_z, const T &  inv_extent_x, const T &  inv_extent_y, const T &  inv_extent_z, const T &  offset_x, const T &  offset_y, const T &  offset_z  )

inline

Computes the filter coordinates. See CoordinateTransformation.cuh for the full parameter documentation; semantics are identical (SYCL port).

ComputeVoxelIndex() [1/2]

template<class TDerived >

Eigen::Vector3i open3d::ml::impl::ComputeVoxelIndex	(	const Eigen::ArrayBase< TDerived > &	pos,
		const typename TDerived::Scalar &	inv_voxel_size
	)

Computes an integer voxel index for a 3D position.

Parameters

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

ComputeVoxelIndex() [2/2]

template<class TVecf >

HOST_DEVICE utility::MiniVec< int, 3 > open3d::ml::impl::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

CountIndexOccurrencesSYCL()

template<class T >

sycl::event open3d::ml::impl::CountIndexOccurrencesSYCL	(	sycl::queue &	queue,
		uint32_t *	count,
		size_t	count_size,
		const T *	indices,
		size_t	indices_size,
		const std::vector< sycl::event > &	deps = {}
	)

Count the number of occurrences of each index value using atomic add. count[] is zero-initialized by this call. Named CountIndexOccurrencesSYCL to distinguish from the unrelated CountNeighborsSYCL in open3d::core::nns (which counts FRS neighbors).

FillColumn() [1/2]

template<class TFeat , class TReal , class TIndex >

void open3d::ml::impl::FillColumn	(	const cudaStream_t &	stream,
		TFeat *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		const TReal *const __restrict__	out_positions,
		TIndex	num_inp,
		const TReal *const __restrict__	inp_positions,
		const TFeat *const __restrict__	inp_features,
		const TFeat *const __restrict__	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *const __restrict__	neighbors_index,
		const TFeat *const __restrict__	neighbors_importance,
		const int64_t *const __restrict__	neighbors_row_splits,
		const TReal *const __restrict__	extents,
		const TReal *const __restrict__	offsets,
		const std::vector< int > &	filter_dims,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize
	)

Copies and transforms the features to a column, which can be multiplied with the filter matrix.

Template Parameters

TFeat	Type for the features
TOut	Type for the output features
TReal	Type for point positions and extents
TIndex	Type for neighbor indexing

Parameters

columns	Output array with shape [num_out, spatial filter dims, in_channels].
in_channels	Number of input channels.
begin_idx	Index of the first output point to process.
end_idx	Index after the last output point to process.
num_out	The number of output points.
out_positions	The positions of the output points. The shape is [num_out, 3].
num_inp	The number of input points.
inp_positions	The positions of the input points. The shape is [num_inp, 3].
inp_features	The input features with shape [num_inp, in_channels].
inp_importance	Optional importance for each input point with shape [num_inp]. Set to null to disable.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
extents	The spatial extents of the filter in coordinate units. extents can be a scalar or a 1D array of shape [num_out] or a 2D array of shape [num_out,3]. The shape depends on individual_extent and isotropic_extent.
offsets	A single 3D vector used in the filter coordinate computation. The shape is [3].
filter_dims	The spatial filter size in voxels. (The filter resolution) with shape [3].
interpolation	The interpolation mode. Either LINEAR or NEAREST_NEIGHBOR.
coordinate_mapping	The coordinate mapping function. One of IDENTITY, BALL_TO_CUBE_RADIAL, BALL_TO_CUBE_VOLUME_PRESERVING.
align_corners	If true then the voxel centers of the outer voxels of the filter array are mapped to the boundary of the filter shape. If false then the boundary of the filter array is mapped to the boundary of the filter shape.
individual_extent	If true each output point has an individual extent.
isotropic_extent	If true each then the extent is isotropic for each output point.
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

FillColumn() [2/2]

template<class TReal , class TIndex , class TKernelIndex >

void open3d::ml::impl::FillColumn	(	const cudaStream_t &	stream,
		TReal *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		TIndex	num_inp,
		const TReal *const __restrict__	inp_features,
		const TReal *const __restrict__	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *const __restrict__	neighbors_index,
		const TKernelIndex *const __restrict__	neighbors_kernel_index,
		const TReal *const __restrict__	neighbors_importance,
		const int64_t *const __restrict__	neighbors_row_splits,
		const int	num_kernel_elements,
		bool	normalize
	)

Copies and transforms the features to a column, which can be multiplied with the filter matrix.

Template Parameters

TReal	Type for positions and features.
TIndex	Type for addressing neighbors.
TKernelIndex	Type for addressing the spatial dimension in a kernel.

Parameters

columns	Output array with shape [num_out, spatial filter dims, in_channels].
in_channels	Number of input channels.
begin_idx	Index of the first output point to process.
end_idx	Index after the last output point to process.
num_out	The number of output points.
num_inp	The number of input points.
inp_features	The input features with shape [num_inp, in_channels].
inp_importance	Optional importance for each input point with shape [num_inp]. Set to null to disable.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_kernel_index	Defines which kernel element to use for each neighbor. This array has the same length as neighbors_index.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
num_kernel_elements	The number of kernel elements. This is the size of the first dimension of the filter, i.e., The filter shape is [num_kernel_elements, in_channels, out_channels]
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

FillColumnSYCL() [1/2]

template<class TFeat , class TReal , class TIndex >

sycl::event open3d::ml::impl::FillColumnSYCL	(	sycl::queue &	queue,
		TFeat *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		const TReal *const	out_positions,
		TIndex	num_inp,
		const TReal *const	inp_positions,
		const TFeat *const	inp_features,
		const TFeat *const	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *const	neighbors_index,
		const TFeat *const	neighbors_importance,
		const int64_t *const	neighbors_row_splits,
		const TReal *const	extents,
		const TReal *const	offsets,
		const std::vector< int > &	filter_dims,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		const std::vector< sycl::event > &	deps = {}
	)

Copies and transforms the features to a column, which can be multiplied with the filter matrix. SYCL port of FillColumn (see ContinuousConvCUDAKernels.h for the full parameter documentation; semantics are identical, only cudaStream_t -> sycl::queue&). Depends on deps and returns the completion event of its last kernel; callers thread this event into e.g. the following GEMM's dependencies instead of blocking-waiting on it.

FillColumnSYCL() [2/2]

template<class TReal , class TIndex , class TKernelIndex >

sycl::event open3d::ml::impl::FillColumnSYCL	(	sycl::queue &	queue,
		TReal *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		TIndex	num_inp,
		const TReal *const	inp_features,
		const TReal *const	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *const	neighbors_index,
		const TKernelIndex *const	neighbors_kernel_index,
		const TReal *const	neighbors_importance,
		const int64_t *const	neighbors_row_splits,
		const int	num_kernel_elements,
		bool	normalize,
		const std::vector< sycl::event > &	deps = {}
	)

SYCL port of FillColumn: copies/gathers neighbor features (optionally normalized) into the im2col-style patch matrix used as the GEMM's B operand. One work-group per output point; work-items in the group stride over input channels (and, for the normalizer sum, over neighbors). Depends on deps and returns the completion event of its last kernel (the zero-init fill and the fill-column kernel are chained internally); callers thread this event into e.g. the following GEMM's dependencies instead of blocking-waiting on it.

FillColumnSYCL< float, float, int32_t >()

template sycl::event open3d::ml::impl::FillColumnSYCL< float, float, int32_t >	(	sycl::queue &	queue,
		float *	columns,
		int	in_channels,
		int32_t	begin_idx,
		int32_t	end_idx,
		int32_t	num_out,
		const float *const	out_positions,
		int32_t	num_inp,
		const float *const	inp_positions,
		const float *const	inp_features,
		const float *const	inp_importance,
		size_t	neighbors_index_size,
		const int32_t *const	neighbors_index,
		const float *const	neighbors_importance,
		const int64_t *const	neighbors_row_splits,
		const float *const	extents,
		const float *const	offsets,
		const std::vector< int > &	filter_dims,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		const std::vector< sycl::event > &	deps
	)

FillColumnTranspose() [1/2]

template<class TFeat , class TReal , class TIndex >

void open3d::ml::impl::FillColumnTranspose	(	const cudaStream_t &	stream,
		TFeat *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		const TReal *const __restrict__	out_positions,
		TIndex	num_inp,
		const TReal *const __restrict__	inp_positions,
		const TFeat *const __restrict__	inp_features,
		const TFeat *const __restrict__	inp_neighbors_importance_sum,
		const int64_t *const __restrict__	inp_neighbors_prefix_sum,
		size_t	neighbors_index_size,
		const TIndex *const __restrict__	neighbors_index,
		const TFeat *const __restrict__	neighbors_importance,
		const int64_t *const __restrict__	neighbors_row_splits,
		const TReal *const __restrict__	extents,
		const TReal *const __restrict__	offsets,
		const std::vector< int > &	filter_dims,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize
	)

FillColumnTranspose() [2/2]

template<class TReal , class TIndex , class TKernelIndex >

void open3d::ml::impl::FillColumnTranspose	(	const cudaStream_t &	stream,
		TReal *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		TIndex	num_inp,
		const TReal *const __restrict__	inp_features,
		const TReal *const __restrict__	inp_neighbors_importance_sum,
		const int64_t *const __restrict__	inp_neighbors_prefix_sum,
		size_t	neighbors_index_size,
		const TIndex *const __restrict__	neighbors_index,
		const TKernelIndex *const __restrict__	neighbors_kernel_index,
		const TReal *const __restrict__	neighbors_importance,
		const int64_t *const __restrict__	neighbors_row_splits,
		const int	num_kernel_elements,
		bool	normalize
	)

Similar as FillColumn but used in the transpose convolution to create the patch matrix.

FillColumnTransposeSYCL() [1/2]

template<class TFeat , class TReal , class TIndex >

sycl::event open3d::ml::impl::FillColumnTransposeSYCL	(	sycl::queue &	queue,
		TFeat *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		const TReal *const	out_positions,
		TIndex	num_inp,
		const TReal *const	inp_positions,
		const TFeat *const	inp_features,
		const TFeat *const	inp_neighbors_importance_sum,
		const int64_t *const	inp_neighbors_prefix_sum,
		size_t	neighbors_index_size,
		const TIndex *const	neighbors_index,
		const TFeat *const	neighbors_importance,
		const int64_t *const	neighbors_row_splits,
		const TReal *const	extents,
		const TReal *const	offsets,
		const std::vector< int > &	filter_dims,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		const std::vector< sycl::event > &	deps = {}
	)

SYCL port of FillColumnTranspose. See ContinuousConvCUDAKernels.h for the full parameter documentation; semantics are identical. Depends on deps and returns the completion event of its last kernel.

FillColumnTransposeSYCL() [2/2]

template<class TReal , class TIndex , class TKernelIndex >

sycl::event open3d::ml::impl::FillColumnTransposeSYCL	(	sycl::queue &	queue,
		TReal *	columns,
		int	in_channels,
		TIndex	begin_idx,
		TIndex	end_idx,
		TIndex	num_out,
		TIndex	num_inp,
		const TReal *const	inp_features,
		const TReal *const	inp_neighbors_importance_sum,
		const int64_t *const	inp_neighbors_prefix_sum,
		size_t	neighbors_index_size,
		const TIndex *const	neighbors_index,
		const TKernelIndex *const	neighbors_kernel_index,
		const TReal *const	neighbors_importance,
		const int64_t *const	neighbors_row_splits,
		const int	num_kernel_elements,
		bool	normalize,
		const std::vector< sycl::event > &	deps = {}
	)

SYCL port of FillColumnTranspose: same as FillColumnSYCL but accumulates (+=) into the patch matrix and normalizes by the input point's neighbor count/importance sum instead of the output point's, matching the transpose-convolution's scatter-style column construction. Depends on deps and returns the completion event of its last kernel.

FillColumnTransposeSYCL< float, float, int32_t >()

template sycl::event open3d::ml::impl::FillColumnTransposeSYCL< float, float, int32_t >	(	sycl::queue &	queue,
		float *	columns,
		int	in_channels,
		int32_t	begin_idx,
		int32_t	end_idx,
		int32_t	num_out,
		const float *const	out_positions,
		int32_t	num_inp,
		const float *const	inp_positions,
		const float *const	inp_features,
		const float *const	inp_neighbors_importance_sum,
		const int64_t *const	inp_neighbors_prefix_sum,
		size_t	neighbors_index_size,
		const int32_t *const	neighbors_index,
		const float *const	neighbors_importance,
		const int64_t *const	neighbors_row_splits,
		const float *const	extents,
		const float *const	offsets,
		const std::vector< int > &	filter_dims,
		InterpolationMode	interpolation,
		CoordinateMapping	coordinate_mapping,
		bool	align_corners,
		bool	individual_extent,
		bool	isotropic_extent,
		bool	normalize,
		const std::vector< sycl::event > &	deps
	)

FillNeighborsIndexAndAttributesSYCL()

template<class TIndex , class TAttr >

sycl::event open3d::ml::impl::FillNeighborsIndexAndAttributesSYCL	(	sycl::queue &	queue,
		uint32_t *	count,
		size_t	count_size,
		TIndex *	out_neighbors_index,
		TAttr *	out_neighbors_attributes,
		const TIndex *	inp_neighbors_index,
		const TAttr *	inp_neighbors_attributes,
		int	num_attributes_per_neighbor,
		size_t	index_size,
		const int64_t *	inp_neighbors_row_splits,
		size_t	inp_num_queries,
		const int64_t *	out_neighbors_row_splits,
		size_t	out_num_queries,
		const std::vector< sycl::event > &	deps = {}
	)

Fill output index and attribute arrays for the inverted list. Uses atomic_ref to claim positions within each output row.

GemmColumnMajorCUDA()

template<class LayoutA = cutlass::layout::ColumnMajor, class LayoutB = cutlass::layout::ColumnMajor>

void open3d::ml::impl::GemmColumnMajorCUDA	(	const cudaStream_t &	stream,
		int	m,
		int	n,
		int	k,
		float	alpha,
		const float *	A,
		int	lda,
		const float *	B,
		int	ldb,
		float	beta,
		float *	C,
		int	ldc
	)

Runs alpha * A * B + beta * C on CUDA via CUTLASS device::Gemm. A is (M x K, LayoutA), B is (K x N, LayoutB), C/D is (M x N, ColumnMajor). C and D are the same buffer (in-place accumulation).

GemmColumnMajorSYCL()

template<class LayoutA , class LayoutB >

sycl::event open3d::ml::impl::GemmColumnMajorSYCL	(	sycl::queue &	queue,
		int	m,
		int	n,
		int	k,
		float	alpha,
		const float *	A,
		int64_t	lda,
		const float *	B,
		int64_t	ldb,
		float	beta,
		float *	C,
		int64_t	ldc,
		bool	allow_tf32 = false,
		const std::vector< sycl::event > &	deps = {}
	)

Column-major GEMM matching the CUDA conv ops' calling convention: C = alpha * A * B + beta * C with A (M x K), B (K x N) and C (M x N) all column-major with the given leading dimensions; C is both read and accumulated into in place.

LayoutA / LayoutB select whether the corresponding operand is stored column-major (cutlass::layout::ColumnMajor) or row-major (cutlass::layout::RowMajor). Open3D's conv kernels use <ColumnMajor, ColumnMajor> for the forward ops (see ContinuousConv.cuh) and <ColumnMajor, RowMajor> for the backprop-filter ops (see ContinuousConvBackpropFilter.cuh); only those two combinations are instantiated, so any other combination is a link error rather than a silent fallback.

allow_tf32 selects the Intel XMX TF32 path; false selects the device-agnostic IEEE float32 path (also used as a fallback if the requested shape cannot use the TF32 path – see GemmSYCL.cpp). deps lists events the GEMM must wait on before reading its operands.

Does NOT block: returns a completion event the caller must depend on before reading/writing any of A/B/C, since sycl-tla's run() exposes no per-kernel event (see GemmSYCL.cpp for why the returned event is a barrier over all prior submits on queue, not a tighter one).

Throws std::runtime_error if no supported tile configuration can implement the requested problem shape.

// Accumulate a chunk's contribution into filter_backprop (B row-major):
sycl::event gemm_event =
        GemmColumnMajorSYCL<cutlass::layout::ColumnMajor,
                            cutlass::layout::RowMajor>(
                queue, out_channels, spatial_filter_size * in_channels,
                num_cols_this_run, /*alpha*/ 1.f, out_features_gradient,
                lda, columns, ldb, /*beta*/ 1.f, filter_backprop, ldc,
                allow_tf32, {fill_column_event});

GemmColumnMajorSYCL< cutlass::layout::ColumnMajor, cutlass::layout::ColumnMajor >()

template sycl::event open3d::ml::impl::GemmColumnMajorSYCL< cutlass::layout::ColumnMajor, cutlass::layout::ColumnMajor >	(	sycl::queue &	,
		int	,
		int	,
		int	,
		float	,
		const float *	,
		int64_t	,
		const float *	,
		int64_t	,
		float	,
		float *	,
		int64_t	,
		bool	,
		const std::vector< sycl::event > &
	)

GemmColumnMajorSYCL< cutlass::layout::ColumnMajor, cutlass::layout::RowMajor >()

template sycl::event open3d::ml::impl::GemmColumnMajorSYCL< cutlass::layout::ColumnMajor, cutlass::layout::RowMajor >	(	sycl::queue &	,
		int	,
		int	,
		int	,
		float	,
		const float *	,
		int64_t	,
		const float *	,
		int64_t	,
		float	,
		float *	,
		int64_t	,
		bool	,
		const std::vector< sycl::event > &
	)

InterpolateSYCL()

template<InterpolationMode INTERPOLATION, class T >

void open3d::ml::impl::InterpolateSYCL ( T *  w, int *  idx, const T &  x, const T &  y, const T &  z, const int &  filter_size_x, const int &  filter_size_y, const int &  filter_size_z, int  num_channels = 1  )

inline

Computes interpolation weights and indices. See CoordinateTransformation .cuh for the full parameter documentation; semantics are identical (SYCL port).

InvertNeighborsListCPU()

template<class TIndex , class TAttr >

void open3d::ml::impl::InvertNeighborsListCPU	(	const TIndex *const	inp_neighbors_index,
		const TAttr *const	inp_neighbors_attributes,
		const int	num_attributes_per_neighbor,
		const int64_t *const	inp_neighbors_row_splits,
		const size_t	inp_num_queries,
		TIndex *	out_neighbors_index,
		TAttr *	out_neighbors_attributes,
		const size_t	index_size,
		int64_t *	out_neighbors_row_splits,
		const size_t	out_num_queries
	)

Inverts a neighbors list, which is a tuple of the form (neighbors_index, neighbors_row_splits, neighbors_attributes). neighbors_index is a nested list of indices to the neighbors. Each entry defines an edge between two indices (points). The neighbors_row_splits defines the start and end of each sublist. neighbors_attributes is an optional array of attributes for each entry in neighbors_index.

Example: The neighbors for point cloud A (3 points) in point cloud B (2 points) is defined by:

• neighbors_index [0 1 0 0]
• neighbors_row_splits [0 2 3 4]
• optional neighbors_attributes [0.1 0.2 0.3 0.4] (1 scalar attribute)

The inverted neighbors list is then the neighbors for point cloud B in A

• neighbors_index [0 1 2 0]
• neighbors_row_splits [0 3 4]
• optional neighbors_attributes [0.1 0.3 0.4 0.2]

Parameters

inp_neighbors_index	The nested list of neighbor indices.
inp_neighbors_attributes	The array of attributes for each entry in inp_neighbors_index. This is optional and can be set to null.
num_attributes_per_neighbor	The number of scalar attributes for each entry in inp_neighbors_index.
inp_neighbors_row_splits	Defines the start and end of the sublists in inp_neighbors_index. This is an exclusive prefix sum with 0 as the first element and the length of inp_neighbors_index as last element. The size is inp_num_queries + 1
inp_num_queries	The number of queries.
out_neighbors_index	The inverted neighbors_index list with the same size as inp_neighbors_index .
out_neighbors_attributes	The inverted array of attributes with the same size as inp_neighbors_attributes .
index_size	This is the size of inp_neighbors_index and out_neighbors_index, both have the same size.
out_neighbors_row_splits	The prefix sum which defines the start and end of the sublists in out_neighbors_index.
out_num_queries	The number of queries with respect to the inverted neighbors list.

InvertNeighborsListSYCL()

template<class TIndex , class TAttr >

sycl::event open3d::ml::impl::InvertNeighborsListSYCL	(	sycl::queue &	queue,
		uint32_t *	count_buf,
		const TIndex *	inp_neighbors_index,
		const TAttr *	inp_neighbors_attributes,
		int	num_attributes_per_neighbor,
		const int64_t *	inp_neighbors_row_splits,
		size_t	inp_num_queries,
		TIndex *	out_neighbors_index,
		TAttr *	out_neighbors_attributes,
		size_t	index_size,
		int64_t *	out_neighbors_row_splits,
		size_t	out_num_queries
	)

Inverts a neighbors list.

All pointer arguments must point to SYCL device-accessible memory (USM or PyTorch XPU tensors).

count_buf must be device memory of at least out_num_queries uint32_t elements (used as a scratch buffer, zero-initialized internally).

Returns

The completion event of the last stage; the caller only needs to wait on it if a host-visible completion signal is required (e.g. before returning a torch::Tensor built from the output buffers to Python, PyTorch's own queue-based dependency tracking makes an explicit wait unnecessary).

MapCylinderToCube()

template<class T , int VECSIZE>

void open3d::ml::impl::MapCylinderToCube ( Eigen::Array< T, VECSIZE, 1 > &  x, Eigen::Array< T, VECSIZE, 1 > &  y, Eigen::Array< T, VECSIZE, 1 > &  z  )

inline

Maps coordinates in a cylinder with radius 1 to a cube. The input and output range of the coordinates is [-1,1]. The cylinder axis is along z.

MapCylinderToCubeSYCL()

template<class T >

void open3d::ml::impl::MapCylinderToCubeSYCL ( T &  x, T &  y, T &  z  )

inline

Maps coordinates in a cylinder with radius 1 to a cube. The input and output range of the coordinates is [-1,1]. The cylinder axis is along z.

MapSphereToCylinder()

template<class T , int VECSIZE>

void open3d::ml::impl::MapSphereToCylinder ( Eigen::Array< T, VECSIZE, 1 > &  x, Eigen::Array< T, VECSIZE, 1 > &  y, Eigen::Array< T, VECSIZE, 1 > &  z  )

inline

Maps coordinates in a sphere with radius 1 to a cylinder. The input and output range of the coordinates is [-1,1]. The cylinder axis is along z.

MapSphereToCylinderSYCL()

template<class T >

void open3d::ml::impl::MapSphereToCylinderSYCL ( T &  x, T &  y, T &  z  )

inline

Maps coordinates in a sphere with radius 1 to a cylinder. The input and output range of the coordinates is [-1,1]. The cylinder axis is along z.

MultiplyAndCopyColumns()

template<class T >

void open3d::ml::impl::MultiplyAndCopyColumns	(	const cudaStream_t &	stream,
		size_t	rows,
		size_t	cols,
		T *__restrict__	out_ptr,
		const T *const __restrict__	col_major_matrix,
		const T *const __restrict__	vector
	)

Multiplies each column with a scalar.

Parameters

out_ptr	Output pointer
col_major_matrix	Matrix with shape [rows, cols] in column major storage order.
vector	A vector with shape [cols].

MultiplyAndCopyColumnsSYCL()

template<class T >

sycl::event open3d::ml::impl::MultiplyAndCopyColumnsSYCL	(	sycl::queue &	queue,
		size_t	rows,
		size_t	cols,
		T *	out_ptr,
		const T *const	col_major_matrix,
		const T *const	vector,
		const std::vector< sycl::event > &	deps = {}
	)

Same as MultiplyColumnsSYCL but writes the result to out_ptr instead of modifying col_major_matrix in place. SYCL port of MultiplyAndCopyColumns. Depends on deps and returns the completion event.

MultiplyAndCopyColumnsSYCL< float >()

template sycl::event open3d::ml::impl::MultiplyAndCopyColumnsSYCL< float >	(	sycl::queue &	queue,
		size_t	rows,
		size_t	cols,
		float *	out_ptr,
		const float *const	col_major_matrix,
		const float *const	vector,
		const std::vector< sycl::event > &	deps
	)

MultiplyColumns()

template<class T >

void open3d::ml::impl::MultiplyColumns	(	const cudaStream_t &	stream,
		size_t	rows,
		size_t	cols,
		T *__restrict__	col_major_matrix,
		const T *const __restrict__	vector
	)

Multiplies each column with a scalar in-place.

Parameters

col_major_matrix	Matrix with shape [rows, cols] in column major storage order.
vector	A vector with shape [cols].

MultiplyColumnsSYCL()

template<class T >

sycl::event open3d::ml::impl::MultiplyColumnsSYCL	(	sycl::queue &	queue,
		size_t	rows,
		size_t	cols,
		T *	col_major_matrix,
		const T *const	vector,
		const std::vector< sycl::event > &	deps = {}
	)

Multiplies each column of a column-major [rows, cols] matrix in place by the corresponding entry of vector (length cols). SYCL port of MultiplyColumns. Depends on deps and returns the completion event.

MultiplyColumnsSYCL< float >()

template sycl::event open3d::ml::impl::MultiplyColumnsSYCL< float >	(	sycl::queue &	queue,
		size_t	rows,
		size_t	cols,
		float *	col_major_matrix,
		const float *const	vector,
		const std::vector< sycl::event > &	deps
	)

RaggedToDenseCPU()

template<class T >

void open3d::ml::impl::RaggedToDenseCPU	(	const T *const	values,
		const int64_t *const	row_splits,
		const size_t	row_splits_size,
		const size_t	out_col_size,
		const T *const	default_value,
		const size_t	default_value_size,
		T *	out_values
	)

Creates a dense tensor from a ragged tensor.

Example where each value has size 2: values = [[0,0],[1,1],[2,2],[3,3],[4,4]] row_splits = [0,2,5] out_col_size=3 default_value=[-1,-1] default_value_size = 2

will return

out_values = [[[0,0],[1,1],[-1,-1]], [[2,2],[3,3],[4,4]]]

Parameters

values	Linear memory with all values.
row_splits	Defines the start and end of each entry in the ragged tensor. This is an exclusive prefix sum with 0 as the first element and the length of all values as the last element.
row_splits_size	The length of the row_splits vector.
out_col_size	The output column size. This is the second dim of the dense output tensor.
default_value	The default value to use if there are not enough values for filling the row.
default_value_size	The size of the default value.
out_values	This is the output array. The size of the array must be [row_splits_size-1, out_col_size, default_value_size].

RaggedToDenseSYCL()

template<class T >

void open3d::ml::impl::RaggedToDenseSYCL	(	sycl::queue &	queue,
		const T *const	values,
		const int64_t *const	row_splits,
		const size_t	row_splits_size,
		const size_t	out_col_size,
		const T *const	default_value,
		const size_t	default_value_size,
		T *	out_values
	)

SYCL kernel body for RaggedToDense. Each work-item handles one row of the ragged tensor.

ReduceSubarraysSumCPU()

template<class T >

void open3d::ml::impl::ReduceSubarraysSumCPU	(	const T *const	values,
		const size_t	values_size,
		const int64_t *const	row_splits,
		const size_t	num_arrays,
		T *	out_sums
	)

Reduces subarrays in linear memory with the sum operation. The sum for empty subarrays is 0.

Parameters

values	The linear array with all values
values_size	Number of elements of values
row_splits	Defines the start and end of each subarray. This is an exclusive prefix sum with 0 as the first element and the length of values as last element. The size is num_arrays + 1
num_arrays	The number of subarrays
out_sums	The preallocated output array with size num_arrays

ReduceSubarraysSumSYCL()

template<class T >

void open3d::ml::impl::ReduceSubarraysSumSYCL	(	sycl::queue &	queue,
		const T *const	values,
		const size_t	values_size,
		const int64_t *const	row_splits,
		const size_t	num_arrays,
		T *	out_sums
	)

Each work-group i sums values[row_splits[i]..row_splits[i+1]) into out_sums[i].

SparseConvBackpropFilterCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvBackpropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TFeat *	out_features_gradient,
		bool	normalize
	)

Computes the backprop for the filter of a sparse convolution.

Parameters

temp	Pointer to temporary memory. If nullptr then the required size of temporary memory will be written to temp_size and no work is done. This function can make use of more memory and returns the maximum size that can be used in max_temp_size.
temp_size	The size of the temporary memory in bytes. This is used as an output if temp is nullptr and returns the minimum temp size required.
max_temp_size	This is used as an output if temp is nullptr and returns the maximum temp size that can be used.
texture_alignment	The texture alignment in bytes. This is used for allocating segments within the temporary memory.
filter_backrop	Output array for the computed filter gradient with shape [depth,height,width, inp channels, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be >=3. The order is [num kernel elements, inp channels, out channels].
num_out	The number of output points.
num_inp	The number of input points.
inp_features	The input features with shape [num_inp, in_channels].
inp_importance	Optional importance for each input point with shape [num_inp]. Set to null to disable.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_kernel_index	Defines which kernel element to use for each neighbor. This array has the same length as neighbors_index.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
out_features_gradient	The gradient signal from the features.
normalize	If true then the output features are normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

SparseConvBackpropFilterSYCL()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvBackpropFilterSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		TIndex	num_out,
		TIndex	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TFeat *	out_features_gradient,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of SparseConvBackpropFilterCUDA. See SparseConvBackpropFilter .cuh for the full parameter documentation; semantics are identical.

SparseConvComputeFeaturesCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		bool	normalize
	)

Computes the output features of a sparse convolution.

Parameters

out_features	Output array for the computed features with shape [num_out, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be >=3. The order is [num kernel elements, inp channels, out channels].
filter	Pointer to the filter values.
num_out	The number of output points.
num_inp	The number of input points.
inp_features	The input features with shape [num_inp, in_channels].
inp_importance	Optional importance for each input point with shape [num_inp]. Set to null to disable.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_kernel_index	Defines which kernel element to use for each neighbor. This array has the same length as neighbors_index.
neighbors_importance	Array of the same length as neighbors_importance. Defines which of the kernel elements to use in the matrix multiplication.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

SparseConvComputeFeaturesSYCL()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvComputeFeaturesSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		TIndex	num_out,
		TIndex	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_importance,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of SparseConvComputeFeaturesCUDA. See SparseConv.cuh for the full parameter documentation; semantics are identical.

SparseConvTransposeBackpropFilterCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvTransposeBackpropFilterCPU	(	TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		size_t	num_out,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TFeat *	out_features_gradient,
		bool	normalize
	)

Computes the backprop for the filter of a transpose sparse convolution.

Parameters

filter_backrop	Output array for the computed filter gradient with shape [depth,height,width, inp channels, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be >=3. The order is [num kernel elements, inp channels, out channels].
filter	Pointer to the filter values.
num_out	The number of output points.
out_importance	Optional importance for each output point with shape [num_out]. Set to null to disable.
num_inp	The number of input points.
inp_features	The input features with shape [num_inp, in_channels].
inp_neighbors_importance_sum	The sum of the neighbors_importance values for each input with shape [num_inp].
inp_neighbors_row_splits	The prefix sum which defines the start and end of the sublists in inp_neighbors_index. The size of the array is num_inp + 1.
neighbors_index_size	The size of the neighbors_index array.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_kernel_index	Defines which kernel element to use for each neighbor. This array has the same length as neighbors_index.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
out_features_gradient	The gradient from the features with shape [num_out, out_channels]
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

SparseConvTransposeBackpropFilterSYCL()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvTransposeBackpropFilterSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	filter_backprop,
		const std::vector< int > &	filter_dims,
		TIndex	num_out,
		const TFeat *	out_importance,
		TIndex	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		const TFeat *	out_features_gradient,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of SparseConvTransposeBackpropFilterCUDA. See SparseConvTransposeBackpropFilter.cuh for the full parameter documentation; semantics are identical.

SparseConvTransposeComputeFeaturesCPU()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvTransposeComputeFeaturesCPU	(	TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		size_t	num_out,
		const TFeat *	out_importance,
		size_t	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_row_splits,
		const TIndex *	neighbor_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbor_importance,
		const int64_t *	neighbors_row_splits,
		bool	normalize
	)

Computes the output features of a transpose sparse convolution.

Parameters

out_features	Output array for the computed features with shape [num_out, out channels]
filter_dims	The sizes of the filter dimensions. The size of filter_dims must be >=3. The order is [num kernel elements, inp channels, out channels].
filter	Pointer to the filter values.
num_out	The number of output points.
out_importance	Optional importance for each output point with shape [num_out]. Set to null to disable.
num_inp	The number of input points.
inp_features	The input features with shape [num_inp, in_channels].
inp_neighbors_importance_sum	The sum of the neighbors_importance values for each input with shape [num_inp].
inp_neighbors_row_splits	The prefix sum which defines the start and end of the sublists in inp_neighbors_index. The size of the array is num_inp + 1.
neighbors_index	The array with lists of neighbors for each output point. The start and end of each sublist is defined by neighbors_row_splits.
neighbors_kernel_index	Defines which kernel element to use for each neighbor. This array has the same length as neighbors_index.
neighbors_importance	Optional importance for each entry in neighbors_index. Set to null to disable.
neighbors_row_splits	The prefix sum which defines the start and end of the sublists in neighbors_index. The size of the array is num_out + 1.
normalize	If true then the result is normalized either by the number of points (neighbors_importance is null) or by the sum of the respective values in neighbors_importance.

SparseConvTransposeComputeFeaturesSYCL()

template<class TFeat , class TOut , class TIndex , class TKernelIndex >

void open3d::ml::impl::SparseConvTransposeComputeFeaturesSYCL	(	sycl::queue &	queue,
		void *	temp,
		size_t &	temp_size,
		size_t &	max_temp_size,
		int	texture_alignment,
		TOut *	out_features,
		const std::vector< int > &	filter_dims,
		const TFeat *	filter,
		TIndex	num_out,
		const TFeat *	out_importance,
		TIndex	num_inp,
		const TFeat *	inp_features,
		const TFeat *	inp_neighbors_importance_sum,
		const int64_t *	inp_neighbors_prefix_sum,
		size_t	neighbors_index_size,
		const TIndex *	neighbors_index,
		const TKernelIndex *	neighbors_kernel_index,
		const TFeat *	neighbors_importance,
		const int64_t *	neighbors_row_splits,
		bool	normalize,
		bool	allow_tf32
	)

SYCL version of SparseConvTransposeComputeFeaturesCUDA. See SparseConvTranspose.cuh for the full parameter documentation; semantics are identical.

SpatialHash() [1/2]

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

inline

SpatialHash() [2/2]

HOST_DEVICE size_t open3d::ml::impl::SpatialHash ( int  x, int  y, int  z  )

inline

Spatial hashing function for integer coordinates.

VoxelizeCPU()

template<class T , int NDIM, class OUTPUT_ALLOCATOR >

void open3d::ml::impl::VoxelizeCPU	(	const size_t	num_points,
		const T *const	points,
		const size_t	batch_size,
		const int64_t *const	row_splits,
		const T *const	voxel_size,
		const T *const	points_range_min,
		const T *const	points_range_max,
		const int64_t	max_points_per_voxel,
		const int64_t	max_voxels,
		OUTPUT_ALLOCATOR &	output_allocator
	)

This function voxelizes a point cloud. The function returns the integer coordinates of the voxels that contain points and a compact list of the indices that associate the voxels to the points.

Template Parameters

T	Floating-point data type for the point positions.
NDIM	The number of dimensions of the points.
OUTPUT_ALLOCATOR	Type of the output_allocator. See output_allocator for more information.

Parameters

num_points	The number of points.
points	Array with the point positions. The shape is [num_points,NDIM].
batch_size	The batch size of points.
row_splits	row_splits for defining batches.
voxel_size	The edge lengths of the voxel. The shape is [NDIM]. This pointer points to host memory!
points_range_min	The lower bound of the domain to be voxelized. The shape is [NDIM]. This pointer points to host memory!
points_range_max	The upper bound of the domain to be voxelized. The shape is [NDIM]. This pointer points to host memory!
max_points_per_voxel	This parameter limits the number of points that are recorderd for each voxel.
max_voxels	This parameter limits the number of voxels that will be generated.
output_allocator	An object that implements functions for allocating the output arrays. The object must implement functions AllocVoxelCoords(int32_t** ptr, int64_t rows, int64_t cols), AllocVoxelPointIndices(int64_t** ptr, int64_t size), AllocVoxelPointRowSplits(int64_t** ptr, int64_t size) and AllocVoxelBatchSplits(int64_t** ptr, int64_t size). All functions should allocate memory and return a pointer to that memory in ptr. The arguments size, rows, and cols define the size of the array as the number of elements. All functions must accept zero size arguments. In this case ptr does not need to be set.

VoxelizeSYCL()

template<class T , int NDIM, class OUTPUT_ALLOCATOR >

void open3d::ml::impl::VoxelizeSYCL	(	sycl::queue &	queue,
		size_t	num_points,
		const T *const	points,
		const size_t	batch_size,
		const int64_t *const	row_splits,
		const T *const	voxel_size,
		const T *const	points_range_min,
		const T *const	points_range_max,
		const int64_t	max_points_per_voxel,
		const int64_t	max_voxels,
		OUTPUT_ALLOCATOR &	output_allocator
	)

Voxelizes a point cloud (SYCL). See Voxelize.cuh for the full parameter documentation; the signature matches the CUDA version except that this function does the work in one call (no temp-size query pass — see the file-level comment on design differences) and takes a SYCL queue directly (kernels run on the caller-supplied queue).

voxel_size, points_range_min, points_range_max point to host memory (matching the CUDA convention; the PyTorch dispatch layer already keeps these small per-dimension arrays on the CPU).

VoxelPooling()

template<class TReal , class TFeat , class OUTPUT_ALLOCATOR >

void open3d::ml::impl::VoxelPooling	(	size_t	num_inp,
		const TReal *const	inp_positions,
		int	in_channels,
		const TFeat *	inp_features,
		TReal	voxel_size,
		OUTPUT_ALLOCATOR &	output_allocator,
		AccumulationFn	position_fn,
		AccumulationFn	feature_fn
	)

Pooling operation for point clouds. Aggregates points that are inside the same voxel.

Template Parameters

TReal	Scalar type for point positions.
TFeat	Scalar type for the features.
OUTPUT_ALLOCATOR	Type of the output_allocator. See output_allocator for more information.

Parameters

num_inp	The number of input points.
inp_positions	Array with 3D point positions.
in_channels	The number of input feature channels.
inp_features	The array with the point features. The shape is [num_inp, in_channels].
voxel_size	The voxel size (voxel edge length) used for pooling.
output_allocator	An object that implements functions for allocating the output arrays. The object must implement functions AllocPooledPositions(TReal** ptr, size_t num) and AllocPooledFeatures(TFeat** ptr, size_t num, channels), with 'num' as the number of output points and 'channels' as in_channels . Both functions should allocate memory and return a pointer to that memory in ptr. Both functions must accept the argument num==0. In this case ptr does not need to be set. Note that AllocPooledPositions must allocate memory for num*3*sizeof(TReal) bytes.
position_fn	Defines how the new point positions will be computed. AVERAGE computes the center of gravity for the points within one voxel. NEAREST_NEIGHBOR selects the point closest to the voxel center. CENTER uses the voxel center for the position of the generated point.
feature_fn	Defines how the new features will be computed. AVERAGE computes the average feature vector. NEAREST_NEIGHBOR selects the feature vector of the point closest to the voxel center. MAX uses the maximum feature among all points within the voxel.

VoxelPoolingBackprop()

template<class TReal , class TFeat >

void open3d::ml::impl::VoxelPoolingBackprop	(	TFeat *	features_backprop,
		size_t	num_inp,
		const TReal *const	inp_positions,
		int	in_channels,
		const TFeat *const	inp_features,
		size_t	num_pooled,
		const TReal *const	pooled_positions,
		const TFeat *const	pooled_features_gradient,
		TReal	voxel_size,
		AccumulationFn	position_fn,
		AccumulationFn	feature_fn
	)

Backpropagation to the features for VoxelPooling.

Parameters

features_backprop	The output array with the gradients for the features.
num_pooled	The number of points after pooling with VoxelPooling.
pooled_positions	Array with the 3D positions after pooling.
pooled_features_gradient	Array with the point features after pooling.

See VoxelPooling for the description of the remaining parameters.