diff options
author | Patrick Mours <pmours@nvidia.com> | 2021-01-08 15:38:26 +0300 |
---|---|---|
committer | Patrick Mours <pmours@nvidia.com> | 2021-01-08 15:38:26 +0300 |
commit | c66f00dc26b08d5f7be6aef080c1a0ec2de19cd7 (patch) | |
tree | 18a26baabf07b5042a246da1e1181eefeabc5b38 | |
parent | e3ae7d1f2f4592a0cc9032e2056b2236a39795f8 (diff) |
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.
-rw-r--r-- | intern/cycles/device/device_optix.cpp | 27 | ||||
-rw-r--r-- | 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<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() |