CollisionDetection.cu File Reference

#include <cstdio>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <vector>
#include "helios_vector_types.h"

Go to the source code of this file.

Data Structures
struct	GPUBVHNode
	GPU-friendly BVH node structure (Legacy AoS format) More...
struct	GPUBVHNodesSoA
	GPU-optimized SoA BVH structure for warp-efficient traversal. More...

Functions
__device__ bool	d_aabbIntersect (const float3 &min1, const float3 &max1, const float3 &min2, const float3 &max2)
	CUDA device function to test AABB intersection.
__device__ __forceinline__ float3	cross (const float3 &a, const float3 &b)
	CUDA device helper functions for vector operations.
__device__ __forceinline__ float	dot (const float3 &a, const float3 &b)
__device__ __forceinline__ float3	normalize (const float3 &v)
__device__ __forceinline__ float3	operator+ (const float3 &a, const float3 &b)
__device__ __forceinline__ float3	operator- (const float3 &a, const float3 &b)
__device__ __forceinline__ float3	operator* (const float3 &a, float scalar)
__device__ __forceinline__ bool	rayTriangleIntersect (const float3 &ray_origin, const float3 &ray_direction, const float3 &v0, const float3 &v1, const float3 &v2, float max_distance, float &hit_distance)
	Fast ray-triangle intersection using Möller-Trumbore algorithm.
__device__ __forceinline__ bool	warpRayAABBIntersect (const float3 &ray_origin, const float3 &ray_dir, const float3 &aabb_min, const float3 &aabb_max, float max_dist)
	Warp-efficient ray-AABB intersection for GPU optimization.
__device__ __forceinline__ bool	rayTriangleIntersectCPU (const float3 &origin, const float3 &direction, const float3 &v0, const float3 &v1, const float3 &v2, float &distance)
	High-performance GPU ray-triangle intersection kernel using BVH traversal.
__device__ __forceinline__ bool	rayPatchIntersect (const float3 &origin, const float3 &direction, const float3 &v0, const float3 &v1, const float3 &v2, const float3 &v3, float &distance)
__device__ bool	rayVoxelIntersect (const float3 &ray_origin, const float3 &ray_direction, const float3 &aabb_min, const float3 &aabb_max, float &distance)
__global__ void	rayPrimitiveBVHKernel (GPUBVHNode *d_bvh_nodes, unsigned int *d_primitive_indices, int *d_primitive_types, float3 *d_primitive_vertices, unsigned int *d_vertex_offsets, float3 *d_ray_origins, float3 *d_ray_directions, float *d_ray_max_distances, int num_rays, int primitive_count, int total_vertex_count, float *d_hit_distances, unsigned int *d_hit_primitive_ids, unsigned int *d_hit_counts, bool find_closest_hit)
void	launchRayPrimitiveIntersection (void *h_bvh_nodes, int node_count, unsigned int *h_primitive_indices, int primitive_count, int *h_primitive_types, float3 *h_primitive_vertices, unsigned int *h_vertex_offsets, int total_vertex_count, float *h_ray_origins, float *h_ray_directions, float *h_ray_max_distances, int num_rays, float *h_hit_distances, unsigned int *h_hit_primitive_ids, unsigned int *h_hit_counts, bool find_closest_hit)
	Launch optimized ray-triangle intersection kernel for true ray tracing.
__global__ void	bvhTraversalKernel (GPUBVHNode *d_nodes, unsigned int *d_primitive_indices, float3 *d_primitive_aabb_min, float3 *d_primitive_aabb_max, float3 *d_query_aabb_min, float3 *d_query_aabb_max, unsigned int *d_results, unsigned int *d_result_counts, int num_queries, int max_results_per_query)
	CUDA kernel for BVH traversal collision detection.
void	launchBVHTraversal (void *h_nodes, int node_count, unsigned int *h_primitive_indices, int primitive_count, float *h_primitive_aabb_min, float *h_primitive_aabb_max, float *h_query_aabb_min, float *h_query_aabb_max, int num_queries, unsigned int *h_results, unsigned int *h_result_counts, int max_results_per_query)
	Launch BVH traversal kernel from C++ code.
__global__ void	intersectRegularGridKernel (const size_t num_rays, float3 *d_ray_origins, float3 *d_ray_directions, float3 grid_center, float3 grid_size, int3 grid_divisions, int primitive_count, int *d_voxel_ray_counts, float *d_voxel_path_lengths, int *d_voxel_transmitted, int *d_voxel_hit_before, int *d_voxel_hit_after, int *d_voxel_hit_inside)
	CUDA kernel to calculate ray path lengths through a regular voxel grid.
void	launchVoxelRayPathLengths (int num_rays, float *h_ray_origins, float *h_ray_directions, float grid_center_x, float grid_center_y, float grid_center_z, float grid_size_x, float grid_size_y, float grid_size_z, int grid_divisions_x, int grid_divisions_y, int grid_divisions_z, int primitive_count, int *h_voxel_ray_counts, float *h_voxel_path_lengths, int *h_voxel_transmitted, int *h_voxel_hit_before, int *h_voxel_hit_after, int *h_voxel_hit_inside)
	Launch CUDA kernel for regular grid voxel ray path length calculation.

Detailed Description

CUDA source file for GPU-accelerated collision detection

Copyright (C) 2016-2025 Brian Bailey

This program is free software: you can redistribute it and/or modify it under the terms of the GNU General Public License as published by the Free Software Foundation, version 2.

This program is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License for more details.

Definition in file CollisionDetection.cu.

Function Documentation

bvhTraversalKernel()

__global__ void bvhTraversalKernel	(	GPUBVHNode *	d_nodes,
		unsigned int *	d_primitive_indices,
		float3 *	d_primitive_aabb_min,
		float3 *	d_primitive_aabb_max,
		float3 *	d_query_aabb_min,
		float3 *	d_query_aabb_max,
		unsigned int *	d_results,
		unsigned int *	d_result_counts,
		int	num_queries,
		int	max_results_per_query
	)

CUDA kernel for BVH traversal collision detection.

Each thread processes one query AABB and traverses the BVH to find collisions.

Parameters

[in]	d_nodes	Array of BVH nodes on GPU
[in]	d_primitive_indices	Array of primitive indices on GPU
[in]	d_query_aabb_min	Array of query AABB minimum corners
[in]	d_query_aabb_max	Array of query AABB maximum corners
[out]	d_results	Array to store collision results
[out]	d_result_counts	Array to store number of results per query
[in]	num_queries	Number of queries to process
[in]	max_results_per_query	Maximum results to store per query

Definition at line 753 of file CollisionDetection.cu.

cross()

__device__ __forceinline__ float3 cross	(	const float3 &	a,
		const float3 &	b
	)

CUDA device helper functions for vector operations.

Definition at line 77 of file CollisionDetection.cu.

d_aabbIntersect()

__device__ bool d_aabbIntersect	(	const float3 &	min1,
		const float3 &	max1,
		const float3 &	min2,
		const float3 &	max2
	)

CUDA device function to test AABB intersection.

Parameters

[in]	min1	Minimum corner of first AABB
[in]	max1	Maximum corner of first AABB
[in]	min2	Minimum corner of second AABB
[in]	max2	Maximum corner of second AABB

Returns

True if AABBs intersect

Definition at line 70 of file CollisionDetection.cu.

dot()

__device__ __forceinline__ float dot	(	const float3 &	a,
		const float3 &	b
	)

Definition at line 81 of file CollisionDetection.cu.

intersectRegularGridKernel()

__global__ void intersectRegularGridKernel	(	const size_t	num_rays,
		float3 *	d_ray_origins,
		float3 *	d_ray_directions,
		float3	grid_center,
		float3	grid_size,
		int3	grid_divisions,
		int	primitive_count,
		int *	d_voxel_ray_counts,
		float *	d_voxel_path_lengths,
		int *	d_voxel_transmitted,
		int *	d_voxel_hit_before,
		int *	d_voxel_hit_after,
		int *	d_voxel_hit_inside
	)

CUDA kernel to calculate ray path lengths through a regular voxel grid.

This kernel computes the intersection of rays with voxels and accumulates path length statistics for integration with aeriallidar and lidar plugins.

Parameters

[in]	num_rays	Number of rays to process
[in]	d_ray_origins	Array of ray origin points
[in]	d_ray_directions	Array of ray direction vectors (normalized)
[in]	grid_center	Center of the voxel grid
[in]	grid_size	Total size of the voxel grid
[in]	grid_divisions	Number of divisions in x, y, z
[out]	d_voxel_ray_counts	Flattened array to accumulate ray counts per voxel
[out]	d_voxel_path_lengths	Flattened array to accumulate path lengths per voxel
[out]	d_voxel_transmitted	Flattened array to count transmitted rays per voxel

Definition at line 941 of file CollisionDetection.cu.

launchBVHTraversal()

void launchBVHTraversal	(	void *	h_nodes,
		int	node_count,
		unsigned int *	h_primitive_indices,
		int	primitive_count,
		float *	h_primitive_aabb_min,
		float *	h_primitive_aabb_max,
		float *	h_query_aabb_min,
		float *	h_query_aabb_max,
		int	num_queries,
		unsigned int *	h_results,
		unsigned int *	h_result_counts,
		int	max_results_per_query
	)

Launch BVH traversal kernel from C++ code.

Parameters

[in]	h_nodes	Host array of BVH nodes
[in]	node_count	Number of BVH nodes
[in]	h_primitive_indices	Host array of primitive indices
[in]	primitive_count	Number of primitive indices
[in]	h_query_aabb_min	Host array of query AABB minimum corners
[in]	h_query_aabb_max	Host array of query AABB maximum corners
[in]	num_queries	Number of queries
[out]	h_results	Host array for results
[out]	h_result_counts	Host array for result counts
[in]	max_results_per_query	Maximum results per query

Definition at line 842 of file CollisionDetection.cu.

launchRayPrimitiveIntersection()

void launchRayPrimitiveIntersection	(	void *	h_bvh_nodes,
		int	node_count,
		unsigned int *	h_primitive_indices,
		int	primitive_count,
		int *	h_primitive_types,
		float3 *	h_primitive_vertices,
		unsigned int *	h_vertex_offsets,
		int	total_vertex_count,
		float *	h_ray_origins,
		float *	h_ray_directions,
		float *	h_ray_max_distances,
		int	num_rays,
		float *	h_hit_distances,
		unsigned int *	h_hit_primitive_ids,
		unsigned int *	h_hit_counts,
		bool	find_closest_hit
	)

Launch optimized ray-triangle intersection kernel for true ray tracing.

This kernel implements proper ray-triangle intersection using the Möller-Trumbore algorithm with GPU optimizations for high-performance ray tracing.

Parameters

[in]	h_bvh_nodes	Host array of BVH nodes
[in]	node_count	Number of BVH nodes
[in]	h_primitive_indices	Host array of primitive indices
[in]	primitive_count	Number of primitives
[in]	h_triangle_vertices	Host array of triangle vertex data (9 floats per triangle: v0,v1,v2)
[in]	h_ray_origins	Host array of ray origins (3 floats per ray)
[in]	h_ray_directions	Host array of ray directions (3 floats per ray)
[in]	h_ray_max_distances	Host array of ray maximum distances
[in]	num_rays	Number of rays to process
[out]	h_hit_distances	Host array for hit distances (closest hit per ray)
[out]	h_hit_primitive_ids	Host array for hit primitive IDs
[out]	h_hit_counts	Host array for number of hits per ray
[in]	find_closest_hit	If true, return only closest hit; if false, return all hits within distance

Definition at line 522 of file CollisionDetection.cu.

launchVoxelRayPathLengths()

void launchVoxelRayPathLengths	(	int	num_rays,
		float *	h_ray_origins,
		float *	h_ray_directions,
		float	grid_center_x,
		float	grid_center_y,
		float	grid_center_z,
		float	grid_size_x,
		float	grid_size_y,
		float	grid_size_z,
		int	grid_divisions_x,
		int	grid_divisions_y,
		int	grid_divisions_z,
		int	primitive_count,
		int *	h_voxel_ray_counts,
		float *	h_voxel_path_lengths,
		int *	h_voxel_transmitted,
		int *	h_voxel_hit_before,
		int *	h_voxel_hit_after,
		int *	h_voxel_hit_inside
	)

Launch CUDA kernel for regular grid voxel ray path length calculation.

Definition at line 1142 of file CollisionDetection.cu.

normalize()

__device__ __forceinline__ float3 normalize

(

const float3 &

v

)

Definition at line 85 of file CollisionDetection.cu.

operator*()

__device__ __forceinline__ float3 operator*	(	const float3 &	a,
		float	scalar
	)

Definition at line 101 of file CollisionDetection.cu.

operator+()

__device__ __forceinline__ float3 operator+	(	const float3 &	a,
		const float3 &	b
	)

Definition at line 93 of file CollisionDetection.cu.

operator-()

__device__ __forceinline__ float3 operator-	(	const float3 &	a,
		const float3 &	b
	)

Definition at line 97 of file CollisionDetection.cu.

rayPatchIntersect()

__device__ __forceinline__ bool rayPatchIntersect	(	const float3 &	origin,
		const float3 &	direction,
		const float3 &	v0,
		const float3 &	v1,
		const float3 &	v2,
		const float3 &	v3,
		float &	distance
	)

Definition at line 260 of file CollisionDetection.cu.

rayPrimitiveBVHKernel()

__global__ void rayPrimitiveBVHKernel	(	GPUBVHNode *	d_bvh_nodes,
		unsigned int *	d_primitive_indices,
		int *	d_primitive_types,
		float3 *	d_primitive_vertices,
		unsigned int *	d_vertex_offsets,
		float3 *	d_ray_origins,
		float3 *	d_ray_directions,
		float *	d_ray_max_distances,
		int	num_rays,
		int	primitive_count,
		int	total_vertex_count,
		float *	d_hit_distances,
		unsigned int *	d_hit_primitive_ids,
		unsigned int *	d_hit_counts,
		bool	find_closest_hit
	)

Definition at line 338 of file CollisionDetection.cu.

rayTriangleIntersect()

__device__ __forceinline__ bool rayTriangleIntersect	(	const float3 &	ray_origin,
		const float3 &	ray_direction,
		const float3 &	v0,
		const float3 &	v1,
		const float3 &	v2,
		float	max_distance,
		float &	hit_distance
	)

Fast ray-triangle intersection using Möller-Trumbore algorithm.

Parameters

[in]	ray_origin	Ray starting point
[in]	ray_direction	Ray direction vector (normalized)
[in]	v0	First triangle vertex
[in]	v1	Second triangle vertex
[in]	v2	Third triangle vertex
[in]	max_distance	Maximum ray distance
[out]	hit_distance	Distance to intersection (if hit)

Returns

True if ray intersects triangle within max_distance

Definition at line 116 of file CollisionDetection.cu.

rayTriangleIntersectCPU()

__device__ __forceinline__ bool rayTriangleIntersectCPU	(	const float3 &	origin,
		const float3 &	direction,
		const float3 &	v0,
		const float3 &	v1,
		const float3 &	v2,
		float &	distance
	)

High-performance GPU ray-triangle intersection kernel using BVH traversal.

This kernel implements proper ray-triangle intersection with BVH acceleration, using the Möller-Trumbore algorithm optimized for GPU warp efficiency.

Parameters

[in]	d_bvh_nodes	BVH nodes on GPU
[in]	d_primitive_indices	Primitive indices on GPU
[in]	d_triangle_vertices	Triangle vertex data on GPU (3 vertices per triangle)
[in]	d_ray_origins	Ray origins on GPU
[in]	d_ray_directions	Ray directions on GPU
[in]	d_ray_max_distances	Ray maximum distances on GPU
[in]	num_rays	Number of rays to process
[out]	d_hit_distances	Closest hit distances per ray
[out]	d_hit_primitive_ids	Hit primitive IDs per ray
[out]	d_hit_counts	Number of hits per ray
[in]	find_closest_hit	If true, return only closest hit

Definition at line 221 of file CollisionDetection.cu.

rayVoxelIntersect()

__device__ bool rayVoxelIntersect	(	const float3 &	ray_origin,
		const float3 &	ray_direction,
		const float3 &	aabb_min,
		const float3 &	aabb_max,
		float &	distance
	)

Definition at line 296 of file CollisionDetection.cu.

warpRayAABBIntersect()

__device__ __forceinline__ bool warpRayAABBIntersect	(	const float3 &	ray_origin,
		const float3 &	ray_dir,
		const float3 &	aabb_min,
		const float3 &	aabb_max,
		float	max_dist
	)

Warp-efficient ray-AABB intersection for GPU optimization.

Parameters

[in]	ray_origin	Ray starting point
[in]	ray_dir	Ray direction vector (normalized)
[in]	aabb_min	AABB minimum corner
[in]	aabb_max	AABB maximum corner
[in]	max_dist	Maximum ray distance

Returns

True if ray intersects AABB within max_dist

Definition at line 168 of file CollisionDetection.cu.