Kernels for ray tracing. More...

#include "KernelDefs.h"
#include "kd-tree/KDKernelDefs.h"
#include "RayPool.h"
#include "MNCudaUtil.h"
#include "MNCudaMemPool.h"
#include "MNCudaMT.h"
#include "MNStatContainer.h"
#include "intersect_dev.h"
#include "photon_dev.h"
#include "sample_dev.h"

Defines
#define	INTERSECT_BLOCKSIZE 128
	Thread block size used for intersection search.
#define	MAX_DIFF_TEX_COUNT 20
	Maximum number of diffuse textures supported.
#define	GENERATETEXVAR(idx) tex_diffTex ## idx
	Macro to generate diffuse texture reference names.
#define	BINDTEX2ARRAY(numTex, idx)
	Binds the given texture reference to a CUDA texture array.
#define	UNBINDTEX(numTex, idx)
	Unbinds the given texture reference.
Functions
void	RTInitializeKernels ()
	Sets cache configurations for ray tracing kernels.
void	BindTextureTextures (const MaterialData &mats)
	Binds diffuse texture arrays to texture references.
void	RTUpdateKernelData (const LightData &lights, const TriangleData &tris, const MaterialData &mats, const KDTreeData &kdTree, float fRayEpsilon)
	Moves scene data to constant memory and texture memory.
void	RTCleanupKernelData ()
	Unbinds textures used for ray tracing kernels.
CUDA device functions
__device__ float4	dev_FetchDiffTex (int idxTex, float u, float v)
	Fetches data from a given diffuse texture.
__device__ float3	dev_GetColorDiffuse (uint idxTri, uint idxMaterial, char4 matFlags, float2 baryHit)
	Determines the diffuse material color for given triangle hit.
__device__ float3	dev_GetColorSpec (uint idxTri, uint idxMaterial, char4 matFlags, float2 baryHit, float *outSpecExp)
	Determines the diffuse material color for given triangle hit.
__device__ float3	dev_SampleLightL (float3 pt, float3 ptLightSample, float3 *outW_i)
	Samples the only primary light source.
__device__ float3	dev_GetReflectedDirectLight (float3 ptEye, float3 pt, float3 nG, float3 nS, float3 ptLightSample, uint idxTri, uint idxMat, char4 matFlags, float2 baryHit)
	Computes the reflected direct light from pt to ptEye.
template<bool needClosest>
__device__ int	dev_FindNextIntersectionKDWhileWhile (const float3 rayO, const float3 rayD, const float tMinRay, const float tMaxRay, float &outLambda, float2 &outBaryHit)
	Looks for the next ray triangle intersection using a while-while standard traversal.
template<bool getMinDist>
__device__ float4	dev_ApproximateNormalAt (float3 ptQuery, float queryRadiusSqr)
	Approximates normal at a given query point.
__device__ void	dev_AddPixelRadiance (const RayChunk &rayChunk, const ShadingPoints &shadingPts, const uint tid, const float3 L_sample, float4 *d_ioRadiance)
	Updates the radiance accumulator by adding a scaled radiance value.
CUDA kernels
__global__ void	kernel_FindIntersections (RayChunk rayChunk, ShadingPoints shadingPts, uint *d_outIsValid)
	Searches ray hit points for given ray chunk.
__global__ void	kernel_SampleAreaLight (uint numShadingPts, float idxSampleX, float invNumSamplesX, float idxSampleY, float invNumSamplesY, float d_randoms1, float d_randoms2, float4 *d_outSamplePts)
	Samples points on area light.
__global__ void	kernel_TraceShadowRaysDelta (ShadingPoints shadingPts, uint *d_outShadowRayResult)
	Traces shadow rays to delta light sources.
__global__ void	kernel_TraceShadowRaysArea (ShadingPoints shadingPts, float4 d_samplePts, uint d_outShadowRayResult)
	Traces shadow rays to area light sources.
__global__ void	kernel_TracePhotons (PhotonData photons, uint d_outIsValid, uint d_outHasNonSpecular, int d_outTriHitIndex, float2 d_outHitBary, float4 d_outHitDiffClr, float4 d_outHitSpecClr)
	Traces photons into the scene.
__global__ void	kernel_AddDirectRadiance (RayChunk rayChunk, ShadingPoints shadingPts, uint d_ShadowRayResult, float4 d_lightSamples, float fScale, float4 *d_ioRadiance)
	Adds direct radiance from primary light source.
__global__ void	kernel_GetDiffuseColors (int d_triHitIndices, float2 d_baryHit, uint numPoints, float4 *d_outClrDiffHit)
	Generates array of diffuse material colors for given hits.
__global__ void	kernel_AddEmittedAndIndirect (float4 d_indirectIllum, RayChunk rayChunk, ShadingPoints shadingPts, float4 d_ioRadiance)
	Adds emitting and indirect component of light transport equation.
__global__ void	kernel_ApproximateNormalAt (KDFinalNodeList lstFinal, float queryRadiusMax, float4 *d_outNormals)
	Approximates normals at the center of each tree node.
Kernel wrappers
void	KernelRTTraceRays (const RayChunk &rayChunk, ShadingPoints &outInters, uint *d_outIsValid)
	Wraps kernel_FindIntersections() kernel call.
void	KernelRTEvaluateLTE (const RayChunk &rayChunk, const ShadingPoints &shadingPts, const LightData &lights, float4 d_radianceIndirect, bool bDirectRT, bool bTraceShadowRays, uint2 areaLightSamples, float4 d_ioRadiance)
	Evaluates the light transport equation.
void	KernelRTTracePhotons (PhotonData &photons, uint d_outIsValid, uint d_outHasNonSpecular, int d_outTriHitIndex, float2 d_outHitBary, float4 d_outHitDiffClr, float4 d_outHitSpecClr)
	Wraps kernel_TracePhotons() kernel call.
void	KernelRTGetDiffuseColors (int d_triHitIndices, float2 d_baryHit, uint numPoints, float4 *d_outClrDiffHit)
	Wraps kernel_GetDiffuseColors() kernel call.
void	KernelRTApproximateNormalAt (const KDFinalNodeList &lstFinal, float queryRadiusMax, float4 *d_outNormals)
	Wraps kernel_ApproximateNormalAt() kernel call.
Variables
__constant__ LightData	c_LightData
	Light data constant memory variable.
__constant__ MaterialProperties	c_MatProps
	Material properties constant memory variable.
__constant__ TriangleData	c_TriData
	Traingle data constant memory variable.
__constant__ KDTreeData	c_KDTree
	Object kd-tree data constant memory variable.
__constant__ float	c_fRayEpsilon = 1e-3f
	Ray epsilon.
texture< float4, 1, cudaReadModeElementType >	tex_TriV0
	First triangle vertices.
texture< float4, 1, cudaReadModeElementType >	tex_TriV1
	Second triangle vertices.
texture< float4, 1, cudaReadModeElementType >	tex_TriV2
	Third triangle vertices.
texture< float4, 1, cudaReadModeElementType >	tex_TriN0
	First vertex triangle normals texture.
texture< uint, 1, cudaReadModeElementType >	tex_TriMatIdx
	Triangle material indices. One per triangle.
texture< float2, 1, cudaReadModeElementType >	tex_TriTexCoordA
	UV-texture coordinate at first triangle vertices.
texture< float2, 1, cudaReadModeElementType >	tex_TriTexCoordB
	UV-texture coordinate at second triangle vertices.
texture< float2, 1, cudaReadModeElementType >	tex_TriTexCoordC
	UV-texture coordinate at third triangle vertices.
texture< uint, 1, cudaReadModeElementType >	tex_kdTree
	Object kd-tree texture for KDTreeData::d_preorderTree.
uint	f_numDiffTextures = 0
	Current number of loaded diffuse textures.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex0
	Diffuse texture 1.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex1
	Diffuse texture 2. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex2
	Diffuse texture 3. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex3
	Diffuse texture 4. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex4
	Diffuse texture 5. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex5
	Diffuse texture 6. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex6
	Diffuse texture 7. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex7
	Diffuse texture 8. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex8
	Diffuse texture 9. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex9
	Diffuse texture 10. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex10
	Diffuse texture 11. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex11
	Diffuse texture 12. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex12
	Diffuse texture 13. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex13
	Diffuse texture 14. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex14
	Diffuse texture 15. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex15
	Diffuse texture 16. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex16
	Diffuse texture 17. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex17
	Diffuse texture 18. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex18
	Diffuse texture 19. See tex_diffTex0.
texture< uchar4, 2, cudaReadModeNormalizedFloat >	tex_diffTex19
	Diffuse texture 20. See tex_diffTex0.

Detailed Description

Kernels for ray tracing.

Todo:: Evaluate shared memory usage and compability with older GPUs.

Author:: Mathias Neumann

Date:: 30.01.2010

Define Documentation

#define BINDTEX2ARRAY	(	numTex,
		idx
	)

Value:

if(numTex > idx) \
    { \
        MNAssert(idx >= 0 && idx < MAX_DIFF_TEX_COUNT); \
        GENERATETEXVAR(idx).addressMode[0] = cudaAddressModeWrap; \
        GENERATETEXVAR(idx).addressMode[1] = cudaAddressModeWrap; \
        GENERATETEXVAR(idx).filterMode = cudaFilterModeLinear; \
        GENERATETEXVAR(idx).normalized = true; \
        mncudaSafeCallNoSync(cudaBindTextureToArray(GENERATETEXVAR(idx), mats.vecTexArrays[Tex_Diffuse][idx], cdClrs)); \
    }

Binds the given texture reference to a CUDA texture array.

This macro was introduced for convenience.

Parameters:

numTex	Total number of textures. Used to avoid illegal bindings.
idx	Zero-based index the texture to bind.

#define GENERATETEXVAR ( idx ) tex_diffTex ## idx

Macro to generate diffuse texture reference names.

This macro was introduced for convenience.

Parameters:

idx	Zero-based index the texture reference.

#define INTERSECT_BLOCKSIZE 128

Thread block size used for intersection search.

For kernels calling dev_FindNextIntersectionKDWhileWhile().

#define MAX_DIFF_TEX_COUNT 20

Maximum number of diffuse textures supported.

The number of diffuse textures has to be limited to a small amount due to the problem that arrays of textures are not possible, yet.

Todo:: Right now, arrays of texture references are not possible in CUDA. Hence I currently use many single texture references to reach diffuse or bump map texture support in MNRT. This led to quite a lot of redundant code. I tried to reduce this using macros, but that only worked to some extent. Accordingly it would be relieving to find a better solution for texture memory variables for real textures (e.g. bump maps). As discussed for tex_diffTex0, the solution of moving the real textures into slices of a 3D CUDA array worked, but also had several major drawbacks.

#define UNBINDTEX	(	numTex,
		idx
	)

Value:

if(numTex > idx) \
        mncudaSafeCallNoSync(cudaUnbindTexture(GENERATETEXVAR(idx)));

Unbinds the given texture reference.

This macro was introduced for convenience.

Parameters:

numTex	Total number of textures. Used to avoid illegal bindings.
idx	Zero-based index the texture to bind.

Function Documentation

__device__ void dev_AddPixelRadiance	(	const RayChunk &	rayChunk,
		const ShadingPoints &	shadingPts,
		const uint	tid,
		const float3	L_sample,
		float4 *	d_ioRadiance
	)

Updates the radiance accumulator by adding a scaled radiance value.

Author:: Mathias Neumann

Date:: 27.06.2010

Parameters:

	rayChunk	Source ray chunk. The influence component determines the scale factor.
	shadingPts	The shading point data. Pixel component gives the index of the radiance accumulator.
	tid	Index of source ray and shading point.
	L_sample	Unscaled radiance value to add.
[in,out]	d_ioRadiance	Radiance accumulator (screen buffer).

template<bool getMinDist>

template< bool getMinDist > __device__ float4 dev_ApproximateNormalAt	(	float3	ptQuery,
		float	queryRadiusSqr
	)

Approximates normal at a given query point.

A range search is performed in the object kd-tree.

Todo:: Improve normal approximation.

Author:: Mathias Neumann

Date:: 23.07.2010

Template Parameters:

getMinDist If true is passed, the minimum "distance" to triangles in the environment is computed and provided in the w-component of the returned value. The xyz-components contain the geometric normal of the "closest" triangle. If false is passed, a weighted interpolation of normals is performed.

Parameters:

ptQuery	The query point.
queryRadiusSqr	The query radius (squared).

Returns:: xyz: Approximated normal at ptQuery. w: Minimum "distance" to a triangle if getMinDist = true.

__device__ float4 dev_FetchDiffTex	(	int	idxTex,
		float	u,
		float	v
	)

Fetches data from a given diffuse texture.

Author:: Mathias Neumann

Date:: April 2010

See also:: tex_diffTex0

Parameters:

idxTex	Zero-based texture index.
u	The u texture coordinate value.
v	The v texture coordinate value.

Returns:: The fetched texture value at given coordinates.

template<bool needClosest>

__device__ int dev_FindNextIntersectionKDWhileWhile	(	const float3	rayO,
		const float3	rayD,
		const float	tMinRay,
		const float	tMaxRay,
		float &	outLambda,
		float2 &	outBaryHit
	)		`[inline]`

Looks for the next ray triangle intersection using a while-while standard traversal.

Note:: Required shared memory per thread block: 16 * INTERSECT_BLOCKSIZE bytes.

Author:: Mathias Neumann

Date:: 04.03.2010

Parameters:

	rayO	The ray origin.
	rayD	The ray direction (normalized).
	tMinRay	The ray segment minimum.
	tMaxRay	The ray segment maximum.
[out]	outLambda	Intersection parameter. Only valid if returned value is not -1.
[out]	outBaryHit	Barycentric hit coordinate. Only valid if returned value is not -1. Extracting the hit coordinate computation by returning just the hit index worsens performance, even if registers are saved.

Returns:: Returns index of the first intersected triangle or -1, if no intersection found.

__device__ float3 dev_GetColorDiffuse	(	uint	idxTri,
		uint	idxMaterial,
		char4	matFlags,
		float2	baryHit
	)

Determines the diffuse material color for given triangle hit.

When there is a diffuse texture for the given material, the color is fetched from that texture. Else the material's diffuse color is used. For area light materials, a white color is returned right now.

Author:: Mathias Neumann

Date:: April 2010

Parameters:

idxTri	The triangle index.
idxMaterial	The material index.
matFlags	Material flags array. See MaterialProperties::flags.
baryHit	Barycentric hit coordinates.

Returns:: Diffuse material color (reflectance).

__device__ float3 dev_GetColorSpec	(	uint	idxTri,
		uint	idxMaterial,
		char4	matFlags,
		float2	baryHit,
		float *	outSpecExp
	)

Determines the diffuse material color for given triangle hit.

Right now, only the material's color is used. There is no support for specular textures.

Author:: Mathias Neumann

Date:: April 2010

Parameters:

	idxTri	The triangle index.
	idxMaterial	The material index.
	matFlags	Material flags array. See MaterialProperties::flags.
	baryHit	Barycentric hit coordinates.
[out]	outSpecExp	Materials specular exponent (shininess).

Returns:: Specular material color (reflectance).

__device__ float3 dev_GetReflectedDirectLight	(	float3	ptEye,
		float3	pt,
		float3	nG,
		float3	nS,
		float3	ptLightSample,
		uint	idxTri,
		uint	idxMat,
		char4	matFlags,
		float2	baryHit
	)

Computes the reflected direct light from pt to ptEye.

Todo:: Currently only diffuse materials are considered. Implement support for other BSDFs.

Author:: Mathias Neumann

Date:: 25.10.2010

Parameters:

ptEye	Eye position.
pt	Position from which the reflected direct light shall be evaluated.
nG	Geometric normal at surface in pt.
nS	Shading normal at surface in pt.
ptLightSample	Sampled point on light source.
idxTri	Index of the triangle in pt.
idxMat	Index of the material in pt.
matFlags	Material flags array. See MaterialProperties::flags.
baryHit	Barycentric hit coordinates.

Returns:: Reflected direct radiance from primary light source.

__device__ float3 dev_SampleLightL	(	float3	pt,
		float3	ptLightSample,
		float3 *	outW_i
	)

Samples the only primary light source.

The light source is defined by the constant memory variable c_LightData.

Author:: Mathias Neumann

Date:: March 2010

Parameters:

	pt	The point for which the incident radiance should be calculated.
	ptLightSample	Sampled point on light source.
[out]	outW_i	Incident light direction (pointing away from pt).

Returns:: Incident, direct radiance from primary light source.

__global__ void kernel_AddDirectRadiance	(	RayChunk	rayChunk,
		ShadingPoints	shadingPts,
		uint *	d_ShadowRayResult,
		float4 *	d_lightSamples,
		float	fScale,
		float4 *	d_ioRadiance
	)

Adds direct radiance from primary light source.

The function dev_GetReflectedDirectLight() is used to evaluate the direct radiance reflected from each shading point into the direction given by the source ray.

Author:: Mathias Neumann

Date:: February 2010

Parameters:

	rayChunk	Source ray chunk. Compacted, so that rays hitting nothing are removed.
	shadingPts	The shading points. Contains corresponding hits for ray chunk. Compacted, so that invalid hits are removed.
[in]	d_ShadowRayResult	The shadow ray result. Binary 0/1 buffer. Can be generated using kernel_TraceShadowRaysArea() for area lights and kernel_TraceShadowRaysDelta() for delta lights.
[in]	d_lightSamples	Sample point on area light sources for each shading point. Set to `NULL` for delta light sources.
	fScale	The scale factor. Radiance will be scaled by this factor, before it is added to the accumulator. Can be used for Monte-Carlo integration.
[in,out]	d_ioRadiance	Radiance accumulator screen buffer, i.e. elements are associated to screen's pixels.

__global__ void kernel_AddEmittedAndIndirect	(	float4 *	d_indirectIllum,
		RayChunk	rayChunk,
		ShadingPoints	shadingPts,
		float4 *	d_ioRadiance
	)

Adds emitting and indirect component of light transport equation.

The emitting component is handled by checking if the material at the hit points is an area light material. The indirect component is assumed to be computed by some other function and is only added in by this kernel.

Author:: Mathias Neumann

Date:: 26.06.2010

Parameters:

[in]	d_indirectIllum	Computed indirect illumination for each source ray.
	rayChunk	Source ray chunk. Compacted, so that rays hitting nothing are removed.
	shadingPts	The shading points. Contains corresponding hits for ray chunk. Compacted, so that invalid hits are removed.
[in,out]	d_ioRadiance	Radiance accumulator screen buffer, i.e. elements are associated to screen's pixels.

__global__ void kernel_ApproximateNormalAt	(	KDFinalNodeList	lstFinal,
		float	queryRadiusMax,
		float4 *	d_outNormals
	)

Approximates normals at the center of each tree node.

Calls dev_ApproximateNormalAt() twice. First, the closest triangle intersection is determined. It that intersection is very close to the node center. The triangle's geometric normal is used as approximated normal. Else the device function is called again, now to perform a weighted interpolation of normals of nearby triangles.

Author:: Mathias Neumann

Date:: 23.07.2010

Parameters:

	lstFinal	The final kd-tree node list.
	queryRadiusMax	The query radius maximum to use for queries in the object kd-tree.
[out]	d_outNormals	Approximated normal for each node. Can be zero, if approximation failed.

__global__ void kernel_FindIntersections	(	RayChunk	rayChunk,
		ShadingPoints	shadingPts,
		uint *	d_outIsValid
	)

Searches ray hit points for given ray chunk.

Calls dev_FindNextIntersectionKDWhileWhile() for intersection search.

Author:: Mathias Neumann

Date:: February 2010

Parameters:

	rayChunk	The ray chunk.
	shadingPts	Target data structure for hit points. Is assumed to be empty. For the i-th ray, the i-th entry in this structure is used.
[out]	d_outIsValid	Binary 0/1 array. Will contain 1 for valid hit points and 0 for invalid hit points.

__global__ void kernel_GetDiffuseColors	(	int *	d_triHitIndices,
		float2 *	d_baryHit,
		uint	numPoints,
		float4 *	d_outClrDiffHit
	)

Generates array of diffuse material colors for given hits.

This method was added as I wanted to reduce redundancy. Material properties including diffuse textures are available in this file, but not in other cu-files.

Author:: Mathias Neumann

Date:: 25.06.2010

Parameters:

[in]	d_triHitIndices	Triangle index for each hit. Should contain -1 for invalid hits.
[in]	d_baryHit	Barycentric coordinates for ecah hit.
	numPoints	Number of hits.
[out]	d_outClrDiffHit	Diffuse colors at hits, will be invalid if no hit. `xyz` contains color and `w` transparency alpha.

__global__ void kernel_SampleAreaLight	(	uint	numShadingPts,
		float	idxSampleX,
		float	invNumSamplesX,
		float	idxSampleY,
		float	invNumSamplesY,
		float *	d_randoms1,
		float *	d_randoms2,
		float4 *	d_outSamplePts
	)

Samples points on area light.

Assumes that c_LightData represents an area light source.

Author:: Mathias Neumann

Date:: June 2010

Parameters:

	numShadingPts	Number of shading points.
	idxSampleX	Index of sample X (for stratified sampling).
	invNumSamplesX	Inverse number of samples X.
	idxSampleY	Index of sample Y (for stratified sampling).
	invNumSamplesY	Inverse number of samples Y.
[in]	d_randoms1	First array of uniform random numbers, one per shading point.
[in]	d_randoms2	Second array of uniform random numbers, one per shading point.
[out]	d_outSamplePts	Sampled point on light source area for each shading point.

__global__ void kernel_TracePhotons	(	PhotonData	photons,
		uint *	d_outIsValid,
		uint *	d_outHasNonSpecular,
		int *	d_outTriHitIndex,
		float2 *	d_outHitBary,
		float4 *	d_outHitDiffClr,
		float4 *	d_outHitSpecClr
	)

Traces photons into the scene.

Uses dev_FindNextIntersectionKDWhileWhile() to find the next intersection.

Author:: Mathias Neumann

Date:: 07.04.2010

Parameters:

	photons	The photons to trace. Works inplace, that is, the new photons replace the old ones.
[out]	d_outIsValid	Binary 0/1 array. Contains 1 iff the corresponding photon is valid. A photon is considered as invalid if it found no intersection or if its flux is zero.
[out]	d_outHasNonSpecular	Binary 0/1 array. If 1, the photon intersected a surface that has non- specular components. Else, even for no intersections, the element is set to 0.
[out]	d_outTriHitIndex	Triangle index of hit triangle for each photon. Will be -1 if no hit.
[out]	d_outHitBary	Barycentric hit coordinates. Will be invalid if no hit.
[out]	d_outHitDiffClr	Diffuse colors at hits, will be invalid if no hit. Used for reflected photon flux. I decided to move this to this file because here I have access to texture and color data. `xyz` contains color and `w` transparency alpha.
[out]	d_outHitSpecClr	Specular colors at hits, will be invalid if no hit. Color in `xyz` and index of refraction in `w`.

__global__ void kernel_TraceShadowRaysArea	(	ShadingPoints	shadingPts,
		float4 *	d_samplePts,
		uint *	d_outShadowRayResult
	)

Traces shadow rays to area light sources.

Assumes that c_LightData represents an area light source. Shadow rays are traced using dev_FindNextIntersectionKDWhileWhile().

Author:: Mathias Neumann

Date:: June 2010

Parameters:

	shadingPts	The shading points.
[in]	d_samplePts	Contains the sample point on the area of the light source (for each shading point). Can be generated using kernel_SampleAreaLight().
[out]	d_outShadowRayResult	The result of shadow ray tracing. Binary 0/1 array. Contains 1 iff the light source is unoccluded for a given shading point.

__global__ void kernel_TraceShadowRaysDelta	(	ShadingPoints	shadingPts,
		uint *	d_outShadowRayResult
	)

Traces shadow rays to delta light sources.

Assumes that c_LightData represents a delta light source, i.e. a light source where only one direction to the light source is possible (point light, directional light). Shadow rays are traced using dev_FindNextIntersectionKDWhileWhile().

Author:: Mathias Neumann

Date:: June 2010

Parameters:

	shadingPts	The shading points.
[out]	d_outShadowRayResult	Contains the result of shadow ray tracing. Binary 0/1 array. Contains 1 iff the light source is unoccluded for a given shading point.

void KernelRTEvaluateLTE	(	const RayChunk &	rayChunk,
		const ShadingPoints &	shadingPts,
		const LightData &	lights,
		float4 *	d_radianceIndirect,
		bool	bDirectRT,
		bool	bTraceShadowRays,
		uint2	areaLightSamples,
		float4 *	d_ioRadiance
	)

Evaluates the light transport equation.

Actually, only emitted and direct illumination are evaluated. The indirect illumination is assumed to be computed elsewhere and only passed as parameter. For the direct part, evaluation depends on light source type. For delta light sources, one shadow ray is enough. For area light sources, Monte-Carlo integration is performed to evaluate the illumination integral.

Author:: Mathias Neumann

Date:: 25.10.2010

See also:: kernel_TraceShadowRaysDelta(), kernel_SampleAreaLight(), kernel_TraceShadowRaysArea(), kernel_AddDirectRadiance(), kernel_AddEmittedAndIndirect()

Parameters:

	rayChunk	Source ray chunk. Compacted, so that rays hitting nothing are removed.
	shadingPts	The shading points. Contains corresponding hits for ray chunk. Compacted, so that invalid hits are removed.
	lights	Current scene's light data.
[in]	d_radianceIndirect	Computed indirect illumination for each source ray.
	bDirectRT	Whether to evaluate direct illumination using ray tracing. If false is passed, only indirect and emitted illumination are computed.
	bTraceShadowRays	Whether to trace shadow rays for direct lighting. If `false` is passed, lights are assumed to be unoccluded.
	areaLightSamples	Number of area light samples to take. Value is two- dimensional to allow stratification of samples.
[in,out]	d_ioRadiance	Radiance accumulator screen buffer, i.e. elements are associated to screen's pixels.

void RTUpdateKernelData	(	const LightData &	lights,
		const TriangleData &	tris,
		const MaterialData &	mats,
		const KDTreeData &	kdTree,
		float	fRayEpsilon
	)

Moves scene data to constant memory and texture memory.

Author:: Mathias Neumann

Date:: February 2010

Parameters:

lights	Current scene's light data.
tris	Current scene's triangle data.
mats	Current scene's material data.
kdTree	Current scene's object kd-tree data.
fRayEpsilon	Ray epsilon for current scene.

Variable Documentation

__constant__ float c_fRayEpsilon = 1e-3f

Ray epsilon.

See SceneConfig::SetRayEpsilon().

uint f_numDiffTextures = 0

Current number of loaded diffuse textures.

Limited by MAX_DIFF_TEX_COUNT.

texture<uchar4, 2, cudaReadModeNormalizedFloat> tex_diffTex0

Diffuse texture 1.

I tried to move textures into some kind of array, but this failed since CUDA doesn't support arrays of texture samplers, check the CUDA FAQ. This FAQ proposed two ways to eliminate this problem:

First: 3D array, 2D textures into slices

I tried this and it worked. But it has some serious drawbacks:

Forced to use cudaFilterModePoint to avoid fetching from different textures.
No normalized addressing possbile.
Texture dimension limited to 2k x 2k.
Wasting of memory: The 3D tex has to as large as the largest texture. Example: One texture 2048x2048 (16 MByte), 10 others 512x512 (10 MByte) would lead to 16 * 11 = 176 MByte instead of 16 + 10 = 26 MByte.

Hence I dropped this.

Second: One texture reference for each texture and switch-instruction when fetching.

Also some horrible drawbacks:

Fixed maximum number of textures.
A lot of code redundancy.

Despite these problems I'll give the second approach a try.

texture<float4, 1, cudaReadModeElementType> tex_TriN0

First vertex triangle normals texture.

Required for orientation of geometric normals in dev_ApproximateNormalAt().

GPU/raytracing.cu File Reference

Defines

Functions

Variables

Detailed Description

Define Documentation

Function Documentation

Variable Documentation