Fix T99201: Cycles render difference with 3D hair curves between OptiX and Emrbee

It should consistently use the Cycles pirmitive ID for self intersection detection,
not the one from the OptiX or Embree acceleration structure.

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

3 files changed, 54 insertions, 23 deletions

diff --git a/intern/cycles/kernel/device/cpu/bvh.h b/intern/cycles/kernel/device/cpu/bvh.h
index 6c06232a692..d9267e1cd6d 100644
--- a/intern/cycles/kernel/device/cpu/bvh.h
+++ b/intern/cycles/kernel/device/cpu/bvh.h
@@ -114,27 +114,37 @@ ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg
                                                           const RTCHit *hit,
                                                           const Ray *ray)
 {
-  bool status = false;
+  int object, prim;
+
   if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) {
-    const int oID = hit->instID[0] / 2;
-    if ((ray->self.object == oID) || (ray->self.light_object == oID)) {
+    object = hit->instID[0] / 2;
+    if ((ray->self.object == object) || (ray->self.light_object == object)) {
       RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData(
           rtcGetGeometry(kernel_data.device_bvh, hit->instID[0]));
-      const int pID = hit->primID +
-                      (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
-      status = intersection_skip_self_shadow(ray->self, oID, pID);
+      prim = hit->primID +
+             (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID));
+    }
+    else {
+      return false;
     }
   }
   else {
-    const int oID = hit->geomID / 2;
-    if ((ray->self.object == oID) || (ray->self.light_object == oID)) {
-      const int pID = hit->primID + (intptr_t)rtcGetGeometryUserData(
-                                        rtcGetGeometry(kernel_data.device_bvh, hit->geomID));
-      status = intersection_skip_self_shadow(ray->self, oID, pID);
+    object = hit->geomID / 2;
+    if ((ray->self.object == object) || (ray->self.light_object == object)) {
+      prim = hit->primID +
+             (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(kernel_data.device_bvh, hit->geomID));
     }
+    else {
+      return false;
+    }
+  }
+
+  const bool is_hair = hit->geomID & 1;
+  if (is_hair) {
+    prim = kernel_data_fetch(curve_segments, prim).prim;
   }
 
-  return status;
+  return intersection_skip_self_shadow(ray->self, object, prim);
 }
 
 ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,
diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal
index 3d173b0d601..3de8c069c30 100644
--- a/intern/cycles/kernel/device/metal/kernel.metal
+++ b/intern/cycles/kernel/device/metal/kernel.metal
@@ -180,11 +180,6 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
   }
 #  endif
 
-  if (intersection_skip_self_shadow(payload.self, object, prim)) {
-    /* continue search */
-    return true;
-  }
-
   const float u = barycentrics.x;
   const float v = barycentrics.y;
   int type = 0;
@@ -205,6 +200,11 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
   }
 #  endif
 
+  if (intersection_skip_self_shadow(payload.self, object, prim)) {
+    /* continue search */
+    return true;
+  }
+
 #  ifndef __TRANSPARENT_SHADOWS__
   /* No transparent shadows support compiled in, make opaque. */
   payload.result = true;
@@ -346,6 +346,14 @@ inline TReturnType metalrt_visibility_test(
   }
 #endif
 
+  if (intersection_type == METALRT_HIT_TRIANGLE) {
+  }
+#  ifdef __HAIR__
+  else {
+    prim = kernel_data_fetch(curve_segments, prim).prim;
+  }
+#  endif
+
   /* Shadow ray early termination. */
   if (visibility & PATH_RAY_SHADOW_OPAQUE) {
     if (intersection_skip_self_shadow(payload.self, object, prim)) {
diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h
index d1d342f6034..fb9907709ce 100644
--- a/intern/cycles/kernel/device/optix/bvh.h
+++ b/intern/cycles/kernel/device/optix/bvh.h
@@ -143,14 +143,10 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
   }
 #  endif
 
-  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
-  if (intersection_skip_self_shadow(ray->self, object, prim)) {
-    return optixIgnoreIntersection();
-  }
-
   float u = 0.0f, v = 0.0f;
   int type = 0;
   if (optixIsTriangleHit()) {
+    /* Triangle. */
     const float2 barycentrics = optixGetTriangleBarycentrics();
     u = barycentrics.x;
     v = barycentrics.y;
@@ -158,6 +154,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
   }
 #  ifdef __HAIR__
   else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
+    /* Curve. */
     u = __uint_as_float(optixGetAttribute_0());
     v = __uint_as_float(optixGetAttribute_1());
 
@@ -174,11 +171,17 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
   }
 #  endif
   else {
+    /* Point. */
     type = kernel_data_fetch(objects, object).primitive_type;
     u = 0.0f;
     v = 0.0f;
   }
 
+  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
+  if (intersection_skip_self_shadow(ray->self, object, prim)) {
+    return optixIgnoreIntersection();
+  }
+
 #  ifndef __TRANSPARENT_SHADOWS__
   /* No transparent shadows support compiled in, make opaque. */
   optixSetPayload_5(true);
@@ -307,7 +310,17 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
   }
 #endif
 
-  const int prim = optixGetPrimitiveIndex();
+  int prim = optixGetPrimitiveIndex();
+  if (optixIsTriangleHit()) {
+    /* Triangle. */
+  }
+#ifdef __HAIR__
+  else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
+    /* Curve. */
+    prim = kernel_data_fetch(curve_segments, prim).prim;
+  }
+#endif
+
   ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
 
   if (visibility & PATH_RAY_SHADOW_OPAQUE) {