diff options
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 | ||||
-rw-r--r-- | intern/cycles/kernel/integrator/shade_surface.h | 8 |
4 files changed, 57 insertions, 28 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) { diff --git a/intern/cycles/kernel/integrator/shade_surface.h b/intern/cycles/kernel/integrator/shade_surface.h index 1514b3956ad..70b20a93b6a 100644 --- a/intern/cycles/kernel/integrator/shade_surface.h +++ b/intern/cycles/kernel/integrator/shade_surface.h @@ -43,11 +43,9 @@ ccl_device_forceinline bool integrate_surface_holdout(KernelGlobals kg, if (((sd->flag & SD_HOLDOUT) || (sd->object_flag & SD_OBJECT_HOLDOUT_MASK)) && (path_flag & PATH_RAY_TRANSPARENT_BACKGROUND)) { const float3 holdout_weight = shader_holdout_apply(kg, sd); - if (kernel_data.background.transparent) { - const float3 throughput = INTEGRATOR_STATE(state, path, throughput); - const float transparent = average(holdout_weight * throughput); - kernel_accum_holdout(kg, state, path_flag, transparent, render_buffer); - } + const float3 throughput = INTEGRATOR_STATE(state, path, throughput); + const float transparent = average(holdout_weight * throughput); + kernel_accum_holdout(kg, state, path_flag, transparent, render_buffer); if (isequal(holdout_weight, one_float3())) { return false; } |