summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorStan Seibert <stan@mtrr.org>2013-01-31 11:17:16 -0500
committertlatorre <tlatorre@uchicago.edu>2021-05-09 08:42:39 -0700
commitf91c0174046bffa95c7d42ea945c47ced48ed892 (patch)
treead732e8e41a56e81b3b60eb66ccc6eeb333ec3c2
parente8f7f1cb5bba52600c05a110b3605232a28744ab (diff)
downloadchroma-f91c0174046bffa95c7d42ea945c47ced48ed892.tar.gz
chroma-f91c0174046bffa95c7d42ea945c47ced48ed892.tar.bz2
chroma-f91c0174046bffa95c7d42ea945c47ced48ed892.zip
Fix bug in intersection calculation that prevented axis-aligned photons from hitting triangles.
-rw-r--r--chroma/cuda/intersect.h36
1 files changed, 22 insertions, 14 deletions
diff --git a/chroma/cuda/intersect.h b/chroma/cuda/intersect.h
index 132b79f..0d4ba8a 100644
--- a/chroma/cuda/intersect.h
+++ b/chroma/cuda/intersect.h
@@ -64,34 +64,42 @@ intersect_triangle(const float3 &origin, const float3 &direction,
Source: Optimizing ray tracing for CUDA by Hannu Saransaari
https://wiki.aalto.fi/download/attachments/40023967/gpgpu.pdf
*/
+#define INFINITY __int_as_float(0x7f800000)
+
__device__ bool
intersect_box(const float3 &neg_origin_inv_dir, const float3 &inv_dir,
const float3 &lower_bound, const float3 &upper_bound,
float& distance_to_box)
{
- float tmin = 0.0f, tmax = 1e30f;
+ float tmin = 0.0f, tmax = INFINITY;
float t0, t1;
// X
- t0 = lower_bound.x * inv_dir.x + neg_origin_inv_dir.x;
- t1 = upper_bound.x * inv_dir.x + neg_origin_inv_dir.x;
-
- tmin = max(tmin, min(t0, t1));
- tmax = min(tmax, max(t0, t1));
+ if (isfinite(inv_dir.x)) {
+ t0 = lower_bound.x * inv_dir.x + neg_origin_inv_dir.x;
+ t1 = upper_bound.x * inv_dir.x + neg_origin_inv_dir.x;
+ tmin = max(tmin, min(t0, t1));
+ tmax = min(tmax, max(t0, t1));
+ }
+
// Y
- t0 = lower_bound.y * inv_dir.y + neg_origin_inv_dir.y;
- t1 = upper_bound.y * inv_dir.y + neg_origin_inv_dir.y;
+ if (isfinite(inv_dir.y)) {
+ t0 = lower_bound.y * inv_dir.y + neg_origin_inv_dir.y;
+ t1 = upper_bound.y * inv_dir.y + neg_origin_inv_dir.y;
- tmin = max(tmin, min(t0, t1));
- tmax = min(tmax, max(t0, t1));
+ tmin = max(tmin, min(t0, t1));
+ tmax = min(tmax, max(t0, t1));
+ }
// Z
- t0 = lower_bound.z * inv_dir.z + neg_origin_inv_dir.z;
- t1 = upper_bound.z * inv_dir.z + neg_origin_inv_dir.z;
+ if (isfinite(inv_dir.z)) {
+ t0 = lower_bound.z * inv_dir.z + neg_origin_inv_dir.z;
+ t1 = upper_bound.z * inv_dir.z + neg_origin_inv_dir.z;
- tmin = max(tmin, min(t0, t1));
- tmax = min(tmax, max(t0, t1));
+ tmin = max(tmin, min(t0, t1));
+ tmax = min(tmax, max(t0, t1));
+ }
if (tmin > tmax)
return false;