Skip to content

Commit

Permalink
[cl] Fix incorrect filtering of ray max dist in intersection code
Browse files Browse the repository at this point in the history
  • Loading branch information
achilleasa committed Aug 14, 2016
1 parent ba033fa commit c21b836
Show file tree
Hide file tree
Showing 2 changed files with 12 additions and 10 deletions.
2 changes: 1 addition & 1 deletion tracer/opencl/CL/constants.cl
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,6 @@

// Intersection constants
#define INTERSECTION_EPSILON 0.00001f
#define INTERSECTION_EPSILON_X2 (INTERSECTION_EPSILON * 2.0f)
#define INTERSECTION_WITH_LIGHT_EPSILON (INTERSECTION_EPSILON * 1e3f)

#endif
20 changes: 11 additions & 9 deletions tracer/opencl/CL/kernels/intersect.cl
Original file line number Diff line number Diff line change
Expand Up @@ -117,7 +117,7 @@ __kernel void rayIntersectionTest(
}

float t = dot(edge02, qVec) * invDet;
if (t > INTERSECTION_EPSILON){
if (t > INTERSECTION_EPSILON && t < ray.origin.w){
gotHit = 1;
stackIndex = -1;
break;
Expand Down Expand Up @@ -191,9 +191,6 @@ __kernel void rayIntersectionQuery(
__global Intersection* intersections
){

Intersection intersection;
intersection.wuvt.w = FLT_MAX;

int globalId = get_global_id(0);
if(globalId >= *numRays){
return;
Expand All @@ -220,6 +217,10 @@ __kernel void rayIntersectionQuery(
float3 origRayOrigin = ray.origin.xyz;
float3 origRayDir = ray.dir.xyz;

// Set initial intersection to the ray max dist
Intersection intersection;
intersection.wuvt.w = ray.origin.w;

// Setup stack
stackIndex = 0;
meshBvhStackStartIndex = -1;
Expand Down Expand Up @@ -341,7 +342,7 @@ __kernel void rayIntersectionQuery(
}

// Update hit flag
hitFlag[globalId] = intersection.wuvt.w < FLT_MAX ? 1 : 0;
hitFlag[globalId] = intersection.wuvt.w < ray.origin.w ? 1 : 0;
intersections[globalId] = intersection;
}

Expand All @@ -359,9 +360,6 @@ __kernel void rayPacketIntersectionQuery(
__global Intersection* intersections
){

Intersection intersection;
intersection.wuvt.w = FLT_MAX;

int globalId = get_global_id(0);
if (globalId >= *numRays){
return;
Expand Down Expand Up @@ -391,6 +389,10 @@ __kernel void rayPacketIntersectionQuery(
float3 origRayOrigin = ray.origin.xyz;
float3 origRayDir = ray.dir.xyz;

// Set initial intersection to the ray max dist
Intersection intersection;
intersection.wuvt.w = ray.origin.w;

// Traversal preferences
int packetWantsLeft, packetWantsRight;

Expand Down Expand Up @@ -568,7 +570,7 @@ __kernel void rayPacketIntersectionQuery(
}

// Update hit flag
hitFlag[globalId] = intersection.wuvt.w < FLT_MAX ? 1 : 0;
hitFlag[globalId] = intersection.wuvt.w < ray.origin.w ? 1 : 0;
intersections[globalId] = intersection;
}

Expand Down

0 comments on commit c21b836

Please sign in to comment.