CudaTracerKernels.hpp File Reference

#include <cuda.h>

Go to the source code of this file.

Classes
struct	KernelConfig
struct	OtraceInput
struct	RayStruct

Macros
#define	KDTREE_MASK   0xF0000000
#define	KDTREE_UNMASK   0x0FFFFFFF
#define	KDTREE_LEAF   0xC0000000
#define	KDTREE_EMPTYLEAF   0x80000000
#define	KDTREE_DIMPOS   28
#define	TRACE_FUNC_KDTREE
#define	TRACE_FUNC_BVH
#define	FETCH_GLOBAL(NAME, IDX, TYPE)   ((const TYPE*)NAME)[IDX]
#define	FETCH_TEXTURE(NAME, IDX, TYPE)   tex1Dfetch(t_ ## NAME, IDX)
#define	STORE_RESULT(RAY, TRI, T, U, V)   results[RAY] = make_int4(TRI, __float_as_int(T), __float_as_int(U), __float_as_int(V))

Enumerations
enum	{ MaxBlockHeight = 6, EntrypointSentinel = 0x76543210 }
enum	BVHLayout {   BVHLayout_AOS_AOS = 0, BVHLayout_AOS_SOA, BVHLayout_SOA_AOS, BVHLayout_SOA_SOA,   BVHLayout_Compact, BVHLayout_Compact2, BVHLayout_CPU, BVHLayout_Max }

Macro Definition Documentation

#define FETCH_GLOBAL	(		NAME,
			IDX,
			TYPE
	)		((const TYPE*)NAME)[IDX]

Definition at line 195 of file CudaTracerKernels.hpp.

#define FETCH_TEXTURE	(		NAME,
			IDX,
			TYPE
	)		tex1Dfetch(t_ ## NAME, IDX)

Definition at line 196 of file CudaTracerKernels.hpp.

#define KDTREE_DIMPOS   28

Definition at line 46 of file CudaTracerKernels.hpp.

#define KDTREE_EMPTYLEAF   0x80000000

Definition at line 45 of file CudaTracerKernels.hpp.

#define KDTREE_LEAF   0xC0000000

Definition at line 44 of file CudaTracerKernels.hpp.

#define KDTREE_MASK   0xF0000000

Definition at line 42 of file CudaTracerKernels.hpp.

#define KDTREE_UNMASK   0x0FFFFFFF

Definition at line 43 of file CudaTracerKernels.hpp.

#define STORE_RESULT	(		RAY,
			TRI,
			T,
			U,
			V
	)		results[RAY] = make_int4(TRI, __float_as_int(T), __float_as_int(U), __float_as_int(V))

Definition at line 198 of file CudaTracerKernels.hpp.

#define TRACE_FUNC_BVH

Value:

extern "C" __global__ void trace_bvh( \
        int             numRays,        /* Total number of rays in the batch. */ \
        bool            anyHit,         /* False if rays need to find the closest hit. */ \
        float4*         rays,           /* Ray input: float3 origin, float tmin, float3 direction, float tmax. */ \
        int4*           results,        /* Ray output: int triangleID, float hitT, int2 padding. */ \
        float4*         nodesA,         /* SOA: bytes 0-15 of each node, AOS/Compact: 64 bytes per node. */ \
        float4*         nodesB,         /* SOA: bytes 16-31 of each node, AOS/Compact: unused. */ \
        float4*         nodesC,         /* SOA: bytes 32-47 of each node, AOS/Compact: unused. */ \
        float4*         nodesD,         /* SOA: bytes 48-63 of each node, AOS/Compact: unused. */ \
        float4*         trisA,          /* SOA: bytes 0-15 of each triangle, AOS: 64 bytes per triangle, Compact: 48 bytes per triangle. */ \
        float4*         trisB,          /* SOA: bytes 16-31 of each triangle, AOS/Compact: unused. */ \
        float4*         trisC,          /* SOA: bytes 32-47 of each triangle, AOS/Compact: unused. */ \
        int*            triIndices)     /* Triangle index remapping table. */

Definition at line 99 of file CudaTracerKernels.hpp.

#define TRACE_FUNC_KDTREE

Value:

extern "C" __global__ void trace_kdtree( \
        int             numRays,        /* Total number of rays in the batch. */ \
        bool            anyHit,         /* False if rays need to find the closest hit. */ \
        float*          bmin,   \
        float*          bmax,   \
        float           delta,  \
        float4*         rays,           /* Ray input: float3 origin, float tmin, float3 direction, float tmax. */ \
        int4*           results,        /* Ray output: int triangleID, float hitT, int2 padding. */ \
        float4*         nodesA,         /* SOA: bytes 0-15 of each node, AOS/Compact: 64 bytes per node. */ \
        float4*         nodesB,         /* SOA: bytes 16-31 of each node, AOS/Compact: unused. */ \
        float4*         nodesC,         /* SOA: bytes 32-47 of each node, AOS/Compact: unused. */ \
        float4*         nodesD,         /* SOA: bytes 48-63 of each node, AOS/Compact: unused. */ \
        float4*         trisA,          /* SOA: bytes 0-15 of each triangle, AOS: 64 bytes per triangle, Compact: 48 bytes per triangle. */ \
        float4*         trisB,          /* SOA: bytes 16-31 of each triangle, AOS/Compact: unused. */ \
        float4*         trisC,          /* SOA: bytes 32-47 of each triangle, AOS/Compact: unused. */ \
        int*            triIndices)     /* Triangle index remapping table. */

Definition at line 81 of file CudaTracerKernels.hpp.

Enumeration Type Documentation

anonymous enum

Enumerator
MaxBlockHeight
EntrypointSentinel

Definition at line 35 of file CudaTracerKernels.hpp.

enum BVHLayout

Enumerator
BVHLayout_AOS_AOS
BVHLayout_AOS_SOA
BVHLayout_SOA_AOS
BVHLayout_SOA_SOA
BVHLayout_Compact
BVHLayout_Compact2
BVHLayout_CPU
BVHLayout_Max

Definition at line 52 of file CudaTracerKernels.hpp.