Defines | Functions | Variables

GPU/raytracing.cu File Reference

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:
numTexTotal number of textures. Used to avoid illegal bindings.
idxZero-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:
idxZero-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:
numTexTotal number of textures. Used to avoid illegal bindings.
idxZero-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:
rayChunkSource ray chunk. The influence component determines the scale factor.
shadingPtsThe shading point data. Pixel component gives the index of the radiance accumulator.
tidIndex of source ray and shading point.
L_sampleUnscaled radiance value to add.
[in,out]d_ioRadianceRadiance 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:
getMinDistIf 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:
ptQueryThe query point.
queryRadiusSqrThe 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:
idxTexZero-based texture index.
uThe u texture coordinate value.
vThe 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:
rayOThe ray origin.
rayDThe ray direction (normalized).
tMinRayThe ray segment minimum.
tMaxRayThe ray segment maximum.
[out]outLambdaIntersection parameter. Only valid if returned value is not -1.
[out]outBaryHitBarycentric 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:
idxTriThe triangle index.
idxMaterialThe material index.
matFlagsMaterial flags array. See MaterialProperties::flags.
baryHitBarycentric 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:
idxTriThe triangle index.
idxMaterialThe material index.
matFlagsMaterial flags array. See MaterialProperties::flags.
baryHitBarycentric hit coordinates.
[out]outSpecExpMaterials 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:
ptEyeEye position.
ptPosition from which the reflected direct light shall be evaluated.
nGGeometric normal at surface in pt.
nSShading normal at surface in pt.
ptLightSampleSampled point on light source.
idxTriIndex of the triangle in pt.
idxMatIndex of the material in pt.
matFlagsMaterial flags array. See MaterialProperties::flags.
baryHitBarycentric 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:
ptThe point for which the incident radiance should be calculated.
ptLightSampleSampled point on light source.
[out]outW_iIncident 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:
rayChunkSource ray chunk. Compacted, so that rays hitting nothing are removed.
shadingPtsThe shading points. Contains corresponding hits for ray chunk. Compacted, so that invalid hits are removed.
[in]d_ShadowRayResultThe shadow ray result. Binary 0/1 buffer. Can be generated using kernel_TraceShadowRaysArea() for area lights and kernel_TraceShadowRaysDelta() for delta lights.
[in]d_lightSamplesSample point on area light sources for each shading point. Set to NULL for delta light sources.
fScaleThe 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_ioRadianceRadiance 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_indirectIllumComputed indirect illumination for each source ray.
rayChunkSource ray chunk. Compacted, so that rays hitting nothing are removed.
shadingPtsThe shading points. Contains corresponding hits for ray chunk. Compacted, so that invalid hits are removed.
[in,out]d_ioRadianceRadiance 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:
lstFinalThe final kd-tree node list.
queryRadiusMaxThe query radius maximum to use for queries in the object kd-tree.
[out]d_outNormalsApproximated 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:
rayChunkThe ray chunk.
shadingPtsTarget 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_outIsValidBinary 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_triHitIndicesTriangle index for each hit. Should contain -1 for invalid hits.
[in]d_baryHitBarycentric coordinates for ecah hit.
numPointsNumber of hits.
[out]d_outClrDiffHitDiffuse 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:
numShadingPtsNumber of shading points.
idxSampleXIndex of sample X (for stratified sampling).
invNumSamplesXInverse number of samples X.
idxSampleYIndex of sample Y (for stratified sampling).
invNumSamplesYInverse number of samples Y.
[in]d_randoms1First array of uniform random numbers, one per shading point.
[in]d_randoms2Second array of uniform random numbers, one per shading point.
[out]d_outSamplePtsSampled 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:
photonsThe photons to trace. Works inplace, that is, the new photons replace the old ones.
[out]d_outIsValidBinary 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_outHasNonSpecularBinary 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_outTriHitIndexTriangle index of hit triangle for each photon. Will be -1 if no hit.
[out]d_outHitBaryBarycentric hit coordinates. Will be invalid if no hit.
[out]d_outHitDiffClrDiffuse 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_outHitSpecClrSpecular 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:
shadingPtsThe shading points.
[in]d_samplePtsContains the sample point on the area of the light source (for each shading point). Can be generated using kernel_SampleAreaLight().
[out]d_outShadowRayResultThe 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:
shadingPtsThe shading points.
[out]d_outShadowRayResultContains 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:
rayChunkSource ray chunk. Compacted, so that rays hitting nothing are removed.
shadingPtsThe shading points. Contains corresponding hits for ray chunk. Compacted, so that invalid hits are removed.
lightsCurrent scene's light data.
[in]d_radianceIndirectComputed indirect illumination for each source ray.
bDirectRTWhether to evaluate direct illumination using ray tracing. If false is passed, only indirect and emitted illumination are computed.
bTraceShadowRaysWhether to trace shadow rays for direct lighting. If false is passed, lights are assumed to be unoccluded.
areaLightSamplesNumber of area light samples to take. Value is two- dimensional to allow stratification of samples.
[in,out]d_ioRadianceRadiance 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:
lightsCurrent scene's light data.
trisCurrent scene's triangle data.
matsCurrent scene's material data.
kdTreeCurrent scene's object kd-tree data.
fRayEpsilonRay epsilon for current scene.

Variable Documentation

__constant__ float c_fRayEpsilon = 1e-3f

Ray epsilon.

See SceneConfig::SetRayEpsilon().

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

MNRT Source Code Documentation (Version 1.0) - Copyright © Mathias Neumann 2010