[Bf-blender-cvs] [d5acb2b2fea] blender-v3.3-release: Fix T94136: Cycles: No Hair Shadows with Transparent BSDF

[Lukas Stockner]

Commit: d5acb2b2fea613a36943fdd40f8af919582fb5d0
Author: Lukas Stockner
Date:   Thu Oct 20 04:38:50 2022 +0200
Branches: blender-v3.3-release
https://developer.blender.org/rBd5acb2b2fea613a36943fdd40f8af919582fb5d0

Fix T94136: Cycles: No Hair Shadows with Transparent BSDF

===================================================================

M	intern/cycles/kernel/bvh/shadow_all.h
M	intern/cycles/kernel/bvh/util.h
M	intern/cycles/kernel/device/cpu/bvh.h
M	intern/cycles/kernel/device/metal/kernel.metal
M	intern/cycles/kernel/device/optix/bvh.h

===================================================================

diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h
index 2ffe1496c72..b31ba479e4f 100644
--- a/intern/cycles/kernel/bvh/shadow_all.h
+++ b/intern/cycles/kernel/bvh/shadow_all.h
@@ -229,7 +229,7 @@ ccl_device_inline
               /* Always use baked shadow transparency for curves. */
               if (isect.type & PRIMITIVE_CURVE) {
                 *r_throughput *= intersection_curve_shadow_transparency(
-                    kg, isect.object, isect.prim, isect.u);
+                    kg, isect.object, isect.prim, isect.type, isect.u);
 
                 if (*r_throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) {
                   return true;
diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h
index a57703a8b8c..9ba787550c5 100644
--- a/intern/cycles/kernel/bvh/util.h
+++ b/intern/cycles/kernel/bvh/util.h
@@ -190,10 +190,8 @@ ccl_device_inline int intersection_find_attribute(KernelGlobals kg,
 /* Cut-off value to stop transparent shadow tracing when practically opaque. */
 #define CURVE_SHADOW_TRANSPARENCY_CUTOFF 0.001f
 
-ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg,
-                                                               const int object,
-                                                               const int prim,
-                                                               const float u)
+ccl_device_inline float intersection_curve_shadow_transparency(
+    KernelGlobals kg, const int object, const int prim, const int type, const float u)
 {
   /* Find attribute. */
   const int offset = intersection_find_attribute(kg, object, ATTR_STD_SHADOW_TRANSPARENCY);
@@ -204,7 +202,7 @@ ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg,
 
   /* Interpolate transparency between curve keys. */
   const KernelCurve kcurve = kernel_data_fetch(curves, prim);
-  const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(kcurve.type);
+  const int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
   const int k1 = k0 + 1;
 
   const float f0 = kernel_data_fetch(attributes_float, offset + k0);
diff --git a/intern/cycles/kernel/device/cpu/bvh.h b/intern/cycles/kernel/device/cpu/bvh.h
index d9267e1cd6d..2d7d8c2d704 100644
--- a/intern/cycles/kernel/device/cpu/bvh.h
+++ b/intern/cycles/kernel/device/cpu/bvh.h
@@ -252,7 +252,7 @@ ccl_device void kernel_embree_filter_occluded_func(const RTCFilterFunctionNArgum
       /* Always use baked shadow transparency for curves. */
       if (current_isect.type & PRIMITIVE_CURVE) {
         ctx->throughput *= intersection_curve_shadow_transparency(
-            kg, current_isect.object, current_isect.prim, current_isect.u);
+            kg, current_isect.object, current_isect.prim, current_isect.type, current_isect.u);
 
         if (ctx->throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) {
           ctx->opaque_hit = true;
diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal
index 5646c7446db..8b69ee025cd 100644
--- a/intern/cycles/kernel/device/metal/kernel.metal
+++ b/intern/cycles/kernel/device/metal/kernel.metal
@@ -228,7 +228,7 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
   /* Always use baked shadow transparency for curves. */
   if (type & PRIMITIVE_CURVE) {
     float throughput = payload.throughput;
-    throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, u);
+    throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, type, u);
     payload.throughput = throughput;
     payload.num_hits += 1;
 
diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h
index fb9907709ce..6d81b44660c 100644
--- a/intern/cycles/kernel/device/optix/bvh.h
+++ b/intern/cycles/kernel/device/optix/bvh.h
@@ -202,7 +202,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
   /* Always use baked shadow transparency for curves. */
   if (type & PRIMITIVE_CURVE) {
     float throughput = __uint_as_float(optixGetPayload_1());
-    throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u);
+    throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, type, u);
     optixSetPayload_1(__float_as_uint(throughput));
     optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits, num_hits + 1));