diff options
author | Brecht Van Lommel <brecht@blender.org> | 2022-08-05 20:49:12 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2022-08-05 20:49:12 +0300 |
commit | fafd1ab9d3a6ea59734288cacb40607230be826a (patch) | |
tree | d49db404b57a425c6976e6c24124c134bc1327a8 /intern/cycles/kernel/device | |
parent | 91250022d003ec8909e1a602445a637d51124dc0 (diff) | |
parent | 4181d82ad1e3cbe9f3774e6c60767e1915201548 (diff) |
Merge branch 'blender-v3.3-release'
Diffstat (limited to 'intern/cycles/kernel/device')
-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) { |