Fix T71255: Particle hair not showing in viewport with OptiX after scaling

The OptiX intersection program for curves uses "optixGetObjectRayDirection"
to get the ray direction in object space (which was inverse transformed
with the current transformation matrix). OptiX does no additional operations
on it, so if there is a scaling transform, the direction is not normalized.
But the curve intersection routine expects that. In addition, the distances
used in "optixGetRayTmax()" and "optixReportIntersection()" are in world
space, so need to adjust them accordingly.

1 files changed, 11 insertions, 4 deletions

diff --git a/intern/cycles/kernel/kernels/optix/kernel_optix.cu b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
index c7223a49d79..e03504316ad 100644
--- a/intern/cycles/kernel/kernels/optix/kernel_optix.cu
+++ b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
@@ -44,7 +44,7 @@ template<bool always = false> ccl_device_forceinline uint get_object_id()
 #endif
   // Choose between always returning object ID or only for instances
   if (always)
-    // Can just remove the high bit since instace always contains object ID
+    // Can just remove the high bit since instance always contains object ID
     return object & 0x7FFFFF;
   // Set to OBJECT_NONE if this is not an instanced object
   else if (object & 0x800000)
@@ -263,8 +263,12 @@ extern "C" __global__ void __intersection__curve()
   const uint type = kernel_tex_fetch(__prim_type, prim);
   const uint visibility = optixGetPayload_4();
 
-  const float3 P = optixGetObjectRayOrigin();
-  const float3 dir = optixGetObjectRayDirection();
+  float3 P = optixGetObjectRayOrigin();
+  float3 dir = optixGetObjectRayDirection();
+
+  // The direction is not normalized by default, but the curve intersection routine expects that
+  float len;
+  dir = normalize_len(dir, &len);
 
 #  ifdef __OBJECT_MOTION__
   const float time = optixGetRayTime();
@@ -274,11 +278,14 @@ extern "C" __global__ void __intersection__curve()
 
   Intersection isect;
   isect.t = optixGetRayTmax();
+  // Transform maximum distance into object space
+  if (isect.t != FLT_MAX)
+    isect.t *= len;
 
   if (!(kernel_data.curve.curveflags & CURVE_KN_INTERPOLATE) ?
           curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type) :
           cardinal_curve_intersect(NULL, &isect, P, dir, visibility, object, prim, time, type)) {
-    optixReportIntersection(isect.t,
+    optixReportIntersection(isect.t / len,
                             type & PRIMITIVE_ALL,
                             __float_as_int(isect.u),   // Attribute_0
                             __float_as_int(isect.v));  // Attribute_1