Cycles: refactor rays to have start and end distance, fix precision issues

For transparency, volume and light intersection rays, adjust these distances
rather than the ray start position. This way we increment the start distance
by the smallest possible float increment to avoid self intersections, and be
sure it works as the distance compared to be will be exactly the same as
before, due to the ray start position and direction remaining the same.

Fix T98764, T96537, hair ray tracing precision issues.

Differential Revision: https://developer.blender.org/D15455

9 files changed, 158 insertions, 65 deletions

diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h
index f375529a6f6..9972de86c47 100644
--- a/intern/cycles/kernel/bvh/bvh.h
+++ b/intern/cycles/kernel/bvh/bvh.h
@@ -175,8 +175,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
   optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
              ray->P,
              ray->D,
-             0.0f,
-             ray->t,
+             ray->tmin,
+             ray->tmax,
              ray->time,
              ray_mask,
              ray_flags,
@@ -203,28 +203,28 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
 #elif defined(__METALRT__)
 
   if (!scene_intersect_valid(ray)) {
-    isect->t = ray->t;
+    isect->t = ray->tmax;
     isect->type = PRIMITIVE_NONE;
     return false;
   }
 
 #  if defined(__KERNEL_DEBUG__)
   if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
-    isect->t = ray->t;
+    isect->t = ray->tmax;
     isect->type = PRIMITIVE_NONE;
     kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
     return false;
   }
 
   if (is_null_intersection_function_table(metal_ancillaries->ift_default)) {
-    isect->t = ray->t;
+    isect->t = ray->tmax;
     isect->type = PRIMITIVE_NONE;
     kernel_assert(!"Invalid ift_default");
     return false;
   }
 #  endif
 
-  metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
   metalrt_intersector_type metalrt_intersect;
 
   if (!kernel_data.bvh.have_curves) {
@@ -263,7 +263,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
 #  endif
 
   if (intersection.type == intersection_type::none) {
-    isect->t = ray->t;
+    isect->t = ray->tmax;
     isect->type = PRIMITIVE_NONE;
 
     return false;
@@ -296,7 +296,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
 
 #  ifdef __EMBREE__
   if (kernel_data.device_bvh) {
-    isect->t = ray->t;
+    isect->t = ray->tmax;
     CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
     IntersectContext rtc_ctx(&ctx);
     RTCRayHit ray_hit;
@@ -360,8 +360,8 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
   optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
              ray->P,
              ray->D,
-             0.0f,
-             ray->t,
+             ray->tmin,
+             ray->tmax,
              ray->time,
              0xFF,
              /* Need to always call into __anyhit__kernel_optix_local_hit. */
@@ -405,7 +405,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
   }
 #    endif
 
-  metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
   metalrt_intersector_type metalrt_intersect;
 
   metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
@@ -476,7 +476,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
         float3 dir = ray->D;
         float3 idir = ray->D;
         Transform ob_itfm;
-        rtc_ray.tfar = ray->t *
+        rtc_ray.tfar = ray->tmax *
                        bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
         /* bvh_instance_motion_push() returns the inverse transform but
          * it's not needed here. */
@@ -542,8 +542,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
   optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
              ray->P,
              ray->D,
-             0.0f,
-             ray->t,
+             ray->tmin,
+             ray->tmax,
              ray->time,
              ray_mask,
              /* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */
@@ -582,7 +582,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
   }
 #    endif
 
-  metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
   metalrt_intersector_type metalrt_intersect;
 
   metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
@@ -701,8 +701,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
   optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
              ray->P,
              ray->D,
-             0.0f,
-             ray->t,
+             ray->tmin,
+             ray->tmax,
              ray->time,
              ray_mask,
              /* Need to always call into __anyhit__kernel_optix_volume_test. */
@@ -744,7 +744,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
   }
 #    endif
 
-  metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
   metalrt_intersector_type metalrt_intersect;
 
   metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
diff --git a/intern/cycles/kernel/bvh/embree.h b/intern/cycles/kernel/bvh/embree.h
index 77eec2468f4..fecbccac2f8 100644
--- a/intern/cycles/kernel/bvh/embree.h
+++ b/intern/cycles/kernel/bvh/embree.h
@@ -83,8 +83,8 @@ ccl_device_inline void kernel_embree_setup_ray(const Ray &ray,
   rtc_ray.dir_x = ray.D.x;
   rtc_ray.dir_y = ray.D.y;
   rtc_ray.dir_z = ray.D.z;
-  rtc_ray.tnear = 0.0f;
-  rtc_ray.tfar = ray.t;
+  rtc_ray.tnear = ray.tmin;
+  rtc_ray.tfar = ray.tmax;
   rtc_ray.time = ray.time;
   rtc_ray.mask = visibility;
 }
diff --git a/intern/cycles/kernel/bvh/local.h b/intern/cycles/kernel/bvh/local.h
index 3b6b30ea93d..017a241ef4a 100644
--- a/intern/cycles/kernel/bvh/local.h
+++ b/intern/cycles/kernel/bvh/local.h
@@ -47,8 +47,9 @@ ccl_device_inline
   float3 P = ray->P;
   float3 dir = bvh_clamp_direction(ray->D);
   float3 idir = bvh_inverse_direction(dir);
+  float tmin = ray->tmin;
   int object = OBJECT_NONE;
-  float isect_t = ray->t;
+  float isect_t = ray->tmax;
 
   if (local_isect != NULL) {
     local_isect->num_hits = 0;
@@ -59,10 +60,13 @@ ccl_device_inline
   if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
 #if BVH_FEATURE(BVH_MOTION)
     Transform ob_itfm;
-    isect_t *= bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
+    const float t_world_to_instance = bvh_instance_motion_push(
+        kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
 #else
-    isect_t *= bvh_instance_push(kg, local_object, ray, &P, &dir, &idir);
+    const float t_world_to_instance = bvh_instance_push(kg, local_object, ray, &P, &dir, &idir);
 #endif
+    isect_t *= t_world_to_instance;
+    tmin *= t_world_to_instance;
     object = local_object;
   }
 
@@ -81,6 +85,7 @@ ccl_device_inline
                                        dir,
 #endif
                                        idir,
+                                       tmin,
                                        isect_t,
                                        node_addr,
                                        PATH_RAY_ALL_VISIBILITY,
@@ -155,6 +160,7 @@ ccl_device_inline
                                            local_object,
                                            prim,
                                            prim_addr,
+                                           tmin,
                                            isect_t,
                                            lcg_state,
                                            max_hits)) {
@@ -191,6 +197,7 @@ ccl_device_inline
                                                   local_object,
                                                   prim,
                                                   prim_addr,
+                                                  tmin,
                                                   isect_t,
                                                   lcg_state,
                                                   max_hits)) {
diff --git a/intern/cycles/kernel/bvh/nodes.h b/intern/cycles/kernel/bvh/nodes.h
index c19dea9223b..e02841fad16 100644
--- a/intern/cycles/kernel/bvh/nodes.h
+++ b/intern/cycles/kernel/bvh/nodes.h
@@ -18,7 +18,8 @@ ccl_device_forceinline Transform bvh_unaligned_node_fetch_space(KernelGlobals kg
 ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
                                                       const float3 P,
                                                       const float3 idir,
-                                                      const float t,
+                                                      const float tmin,
+                                                      const float tmax,
                                                       const int node_addr,
                                                       const uint visibility,
                                                       float dist[2])
@@ -39,8 +40,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
   float c0hiy = (node1.z - P.y) * idir.y;
   float c0loz = (node2.x - P.z) * idir.z;
   float c0hiz = (node2.z - P.z) * idir.z;
-  float c0min = max4(0.0f, min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz));
-  float c0max = min4(t, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz));
+  float c0min = max4(tmin, min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz));
+  float c0max = min4(tmax, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz));
 
   float c1lox = (node0.y - P.x) * idir.x;
   float c1hix = (node0.w - P.x) * idir.x;
@@ -48,8 +49,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
   float c1hiy = (node1.w - P.y) * idir.y;
   float c1loz = (node2.y - P.z) * idir.z;
   float c1hiz = (node2.w - P.z) * idir.z;
-  float c1min = max4(0.0f, min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz));
-  float c1max = min4(t, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz));
+  float c1min = max4(tmin, min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz));
+  float c1max = min4(tmax, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz));
 
   dist[0] = c0min;
   dist[1] = c1min;
@@ -66,7 +67,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
 ccl_device_forceinline bool bvh_unaligned_node_intersect_child(KernelGlobals kg,
                                                                const float3 P,
                                                                const float3 dir,
-                                                               const float t,
+                                                               const float tmin,
+                                                               const float tmax,
                                                                int node_addr,
                                                                int child,
                                                                float dist[2])
@@ -83,8 +85,8 @@ ccl_device_forceinline bool bvh_unaligned_node_intersect_child(KernelGlobals kg,
   const float far_x = max(lower_xyz.x, upper_xyz.x);
   const float far_y = max(lower_xyz.y, upper_xyz.y);
   const float far_z = max(lower_xyz.z, upper_xyz.z);
-  const float tnear = max4(0.0f, near_x, near_y, near_z);
-  const float tfar = min4(t, far_x, far_y, far_z);
+  const float tnear = max4(tmin, near_x, near_y, near_z);
+  const float tfar = min4(tmax, far_x, far_y, far_z);
   *dist = tnear;
   return tnear <= tfar;
 }
@@ -93,7 +95,8 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg,
                                                         const float3 P,
                                                         const float3 dir,
                                                         const float3 idir,
-                                                        const float t,
+                                                        const float tmin,
+                                                        const float tmax,
                                                         const int node_addr,
                                                         const uint visibility,
                                                         float dist[2])
@@ -102,7 +105,7 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg,
 #ifdef __VISIBILITY_FLAG__
   float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
 #endif
-  if (bvh_unaligned_node_intersect_child(kg, P, dir, t, node_addr, 0, &dist[0])) {
+  if (bvh_unaligned_node_intersect_child(kg, P, dir, tmin, tmax, node_addr, 0, &dist[0])) {
 #ifdef __VISIBILITY_FLAG__
     if ((__float_as_uint(cnodes.x) & visibility))
 #endif
@@ -110,7 +113,7 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg,
       mask |= 1;
     }
   }
-  if (bvh_unaligned_node_intersect_child(kg, P, dir, t, node_addr, 1, &dist[1])) {
+  if (bvh_unaligned_node_intersect_child(kg, P, dir, tmin, tmax, node_addr, 1, &dist[1])) {
 #ifdef __VISIBILITY_FLAG__
     if ((__float_as_uint(cnodes.y) & visibility))
 #endif
@@ -125,16 +128,17 @@ ccl_device_forceinline int bvh_node_intersect(KernelGlobals kg,
                                               const float3 P,
                                               const float3 dir,
                                               const float3 idir,
-                                              const float t,
+                                              const float tmin,
+                                              const float tmax,
                                               const int node_addr,
                                               const uint visibility,
                                               float dist[2])
 {
   float4 node = kernel_data_fetch(bvh_nodes, node_addr);
   if (__float_as_uint(node.x) & PATH_RAY_NODE_UNALIGNED) {
-    return bvh_unaligned_node_intersect(kg, P, dir, idir, t, node_addr, visibility, dist);
+    return bvh_unaligned_node_intersect(kg, P, dir, idir, tmin, tmax, node_addr, visibility, dist);
   }
   else {
-    return bvh_aligned_node_intersect(kg, P, idir, t, node_addr, visibility, dist);
+    return bvh_aligned_node_intersect(kg, P, idir, tmin, tmax, node_addr, visibility, dist);
   }
 }
diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h
index e86fe867eac..db3c91569aa 100644
--- a/intern/cycles/kernel/bvh/shadow_all.h
+++ b/intern/cycles/kernel/bvh/shadow_all.h
@@ -49,6 +49,7 @@ ccl_device_inline
   float3 P = ray->P;
   float3 dir = bvh_clamp_direction(ray->D);
   float3 idir = bvh_inverse_direction(dir);
+  float tmin = ray->tmin;
   int object = OBJECT_NONE;
   uint num_hits = 0;
 
@@ -59,12 +60,12 @@ ccl_device_inline
   /* Max distance in world space. May be dynamically reduced when max number of
    * recorded hits is exceeded and we no longer need to find hits beyond the max
    * distance found. */
-  float t_max_world = ray->t;
+  float t_max_world = ray->tmax;
 
   /* Current maximum distance to the intersection.
    * Is calculated as a ray length, transformed to an object space when entering
    * instance node. */
-  float t_max_current = ray->t;
+  float t_max_current = ray->tmax;
 
   /* Conversion from world to local space for the current instance if any, 1.0
    * otherwise. */
@@ -88,6 +89,7 @@ ccl_device_inline
                                        dir,
 #endif
                                        idir,
+                                       tmin,
                                        t_max_current,
                                        node_addr,
                                        visibility,
@@ -156,8 +158,16 @@ ccl_device_inline
 
             switch (type & PRIMITIVE_ALL) {
               case PRIMITIVE_TRIANGLE: {
-                hit = triangle_intersect(
-                    kg, &isect, P, dir, t_max_current, visibility, prim_object, prim, prim_addr);
+                hit = triangle_intersect(kg,
+                                         &isect,
+                                         P,
+                                         dir,
+                                         tmin,
+                                         t_max_current,
+                                         visibility,
+                                         prim_object,
+                                         prim,
+                                         prim_addr);
                 break;
               }
 #if BVH_FEATURE(BVH_MOTION)
@@ -166,6 +176,7 @@ ccl_device_inline
                                                 &isect,
                                                 P,
                                                 dir,
+                                                tmin,
                                                 t_max_current,
                                                 ray->time,
                                                 visibility,
@@ -189,8 +200,16 @@ ccl_device_inline
                 }
 
                 const int curve_type = kernel_data_fetch(prim_type, prim_addr);
-                hit = curve_intersect(
-                    kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, curve_type);
+                hit = curve_intersect(kg,
+                                      &isect,
+                                      P,
+                                      dir,
+                                      tmin,
+                                      t_max_current,
+                                      prim_object,
+                                      prim,
+                                      ray->time,
+                                      curve_type);
 
                 break;
               }
@@ -207,8 +226,16 @@ ccl_device_inline
                 }
 
                 const int point_type = kernel_data_fetch(prim_type, prim_addr);
-                hit = point_intersect(
-                    kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, point_type);
+                hit = point_intersect(kg,
+                                      &isect,
+                                      P,
+                                      dir,
+                                      tmin,
+                                      t_max_current,
+                                      prim_object,
+                                      prim,
+                                      ray->time,
+                                      point_type);
                 break;
               }
 #endif /* BVH_FEATURE(BVH_POINTCLOUD) */
@@ -302,6 +329,7 @@ ccl_device_inline
 
           /* Convert intersection to object space. */
           t_max_current *= t_world_to_instance;
+          tmin *= t_world_to_instance;
 
           ++stack_ptr;
           kernel_assert(stack_ptr < BVH_STACK_SIZE);
@@ -323,7 +351,8 @@ ccl_device_inline
 #endif
 
       /* Restore world space ray length. */
-      t_max_current = ray->t;
+      tmin = ray->tmin;
+      t_max_current = ray->tmax;
 
       object = OBJECT_NONE;
       t_world_to_instance = 1.0f;
diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h
index 784fbf4fd11..0ff38bf02de 100644
--- a/intern/cycles/kernel/bvh/traversal.h
+++ b/intern/cycles/kernel/bvh/traversal.h
@@ -43,13 +43,14 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
   float3 P = ray->P;
   float3 dir = bvh_clamp_direction(ray->D);
   float3 idir = bvh_inverse_direction(dir);
+  float tmin = ray->tmin;
   int object = OBJECT_NONE;
 
 #if BVH_FEATURE(BVH_MOTION)
   Transform ob_itfm;
 #endif
 
-  isect->t = ray->t;
+  isect->t = ray->tmax;
   isect->u = 0.0f;
   isect->v = 0.0f;
   isect->prim = PRIM_NONE;
@@ -71,6 +72,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                                          dir,
 #endif
                                          idir,
+                                         tmin,
                                          isect->t,
                                          node_addr,
                                          visibility,
@@ -133,8 +135,16 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
 
             switch (type & PRIMITIVE_ALL) {
               case PRIMITIVE_TRIANGLE: {
-                if (triangle_intersect(
-                        kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr)) {
+                if (triangle_intersect(kg,
+                                       isect,
+                                       P,
+                                       dir,
+                                       tmin,
+                                       isect->t,
+                                       visibility,
+                                       prim_object,
+                                       prim,
+                                       prim_addr)) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
                     return true;
@@ -147,6 +157,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
                                               isect,
                                               P,
                                               dir,
+                                              tmin,
                                               isect->t,
                                               ray->time,
                                               visibility,
@@ -174,7 +185,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
 
                 const int curve_type = kernel_data_fetch(prim_type, prim_addr);
                 const bool hit = curve_intersect(
-                    kg, isect, P, dir, isect->t, prim_object, prim, ray->time, curve_type);
+                    kg, isect, P, dir, tmin, isect->t, prim_object, prim, ray->time, curve_type);
                 if (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
@@ -195,7 +206,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
 
                 const int point_type = kernel_data_fetch(prim_type, prim_addr);
                 const bool hit = point_intersect(
-                    kg, isect, P, dir, isect->t, prim_object, prim, ray->time, point_type);
+                    kg, isect, P, dir, tmin, isect->t, prim_object, prim, ray->time, point_type);
                 if (hit) {
                   /* shadow ray early termination */
                   if (visibility & PATH_RAY_SHADOW_OPAQUE)
@@ -212,11 +223,15 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
           object = kernel_data_fetch(prim_object, -prim_addr - 1);
 
 #if BVH_FEATURE(BVH_MOTION)
-          isect->t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm);
+          const float t_world_to_instance = bvh_instance_motion_push(
+              kg, object, ray, &P, &dir, &idir, &ob_itfm);
 #else
-          isect->t *= bvh_instance_push(kg, object, ray, &P, &dir, &idir);
+          const float t_world_to_instance = bvh_instance_push(kg, object, ray, &P, &dir, &idir);
 #endif
 
+          isect->t *= t_world_to_instance;
+          tmin *= t_world_to_instance;
+
           ++stack_ptr;
           kernel_assert(stack_ptr < BVH_STACK_SIZE);
           traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
@@ -235,6 +250,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
 #else
       isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
 #endif
+      tmin = ray->tmin;
 
       object = OBJECT_NONE;
       node_addr = traversal_stack[stack_ptr];
diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h
index 572e023db25..1795ae4c790 100644
--- a/intern/cycles/kernel/bvh/util.h
+++ b/intern/cycles/kernel/bvh/util.h
@@ -5,6 +5,19 @@
 
 CCL_NAMESPACE_BEGIN
 
+/* Offset intersection distance by the smallest possible amount, to skip
+ * intersections at this distance. This works in cases where the ray start
+ * position is unchanged and only tmin is updated, since for self
+ * intersection we'll be comparing against the exact same distances. */
+ccl_device_forceinline float intersection_t_offset(const float t)
+{
+  /* This is a simplified version of nextafterf(t, FLT_MAX), only dealing with
+   * non-negative and finite t. */
+  kernel_assert(t >= 0.0f && isfinite_safe(t));
+  const uint32_t bits = (t == 0.0f) ? 1 : __float_as_uint(t) + 1;
+  return __uint_as_float(bits);
+}
+
 #if defined(__KERNEL_CPU__)
 ccl_device int intersections_compare(const void *a, const void *b)
 {
diff --git a/intern/cycles/kernel/bvh/volume.h b/intern/cycles/kernel/bvh/volume.h
index 9715712a8f2..bd4e508ecac 100644
--- a/intern/cycles/kernel/bvh/volume.h
+++ b/intern/cycles/kernel/bvh/volume.h
@@ -46,13 +46,14 @@ ccl_device_inline
   float3 P = ray->P;
   float3 dir = bvh_clamp_direction(ray->D);
   float3 idir = bvh_inverse_direction(dir);
+  float tmin = ray->tmin;
   int object = OBJECT_NONE;
 
 #if BVH_FEATURE(BVH_MOTION)
   Transform ob_itfm;
 #endif
 
-  isect->t = ray->t;
+  isect->t = ray->tmax;
   isect->u = 0.0f;
   isect->v = 0.0f;
   isect->prim = PRIM_NONE;
@@ -73,6 +74,7 @@ ccl_device_inline
                                        dir,
 #endif
                                        idir,
+                                       tmin,
                                        isect->t,
                                        node_addr,
                                        visibility,
@@ -140,7 +142,7 @@ ccl_device_inline
                   continue;
                 }
                 triangle_intersect(
-                    kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr);
+                    kg, isect, P, dir, tmin, isect->t, visibility, prim_object, prim, prim_addr);
               }
               break;
             }
@@ -165,6 +167,7 @@ ccl_device_inline
                                           isect,
                                           P,
                                           dir,
+                                          tmin,
                                           isect->t,
                                           ray->time,
                                           visibility,
@@ -186,11 +189,15 @@ ccl_device_inline
           int object_flag = kernel_data_fetch(object_flag, object);
           if (object_flag & SD_OBJECT_HAS_VOLUME) {
 #if BVH_FEATURE(BVH_MOTION)
-            isect->t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm);
+            const float t_world_to_instance = bvh_instance_motion_push(
+                kg, object, ray, &P, &dir, &idir, &ob_itfm);
 #else
-            isect->t *= bvh_instance_push(kg, object, ray, &P, &dir, &idir);
+            const float t_world_to_instance = bvh_instance_push(kg, object, ray, &P, &dir, &idir);
 #endif
 
+            isect->t *= t_world_to_instance;
+            tmin *= t_world_to_instance;
+
             ++stack_ptr;
             kernel_assert(stack_ptr < BVH_STACK_SIZE);
             traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
@@ -217,6 +224,8 @@ ccl_device_inline
       isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
 #endif
 
+      tmin = ray->tmin;
+
       object = OBJECT_NONE;
       node_addr = traversal_stack[stack_ptr];
       --stack_ptr;
diff --git a/intern/cycles/kernel/bvh/volume_all.h b/intern/cycles/kernel/bvh/volume_all.h
index d06ea8fe557..c6eeb07a14d 100644
--- a/intern/cycles/kernel/bvh/volume_all.h
+++ b/intern/cycles/kernel/bvh/volume_all.h
@@ -44,12 +44,12 @@ ccl_device_inline
   int node_addr = kernel_data.bvh.root;
 
   /* ray parameters in registers */
-  const float tmax = ray->t;
   float3 P = ray->P;
   float3 dir = bvh_clamp_direction(ray->D);
   float3 idir = bvh_inverse_direction(dir);
+  float tmin = ray->tmin;
   int object = OBJECT_NONE;
-  float isect_t = tmax;
+  float isect_t = ray->tmax;
 
 #if BVH_FEATURE(BVH_MOTION)
   Transform ob_itfm;
@@ -58,7 +58,7 @@ ccl_device_inline
   int num_hits_in_instance = 0;
 
   uint num_hits = 0;
-  isect_array->t = tmax;
+  isect_array->t = ray->tmax;
 
   /* traversal loop */
   do {
@@ -75,6 +75,7 @@ ccl_device_inline
                                        dir,
 #endif
                                        idir,
+                                       tmin,
                                        isect_t,
                                        node_addr,
                                        visibility,
@@ -141,8 +142,16 @@ ccl_device_inline
                 if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
                   continue;
                 }
-                hit = triangle_intersect(
-                    kg, isect_array, P, dir, isect_t, visibility, prim_object, prim, prim_addr);
+                hit = triangle_intersect(kg,
+                                         isect_array,
+                                         P,
+                                         dir,
+                                         tmin,
+                                         isect_t,
+                                         visibility,
+                                         prim_object,
+                                         prim,
+                                         prim_addr);
                 if (hit) {
                   /* Move on to next entry in intersections array. */
                   isect_array++;
@@ -189,6 +198,7 @@ ccl_device_inline
                                                 isect_array,
                                                 P,
                                                 dir,
+                                                tmin,
                                                 isect_t,
                                                 ray->time,
                                                 visibility,
@@ -232,11 +242,15 @@ ccl_device_inline
           int object_flag = kernel_data_fetch(object_flag, object);
           if (object_flag & SD_OBJECT_HAS_VOLUME) {
 #if BVH_FEATURE(BVH_MOTION)
-            isect_t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm);
+            const float t_world_to_instance = bvh_instance_motion_push(
+                kg, object, ray, &P, &dir, &idir, &ob_itfm);
 #else
-            isect_t *= bvh_instance_push(kg, object, ray, &P, &dir, &idir);
+            const float t_world_to_instance = bvh_instance_push(kg, object, ray, &P, &dir, &idir);
 #endif
 
+            isect_t *= t_world_to_instance;
+            tmin *= t_world_to_instance;
+
             num_hits_in_instance = 0;
             isect_array->t = isect_t;
 
@@ -280,7 +294,8 @@ ccl_device_inline
 #endif
       }
 
-      isect_t = tmax;
+      tmin = ray->tmin;
+      isect_t = ray->tmax;
       isect_array->t = isect_t;
 
       object = OBJECT_NONE;