[Bf-blender-cvs] [c66f00dc26b] master: Fix Cycles rendering with OptiX after instance limit increase when building with old SDK

[Patrick Mours]

Commit: c66f00dc26b08d5f7be6aef080c1a0ec2de19cd7
Author: Patrick Mours
Date:   Fri Jan 8 13:38:26 2021 +0100
Branches: master
https://developer.blender.org/rBc66f00dc26b08d5f7be6aef080c1a0ec2de19cd7

Fix Cycles rendering with OptiX after instance limit increase when building with old SDK

Commit d259e7dcfbbd37cec5a45fdfb554f24de10d0268 increased the instance limit, but only provided
a fall back for the host code for older OptiX SDKs, not for kernel code. This caused a mismatch when
an old SDK was used (as is currently the case on buildbot) and subsequent rendering artifacts. This
fixes that by moving the bit that is checked to a common location that works with both old an new
SDK versions.

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

M	intern/cycles/device/device_optix.cpp
M	intern/cycles/kernel/kernels/optix/kernel_optix.cu

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

diff --git a/intern/cycles/device/device_optix.cpp b/intern/cycles/device/device_optix.cpp
index de98e3f3594..f19289f966e 100644
--- a/intern/cycles/device/device_optix.cpp
+++ b/intern/cycles/device/device_optix.cpp
@@ -1514,16 +1514,19 @@ class OptiXDevice : public CUDADevice {
     }
     else {
       unsigned int num_instances = 0;
+      unsigned int max_num_instances = 0xFFFFFFFF;
 
       bvh_optix->as_data.free();
       bvh_optix->traversable_handle = 0;
       bvh_optix->motion_transform_data.free();
 
-#  if OPTIX_ABI_VERSION < 23
-      if (bvh->objects.size() > 0x7FFFFF) {
-#  else
-      if (bvh->objects.size() > 0x7FFFFFF) {
-#  endif
+      optixDeviceContextGetProperty(context,
+                                    OPTIX_DEVICE_PROPERTY_LIMIT_MAX_INSTANCE_ID,
+                                    &max_num_instances,
+                                    sizeof(max_num_instances));
+      // Do not count first bit, which is used to distinguish instanced and non-instanced objects
+      max_num_instances >>= 1;
+      if (bvh->objects.size() > max_num_instances) {
         progress.set_error(
             "Failed to build OptiX acceleration structure because there are too many instances");
         return;
@@ -1582,8 +1585,8 @@ class OptiXDevice : public CUDADevice {
         instance.transform[5] = 1.0f;
         instance.transform[10] = 1.0f;
 
-        // Set user instance ID to object index
-        instance.instanceId = ob->get_device_index();
+        // Set user instance ID to object index (but leave low bit blank)
+        instance.instanceId = ob->get_device_index() << 1;
 
         // Have to have at least one bit in the mask, or else instance would always be culled
         instance.visibilityMask = 1;
@@ -1689,13 +1692,9 @@ class OptiXDevice : public CUDADevice {
           else {
             // Disable instance transform if geometry already has it applied to vertex data
             instance.flags = OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
-            // Non-instanced objects read ID from prim_object, so
-            // distinguish them from instanced objects with high bit set
-#  if OPTIX_ABI_VERSION < 23
-            instance.instanceId |= 0x800000;
-#  else
-            instance.instanceId |= 0x8000000;
-#  endif
+            // Non-instanced objects read ID from 'prim_object', so distinguish
+            // them from instanced objects with the low bit set
+            instance.instanceId |= 1;
           }
         }
       }
diff --git a/intern/cycles/kernel/kernels/optix/kernel_optix.cu b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
index 0c2c84fdbdf..7f609eab474 100644
--- a/intern/cycles/kernel/kernels/optix/kernel_optix.cu
+++ b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
@@ -45,13 +45,12 @@ template<bool always = false> ccl_device_forceinline uint get_object_id()
   uint object = optixGetInstanceId();
 #endif
   // Choose between always returning object ID or only for instances
-  if (always)
-    // Can just remove the high bit since instance always contains object ID
-    return object & 0x7FFFFFF;  // OPTIX_ABI_VERSION >= 23 ? 0x7FFFFFF : 0x7FFFFF
-  // Set to OBJECT_NONE if this is not an instanced object
-  else if (object & 0x8000000)  // OPTIX_ABI_VERSION >= 23 ? 0x8000000 : 0x800000
-    object = OBJECT_NONE;
-  return object;
+  if (always || (object & 1) == 0)
+    // Can just remove the low bit since instance always contains object ID
+    return object >> 1;
+  else
+    // Set to OBJECT_NONE if this is not an instanced object
+    return OBJECT_NONE;
 }
 
 extern "C" __global__ void __raygen__kernel_optix_path_trace()