diff options
| -rw-r--r-- | chroma/cuda/intersect.h | 91 | ||||
| -rw-r--r-- | chroma/cuda/mesh.h | 15 |
2 files changed, 36 insertions, 70 deletions
diff --git a/chroma/cuda/intersect.h b/chroma/cuda/intersect.h index 978bde8..132b79f 100644 --- a/chroma/cuda/intersect.h +++ b/chroma/cuda/intersect.h @@ -61,79 +61,42 @@ intersect_triangle(const float3 &origin, const float3 &direction, `distance_to_box` to the distance from `origin` to the intersection and return true, else return false. `direction` must be normalized to one. - Source: "An Efficient and Robust Ray-Box Intersection Algorithm." - by Williams, et. al. */ + Source: Optimizing ray tracing for CUDA by Hannu Saransaari + https://wiki.aalto.fi/download/attachments/40023967/gpgpu.pdf +*/ __device__ bool -intersect_box(const float3 &origin, const float3 &direction, +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 kmin, kmax, kymin, kymax, kzmin, kzmax; - - float divx = 1.0f/direction.x; - if (divx >= 0.0f) - { - kmin = (lower_bound.x - origin.x)*divx; - kmax = (upper_bound.x - origin.x)*divx; - } - else - { - kmin = (upper_bound.x - origin.x)*divx; - kmax = (lower_bound.x - origin.x)*divx; - } - - if (kmax < 0.0f) - return false; - - float divy = 1.0f/direction.y; - if (divy >= 0.0f) - { - kymin = (lower_bound.y - origin.y)*divy; - kymax = (upper_bound.y - origin.y)*divy; - } - else - { - kymin = (upper_bound.y - origin.y)*divy; - kymax = (lower_bound.y - origin.y)*divy; - } - - if (kymax < 0.0f) - return false; - - if (kymin > kmin) - kmin = kymin; - - if (kymax < kmax) - kmax = kymax; - - if (kmin > kmax) - return false; + float tmin = 0.0f, tmax = 1e30f; + float t0, t1; - float divz = 1.0f/direction.z; - if (divz >= 0.0f) - { - kzmin = (lower_bound.z - origin.z)*divz; - kzmax = (upper_bound.z - origin.z)*divz; - } - else - { - kzmin = (upper_bound.z - origin.z)*divz; - kzmax = (lower_bound.z - origin.z)*divz; - } - - if (kzmax < 0.0f) - return false; - - if (kzmin > kmin) - kmin = kzmin; + // 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; + + tmin = max(tmin, min(t0, t1)); + tmax = min(tmax, max(t0, t1)); - if (kzmax < kmax) - kmax = kzmax; + // 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)); - if (kmin > kmax) + if (tmin > tmax) return false; - distance_to_box = kmin; + distance_to_box = tmin; return true; } diff --git a/chroma/cuda/mesh.h b/chroma/cuda/mesh.h index 3b79768..70ffc23 100644 --- a/chroma/cuda/mesh.h +++ b/chroma/cuda/mesh.h @@ -13,12 +13,12 @@ is less than zero or the distance from `origin` to the intersection is less than `min_distance`, return true, else return false. */ __device__ bool -intersect_node(const float3 &origin, const float3 &direction, +intersect_node(const float3 &neg_origin_inv_dir, const float3 &inv_dir, Geometry *g, const Node &node, const float min_distance=-1.0f) { float distance_to_box; - if (intersect_box(origin, direction, node.lower, node.upper, + if (intersect_box(neg_origin_inv_dir, inv_dir, node.lower, node.upper, distance_to_box)) { if (min_distance < 0.0f) return true; @@ -49,7 +49,10 @@ intersect_mesh(const float3 &origin, const float3& direction, Geometry *g, Node root = get_node(g, 0); - if (!intersect_node(origin, direction, g, root, min_distance)) + float3 neg_origin_inv_dir = -origin / direction; + float3 inv_dir = 1.0f / direction; + + if (!intersect_node(neg_origin_inv_dir, inv_dir, g, root, min_distance)) return -1; unsigned int child_ptr_stack[STACK_SIZE]; @@ -71,7 +74,7 @@ intersect_mesh(const float3 &origin, const float3& direction, Geometry *g, if (node.kind == PADDING_NODE) break; // this node and rest of children are padding - if (intersect_node(origin, direction, g, node, min_distance)) { + if (intersect_node(neg_origin_inv_dir, inv_dir, g, node, min_distance)) { if (node.kind == LEAF_NODE) { // This node wraps a triangle @@ -98,8 +101,8 @@ intersect_mesh(const float3 &origin, const float3& direction, Geometry *g, } // hit node? if (curr >= STACK_SIZE) { - printf("warning: intersect_mesh() aborted; node > tail\n"); - break; + printf("warning: intersect_mesh() aborted; node > tail\n"); + break; } } // loop over children, starting with first_child |