griefith/test
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Fix: Cycles: Do not count volume bounds bounce as transparent

Browse Source 
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
This commit is contained in:
Weizhen Huang
2025-04-24 13:10:33 +02:00
committed by Weizhen Huang
parent 1cd2b34685
commit 23c762e388
15 changed files with 127 additions and 85 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
intern/cycles/kernel/bvh/bvh.h
View File

@@ -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);

15
intern/cycles/kernel/bvh/shadow_all.h
View File

@@ -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

22
intern/cycles/kernel/device/cpu/bvh.h
View File

@@ -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;

6
intern/cycles/kernel/device/hiprt/bvh.h
View File

@@ -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;

47
intern/cycles/kernel/device/hiprt/common.h
View File

@@ -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;

10
intern/cycles/kernel/device/metal/bvh.h
View File

@@ -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;

31
intern/cycles/kernel/device/metal/kernel.metal
View File

@@ -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);

27
intern/cycles/kernel/device/optix/bvh.h
View File

@@ -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;

1
intern/cycles/kernel/geom/shader_data.h
View File

@@ -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);
}

14
intern/cycles/kernel/integrator/intersect_dedicated_light.h
View File

@@ -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;
}

12
intern/cycles/kernel/integrator/intersect_shadow.h
View File

@@ -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);

7
intern/cycles/kernel/integrator/shade_shadow.h
View File

@@ -47,6 +47,9 @@ ccl_device_inline Spectrum integrate_transparent_surface_shadow(KernelGlobals kg
surface_shader_eval<KERNEL_FEATURE_NODE_MASK_SURFACE_SHADOW>(
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. */

4
intern/cycles/kernel/integrator/shade_surface.h
View File

@@ -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) {

2
intern/cycles/kernel/integrator/shade_volume.h
View File

@@ -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. */

2
intern/cycles/kernel/integrator/shadow_state_template.h
View File

@@ -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 */