Merge branch 'blender-v3.3-release'

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) {