open3d::ml::contrib Namespace Reference

Data Structures
struct	Point
struct	PointCloud
class	PointXYZ
class	SampledData

Functions
void	BallQuerySYCL (sycl::queue &queue, int b, int n, int m, float radius, int nsample, const float *const new_xyz, const float *const xyz, int *const idx)
PointXYZ	max_point (std::vector< PointXYZ > points)
PointXYZ	min_point (std::vector< PointXYZ > points)
PointXYZ	operator+ (const PointXYZ A, const PointXYZ B)
PointXYZ	operator- (const PointXYZ A, const PointXYZ B)
PointXYZ	operator* (const PointXYZ P, const float a)
PointXYZ	operator* (const float a, const PointXYZ P)
std::ostream &	operator<< (std::ostream &os, const PointXYZ P)
bool	operator== (const PointXYZ A, const PointXYZ B)
int	OptNumThreads (int work_size)
dim3	OptBlockConfig (int x, int y)
void	grid_subsampling (std::vector< PointXYZ > &original_points, std::vector< PointXYZ > &subsampled_points, std::vector< float > &original_features, std::vector< float > &subsampled_features, std::vector< int > &original_classes, std::vector< int > &subsampled_classes, float sampleDl, int verbose)
void	batch_grid_subsampling (std::vector< PointXYZ > &original_points, std::vector< PointXYZ > &subsampled_points, std::vector< float > &original_features, std::vector< float > &subsampled_features, std::vector< int > &original_classes, std::vector< int > &subsampled_classes, std::vector< int > &original_batches, std::vector< int > &subsampled_batches, float sampleDl, int max_p)
void	ThreeNNSYCL (sycl::queue &queue, int b, int n, int m, const float *const unknown, const float *const known, float *const dist2, int *const idx)
void	ThreeInterpolateSYCL (sycl::queue &queue, int b, int c, int m, int n, const float *const points, const int *const idx, const float *const weight, float *const out)
void	ThreeInterpolateGradSYCL (sycl::queue &queue, int b, int c, int n, int m, const float *const grad_out, const int *const idx, const float *const weight, float *const grad_points)
void	IoUBevCPUKernel (const float *boxes_a, const float *boxes_b, float *iou, int num_a, int num_b)
void	IoU3dCPUKernel (const float *boxes_a, const float *boxes_b, float *iou, int num_a, int num_b)
OPEN3D_HOST_DEVICE float	Cross (const Point &a, const Point &b)
OPEN3D_HOST_DEVICE float	Cross (const Point &p1, const Point &p2, const Point &p0)
OPEN3D_HOST_DEVICE int	CheckRectCross (const Point &p1, const Point &p2, const Point &q1, const Point &q2)
OPEN3D_HOST_DEVICE int	CheckInBox2D (const float *box, const Point &p)
OPEN3D_HOST_DEVICE int	Intersection (const Point &p1, const Point &p0, const Point &q1, const Point &q0, Point &ans)
OPEN3D_HOST_DEVICE void	RotateAroundCenter (const Point &center, const float angle_cos, const float angle_sin, Point &p)
OPEN3D_HOST_DEVICE int	PointCmp (const Point &a, const Point &b, const Point &center)
OPEN3D_HOST_DEVICE float	BoxOverlap (const float *box_a, const float *box_b)
OPEN3D_HOST_DEVICE float	IoUBev2DWithMinAndMax (const float *box_a, const float *box_b, bool intersection_only=false)
	(x_min, z_min, x_max, z_max, y_rotate)
OPEN3D_HOST_DEVICE float	IoUBev2DWithCenterAndSize (const float *box_a, const float *box_b, bool intersection_only=false)
	(x_center, z_center, x_size, z_size, y_rotate)
OPEN3D_HOST_DEVICE float	IoU3DWithCenterAndSize (const float *box_a, const float *box_b)
	(x_center, y_max, z_center, x_size, y_size, z_size, y_rotate)
void	IoUBevSYCLKernel (sycl::queue &queue, const float *boxes_a, const float *boxes_b, float *iou, int num_a, int num_b)
void	IoU3dSYCLKernel (sycl::queue &queue, const float *boxes_a, const float *boxes_b, float *iou, int num_a, int num_b)
void	IoUBevSYCLKernel (const core::Device &device, const float *boxes_a, const float *boxes_b, float *iou, int num_a, int num_b)
void	IoU3dSYCLKernel (const core::Device &device, const float *boxes_a, const float *boxes_b, float *iou, int num_a, int num_b)
std::vector< int64_t >	NmsCPUKernel (const float *boxes, const float *scores, int n, double nms_overlap_thresh)
template<typename T >
std::vector< int64_t >	SortIndexesDescending (const T *values, int64_t num)
OPEN3D_HOST_DEVICE int	NmsGreedyKeepCore (const uint64_t *mask, const int64_t *sort_indices, int n, uint64_t *remv, int64_t *keep_indices)
std::vector< int64_t >	NmsGreedyKeep (const uint64_t *mask, const int64_t *sort_indices, int n)
int	NmsSYCLKernel (sycl::queue &queue, const float *boxes, const float *scores, int n, double nms_overlap_thresh, int64_t *keep_indices_out)
void	roipool3dLauncherCPU (int batch_size, int pts_num, int boxes_num, int feature_in_len, int sampled_pts_num, const float *xyz, const float *boxes3d, const float *pts_feature, float *pooled_features, int *pooled_empty_flag)
OPEN3D_HOST_DEVICE bool	pt_in_box3d (float x, float y, float z, float cx, float bottom_y, float cz, float h, float w, float l, float angle, float max_dis)
void	roipool3dLauncherSYCL (sycl::queue &queue, int batch_size, int pts_num, int boxes_num, int feature_in_len, int sampled_pts_num, const float *xyz, const float *boxes3d, const float *pts_feature, float *pooled_features, int *pooled_empty_flag)
void	TrilinearDevoxelizeSYCL (sycl::queue &queue, int b, int c, int n, int r, int r2, int r3, bool is_training, const float *const coords, const float *const feat, int *const inds, float *const wgts, float *const outs)
void	TrilinearDevoxelizeGradSYCL (sycl::queue &queue, int b, int c, int n, int r3, const int *const inds, const float *const wgts, const float *const grad_y, float *const grad_x)

Variables
constexpr int	NMS_BLOCK_SIZE = sizeof(uint64_t) * 8
constexpr float	EPS = static_cast<float>(1e-8)

Function Documentation

BallQuerySYCL()

void open3d::ml::contrib::BallQuerySYCL ( sycl::queue &  queue, int  b, int  n, int  m, float  radius, int  nsample, const float *const  new_xyz, const float *const  xyz, int *const  idx  )

inline

Finds up to nsample points within radius of each query point.

Parameters

queue	SYCL queue to run the kernel on.
b	Batch size.
n	Number of candidate points per batch (xyz).
m	Number of query points per batch (new_xyz).
radius	Search radius.
nsample	Max number of neighbors to record per query point.
new_xyz	Query point positions, shape [b, m, 3].
xyz	Candidate point positions, shape [b, n, 3].
idx	Output neighbor indices, shape [b, m, nsample]. Must be zero-initialized by the caller (matches the CUDA/original-SYCL contract of leaving index 0 as the implicit match when no candidate is in range).

batch_grid_subsampling()

void open3d::ml::contrib::batch_grid_subsampling	(	std::vector< PointXYZ > &	original_points,
		std::vector< PointXYZ > &	subsampled_points,
		std::vector< float > &	original_features,
		std::vector< float > &	subsampled_features,
		std::vector< int > &	original_classes,
		std::vector< int > &	subsampled_classes,
		std::vector< int > &	original_batches,
		std::vector< int > &	subsampled_batches,
		float	sampleDl,
		int	max_p
	)

BoxOverlap()

OPEN3D_HOST_DEVICE float open3d::ml::contrib::BoxOverlap ( const float *  box_a, const float *  box_b  )

inline

CheckInBox2D()

OPEN3D_HOST_DEVICE int open3d::ml::contrib::CheckInBox2D ( const float *  box, const Point &  p  )

inline

CheckRectCross()

OPEN3D_HOST_DEVICE int open3d::ml::contrib::CheckRectCross ( const Point &  p1, const Point &  p2, const Point &  q1, const Point &  q2  )

inline

Cross() [1/2]

OPEN3D_HOST_DEVICE float open3d::ml::contrib::Cross ( const Point &  a, const Point &  b  )

inline

Cross() [2/2]

OPEN3D_HOST_DEVICE float open3d::ml::contrib::Cross ( const Point &  p1, const Point &  p2, const Point &  p0  )

inline

grid_subsampling()

void open3d::ml::contrib::grid_subsampling	(	std::vector< PointXYZ > &	original_points,
		std::vector< PointXYZ > &	subsampled_points,
		std::vector< float > &	original_features,
		std::vector< float > &	subsampled_features,
		std::vector< int > &	original_classes,
		std::vector< int > &	subsampled_classes,
		float	sampleDl,
		int	verbose
	)

Intersection()

OPEN3D_HOST_DEVICE int open3d::ml::contrib::Intersection ( const Point &  p1, const Point &  p0, const Point &  q1, const Point &  q0, Point &  ans  )

inline

IoU3dCPUKernel()

void open3d::ml::contrib::IoU3dCPUKernel	(	const float *	boxes_a,
		const float *	boxes_b,
		float *	iou,
		int	num_a,
		int	num_b
	)

Parameters

boxes_a	(num_a, 7) float32.
boxes_b	(num_b, 7) float32.
iou	(num_a, num_b) float32, output iou values.
num_a	Number of boxes in boxes_a.
num_b	Number of boxes in boxes_b.

IoU3dSYCLKernel() [1/2]

void open3d::ml::contrib::IoU3dSYCLKernel	(	const core::Device &	device,
		const float *	boxes_a,
		const float *	boxes_b,
		float *	iou,
		int	num_a,
		int	num_b
	)

IoU3dSYCLKernel() [2/2]

void open3d::ml::contrib::IoU3dSYCLKernel	(	sycl::queue &	queue,
		const float *	boxes_a,
		const float *	boxes_b,
		float *	iou,
		int	num_a,
		int	num_b
	)

IoU3DWithCenterAndSize()

OPEN3D_HOST_DEVICE float open3d::ml::contrib::IoU3DWithCenterAndSize ( const float *  box_a, const float *  box_b  )

inline

(x_center, y_max, z_center, x_size, y_size, z_size, y_rotate)

IoUBev2DWithCenterAndSize()

OPEN3D_HOST_DEVICE float open3d::ml::contrib::IoUBev2DWithCenterAndSize ( const float *  box_a, const float *  box_b, bool  intersection_only = false  )

inline

(x_center, z_center, x_size, z_size, y_rotate)

IoUBev2DWithMinAndMax()

OPEN3D_HOST_DEVICE float open3d::ml::contrib::IoUBev2DWithMinAndMax ( const float *  box_a, const float *  box_b, bool  intersection_only = false  )

inline

(x_min, z_min, x_max, z_max, y_rotate)

IoUBevCPUKernel()

void open3d::ml::contrib::IoUBevCPUKernel	(	const float *	boxes_a,
		const float *	boxes_b,
		float *	iou,
		int	num_a,
		int	num_b
	)

Parameters

boxes_a	(num_a, 5) float32.
boxes_b	(num_b, 5) float32.
iou	(num_a, num_b) float32, output iou values.
num_a	Number of boxes in boxes_a.
num_b	Number of boxes in boxes_b. intersection over union.

IoUBevSYCLKernel() [1/2]

void open3d::ml::contrib::IoUBevSYCLKernel	(	const core::Device &	device,
		const float *	boxes_a,
		const float *	boxes_b,
		float *	iou,
		int	num_a,
		int	num_b
	)

IoUBevSYCLKernel() [2/2]

void open3d::ml::contrib::IoUBevSYCLKernel	(	sycl::queue &	queue,
		const float *	boxes_a,
		const float *	boxes_b,
		float *	iou,
		int	num_a,
		int	num_b
	)

max_point()

PointXYZ open3d::ml::contrib::max_point

(

std::vector< PointXYZ >

points

)

min_point()

PointXYZ open3d::ml::contrib::min_point

(

std::vector< PointXYZ >

points

)

NmsCPUKernel()

std::vector< int64_t > open3d::ml::contrib::NmsCPUKernel	(	const float *	boxes,
		const float *	scores,
		int	n,
		double	nms_overlap_thresh
	)

Parameters

boxes	(n, 5) float32.
scores	(n,) float32.
n	Number of boxes.
nms_overlap_thresh	When a high-score box is selected, other remaining boxes with IoU > nms_overlap_thresh will be discarded.

Returns

Selected box indices to keep.

NmsGreedyKeep()

std::vector< int64_t > open3d::ml::contrib::NmsGreedyKeep ( const uint64_t *  mask, const int64_t *  sort_indices, int  n  )

inline

Host convenience wrapper around NmsGreedyKeepCore() for the CPU backend (Nms.cpp), which already has mask and sort_indices in host memory.

NmsGreedyKeepCore()

OPEN3D_HOST_DEVICE int open3d::ml::contrib::NmsGreedyKeepCore ( const uint64_t *  mask, const int64_t *  sort_indices, int  n, uint64_t *  remv, int64_t *  keep_indices  )

inline

Greedily walks the sorted boxes and keeps a box iff it does not overlap (per mask) with any previously-kept box, writing the kept original box indices (in sorted/score-descending order) into the caller-allocated keep_indices (capacity >= n). remv is caller-allocated scratch with capacity >= num_block_cols = ceil(n / NMS_BLOCK_SIZE).

This is the single reduction shared, byte-for-byte, across the CPU/CUDA/SYCL backends: the loop is inherently sequential (each step depends on all previous ones via remv) and mask is tiny (n x num_block_cols bits), so CUDA/SYCL run it as a single-thread device kernel (see Nms.cu / NmsSYCL.cpp) rather than the host, keeping the whole op on-device; CPU runs it directly (see NmsGreedyKeep below).

Parameters

mask	(n, num_block_cols) bitmask, mask[i, j] bit k is 1 iff sorted box i overlaps with sorted box (NMS_BLOCK_SIZE*j + k).
sort_indices	(n,) maps sorted position -> original box index.
n	Number of boxes.
remv	Scratch buffer, capacity >= ceil(n / NMS_BLOCK_SIZE).
keep_indices	Output buffer, capacity >= n.

Returns

Number of boxes kept (i.e. the number of valid entries written to keep_indices).

NmsSYCLKernel()

int open3d::ml::contrib::NmsSYCLKernel	(	sycl::queue &	queue,
		const float *	boxes,
		const float *	scores,
		int	n,
		double	nms_overlap_thresh,
		int64_t *	keep_indices_out
	)

operator*() [1/2]

PointXYZ open3d::ml::contrib::operator* ( const float  a, const PointXYZ  P  )

inline

operator*() [2/2]

PointXYZ open3d::ml::contrib::operator* ( const PointXYZ  P, const float  a  )

inline

operator+()

PointXYZ open3d::ml::contrib::operator+ ( const PointXYZ  A, const PointXYZ  B  )

inline

operator-()

PointXYZ open3d::ml::contrib::operator- ( const PointXYZ  A, const PointXYZ  B  )

inline

operator<<()

std::ostream & open3d::ml::contrib::operator<< ( std::ostream &  os, const PointXYZ  P  )

inline

operator==()

bool open3d::ml::contrib::operator== ( const PointXYZ  A, const PointXYZ  B  )

inline

OptBlockConfig()

dim3 open3d::ml::contrib::OptBlockConfig ( int  x, int  y  )

inline

OptNumThreads()

int open3d::ml::contrib::OptNumThreads ( int  work_size )

inline

PointCmp()

OPEN3D_HOST_DEVICE int open3d::ml::contrib::PointCmp ( const Point &  a, const Point &  b, const Point &  center  )

inline

pt_in_box3d()

OPEN3D_HOST_DEVICE bool open3d::ml::contrib::pt_in_box3d ( float  x, float  y, float  z, float  cx, float  bottom_y, float  cz, float  h, float  w, float  l, float  angle, float  max_dis  )

inline

Checks whether point (x, y, z) lies inside the rotated 3D box centered at (cx, cy=bottom_y-h/2, cz) with size (h, w, l) and heading angle. Shared by the CPU/CUDA/SYCL roi_pool launchers. Returns bool (not the original PointRCNN int) since callers only ever store it into an int flag array or test it for truthiness; bool->int conversion is well defined (true/false -> 1/0) so all three backends are unaffected.

roipool3dLauncherCPU()

void open3d::ml::contrib::roipool3dLauncherCPU	(	int	batch_size,
		int	pts_num,
		int	boxes_num,
		int	feature_in_len,
		int	sampled_pts_num,
		const float *	xyz,
		const float *	boxes3d,
		const float *	pts_feature,
		float *	pooled_features,
		int *	pooled_empty_flag
	)

roipool3dLauncherSYCL()

void open3d::ml::contrib::roipool3dLauncherSYCL	(	sycl::queue &	queue,
		int	batch_size,
		int	pts_num,
		int	boxes_num,
		int	feature_in_len,
		int	sampled_pts_num,
		const float *	xyz,
		const float *	boxes3d,
		const float *	pts_feature,
		float *	pooled_features,
		int *	pooled_empty_flag
	)

RotateAroundCenter()

OPEN3D_HOST_DEVICE void open3d::ml::contrib::RotateAroundCenter ( const Point &  center, const float  angle_cos, const float  angle_sin, Point &  p  )

inline

SortIndexesDescending()

template<typename T >

std::vector< int64_t > open3d::ml::contrib::SortIndexesDescending ( const T *  values, int64_t  num  )

inline

Sorts values in descending (or ascending) order and returns the permutation indices. Shared by the CPU/CUDA/SYCL NMS launchers, which all need a host-side ranking of the (small) scores array before running the pairwise-overlap kernel.

ThreeInterpolateGradSYCL()

void open3d::ml::contrib::ThreeInterpolateGradSYCL ( sycl::queue &  queue, int  b, int  c, int  n, int  m, const float *const  grad_out, const int *const  idx, const float *const  weight, float *const  grad_points  )

inline

Gradient of ThreeInterpolateSYCL w.r.t. points. Scatters (atomically accumulates) grad_out * weight into grad_points at the 3 neighbor indices. grad_points must be zero-initialized by the caller before this call. Ports three_interpolate_grad_kernel.

ThreeInterpolateSYCL()

void open3d::ml::contrib::ThreeInterpolateSYCL ( sycl::queue &  queue, int  b, int  c, int  m, int  n, const float *const  points, const int *const  idx, const float *const  weight, float *const  out  )

inline

Weighted gather of 3 nearest neighbor features. Ports three_interpolate_kernel.

Parameters

points	Source features, shape [b, c, m].
idx	Indices of the 3 nearest, shape [b, n, 3].
weight	Interpolation weights, shape [b, n, 3].
out	Output features, shape [b, c, n].

ThreeNNSYCL()

void open3d::ml::contrib::ThreeNNSYCL ( sycl::queue &  queue, int  b, int  n, int  m, const float *const  unknown, const float *const  known, float *const  dist2, int *const  idx  )

inline

Finds, for each of the n "unknown" points, the 3 nearest of the m "known" points (per batch element). Ports three_nn_kernel.

Parameters

queue	SYCL queue to run the kernel on.
unknown	Query point positions, shape [b, n, 3].
known	Reference point positions, shape [b, m, 3].
dist2	Output squared distances to the 3 nearest, shape [b, n, 3].
idx	Output indices of the 3 nearest, shape [b, n, 3].

TrilinearDevoxelizeGradSYCL()

void open3d::ml::contrib::TrilinearDevoxelizeGradSYCL ( sycl::queue &  queue, int  b, int  c, int  n, int  r3, const int *const  inds, const float *const  wgts, const float *const  grad_y, float *const  grad_x  )

inline

Gradient of TrilinearDevoxelizeSYCL w.r.t. the feature grid. grad_x must be zero-initialized by the caller before this call (uses atomic scatter).

TrilinearDevoxelizeSYCL()

void open3d::ml::contrib::TrilinearDevoxelizeSYCL ( sycl::queue &  queue, int  b, int  c, int  n, int  r, int  r2, int  r3, bool  is_training, const float *const  coords, const float *const  feat, int *const  inds, float *const  wgts, float *const  outs  )

inline

Trilinear devoxelization: for each of the n query points per batch, gathers and interpolates the 8 surrounding voxels of the [c, r, r, r] feature grid. If is_training, also records the 8 voxel indices and weights (needed by the backward pass).

Variable Documentation

EPS

constexpr float open3d::ml::contrib::EPS = static_cast<float>(1e-8)

constexpr

NMS_BLOCK_SIZE

constexpr int open3d::ml::contrib::NMS_BLOCK_SIZE = sizeof(uint64_t) * 8

constexpr