diff options
author | Brecht Van Lommel <brecht@blender.org> | 2022-08-05 15:35:39 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2022-08-05 16:03:47 +0300 |
commit | fa514564b0c3ec4769c814c2bee587ce9e386571 (patch) | |
tree | 825627ca2b4c9ea153cf1bc8edb0be29135fbf4a /intern/cycles/kernel | |
parent | 45f483681fef58bfc0c2099629c940be53e24972 (diff) |
Fix T99201: Cycles render difference with 3D hair curves between OptiX and Emrbee
It should consistently use the Cycles pirmitive ID for self intersection detection,
not the one from the OptiX or Embree acceleration structure.
Differential Revision: https://developer.blender.org/D15632
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r-- | intern/cycles/kernel/device/cpu/bvh.h | 34 | ||||
-rw-r--r-- | intern/cycles/kernel/device/metal/kernel.metal | 18 | ||||
-rw-r--r-- | intern/cycles/kernel/device/optix/bvh.h | 25 |
3 files changed, 54 insertions, 23 deletions
diff --git a/intern/cycles/kernel/device/cpu/bvh.h b/intern/cycles/kernel/device/cpu/bvh.h index 6c06232a692..d9267e1cd6d 100644 --- a/intern/cycles/kernel/device/cpu/bvh.h +++ b/intern/cycles/kernel/device/cpu/bvh.h @@ -114,27 +114,37 @@ ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg const RTCHit *hit, const Ray *ray) { - bool status = false; + int object, prim; + if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) { - const int oID = hit->instID[0] / 2; - if ((ray->self.object == oID) || (ray->self.light_object == oID)) { + object = hit->instID[0] / 2; + if ((ray->self.object == object) || (ray->self.light_object == object)) { RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( rtcGetGeometry(kernel_data.device_bvh, hit->instID[0])); - const int pID = hit->primID + - (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); - status = intersection_skip_self_shadow(ray->self, oID, pID); + prim = hit->primID + + (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); + } + else { + return false; } } else { - const int oID = hit->geomID / 2; - if ((ray->self.object == oID) || (ray->self.light_object == oID)) { - const int pID = hit->primID + (intptr_t)rtcGetGeometryUserData( - rtcGetGeometry(kernel_data.device_bvh, hit->geomID)); - status = intersection_skip_self_shadow(ray->self, oID, pID); + object = hit->geomID / 2; + if ((ray->self.object == object) || (ray->self.light_object == object)) { + prim = hit->primID + + (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(kernel_data.device_bvh, hit->geomID)); } + else { + return false; + } + } + + const bool is_hair = hit->geomID & 1; + if (is_hair) { + prim = kernel_data_fetch(curve_segments, prim).prim; } - return status; + return intersection_skip_self_shadow(ray->self, object, prim); } ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg, diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 3d173b0d601..3de8c069c30 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -180,11 +180,6 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, } # endif - if (intersection_skip_self_shadow(payload.self, object, prim)) { - /* continue search */ - return true; - } - const float u = barycentrics.x; const float v = barycentrics.y; int type = 0; @@ -205,6 +200,11 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, } # endif + if (intersection_skip_self_shadow(payload.self, object, prim)) { + /* continue search */ + return true; + } + # ifndef __TRANSPARENT_SHADOWS__ /* No transparent shadows support compiled in, make opaque. */ payload.result = true; @@ -346,6 +346,14 @@ inline TReturnType metalrt_visibility_test( } #endif + if (intersection_type == METALRT_HIT_TRIANGLE) { + } +# ifdef __HAIR__ + else { + prim = kernel_data_fetch(curve_segments, prim).prim; + } +# endif + /* Shadow ray early termination. */ if (visibility & PATH_RAY_SHADOW_OPAQUE) { if (intersection_skip_self_shadow(payload.self, object, prim)) { diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h index d1d342f6034..fb9907709ce 100644 --- a/intern/cycles/kernel/device/optix/bvh.h +++ b/intern/cycles/kernel/device/optix/bvh.h @@ -143,14 +143,10 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() } # endif - ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); - if (intersection_skip_self_shadow(ray->self, object, prim)) { - return optixIgnoreIntersection(); - } - float u = 0.0f, v = 0.0f; int type = 0; if (optixIsTriangleHit()) { + /* Triangle. */ const float2 barycentrics = optixGetTriangleBarycentrics(); u = barycentrics.x; v = barycentrics.y; @@ -158,6 +154,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() } # ifdef __HAIR__ else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) { + /* Curve. */ u = __uint_as_float(optixGetAttribute_0()); v = __uint_as_float(optixGetAttribute_1()); @@ -174,11 +171,17 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() } # endif else { + /* Point. */ type = kernel_data_fetch(objects, object).primitive_type; u = 0.0f; v = 0.0f; } + ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); + if (intersection_skip_self_shadow(ray->self, object, prim)) { + return optixIgnoreIntersection(); + } + # ifndef __TRANSPARENT_SHADOWS__ /* No transparent shadows support compiled in, make opaque. */ optixSetPayload_5(true); @@ -307,7 +310,17 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test() } #endif - const int prim = optixGetPrimitiveIndex(); + int prim = optixGetPrimitiveIndex(); + if (optixIsTriangleHit()) { + /* Triangle. */ + } +#ifdef __HAIR__ + else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) { + /* Curve. */ + prim = kernel_data_fetch(curve_segments, prim).prim; + } +#endif + ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); if (visibility & PATH_RAY_SHADOW_OPAQUE) { |