From 23c762e388741479f2edf9e0c5a1bfbbebe3287c Mon Sep 17 00:00:00 2001 From: Weizhen Huang Date: Thu, 24 Apr 2025 13:10:33 +0200 Subject: [PATCH] Fix: Cycles: Do not count volume bounds bounce as transparent In forward path tracing, when we pass volume bounding meshes, we accumulate `volume_bounds_bounce`. We should match this behaviour in NEE instead of accumulating `transparent_bounce`. Pull Request: https://projects.blender.org/blender/blender/pulls/137556 --- intern/cycles/kernel/bvh/bvh.h | 12 ++--- intern/cycles/kernel/bvh/shadow_all.h | 15 +++--- intern/cycles/kernel/device/cpu/bvh.h | 22 +++++---- intern/cycles/kernel/device/hiprt/bvh.h | 6 +-- intern/cycles/kernel/device/hiprt/common.h | 47 ++++++++++++------- intern/cycles/kernel/device/metal/bvh.h | 10 ++-- .../cycles/kernel/device/metal/kernel.metal | 31 +++++++----- intern/cycles/kernel/device/optix/bvh.h | 27 +++++++---- intern/cycles/kernel/geom/shader_data.h | 1 + .../integrator/intersect_dedicated_light.h | 14 +++--- .../kernel/integrator/intersect_shadow.h | 12 ++--- .../cycles/kernel/integrator/shade_shadow.h | 7 +++ .../cycles/kernel/integrator/shade_surface.h | 4 ++ .../cycles/kernel/integrator/shade_volume.h | 2 + .../kernel/integrator/shadow_state_template.h | 2 + 15 files changed, 127 insertions(+), 85 deletions(-) diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index 78c050369d8..aa172b713d5 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -229,7 +229,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, IntegratorShadowState state, const ccl_private Ray *ray, const uint visibility, - const uint max_hits, + const uint max_transparent_hits, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { @@ -244,7 +244,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, { if (kernel_data.device_bvh) { return kernel_embree_intersect_shadow_all( - kg, state, ray, visibility, max_hits, num_recorded_hits, throughput); + kg, state, ray, visibility, max_transparent_hits, num_recorded_hits, throughput); } } # endif @@ -256,24 +256,24 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, # ifdef __HAIR__ if (kernel_data.bvh.have_curves) { return bvh_intersect_shadow_all_hair_motion( - kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); + kg, ray, state, visibility, max_transparent_hits, num_recorded_hits, throughput); } # endif /* __HAIR__ */ return bvh_intersect_shadow_all_motion( - kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); + kg, ray, state, visibility, max_transparent_hits, num_recorded_hits, throughput); } # endif /* __OBJECT_MOTION__ */ # ifdef __HAIR__ if (kernel_data.bvh.have_curves) { return bvh_intersect_shadow_all_hair( - kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); + kg, ray, state, visibility, max_transparent_hits, num_recorded_hits, throughput); } # endif /* __HAIR__ */ return bvh_intersect_shadow_all( - kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); + kg, ray, state, visibility, max_transparent_hits, num_recorded_hits, throughput); } kernel_assert(false); diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h index b2978c80975..9ff7d200f31 100644 --- a/intern/cycles/kernel/bvh/shadow_all.h +++ b/intern/cycles/kernel/bvh/shadow_all.h @@ -31,7 +31,7 @@ ccl_device_inline const ccl_private Ray *ray, IntegratorShadowState state, const uint visibility, - const uint max_hits, + const uint max_transparent_hits, ccl_private uint *r_num_recorded_hits, ccl_private float *r_throughput) { @@ -54,7 +54,7 @@ ccl_device_inline float3 idir = bvh_inverse_direction(dir); float tmin = ray->tmin; int object = OBJECT_NONE; - uint num_hits = 0; + uint num_transparent_hits = 0; /* Max distance in world space. May be dynamically reduced when max number of * recorded hits is exceeded and we no longer need to find hits beyond the max @@ -241,8 +241,9 @@ ccl_device_inline continue; } - num_hits++; - if (num_hits > max_hits) { + /* Only count transparent bounces, volume bounds bounces are counted when shading. */ + num_transparent_hits += !(flags & SD_HAS_ONLY_VOLUME); + if (num_transparent_hits > max_transparent_hits) { return true; } @@ -268,7 +269,7 @@ ccl_device_inline * so that we can detect this and trace another ray if needed. */ ++(*r_num_recorded_hits); - const uint max_record_hits = min(max_hits, (uint)INTEGRATOR_SHADOW_ISECT_SIZE); + const uint max_record_hits = (uint)INTEGRATOR_SHADOW_ISECT_SIZE; if (*r_num_recorded_hits <= max_record_hits || isect.t < tmax_hits) { integrator_state_write_shadow_isect(state, &isect, isect_index); @@ -331,12 +332,12 @@ ccl_device_inline bool BVH_FUNCTION_NAME(KernelGlobals kg, const ccl_private Ray *ray, IntegratorShadowState state, const uint visibility, - const uint max_hits, + const uint max_transparent_hits, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { return BVH_FUNCTION_FULL_NAME(BVH)( - kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); + kg, ray, state, visibility, max_transparent_hits, num_recorded_hits, throughput); } #undef BVH_FUNCTION_NAME diff --git a/intern/cycles/kernel/device/cpu/bvh.h b/intern/cycles/kernel/device/cpu/bvh.h index 366b24b0a6e..5142bb1d8b6 100644 --- a/intern/cycles/kernel/device/cpu/bvh.h +++ b/intern/cycles/kernel/device/cpu/bvh.h @@ -80,8 +80,8 @@ struct CCLShadowContext float throughput; float max_t; bool opaque_hit; - numhit_t max_hits; - numhit_t num_hits; + numhit_t max_transparent_hits; + numhit_t num_transparent_hits; numhit_t num_recorded_hits; }; @@ -343,7 +343,7 @@ ccl_device_forceinline void kernel_embree_filter_occluded_shadow_all_func_impl( } #endif - /* If no transparent shadows or max number of hits exceeded, all light is blocked. */ + /* If no transparent shadows, all light is blocked. */ const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type); if ((flags & SD_HAS_TRANSPARENT_SHADOW) == 0) { ctx->opaque_hit = true; @@ -351,14 +351,16 @@ ccl_device_forceinline void kernel_embree_filter_occluded_shadow_all_func_impl( } if (intersection_skip_shadow_already_recoded( - kg, ctx->isect_s, current_isect.object, current_isect.prim, ctx->num_hits)) + kg, ctx->isect_s, current_isect.object, current_isect.prim, ctx->num_recorded_hits)) { *args->valid = 0; return; } - ++ctx->num_hits; - if (ctx->num_hits > ctx->max_hits) { + /* Only count transparent bounces, volume bounds bounces are counted during shading. */ + ctx->num_transparent_hits += !(flags & SD_HAS_ONLY_VOLUME); + if (ctx->num_transparent_hits > ctx->max_transparent_hits) { + /* Max number of hits exceeded. */ ctx->opaque_hit = true; return; } @@ -390,7 +392,7 @@ ccl_device_forceinline void kernel_embree_filter_occluded_shadow_all_func_impl( /* This tells Embree to continue tracing. */ *args->valid = 0; - const numhit_t max_record_hits = min(ctx->max_hits, numhit_t(INTEGRATOR_SHADOW_ISECT_SIZE)); + const numhit_t max_record_hits = numhit_t(INTEGRATOR_SHADOW_ISECT_SIZE); /* If the maximum number of hits was reached, replace the furthest intersection * with a closer one so we get the N closest intersections. */ if (isect_index >= max_record_hits) { @@ -862,7 +864,7 @@ ccl_device_intersect bool kernel_embree_intersect_shadow_all(KernelGlobals kg, IntegratorShadowState state, const ccl_private Ray *ray, const uint visibility, - const uint max_hits, + const uint max_transparent_hits, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { @@ -882,11 +884,11 @@ ccl_device_intersect bool kernel_embree_intersect_shadow_all(KernelGlobals kg, CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_SHADOW_ALL); rtcInitIntersectContext(&ctx); # endif - ctx.num_hits = ctx.num_recorded_hits = numhit_t(0); + ctx.num_transparent_hits = ctx.num_recorded_hits = numhit_t(0); ctx.throughput = 1.0f; ctx.opaque_hit = false; ctx.isect_s = state; - ctx.max_hits = numhit_t(max_hits); + ctx.max_transparent_hits = numhit_t(max_transparent_hits); ctx.max_t = ray->tmax; ctx.ray = ray; RTCRay rtc_ray; diff --git a/intern/cycles/kernel/device/hiprt/bvh.h b/intern/cycles/kernel/device/hiprt/bvh.h index de06c31d1c3..97e9de216b0 100644 --- a/intern/cycles/kernel/device/hiprt/bvh.h +++ b/intern/cycles/kernel/device/hiprt/bvh.h @@ -166,7 +166,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, IntegratorShadowState state, const ccl_private Ray *ray, const uint visibility, - const uint max_hits, + const uint max_transparent_hits, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { @@ -187,8 +187,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, payload.prim_type = PRIMITIVE_NONE; payload.ray_time = ray->time; payload.in_state = state; - payload.max_hits = max_hits; - payload.num_hits = 0; + payload.max_transparent_hits = max_transparent_hits; + payload.num_transparent_hits = 0; payload.r_num_recorded_hits = num_recorded_hits; payload.r_throughput = throughput; diff --git a/intern/cycles/kernel/device/hiprt/common.h b/intern/cycles/kernel/device/hiprt/common.h index 99f294c7e6d..f7ba86137a1 100644 --- a/intern/cycles/kernel/device/hiprt/common.h +++ b/intern/cycles/kernel/device/hiprt/common.h @@ -19,8 +19,8 @@ struct RayPayload { * NOTE: This assumes that reinterpret_cast from void pointer to RayPayload works correctly. */ struct ShadowPayload : RayPayload { int in_state; - uint max_hits; - uint num_hits; + uint max_transparent_hits; + uint num_transparent_hits; uint *r_num_recorded_hits; float *r_throughput; }; @@ -380,8 +380,8 @@ ccl_device_inline bool shadow_intersection_filter(const hiprtRay &ray, { KernelGlobals kg = payload->kg; - const uint num_hits = payload->num_hits; - const uint max_hits = payload->max_hits; + uint num_transparent_hits = payload->num_transparent_hits; + const uint max_transparent_hits = payload->max_transparent_hits; const int state = payload->in_state; const RaySelfPrimitives &self = payload->self; @@ -407,7 +407,9 @@ ccl_device_inline bool shadow_intersection_filter(const hiprtRay &ray, return true; /* No hit -continue traversal. */ } - if (intersection_skip_shadow_already_recoded(kg, state, object, prim, num_hits)) { + if (intersection_skip_shadow_already_recoded( + kg, state, object, prim, *payload->r_num_recorded_hits)) + { return true; } @@ -418,18 +420,22 @@ ccl_device_inline bool shadow_intersection_filter(const hiprtRay &ray, # ifndef __TRANSPARENT_SHADOWS__ return false; # else - if (num_hits >= max_hits || - !(intersection_get_shader_flags(kg, prim, primitive_type) & SD_HAS_TRANSPARENT_SHADOW)) - { + const int flags = intersection_get_shader_flags(kg, prim, primitive_type); + if (!(flags & SD_HAS_TRANSPARENT_SHADOW)) { + return false; + } + + num_transparent_hits += !(flags & SD_HAS_ONLY_VOLUME); + if (num_transparent_hits > max_transparent_hits) { return false; } uint record_index = *payload->r_num_recorded_hits; - payload->num_hits = num_hits + 1; + payload->num_transparent_hits = num_transparent_hits; *(payload->r_num_recorded_hits) += 1; - const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + const uint max_record_hits = INTEGRATOR_SHADOW_ISECT_SIZE; if (record_index >= max_record_hits) { float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); uint max_recorded_hit = 0; @@ -467,9 +473,9 @@ ccl_device_inline bool shadow_intersection_filter_curves(const hiprtRay &ray, { KernelGlobals kg = payload->kg; - const uint num_hits = payload->num_hits; + uint num_transparent_hits = payload->num_transparent_hits; const uint num_recorded_hits = *(payload->r_num_recorded_hits); - const uint max_hits = payload->max_hits; + const uint max_transparent_hits = payload->max_transparent_hits; const RaySelfPrimitives &self = payload->self; const int object = kernel_data_fetch(user_instance_id, hit.instanceID); @@ -494,7 +500,10 @@ ccl_device_inline bool shadow_intersection_filter_curves(const hiprtRay &ray, return true; /* No hit -continue traversal. */ } - if (intersection_skip_shadow_already_recoded(kg, payload->in_state, object, prim, num_hits)) { + /* FIXME: transparent curves are not recorded, this check doesn't work. */ + if (intersection_skip_shadow_already_recoded( + kg, payload->in_state, object, prim, num_recorded_hits)) + { return true; } @@ -510,16 +519,20 @@ ccl_device_inline bool shadow_intersection_filter_curves(const hiprtRay &ray, # ifndef __TRANSPARENT_SHADOWS__ return false; # else - if (num_hits >= max_hits || - !(intersection_get_shader_flags(kg, prim, primitive_type) & SD_HAS_TRANSPARENT_SHADOW)) - { + const int flags = intersection_get_shader_flags(kg, prim, primitive_type); + if (!(flags & SD_HAS_TRANSPARENT_SHADOW)) { + return false; + } + + num_transparent_hits += !(flags & SD_HAS_ONLY_VOLUME); + if (num_transparent_hits > max_transparent_hits) { return false; } float throughput = *payload->r_throughput; throughput *= intersection_curve_shadow_transparency(kg, object, prim, primitive_type, u); *payload->r_throughput = throughput; - payload->num_hits += 1; + payload->num_transparent_hits = num_transparent_hits; if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { return false; diff --git a/intern/cycles/kernel/device/metal/bvh.h b/intern/cycles/kernel/device/metal/bvh.h index 978337578f8..3ee0b31924f 100644 --- a/intern/cycles/kernel/device/metal/bvh.h +++ b/intern/cycles/kernel/device/metal/bvh.h @@ -54,8 +54,8 @@ struct MetalRTIntersectionShadowAllPayload { RaySelfPrimitives self; int state; float throughput; - short max_hits; - short num_hits; + short max_transparent_hits; + short num_transparent_hits; short num_recorded_hits; bool result; }; @@ -466,7 +466,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, IntegratorShadowState state, const ccl_private Ray *ray, const uint visibility, - const uint max_hits, + const uint max_transparent_hits, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { @@ -482,8 +482,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, MetalRTIntersectionShadowAllPayload payload; payload.self = ray->self; - payload.max_hits = max_hits; - payload.num_hits = 0; + payload.max_transparent_hits = max_transparent_hits; + payload.num_transparent_hits = 0; payload.num_recorded_hits = 0; payload.throughput = 1.0f; payload.result = false; diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 28ab53a8720..37c4c9b7493 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -234,7 +234,6 @@ bool metalrt_shadow_all_hit( MetalKernelContext context(launch_params_metal); const IntegratorShadowState state = payload.state; - short num_hits = payload.num_hits; if (context.intersection_skip_self_shadow(payload.self, object, prim)) { /* continue search */ @@ -248,7 +247,10 @@ bool metalrt_shadow_all_hit( } # endif - if (context.intersection_skip_shadow_already_recoded(nullptr, state, object, prim, num_hits)) { + short num_recorded_hits = payload.num_recorded_hits; + if (context.intersection_skip_shadow_already_recoded( + nullptr, state, object, prim, num_recorded_hits)) + { return true; } @@ -258,15 +260,22 @@ bool metalrt_shadow_all_hit( /* terminate ray */ return false; # else - short max_hits = payload.max_hits; - short num_recorded_hits = payload.num_recorded_hits; /* If no transparent shadows, all light is blocked and we can stop immediately. */ - if (num_hits >= max_hits || - !(context.intersection_get_shader_flags(nullptr, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) - { + const int flags = context.intersection_get_shader_flags(nullptr, prim, type); + if (!(flags & SD_HAS_TRANSPARENT_SHADOW)) { + payload.result = true; + /* Terminate ray. */ + return false; + } + + /* Only count transparent bounces, volume bounds bounces are counted during shading. */ + short num_transparent_hits = payload.num_transparent_hits + !(flags & SD_HAS_ONLY_VOLUME); + short max_transparent_hits = payload.max_transparent_hits; + + if (num_transparent_hits > max_transparent_hits) { + /* Max number of hits exceeded. */ payload.result = true; - /* terminate ray */ return false; } @@ -276,7 +285,7 @@ bool metalrt_shadow_all_hit( float throughput = payload.throughput; throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, type, u); payload.throughput = throughput; - payload.num_hits += 1; + payload.num_transparent_hits = num_transparent_hits; if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { /* Accept result and terminate if throughput is sufficiently low */ @@ -289,12 +298,12 @@ bool metalrt_shadow_all_hit( } # endif - payload.num_hits += 1; + payload.num_transparent_hits = num_transparent_hits; payload.num_recorded_hits += 1; uint record_index = num_recorded_hits; - const uint max_record_hits = min(uint(max_hits), INTEGRATOR_SHADOW_ISECT_SIZE); + const uint max_record_hits = INTEGRATOR_SHADOW_ISECT_SIZE; if (record_index >= max_record_hits) { /* If maximum number of hits reached, find a hit to replace. */ float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h index 51cef89ad8b..b301233a8f1 100644 --- a/intern/cycles/kernel/device/optix/bvh.h +++ b/intern/cycles/kernel/device/optix/bvh.h @@ -195,15 +195,22 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() optixSetPayload_5(true); return optixTerminateRay(); # else - const uint max_hits = optixGetPayload_3(); + const uint max_transparent_hits = optixGetPayload_3(); const uint num_hits_packed = optixGetPayload_2(); const uint num_recorded_hits = uint16_unpack_from_uint_0(num_hits_packed); - const uint num_hits = uint16_unpack_from_uint_1(num_hits_packed); + uint num_transparent_hits = uint16_unpack_from_uint_1(num_hits_packed); /* If no transparent shadows, all light is blocked and we can stop immediately. */ - if (num_hits >= max_hits || - !(intersection_get_shader_flags(nullptr, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) - { + const int flags = intersection_get_shader_flags(nullptr, prim, type); + if (!(flags & SD_HAS_TRANSPARENT_SHADOW)) { + optixSetPayload_5(true); + return optixTerminateRay(); + } + + /* Only count transparent bounces, volume bounds bounces are counted during shading. */ + num_transparent_hits += !(flags & SD_HAS_ONLY_VOLUME); + if (num_transparent_hits > max_transparent_hits) { + /* Max number of hits exceeded. */ optixSetPayload_5(true); return optixTerminateRay(); } @@ -213,7 +220,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() float throughput = __uint_as_float(optixGetPayload_1()); throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, type, u); optixSetPayload_1(__float_as_uint(throughput)); - optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits, num_hits + 1)); + optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits, num_transparent_hits)); if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { optixSetPayload_5(true); @@ -227,13 +234,13 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() } /* Record transparent intersection. */ - optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits + 1, num_hits + 1)); + optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits + 1, num_transparent_hits)); uint record_index = num_recorded_hits; const IntegratorShadowState state = optixGetPayload_0(); - const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + const uint max_record_hits = INTEGRATOR_SHADOW_ISECT_SIZE; if (record_index >= max_record_hits) { /* If maximum number of hits reached, find a hit to replace. */ float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); @@ -584,14 +591,14 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, IntegratorShadowState state, const ccl_private Ray *ray, const uint visibility, - const uint max_hits, + const uint max_transparent_hits, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { uint p0 = state; uint p1 = __float_as_uint(1.0f); /* Throughput. */ uint p2 = 0; /* Number of hits. */ - uint p3 = max_hits; + uint p3 = max_transparent_hits; uint p4 = visibility; uint p5 = false; uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; diff --git a/intern/cycles/kernel/geom/shader_data.h b/intern/cycles/kernel/geom/shader_data.h index 99cd7ae04f5..f38eb1974d1 100644 --- a/intern/cycles/kernel/geom/shader_data.h +++ b/intern/cycles/kernel/geom/shader_data.h @@ -86,6 +86,7 @@ ccl_device_inline void shader_setup_from_ray(KernelGlobals kg, triangle_shader_setup(kg, sd); } else { + kernel_assert(sd->type == PRIMITIVE_MOTION_TRIANGLE); /* motion triangle */ motion_triangle_shader_setup(kg, sd); } diff --git a/intern/cycles/kernel/integrator/intersect_dedicated_light.h b/intern/cycles/kernel/integrator/intersect_dedicated_light.h index dd055e6a3bc..1c188cd771c 100644 --- a/intern/cycles/kernel/integrator/intersect_dedicated_light.h +++ b/intern/cycles/kernel/integrator/intersect_dedicated_light.h @@ -43,7 +43,7 @@ ccl_device int shadow_linking_pick_mesh_intersection(KernelGlobals kg, const uint visibility = path_state_ray_visibility(state); int transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce); - int volume_bounce = INTEGRATOR_STATE(state, path, volume_bounce); + int volume_bounds_bounce = INTEGRATOR_STATE(state, path, volume_bounds_bounce); /* TODO: Replace the look with sequential calls to the kernel, similar to the transparent shadow * intersection kernel. */ @@ -88,17 +88,17 @@ ccl_device int shadow_linking_pick_mesh_intersection(KernelGlobals kg, else { /* Lights past the maximum allowed transparency bounce do not contribute any light, so * consider them as fully blocked and only consider lights prior to this intersection. */ - if (shader_flags & SD_HAS_TRANSPARENT_SHADOW) { - ++transparent_bounce; - if (transparent_bounce >= kernel_data.integrator.transparent_max_bounce) { + if (shader_flags & SD_HAS_ONLY_VOLUME) { + ++volume_bounds_bounce; + if (volume_bounds_bounce >= VOLUME_BOUNDS_MAX) { ray->tmax = current_isect.t; break; } } else { - kernel_assert(shader_flags & SD_HAS_ONLY_VOLUME); - ++volume_bounce; - if (volume_bounce >= kernel_data.integrator.max_volume_bounce) { + kernel_assert(shader_flags & SD_HAS_TRANSPARENT_SHADOW); + ++transparent_bounce; + if (transparent_bounce >= kernel_data.integrator.transparent_max_bounce) { ray->tmax = current_isect.t; break; } diff --git a/intern/cycles/kernel/integrator/intersect_shadow.h b/intern/cycles/kernel/integrator/intersect_shadow.h index ee4fadf969f..f3bfe71e17d 100644 --- a/intern/cycles/kernel/integrator/intersect_shadow.h +++ b/intern/cycles/kernel/integrator/intersect_shadow.h @@ -116,11 +116,11 @@ ccl_device bool integrate_intersect_shadow_transparent(KernelGlobals kg, { /* Limit the number hits to the max transparent bounces allowed and the size that we * have available in the integrator state. */ - const uint max_hits = integrate_shadow_max_transparent_hits(kg, state); + const uint max_transparent_hits = integrate_shadow_max_transparent_hits(kg, state); uint num_hits = 0; float throughput = 1.0f; bool opaque_hit = scene_intersect_shadow_all( - kg, state, ray, visibility, max_hits, &num_hits, &throughput); + kg, state, ray, visibility, max_transparent_hits, &num_hits, &throughput); /* Computed throughput from baked shadow transparency, where we can bypass recording * intersections and shader evaluation. */ @@ -128,14 +128,8 @@ ccl_device bool integrate_intersect_shadow_transparent(KernelGlobals kg, INTEGRATOR_STATE_WRITE(state, shadow_path, throughput) *= throughput; } - /* If number of hits exceed the transparent bounces limit, make opaque. */ - if (num_hits > max_hits) { - opaque_hit = true; - } - if (!opaque_hit) { - const uint num_recorded_hits = min(num_hits, - min(max_hits, (uint)INTEGRATOR_SHADOW_ISECT_SIZE)); + const uint num_recorded_hits = min(num_hits, (uint)INTEGRATOR_SHADOW_ISECT_SIZE); if (num_recorded_hits > 0) { sort_shadow_intersections(state, num_recorded_hits); diff --git a/intern/cycles/kernel/integrator/shade_shadow.h b/intern/cycles/kernel/integrator/shade_shadow.h index ea3a4e755dd..45b4d604170 100644 --- a/intern/cycles/kernel/integrator/shade_shadow.h +++ b/intern/cycles/kernel/integrator/shade_shadow.h @@ -47,6 +47,9 @@ ccl_device_inline Spectrum integrate_transparent_surface_shadow(KernelGlobals kg surface_shader_eval( kg, state, shadow_sd, nullptr, PATH_RAY_SHADOW); } + else { + INTEGRATOR_STATE_WRITE(state, shadow_path, volume_bounds_bounce) += 1; + } # ifdef __VOLUME__ /* Exit/enter volume. */ @@ -135,6 +138,10 @@ ccl_device_inline bool integrate_transparent_shadow(KernelGlobals kg, INTEGRATOR_STATE_WRITE(state, shadow_path, rng_offset) += PRNG_BOUNCE_NUM; } + if (INTEGRATOR_STATE(state, shadow_path, volume_bounds_bounce) > VOLUME_BOUNDS_MAX) { + return true; + } + /* Note we do not need to check max_transparent_bounce here, the number * of intersections is already limited and made opaque in the * INTERSECT_SHADOW kernel. */ diff --git a/intern/cycles/kernel/integrator/shade_surface.h b/intern/cycles/kernel/integrator/shade_surface.h index a8571deb353..49674c3c0c6 100644 --- a/intern/cycles/kernel/integrator/shade_surface.h +++ b/intern/cycles/kernel/integrator/shade_surface.h @@ -239,6 +239,8 @@ integrate_direct_light_shadow_init_common(KernelGlobals kg, INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transparent_bounce) = INTEGRATOR_STATE( state, path, transparent_bounce); + INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, volume_bounds_bounce) = INTEGRATOR_STATE( + state, path, volume_bounds_bounce); INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, glossy_bounce) = INTEGRATOR_STATE( state, path, glossy_bounce); @@ -682,6 +684,8 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg, INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, flag) = shadow_flag; INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bounce) = bounce; INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transparent_bounce) = transparent_bounce; + INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, volume_bounds_bounce) = INTEGRATOR_STATE( + state, path, volume_bounds_bounce); INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput; if (kernel_data.kernel_features & KERNEL_FEATURE_AO_ADDITIVE) { diff --git a/intern/cycles/kernel/integrator/shade_volume.h b/intern/cycles/kernel/integrator/shade_volume.h index bbadaaab2b7..89760b67165 100644 --- a/intern/cycles/kernel/integrator/shade_volume.h +++ b/intern/cycles/kernel/integrator/shade_volume.h @@ -858,6 +858,8 @@ ccl_device_forceinline void integrate_volume_direct_light( state, path, glossy_bounce); INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) = INTEGRATOR_STATE( state, path, transmission_bounce); + INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, volume_bounds_bounce) = INTEGRATOR_STATE( + state, path, volume_bounds_bounce); INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput_phase; /* Write Light-group, +1 as light-group is int but we need to encode into a uint8_t. */ diff --git a/intern/cycles/kernel/integrator/shadow_state_template.h b/intern/cycles/kernel/integrator/shadow_state_template.h index a3f9d4a2277..770d8f8c149 100644 --- a/intern/cycles/kernel/integrator/shadow_state_template.h +++ b/intern/cycles/kernel/integrator/shadow_state_template.h @@ -23,6 +23,8 @@ KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, diffuse_bounce, KERNEL_FEATURE_PATH_ KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, glossy_bounce, KERNEL_FEATURE_PATH_TRACING) /* Current transmission ray bounce depth. */ KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, transmission_bounce, KERNEL_FEATURE_PATH_TRACING) +/* Current volume bounds ray bounce depth. */ +KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, volume_bounds_bounce, KERNEL_FEATURE_PATH_TRACING) /* DeviceKernel bit indicating queued kernels. */ KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, queued_kernel, KERNEL_FEATURE_PATH_TRACING) /* enum PathRayFlag */