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

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...

Macros
#define	HELIOS_CUDA_CHECK(call)

#define	BVH_TRAVERSAL_STACK_CAPACITY 128

Functions
void	helios::helios_runtime_error (const std::string &error_message)
	Function to throw a runtime error.

__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__ float3	safeRayInvDir (const float3 &dir)
	Warp-efficient ray-AABB intersection for GPU optimization.

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

__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)

__device__ __forceinline__ float3	computeHitNormal (int ptype, const float3 *d_primitive_vertices, unsigned int vertex_offset, const float3 &ray_origin, const float3 &ray_direction, float hit_distance)
	Compute the surface normal for a hit primitive, matching the CPU intersectPrimitiveThreadSafe convention.

__device__ __forceinline__ bool	sampleMaskOpaqueGPU (int mask_id, float u, float v, const unsigned char d_mask_data, const unsigned int d_mask_offsets, const int *d_mask_sizes)

__device__ __forceinline__ bool	isHitOpaqueGPU (int ptype, const float3 verts, int mask_id, int uv_id, const float uv4, const float3 &hit_point, const unsigned char d_mask_data, const unsigned int d_mask_offsets, const int *d_mask_sizes)

__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, const unsigned char d_mask_data, const unsigned int d_mask_offsets, const int d_mask_sizes, const int d_mask_IDs, const float d_uv_data, const int d_uv_IDs, float3 d_ray_origins, float3 d_ray_directions, float d_ray_max_distances, float uniform_max_distance, 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, float3 d_hit_normals, bool find_closest_hit)

void	launchRaysOnResidentScene (void d_bvh_nodes, int node_count, unsigned int d_primitive_indices, int primitive_count, int d_primitive_types, float3 d_primitive_vertices, unsigned int d_vertex_offsets, const unsigned char d_mask_data, const unsigned int d_mask_offsets, const int d_mask_sizes, const int d_mask_IDs, const float d_uv_data, const int d_uv_IDs, int total_vertex_count, const float h_ray_origins, const float h_ray_directions, const float h_ray_max_distances, float uniform_max_distance, int num_rays, float h_hit_distances, unsigned int h_hit_primitive_ids, unsigned int h_hit_counts, float h_hit_normals, bool find_closest_hit)
	Launch the ray-primitive intersection kernel against scene geometry that is ALREADY resident on the device.

__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.

bool	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.

Variables
__device__ unsigned int	d_bvh_stack_overflow = 0

Detailed Description

CUDA source file for GPU-accelerated collision detection

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.

Macro Definition Documentation

◆ BVH_TRAVERSAL_STACK_CAPACITY

#define BVH_TRAVERSAL_STACK_CAPACITY 128

Per-thread BVH traversal stack capacity. An iterative DFS that pushes both children holds at most ~one node per tree level, so the capacity must exceed the host builder's MAX_DEPTH (64, see buildBVHRecursive). 128 mirrors the CPU STACK_CAPACITY (CollisionDetection_RayTracing.cpp) and leaves generous headroom; the guarded push in each kernel is a defensive backstop that must never actually trigger for a valid tree. The stack lives in per-thread local memory (it is thread-private, so the previous block-shared array gave no sharing benefit while capping occupancy at 32 KiB/block).

Definition at line 64 of file CollisionDetection.cu.

◆ HELIOS_CUDA_CHECK

#define HELIOS_CUDA_CHECK ( call )

Value:

    do {                                                                                                                                                                                                                                        \
        cudaError_t _helios_cuda_err = (call);                                                                                                                                                                                                  \
        if (_helios_cuda_err != cudaSuccess) {                                                                                                                                                                                                  \
            helios::helios_runtime_error(std::string("CUDA error (") + #call + "): " + cudaGetErrorString(_helios_cuda_err));                                                                                                                    \
        }                                                                                                                                                                                                                                       \
    } while (0)

Fail-fast wrapper for CUDA runtime calls. Helios policy forbids silent fallbacks: a failed CUDA call throws a helios_runtime_error rather than the legacy fprintf(stderr)+return, which used to leave output arrays partially written. Used by the resident-scene launch path below.

Definition at line 34 of 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 854 of file CollisionDetection.cu.

◆ computeHitNormal()

__device__ __forceinline__ float3 computeHitNormal	(	int	ptype,
		const float3 *	d_primitive_vertices,
		unsigned int	vertex_offset,
		const float3 &	ray_origin,
		const float3 &	ray_direction,
		float	hit_distance
	)

Compute the surface normal for a hit primitive, matching the CPU intersectPrimitiveThreadSafe convention.

Triangles/patches use normalize(cross(v1-v0, v2-v0)) face-forwarded toward the ray origin (n flipped so it points back at the ray), exactly as the CPU SoA path does, so GPU and CPU normals (and the LiDAR hit_fnorm = dot(dir,n) derived from them) agree. Voxels use the axis-aligned face normal of the hit face (no face-forward, matching CPU). Degenerate/unknown primitives fall back to normalize(-direction).

Definition at line 394 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 107 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 100 of file CollisionDetection.cu.

◆ dot()

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

Definition at line 111 of file CollisionDetection.cu.

◆ helios_runtime_error()

void helios::helios_runtime_error ( const std::string & error_message )

Function to throw a runtime error.

Parameters

[in] error_message Message to be issued when error is thrown.

Definition at line 44 of file global.cpp.

◆ 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 1065 of file CollisionDetection.cu.

◆ isHitOpaqueGPU()

__device__ __forceinline__ bool isHitOpaqueGPU	(	int	ptype,
		const float3 *	verts,
		int	mask_id,
		int	uv_id,
		const float *	uv4,
		const float3 &	hit_point,
		const unsigned char *	d_mask_data,
		const unsigned int *	d_mask_offsets,
		const int *	d_mask_sizes
	)

Decide whether a ray-primitive hit lands on an opaque texel. Computes the (u,v) at the hit point with the SAME interpolation as the CPU CollisionDetection::isHitTexelOpaque() (patch: BL->BR / BL->TL basis projection; triangle: barycentric Cramer solve) so the GPU and CPU synthetic scans reject identical texels. Voxels / mask-less primitives and degenerate/missing-UV configurations are treated as solid, exactly as the CPU path does.

Definition at line 463 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 949 of file CollisionDetection.cu.

◆ launchRaysOnResidentScene()

void launchRaysOnResidentScene	(	void *	d_bvh_nodes,
		int	node_count,
		unsigned int *	d_primitive_indices,
		int	primitive_count,
		int *	d_primitive_types,
		float3 *	d_primitive_vertices,
		unsigned int *	d_vertex_offsets,
		const unsigned char *	d_mask_data,
		const unsigned int *	d_mask_offsets,
		const int *	d_mask_sizes,
		const int *	d_mask_IDs,
		const float *	d_uv_data,
		const int *	d_uv_IDs,
		int	total_vertex_count,
		const float *	h_ray_origins,
		const float *	h_ray_directions,
		const float *	h_ray_max_distances,
		float	uniform_max_distance,
		int	num_rays,
		float *	h_hit_distances,
		unsigned int *	h_hit_primitive_ids,
		unsigned int *	h_hit_counts,
		float *	h_hit_normals,
		bool	find_closest_hit
	)

Launch the ray-primitive intersection kernel against scene geometry that is ALREADY resident on the device.

Only the per-call ray inputs and hit outputs are allocated, uploaded, and freed here; the BVH, primitive indices, types, packed vertices, and vertex offsets are passed in as device pointers owned by the caller (uploaded once per scan by transferBVHToGPU). This is what keeps chunked synthetic scans from re-uploading the whole scene per chunk. Pass h_hit_normals != nullptr to also read back the per-ray face-forwarded surface normal (flat xyz, 3 floats/ray). CUDA failures throw helios_runtime_error (fail-fast) rather than writing partial results.

Parameters

[in]	d_bvh_nodes	Device BVH nodes (resident)
[in]	node_count	Number of BVH nodes
[in]	d_primitive_indices	Device primitive indices (resident)
[in]	primitive_count	Number of primitives
[in]	d_primitive_types	Device per-primitive type codes (resident)
[in]	d_primitive_vertices	Device packed primitive vertices (resident)
[in]	d_vertex_offsets	Device per-primitive vertex offsets (resident)
[in]	d_mask_data	Device texture transparency mask bytes, or null when the scene has no masks (resident)
[in]	d_mask_offsets	Device per-mask start index into d_mask_data (resident)
[in]	d_mask_sizes	Device per-mask width/height (resident)
[in]	d_mask_IDs	Device per-primitive mask index, -1 = none (resident)
[in]	d_uv_data	Device per-primitive UVs, 4 vec2 per primitive (resident)
[in]	d_uv_IDs	Device per-primitive UV flag, -1 = parametric/none (resident)
[in]	total_vertex_count	Number of packed vertices
[in]	h_ray_origins	Host ray origins (3 floats per ray)
[in]	h_ray_directions	Host ray directions (3 floats per ray)
[in]	h_ray_max_distances	Host ray maximum distances
[in]	num_rays	Number of rays to process
[out]	h_hit_distances	Host array for closest-hit distances
[out]	h_hit_primitive_ids	Host array for hit primitive IDs
[out]	h_hit_counts	Host array for per-ray hit counts (0 == miss)
[out]	h_hit_normals	Host array for per-ray surface normals (3 floats per ray), or nullptr to skip
[in]	find_closest_hit	If true, return only the closest hit

Definition at line 757 of file CollisionDetection.cu.

◆ launchVoxelRayPathLengths()

bool 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.

Returns: true if GPU execution succeeded, false if GPU unavailable or error occurred

Definition at line 1267 of file CollisionDetection.cu.

◆ normalize()

__device__ __forceinline__ float3 normalize ( const float3 & v )

Definition at line 115 of file CollisionDetection.cu.

◆ operator*()

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

Definition at line 131 of file CollisionDetection.cu.

◆ operator+()

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

Definition at line 123 of file CollisionDetection.cu.

◆ operator-()

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

Definition at line 127 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 306 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,
		const unsigned char *	d_mask_data,
		const unsigned int *	d_mask_offsets,
		const int *	d_mask_sizes,
		const int *	d_mask_IDs,
		const float *	d_uv_data,
		const int *	d_uv_IDs,
		float3 *	d_ray_origins,
		float3 *	d_ray_directions,
		float *	d_ray_max_distances,
		float	uniform_max_distance,
		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,
		float3 *	d_hit_normals,
		bool	find_closest_hit
	)

Definition at line 514 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 146 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 265 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 344 of file CollisionDetection.cu.

◆ safeRayInvDir()

__device__ __forceinline__ float3 safeRayInvDir ( const float3 & dir )

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 Sign-preserving reciprocal of a ray direction for slab-method AABB tests. Clamps each near-zero component to +/-PARALLEL_EPS before the reciprocal so an axis-aligned ray lying exactly on a box face computes (bound - origin == 0) * huge == 0 (finite) instead of (bound - origin == 0) * inf == NaN, which would corrupt the fminf/fmaxf slab interval and spuriously reject a box the ray actually crosses. Mirrors the CPU aabbIntersectSoA SSE path (CollisionDetection_RayTracing.cpp).

Definition at line 204 of file CollisionDetection.cu.

◆ sampleMaskOpaqueGPU()

__device__ __forceinline__ bool sampleMaskOpaqueGPU	(	int	mask_id,
		float	u,
		float	v,
		const unsigned char *	d_mask_data,
		const unsigned int *	d_mask_offsets,
		const int *	d_mask_sizes
	)

Sample a texture transparency mask at UV (u,v); returns true if the texel is opaque (hit should count), false if transparent (ray passes through). Mirrors the CPU CollisionDetection::isHitTexelOpaque() texel lookup: wrap UV into [0,1), px = u*width, py = (1-v)*height (mask rows are top-to-bottom, UV y=0 is the bottom), clamp, fetch the byte.

Definition at line 437 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
	)

Definition at line 212 of file CollisionDetection.cu.

Variable Documentation

◆ d_bvh_stack_overflow

__device__ unsigned int d_bvh_stack_overflow = 0

Device flag raised (atomically) by a traversal kernel if a child push would exceed BVH_TRAVERSAL_STACK_CAPACITY. The host launch wrappers reset it before launch and throw helios_runtime_error after the kernel if it is set, mirroring the CPU invariant that a valid SAH tree never overflows the stack (fail-fast instead of silently dropping primitives).

Definition at line 69 of file CollisionDetection.cu.

	1.3.77

Data Structures

Macros

Functions

Variables

Detailed Description

Macro Definition Documentation

◆ BVH_TRAVERSAL_STACK_CAPACITY

◆ HELIOS_CUDA_CHECK

Function Documentation

◆ bvhTraversalKernel()

◆ computeHitNormal()

◆ cross()

◆ d_aabbIntersect()

◆ dot()

◆ helios_runtime_error()

◆ intersectRegularGridKernel()

◆ isHitOpaqueGPU()

◆ launchBVHTraversal()

◆ launchRaysOnResidentScene()

◆ launchVoxelRayPathLengths()

◆ normalize()

◆ operator*()

◆ operator+()

◆ operator-()

◆ rayPatchIntersect()

◆ rayPrimitiveBVHKernel()

◆ rayTriangleIntersect()

◆ rayTriangleIntersectCPU()

◆ rayVoxelIntersect()

◆ safeRayInvDir()

◆ sampleMaskOpaqueGPU()

◆ warpRayAABBIntersect()

Variable Documentation

◆ d_bvh_stack_overflow