From c66f00dc26b08d5f7be6aef080c1a0ec2de19cd7 Mon Sep 17 00:00:00 2001 From: Patrick Mours Date: Fri, 8 Jan 2021 13:38:26 +0100 Subject: 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. --- intern/cycles/device/device_optix.cpp | 27 +++++++++++----------- intern/cycles/kernel/kernels/optix/kernel_optix.cu | 13 +++++------ 2 files changed, 19 insertions(+), 21 deletions(-) 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 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() -- cgit v1.2.3