[Bf-blender-cvs] [59aef0ad5d6] master: Fix T71255: Particle hair not showing in viewport with OptiX after scaling

[Patrick Mours]

Commit: 59aef0ad5d6440e855f53db791b28076f50df2b2
Author: Patrick Mours
Date:   Fri Nov 22 17:30:22 2019 +0100
Branches: master
https://developer.blender.org/rB59aef0ad5d6440e855f53db791b28076f50df2b2

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.

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

M	intern/cycles/kernel/kernels/optix/kernel_optix.cu

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

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