From daeb64ad889de3ca0fdd6dc5bb48defaa568cb77 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Thu, 17 Jun 2021 17:22:41 +0200 Subject: [PATCH 1/7] Fix various issues in shadow handling of volumes Properly compute volume shading for all segments before, between and after surface hits. --- .../integrator/integrator_intersect_shadow.h | 7 +- .../integrator/integrator_shade_shadow.h | 64 +++++++++++-------- .../integrator/integrator_shade_surface.h | 4 +- .../kernel/integrator/integrator_state_util.h | 4 +- .../kernel/integrator/integrator_subsurface.h | 2 +- 5 files changed, 47 insertions(+), 34 deletions(-) diff --git a/intern/cycles/kernel/integrator/integrator_intersect_shadow.h b/intern/cycles/kernel/integrator/integrator_intersect_shadow.h index 06830d3..ea25e26 100644 --- a/intern/cycles/kernel/integrator/integrator_intersect_shadow.h +++ b/intern/cycles/kernel/integrator/integrator_intersect_shadow.h @@ -97,6 +97,9 @@ ccl_device bool integrate_intersect_shadow_transparent(INTEGRATOR_STATE_ARGS, INTEGRATOR_STATE_WRITE(shadow_path, num_hits) = num_hits; } + else { + INTEGRATOR_STATE_WRITE(shadow_path, num_hits) = 0; + } return opaque_hit; } @@ -123,8 +126,8 @@ ccl_device void integrator_intersect_shadow(INTEGRATOR_STATE_ARGS) INTEGRATOR_STATE_PASS, &ray, visibility); #endif - if (opaque_hit) { - /* Hit an opaque surface, shadow path ends here. */ + if (opaque_hit && INTEGRATOR_STATE_ARRAY(shadow_volume_stack, 0, shader) == SHADER_NONE) { + /* Hit an opaque surface and no volumes, shadow path ends here. */ INTEGRATOR_SHADOW_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW); return; } diff --git a/intern/cycles/kernel/integrator/integrator_shade_shadow.h b/intern/cycles/kernel/integrator/integrator_shade_shadow.h index 975ad75..15d7229 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_shadow.h +++ b/intern/cycles/kernel/integrator/integrator_shade_shadow.h @@ -59,7 +59,9 @@ ccl_device_inline float3 integrate_transparent_surface_shadow(INTEGRATOR_STATE_A } # ifdef __VOLUME__ -ccl_device_inline float3 integrate_transparent_volume_shadow(INTEGRATOR_STATE_ARGS, const int hit) +ccl_device_inline float3 integrate_transparent_volume_shadow(INTEGRATOR_STATE_ARGS, + const int hit, + const int num_recorded_hits) { /* TODO: deduplicate with surface, or does it not matter for memory usage? */ ShaderDataTinyStorage shadow_sd_storage; @@ -67,10 +69,14 @@ ccl_device_inline float3 integrate_transparent_volume_shadow(INTEGRATOR_STATE_AR /* Setup shader data. */ Ray ray ccl_optional_struct_init; - integrator_state_read_ray(INTEGRATOR_STATE_PASS, &ray); + integrator_state_read_shadow_ray(INTEGRATOR_STATE_PASS, &ray); - Intersection isect ccl_optional_struct_init; - integrator_state_read_shadow_isect(INTEGRATOR_STATE_PASS, &isect, hit); + /* Modify ray position and length to match current segment. */ + const float start_t = (hit == 0) ? 0.0f : INTEGRATOR_STATE_ARRAY(shadow_isect, hit - 1, t); + const float end_t = (hit < num_recorded_hits) ? INTEGRATOR_STATE_ARRAY(shadow_isect, hit, t) : + ray.t; + ray.P += start_t * ray.D; + ray.t = end_t - start_t; shader_setup_from_volume(kg, shadow_sd, &ray); @@ -81,49 +87,55 @@ ccl_device_inline float3 integrate_transparent_volume_shadow(INTEGRATOR_STATE_AR } /* Integrate extinction over segment. */ - const float start_t = (hit == 0) ? 0.0f : INTEGRATOR_STATE_ARRAY(shadow_isect, hit - 1, t); - const float end_t = isect.t; - const float t = end_t - start_t; - - return exp3(-sigma_a * t); + return volume_color_transmittance(sigma_a, ray.t); } # endif +ccl_device_inline bool shadow_intersections_remaining(const int num_hits) +{ + return num_hits >= INTEGRATOR_SHADOW_ISECT_SIZE; +} + ccl_device_inline bool integrate_transparent_shadow(INTEGRATOR_STATE_ARGS, const int num_hits) { /* Accumulate shadow for transparent surfaces. */ const int num_recorded_hits = min(num_hits, INTEGRATOR_SHADOW_ISECT_SIZE); - for (int hit = 0; hit < num_recorded_hits; hit++) { -# ifdef __VOLUME__ + for (int hit = 0; hit < num_recorded_hits + 1; hit++) { /* Volume shaders. */ - if (INTEGRATOR_STATE_ARRAY(shadow_volume_stack, 0, shader) != SHADER_NONE) { - const float3 shadow = integrate_transparent_volume_shadow(INTEGRATOR_STATE_PASS, hit); + if (hit < num_recorded_hits || !shadow_intersections_remaining(num_hits)) { +# ifdef __VOLUME__ + if (INTEGRATOR_STATE_ARRAY(shadow_volume_stack, 0, shader) != SHADER_NONE) { + const float3 shadow = integrate_transparent_volume_shadow( + INTEGRATOR_STATE_PASS, hit, num_recorded_hits); + const float3 throughput = INTEGRATOR_STATE(shadow_path, throughput) * shadow; + if (is_zero(throughput)) { + return true; + } + + INTEGRATOR_STATE_WRITE(shadow_path, throughput) = throughput; + } +# endif + } + + /* Surface shaders. */ + if (hit < num_recorded_hits) { + const float3 shadow = integrate_transparent_surface_shadow(INTEGRATOR_STATE_PASS, hit); const float3 throughput = INTEGRATOR_STATE(shadow_path, throughput) * shadow; if (is_zero(throughput)) { return true; } INTEGRATOR_STATE_WRITE(shadow_path, throughput) = throughput; + INTEGRATOR_STATE_WRITE(shadow_path, transparent_bounce) += 1; } -# endif - - /* Surface shaders. */ - const float3 shadow = integrate_transparent_surface_shadow(INTEGRATOR_STATE_PASS, hit); - const float3 throughput = INTEGRATOR_STATE(shadow_path, throughput) * shadow; - if (is_zero(throughput)) { - return true; - } - - INTEGRATOR_STATE_WRITE(shadow_path, throughput) = throughput; - INTEGRATOR_STATE_WRITE(shadow_path, transparent_bounce) += 1; /* 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. */ } - if (num_hits >= INTEGRATOR_SHADOW_ISECT_SIZE) { + if (shadow_intersections_remaining(num_hits)) { /* There are more hits that we could not recorded due to memory usage, * adjust ray to intersect again from the last hit. */ const float last_hit_t = INTEGRATOR_STATE_ARRAY(shadow_isect, num_recorded_hits - 1, t); @@ -151,7 +163,7 @@ ccl_device void integrator_shade_shadow(INTEGRATOR_STATE_ARGS, } #endif - if (num_hits >= INTEGRATOR_SHADOW_ISECT_SIZE) { + if (shadow_intersections_remaining(num_hits)) { /* More intersections to find, continue shadow ray. */ INTEGRATOR_SHADOW_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW); diff --git a/intern/cycles/kernel/integrator/integrator_shade_surface.h b/intern/cycles/kernel/integrator/integrator_shade_surface.h index d305f20..f357422 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_surface.h +++ b/intern/cycles/kernel/integrator/integrator_shade_surface.h @@ -176,9 +176,7 @@ ccl_device_forceinline void integrate_surface_direct_light(INTEGRATOR_STATE_ARGS INTEGRATOR_STATE_WRITE(shadow_path, diffuse_glossy_ratio) = diffuse_glossy_ratio; INTEGRATOR_STATE_WRITE(shadow_path, throughput) = throughput; - integrator_state_copy_volume_stack_to_shadow(INTEGRATOR_STATE_PASS); - - /* Branch of shadow kernel. */ + /* Branch off shadow kernel. */ INTEGRATOR_SHADOW_PATH_INIT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW); } #endif diff --git a/intern/cycles/kernel/integrator/integrator_state_util.h b/intern/cycles/kernel/integrator/integrator_state_util.h index ec81108..2d38522 100644 --- a/intern/cycles/kernel/integrator/integrator_state_util.h +++ b/intern/cycles/kernel/integrator/integrator_state_util.h @@ -157,8 +157,8 @@ ccl_device_forceinline void integrator_state_copy_volume_stack_to_shadow(INTEGRA ccl_device_forceinline VolumeStack integrator_state_read_shadow_volume_stack(INTEGRATOR_STATE_CONST_ARGS, int i) { - VolumeStack entry = {INTEGRATOR_STATE_ARRAY(volume_stack, i, object), - INTEGRATOR_STATE_ARRAY(volume_stack, i, shader)}; + VolumeStack entry = {INTEGRATOR_STATE_ARRAY(shadow_volume_stack, i, object), + INTEGRATOR_STATE_ARRAY(shadow_volume_stack, i, shader)}; return entry; } diff --git a/intern/cycles/kernel/integrator/integrator_subsurface.h b/intern/cycles/kernel/integrator/integrator_subsurface.h index a25e845..fe7bb1f 100644 --- a/intern/cycles/kernel/integrator/integrator_subsurface.h +++ b/intern/cycles/kernel/integrator/integrator_subsurface.h @@ -124,7 +124,7 @@ ccl_device bool subsurface_bounce(INTEGRATOR_STATE_ARGS, ShaderData *sd, const S INTEGRATOR_STATE_WRITE(subsurface, radius) = bssrdf->radius; INTEGRATOR_STATE_WRITE(subsurface, roughness) = roughness; - return true; + return LABEL_SUBSURFACE_SCATTER; } ccl_device void subsurface_shader_data_setup(INTEGRATOR_STATE_ARGS, ShaderData *sd) -- 2.25.1 From c6f160741ba6e2b2273f1466f29e8f63bfa68b82 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Fri, 25 Jun 2021 20:31:16 +0200 Subject: [PATCH 2/7] Fix incorrect volume stack for shadow rays --- .../integrator/integrator_shade_surface.h | 40 ++++++++++++------- intern/cycles/kernel/kernel_types.h | 1 + 2 files changed, 27 insertions(+), 14 deletions(-) diff --git a/intern/cycles/kernel/integrator/integrator_shade_surface.h b/intern/cycles/kernel/integrator/integrator_shade_surface.h index f357422..c78365f 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_surface.h +++ b/intern/cycles/kernel/integrator/integrator_shade_surface.h @@ -157,6 +157,15 @@ ccl_device_forceinline void integrate_surface_direct_light(INTEGRATOR_STATE_ARGS light_sample_to_shadow_ray(sd, &ls, &ray); const bool is_light = light_sample_is_light(&ls); + /* Copy volume stack and enter/exit volume. */ + integrator_state_copy_volume_stack_to_shadow(INTEGRATOR_STATE_PASS); + + if (is_transmission) { +# ifdef __VOLUME__ + shadow_volume_stack_enter_exit(INTEGRATOR_STATE_PASS, sd); +# endif + } + /* Write shadow ray and associated state to global memory. */ integrator_state_write_shadow_ray(INTEGRATOR_STATE_PASS, &ray); @@ -182,13 +191,13 @@ ccl_device_forceinline void integrate_surface_direct_light(INTEGRATOR_STATE_ARGS #endif /* Path tracing: bounce off or through surface with new direction. */ -ccl_device_forceinline bool integrate_surface_bsdf_bssrdf_bounce(INTEGRATOR_STATE_ARGS, - ShaderData *sd, - const RNGState *rng_state) +ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(INTEGRATOR_STATE_ARGS, + ShaderData *sd, + const RNGState *rng_state) { /* Sample BSDF or BSSRDF. */ if (!(sd->flag & (SD_BSDF | SD_BSSRDF))) { - return false; + return 0; } float bsdf_u, bsdf_v; @@ -213,7 +222,7 @@ ccl_device_forceinline bool integrate_surface_bsdf_bssrdf_bounce(INTEGRATOR_STAT kg, sd, sc, bsdf_u, bsdf_v, &bsdf_eval, &bsdf_omega_in, &bsdf_domega_in, &bsdf_pdf); if (bsdf_pdf == 0.0f || bsdf_eval_is_zero(&bsdf_eval)) { - return false; + return 0; } /* Setup ray. Note that clipping works through transparent bounces. */ @@ -246,7 +255,7 @@ ccl_device_forceinline bool integrate_surface_bsdf_bssrdf_bounce(INTEGRATOR_STAT } path_state_next(INTEGRATOR_STATE_PASS, label); - return true; + return label; } #ifdef __VOLUME__ @@ -254,7 +263,7 @@ ccl_device_forceinline bool integrate_surface_volume_only_bounce(INTEGRATOR_STAT ShaderData *sd) { if (!path_state_volume_next(INTEGRATOR_STATE_PASS)) { - return false; + return 0; } /* Setup ray position, direction stays unchanged. */ @@ -267,7 +276,7 @@ ccl_device_forceinline bool integrate_surface_volume_only_bounce(INTEGRATOR_STAT INTEGRATOR_STATE_WRITE(ray, dP) = differential_make_compact(sd->dP); # endif - return true; + return LABEL_TRANSMIT | LABEL_TRANSPARENT; } #endif @@ -280,7 +289,7 @@ ccl_device bool integrate_surface(INTEGRATOR_STATE_ARGS, ShaderData sd; integrate_surface_shader_setup(INTEGRATOR_STATE_PASS, &sd); - bool continue_path; + int continue_path_label = 0; /* Skip most work for volume bounding surface. */ #ifdef __VOLUME__ @@ -363,20 +372,23 @@ ccl_device bool integrate_surface(INTEGRATOR_STATE_ARGS, # endif /* __AO__ */ #endif - continue_path = integrate_surface_bsdf_bssrdf_bounce(INTEGRATOR_STATE_PASS, &sd, &rng_state); + continue_path_label = integrate_surface_bsdf_bssrdf_bounce( + INTEGRATOR_STATE_PASS, &sd, &rng_state); #ifdef __VOLUME__ } else { - continue_path = integrate_surface_volume_only_bounce(INTEGRATOR_STATE_PASS, &sd); + continue_path_label = integrate_surface_volume_only_bounce(INTEGRATOR_STATE_PASS, &sd); } #endif - /* Enter/Exit volume. */ + if (continue_path_label & LABEL_TRANSMIT) { + /* Enter/Exit volume. */ #ifdef __VOLUME__ - volume_stack_enter_exit(INTEGRATOR_STATE_PASS, &sd); + volume_stack_enter_exit(INTEGRATOR_STATE_PASS, &sd); #endif + } - return continue_path; + return continue_path_label != 0; } template Date: Thu, 17 Jun 2021 17:27:48 +0200 Subject: [PATCH 3/7] Cycles X: tweak a few functions for easier reuse in volume code --- .../integrator/integrator_intersect_closest.h | 39 ++++++++++++------- .../integrator/integrator_shade_background.h | 3 +- .../integrator/integrator_shade_light.h | 3 +- .../integrator/integrator_shade_surface.h | 3 +- .../integrator/integrator_shade_volume.h | 14 +++---- intern/cycles/kernel/kernel_accumulate.h | 3 +- intern/cycles/kernel/kernel_emission.h | 24 ++++-------- intern/cycles/kernel/kernel_shader.h | 10 ++++- 8 files changed, 55 insertions(+), 44 deletions(-) diff --git a/intern/cycles/kernel/integrator/integrator_intersect_closest.h b/intern/cycles/kernel/integrator/integrator_intersect_closest.h index b0fc87f..a4538e8 100644 --- a/intern/cycles/kernel/integrator/integrator_intersect_closest.h +++ b/intern/cycles/kernel/integrator/integrator_intersect_closest.h @@ -29,20 +29,18 @@ CCL_NAMESPACE_BEGIN template -ccl_device_forceinline bool integrator_intersect_shader_next_kernel( - INTEGRATOR_STATE_ARGS, const Intersection *ccl_restrict isect) +ccl_device_forceinline bool integrator_intersect_terminate(INTEGRATOR_STATE_ARGS, + const Intersection *ccl_restrict isect, + const int shader_flags) { - /* Find shader from intersection. */ - const int shader = intersection_get_shader(kg, isect); - const int flags = kernel_tex_fetch(__shaders, shader).flags; /* Optional AO bounce termination. */ if (path_state_ao_bounce(INTEGRATOR_STATE_PASS)) { - if (flags & (SD_HAS_TRANSPARENT_SHADOW | SD_HAS_EMISSION)) { + if (shader_flags & (SD_HAS_TRANSPARENT_SHADOW | SD_HAS_EMISSION)) { INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_TERMINATE_AFTER_TRANSPARENT; } else { - return false; + return true; } } @@ -60,18 +58,28 @@ ccl_device_forceinline bool integrator_intersect_shader_next_kernel( const float terminate = path_state_rng_1D(kg, &rng_state, PRNG_TERMINATE); if (probability == 0.0f || terminate >= probability) { - if (flags & (SD_HAS_TRANSPARENT_SHADOW | SD_HAS_EMISSION)) { + if (shader_flags & (SD_HAS_TRANSPARENT_SHADOW | SD_HAS_EMISSION)) { /* Mark path to be terminated right after shader evaluation. */ INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_TERMINATE_IMMEDIATE; } else { - return false; + return true; } } } + return false; +} + +template +ccl_device_forceinline void integrator_intersect_shader_next_kernel( + INTEGRATOR_STATE_ARGS, + const Intersection *ccl_restrict isect, + const int shader, + const int shader_flags) +{ /* Setup next kernel to execute. */ - if (flags & SD_HAS_RAYTRACE) { + if (shader_flags & SD_HAS_RAYTRACE) { INTEGRATOR_PATH_NEXT_SORTED( current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE, shader); } @@ -82,8 +90,6 @@ ccl_device_forceinline bool integrator_intersect_shader_next_kernel( /* Setup shadow catcher. */ const int object_flags = intersection_get_object_flags(kg, isect); kernel_shadow_catcher_split(INTEGRATOR_STATE_PASS, object_flags); - - return true; } ccl_device void integrator_intersect_closest(INTEGRATOR_STATE_ARGS) @@ -142,8 +148,13 @@ ccl_device void integrator_intersect_closest(INTEGRATOR_STATE_ARGS) } else { /* Hit a surface, continue with surface kernel unless terminated. */ - if (integrator_intersect_shader_next_kernel( - INTEGRATOR_STATE_PASS, &isect)) { + const int shader = intersection_get_shader(kg, &isect); + const int flags = kernel_tex_fetch(__shaders, shader).flags; + + if (!integrator_intersect_terminate( + INTEGRATOR_STATE_PASS, &isect, flags)) { + integrator_intersect_shader_next_kernel( + INTEGRATOR_STATE_PASS, &isect, shader, flags); return; } else { diff --git a/intern/cycles/kernel/integrator/integrator_shade_background.h b/intern/cycles/kernel/integrator/integrator_shade_background.h index 49666fe..3c310de 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_background.h +++ b/intern/cycles/kernel/integrator/integrator_shade_background.h @@ -166,7 +166,8 @@ ccl_device_inline void integrate_distant_lights(INTEGRATOR_STATE_ARGS, } /* Write to render buffer. */ - kernel_accum_emission(INTEGRATOR_STATE_PASS, light_eval, render_buffer); + const float3 throughput = INTEGRATOR_STATE(path, throughput); + kernel_accum_emission(INTEGRATOR_STATE_PASS, throughput, light_eval, render_buffer); } } } diff --git a/intern/cycles/kernel/integrator/integrator_shade_light.h b/intern/cycles/kernel/integrator/integrator_shade_light.h index d22fb06..04c8a25 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_light.h +++ b/intern/cycles/kernel/integrator/integrator_shade_light.h @@ -88,7 +88,8 @@ ccl_device_inline void integrate_light(INTEGRATOR_STATE_ARGS, } /* Write to render buffer. */ - kernel_accum_emission(INTEGRATOR_STATE_PASS, light_eval, render_buffer); + const float3 throughput = INTEGRATOR_STATE(path, throughput); + kernel_accum_emission(INTEGRATOR_STATE_PASS, throughput, light_eval, render_buffer); } ccl_device void integrator_shade_light(INTEGRATOR_STATE_ARGS, diff --git a/intern/cycles/kernel/integrator/integrator_shade_surface.h b/intern/cycles/kernel/integrator/integrator_shade_surface.h index c78365f..2c3da8c 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_surface.h +++ b/intern/cycles/kernel/integrator/integrator_shade_surface.h @@ -94,7 +94,8 @@ ccl_device_forceinline void integrate_surface_emission(INTEGRATOR_STATE_CONST_AR L *= mis_weight; } - kernel_accum_emission(INTEGRATOR_STATE_PASS, L, render_buffer); + const float3 throughput = INTEGRATOR_STATE(path, throughput); + kernel_accum_emission(INTEGRATOR_STATE_PASS, throughput, L, render_buffer); } #endif /* __EMISSION__ */ diff --git a/intern/cycles/kernel/integrator/integrator_shade_volume.h b/intern/cycles/kernel/integrator/integrator_shade_volume.h index 279f692..e7b0c18 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_volume.h +++ b/intern/cycles/kernel/integrator/integrator_shade_volume.h @@ -154,14 +154,12 @@ ccl_device void integrator_shade_volume(INTEGRATOR_STATE_ARGS, } else { /* Hit a surface, continue with surface kernel unless terminated. */ - if (integrator_intersect_shader_next_kernel( - INTEGRATOR_STATE_PASS, &isect)) { - return; - } - else { - INTEGRATOR_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME); - return; - } + const int shader = intersection_get_shader(kg, &isect); + const int flags = kernel_tex_fetch(__shaders, shader).flags; + + integrator_intersect_shader_next_kernel( + INTEGRATOR_STATE_PASS, &isect, shader, flags); + return; } } #endif /* __VOLUME__ */ diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h index 3dff246..54b8ff7 100644 --- a/intern/cycles/kernel/kernel_accumulate.h +++ b/intern/cycles/kernel/kernel_accumulate.h @@ -658,10 +658,11 @@ ccl_device_inline void kernel_accum_background(INTEGRATOR_STATE_CONST_ARGS, /* Write emission to render buffer. */ ccl_device_inline void kernel_accum_emission(INTEGRATOR_STATE_CONST_ARGS, + const float3 throughput, const float3 L, ccl_global float *ccl_restrict render_buffer) { - float3 contribution = INTEGRATOR_STATE(path, throughput) * L; + float3 contribution = throughput * L; kernel_accum_clamp(kg, &contribution, INTEGRATOR_STATE(path, bounce) - 1); ccl_global float *buffer = kernel_accum_pixel_render_buffer(INTEGRATOR_STATE_PASS, diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h index 67469dc..c5173cc 100644 --- a/intern/cycles/kernel/kernel_emission.h +++ b/intern/cycles/kernel/kernel_emission.h @@ -133,14 +133,20 @@ ccl_device_inline bool light_sample_terminate(const KernelGlobals *ccl_restrict } /* Create shadow ray towards light sample. */ +template ccl_device_inline void light_sample_to_shadow_ray(const ShaderData *sd, const LightSample *ls, Ray *ray) { if (ls->shader & SHADER_CAST_SHADOW) { /* setup ray */ - bool transmit = (dot(sd->Ng, ls->D) < 0.0f); - ray->P = ray_offset(sd->P, (transmit) ? -sd->Ng : sd->Ng); + if (is_volume) { + ray->P = sd->P; + } + else { + bool transmit = (dot(sd->Ng, ls->D) < 0.0f); + ray->P = ray_offset(sd->P, (transmit) ? -sd->Ng : sd->Ng); + } if (ls->t == FLT_MAX) { /* distant light */ @@ -164,18 +170,4 @@ ccl_device_inline void light_sample_to_shadow_ray(const ShaderData *sd, ray->time = sd->time; } -/* Volume phase evaluation code - to be moved into volume code. */ -#if 0 -# ifdef __VOLUME__ - float bsdf_pdf; - shader_volume_phase_eval(kg, sd, ls->D, eval, &bsdf_pdf); - if (ls->shader & SHADER_USE_MIS) { - /* Multiple importance sampling. */ - float mis_weight = power_heuristic(ls->pdf, bsdf_pdf); - light_eval *= mis_weight; - } - } -# endif -#endif - CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h index ab9f6e7..ee27bab 100644 --- a/intern/cycles/kernel/kernel_shader.h +++ b/intern/cycles/kernel/kernel_shader.h @@ -1098,13 +1098,19 @@ ccl_device void shader_volume_phase_eval(const KernelGlobals *kg, const ShaderData *sd, const float3 omega_in, BsdfEval *eval, - float *pdf) + const float light_pdf, + const uint light_shader_flags) { PROFILING_INIT(kg, PROFILING_CLOSURE_VOLUME_EVAL); bsdf_eval_init(eval, false, zero_float3(), kernel_data.film.use_light_pass); - _shader_volume_phase_multi_eval(sd, omega_in, pdf, -1, eval, 0.0f, 0.0f); + float pdf; + _shader_volume_phase_multi_eval(sd, omega_in, &pdf, -1, eval, 0.0f, 0.0f); + if (light_shader_flags & SHADER_USE_MIS) { + float weight = power_heuristic(light_pdf, pdf); + bsdf_eval_mul(eval, weight); + } } ccl_device int shader_volume_phase_sample(const KernelGlobals *kg, -- 2.25.1 From 5060f4e551a085f57f9bc173f60459f34ee52101 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Thu, 17 Jun 2021 17:29:25 +0200 Subject: [PATCH 4/7] Cycles X: restore two more volume stack functions --- .../integrator/integrator_volume_stack.h | 77 +++++++++++++++++++ intern/cycles/kernel/kernel_volume.h | 65 ---------------- 2 files changed, 77 insertions(+), 65 deletions(-) diff --git a/intern/cycles/kernel/integrator/integrator_volume_stack.h b/intern/cycles/kernel/integrator/integrator_volume_stack.h index 05855fe..824017b 100644 --- a/intern/cycles/kernel/integrator/integrator_volume_stack.h +++ b/intern/cycles/kernel/integrator/integrator_volume_stack.h @@ -277,4 +277,81 @@ ccl_device_inline void volume_stack_clean(INTEGRATOR_STATE_ARGS) } } +template +ccl_device float volume_stack_step_size(INTEGRATOR_STATE_ARGS, StackReadOp stack_read) +{ + float step_size = FLT_MAX; + + for (int i = 0;; i++) { + VolumeStack entry = stack_read(i); + if (entry.shader == SHADER_NONE) { + break; + } + + int shader_flag = kernel_tex_fetch(__shaders, (entry.shader & SHADER_MASK)).flags; + + bool heterogeneous = false; + + if (shader_flag & SD_HETEROGENEOUS_VOLUME) { + heterogeneous = true; + } + else if (shader_flag & SD_NEED_VOLUME_ATTRIBUTES) { + /* We want to render world or objects without any volume grids + * as homogeneous, but can only verify this at run-time since other + * heterogeneous volume objects may be using the same shader. */ + int object = entry.object; + if (object != OBJECT_NONE) { + int object_flag = kernel_tex_fetch(__object_flag, object); + if (object_flag & SD_OBJECT_HAS_VOLUME_ATTRIBUTES) { + heterogeneous = true; + } + } + } + + if (heterogeneous) { + float object_step_size = object_volume_step_size(kg, entry.object); + object_step_size *= kernel_data.integrator.volume_step_rate; + step_size = fminf(object_step_size, step_size); + } + } + + return step_size; +} + +template +ccl_device int volume_stack_sampling_method(INTEGRATOR_STATE_ARGS, StackReadOp stack_read) +{ + if (kernel_data.integrator.num_all_lights == 0) + return 0; + + int method = -1; + + for (int i = 0;; i++) { + VolumeStack entry = stack_read(i); + if (entry.shader == SHADER_NONE) { + break; + } + + int shader_flag = kernel_tex_fetch(__shaders, (entry.shader & SHADER_MASK)).flags; + + if (shader_flag & SD_VOLUME_MIS) { + return SD_VOLUME_MIS; + } + else if (shader_flag & SD_VOLUME_EQUIANGULAR) { + if (method == 0) + return SD_VOLUME_MIS; + + method = SD_VOLUME_EQUIANGULAR; + } + else { + if (method == SD_VOLUME_EQUIANGULAR) + return SD_VOLUME_MIS; + + method = 0; + } + } + + return method; +} + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_volume.h b/intern/cycles/kernel/kernel_volume.h index 25b8bae..efdc44e 100644 --- a/intern/cycles/kernel/kernel_volume.h +++ b/intern/cycles/kernel/kernel_volume.h @@ -24,71 +24,6 @@ CCL_NAMESPACE_BEGIN * work in volumes and subsurface scattering. */ #define VOLUME_THROUGHPUT_EPSILON 1e-6f -ccl_device float volume_stack_step_size(const KernelGlobals *kg, ccl_addr_space VolumeStack *stack) -{ - float step_size = FLT_MAX; - - for (int i = 0; stack[i].shader != SHADER_NONE; i++) { - int shader_flag = kernel_tex_fetch(__shaders, (stack[i].shader & SHADER_MASK)).flags; - - bool heterogeneous = false; - - if (shader_flag & SD_HETEROGENEOUS_VOLUME) { - heterogeneous = true; - } - else if (shader_flag & SD_NEED_VOLUME_ATTRIBUTES) { - /* We want to render world or objects without any volume grids - * as homogeneous, but can only verify this at run-time since other - * heterogeneous volume objects may be using the same shader. */ - int object = stack[i].object; - if (object != OBJECT_NONE) { - int object_flag = kernel_tex_fetch(__object_flag, object); - if (object_flag & SD_OBJECT_HAS_VOLUME_ATTRIBUTES) { - heterogeneous = true; - } - } - } - - if (heterogeneous) { - float object_step_size = object_volume_step_size(kg, stack[i].object); - object_step_size *= kernel_data.integrator.volume_step_rate; - step_size = fminf(object_step_size, step_size); - } - } - - return step_size; -} - -ccl_device int volume_stack_sampling_method(const KernelGlobals *kg, VolumeStack *stack) -{ - if (kernel_data.integrator.num_all_lights == 0) - return 0; - - int method = -1; - - for (int i = 0; stack[i].shader != SHADER_NONE; i++) { - int shader_flag = kernel_tex_fetch(__shaders, (stack[i].shader & SHADER_MASK)).flags; - - if (shader_flag & SD_VOLUME_MIS) { - return SD_VOLUME_MIS; - } - else if (shader_flag & SD_VOLUME_EQUIANGULAR) { - if (method == 0) - return SD_VOLUME_MIS; - - method = SD_VOLUME_EQUIANGULAR; - } - else { - if (method == SD_VOLUME_EQUIANGULAR) - return SD_VOLUME_MIS; - - method = 0; - } - } - - return method; -} - ccl_device_inline void kernel_volume_step_init(const KernelGlobals *kg, ccl_addr_space PathState *state, const float object_step_size, -- 2.25.1 From 1d5c057f6a18359f28caae0f11af5e4032bcc857 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Thu, 17 Jun 2021 17:32:00 +0200 Subject: [PATCH 5/7] Cycles X: remove old volume code --- intern/cycles/kernel/CMakeLists.txt | 2 - intern/cycles/kernel/kernel_path.h | 142 ---- intern/cycles/kernel/kernel_path_volume.h | 262 ------ intern/cycles/kernel/kernel_volume.h | 989 ---------------------- 4 files changed, 1395 deletions(-) delete mode 100644 intern/cycles/kernel/kernel_path_volume.h delete mode 100644 intern/cycles/kernel/kernel_volume.h diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index 33829ac..e09d190 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -83,7 +83,6 @@ set(SRC_HEADERS kernel_passes.h kernel_path.h kernel_path_state.h - kernel_path_volume.h kernel_profiling.h kernel_projection.h kernel_random.h @@ -93,7 +92,6 @@ set(SRC_HEADERS kernel_subsurface.h kernel_textures.h kernel_types.h - kernel_volume.h kernel_work_stealing.h kernel_write_passes.h ) diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h index 3e98413..c08f61a 100644 --- a/intern/cycles/kernel/kernel_path.h +++ b/intern/cycles/kernel/kernel_path.h @@ -16,150 +16,8 @@ #pragma once -#ifdef __OSL__ -# include "kernel/osl/osl_shader.h" -#endif - -// clang-format off -#include "kernel/kernel_random.h" -#include "kernel/kernel_projection.h" -#include "kernel/kernel_montecarlo.h" -#include "kernel/kernel_differential.h" -#include "kernel/kernel_camera.h" - -#include "kernel/geom/geom.h" -#include "kernel/bvh/bvh.h" - -#include "kernel/kernel_write_passes.h" -#include "kernel/kernel_accumulate.h" -#include "kernel/kernel_shader.h" -#include "kernel/kernel_light.h" -#include "kernel/kernel_adaptive_sampling.h" -#include "kernel/kernel_passes.h" - -#if defined(__VOLUME__) || defined(__SUBSURFACE__) -# include "kernel/kernel_volume.h" -#endif - -#ifdef __SUBSURFACE__ -# include "kernel/kernel_subsurface.h" -#endif - -#include "kernel/kernel_path_state.h" -#include "kernel/kernel_shadow.h" -#include "kernel/kernel_emission.h" -#include "kernel/kernel_path_common.h" -#include "kernel/kernel_path_surface.h" -#include "kernel/kernel_path_volume.h" -#include "kernel/kernel_path_subsurface.h" -// clang-format on - CCL_NAMESPACE_BEGIN -#ifdef __VOLUME__ -ccl_device_forceinline VolumeIntegrateResult kernel_path_volume(const KernelGlobals *kg, - ShaderData *sd, - PathState *state, - Ray *ray, - float3 *throughput, - ccl_addr_space Intersection *isect, - bool hit, - ShaderData *emission_sd, - PathRadiance *L) -{ - PROFILING_INIT(kg, PROFILING_VOLUME); - - /* Sanitize volume stack. */ - if (!hit) { - kernel_volume_clean_stack(kg, state->volume_stack); - } - - if (state->volume_stack[0].shader == SHADER_NONE) { - return VOLUME_PATH_ATTENUATED; - } - - /* volume attenuation, emission, scatter */ - Ray volume_ray = *ray; - volume_ray.t = (hit) ? isect->t : FLT_MAX; - - float step_size = volume_stack_step_size(kg, state->volume_stack); - -# ifdef __VOLUME_DECOUPLED__ - int sampling_method = volume_stack_sampling_method(kg, state->volume_stack); - bool direct = (state->flag & PATH_RAY_CAMERA) != 0; - bool decoupled = kernel_volume_use_decoupled(kg, step_size, direct, sampling_method); - - if (decoupled) { - /* cache steps along volume for repeated sampling */ - VolumeSegment volume_segment; - - shader_setup_from_volume(kg, sd, &volume_ray); - kernel_volume_decoupled_record(kg, state, &volume_ray, sd, &volume_segment, step_size); - - volume_segment.sampling_method = sampling_method; - - /* emission */ - if (volume_segment.closure_flag & SD_EMISSION) - path_radiance_accum_emission(kg, L, state, *throughput, volume_segment.accum_emission); - - /* scattering */ - VolumeIntegrateResult result = VOLUME_PATH_ATTENUATED; - - if (volume_segment.closure_flag & SD_SCATTER) { - int all = kernel_data.integrator.sample_all_lights_indirect; - - /* direct light sampling */ - kernel_branched_path_volume_connect_light( - kg, sd, emission_sd, *throughput, state, L, all, &volume_ray, &volume_segment); - - /* indirect sample. if we use distance sampling and take just - * one sample for direct and indirect light, we could share - * this computation, but makes code a bit complex */ - float rphase = path_state_rng_1D(kg, state, PRNG_PHASE_CHANNEL); - float rscatter = path_state_rng_1D(kg, state, PRNG_SCATTER_DISTANCE); - - result = kernel_volume_decoupled_scatter( - kg, state, &volume_ray, sd, throughput, rphase, rscatter, &volume_segment, NULL, true); - } - - /* free cached steps */ - kernel_volume_decoupled_free(kg, &volume_segment); - - if (result == VOLUME_PATH_SCATTERED) { - if (kernel_path_volume_bounce(kg, sd, throughput, state, &L->state, ray)) - return VOLUME_PATH_SCATTERED; - else - return VOLUME_PATH_MISSED; - } - else { - *throughput *= volume_segment.accum_transmittance; - } - } - else -# endif /* __VOLUME_DECOUPLED__ */ - { - /* integrate along volume segment with distance sampling */ - VolumeIntegrateResult result = kernel_volume_integrate( - kg, state, sd, &volume_ray, L, throughput, step_size); - -# ifdef __VOLUME_SCATTER__ - if (result == VOLUME_PATH_SCATTERED) { - /* direct lighting */ - kernel_path_volume_connect_light(kg, sd, emission_sd, *throughput, state, L); - - /* indirect light bounce */ - if (kernel_path_volume_bounce(kg, sd, throughput, state, &L->state, ray)) - return VOLUME_PATH_SCATTERED; - else - return VOLUME_PATH_MISSED; - } -# endif /* __VOLUME_SCATTER__ */ - } - - return VOLUME_PATH_ATTENUATED; -} -#endif /* __VOLUME__ */ - ccl_device_inline void kernel_path_ao(const KernelGlobals *kg, ShaderData *sd, ShaderData *emission_sd, diff --git a/intern/cycles/kernel/kernel_path_volume.h b/intern/cycles/kernel/kernel_path_volume.h deleted file mode 100644 index 9e628b5..0000000 --- a/intern/cycles/kernel/kernel_path_volume.h +++ /dev/null @@ -1,262 +0,0 @@ -/* - * Copyright 2011-2013 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -CCL_NAMESPACE_BEGIN - -#ifdef __VOLUME_SCATTER__ - -ccl_device_inline void kernel_path_volume_connect_light(const KernelGlobals *kg, - ShaderData *sd, - ShaderData *emission_sd, - float3 throughput, - ccl_addr_space PathState *state, - PathRadiance *L) -{ -# ifdef __EMISSION__ - /* sample illumination from lights to find path contribution */ - Ray light_ray ccl_optional_struct_init; - BsdfEval L_light ccl_optional_struct_init; - bool is_lamp = false; - bool has_emission = false; - - light_ray.t = 0.0f; -# ifdef __OBJECT_MOTION__ - /* connect to light from given point where shader has been evaluated */ - light_ray.time = sd->time; -# endif - - if (kernel_data.integrator.use_direct_light) { - float light_u, light_v; - path_state_rng_2D(kg, state, PRNG_LIGHT_U, &light_u, &light_v); - - LightSample ls ccl_optional_struct_init; - if (light_sample(kg, -1, light_u, light_v, sd->time, sd->P, state->bounce, &ls)) { - float terminate = path_state_rng_light_termination(kg, state); - has_emission = direct_emission( - kg, sd, emission_sd, &ls, state, &light_ray, &L_light, &is_lamp, terminate); - } - } - - /* trace shadow ray */ - float3 shadow; - - const bool blocked = shadow_blocked(kg, sd, emission_sd, state, &light_ray, &shadow); - - if (has_emission && !blocked) { - /* accumulate */ - path_radiance_accum_light(kg, L, state, throughput, &L_light, shadow, 1.0f, is_lamp); - } -# endif /* __EMISSION__ */ -} - -ccl_device_noinline_cpu bool kernel_path_volume_bounce(const KernelGlobals *kg, - ShaderData *sd, - ccl_addr_space float3 *throughput, - ccl_addr_space PathState *state, - PathRadianceState *L_state, - ccl_addr_space Ray *ray) -{ - /* sample phase function */ - float phase_pdf; - BsdfEval phase_eval ccl_optional_struct_init; - float3 phase_omega_in ccl_optional_struct_init; - differential3 phase_domega_in ccl_optional_struct_init; - float phase_u, phase_v; - path_state_rng_2D(kg, state, PRNG_BSDF_U, &phase_u, &phase_v); - int label; - - label = shader_volume_phase_sample( - kg, sd, phase_u, phase_v, &phase_eval, &phase_omega_in, &phase_domega_in, &phase_pdf); - - if (phase_pdf == 0.0f || bsdf_eval_is_zero(&phase_eval)) - return false; - - /* modify throughput */ - path_radiance_bsdf_bounce(kg, L_state, throughput, &phase_eval, phase_pdf, state->bounce, label); - - /* set labels */ - state->ray_pdf = phase_pdf; -# ifdef __LAMP_MIS__ - state->ray_t = 0.0f; -# endif - state->min_ray_pdf = fminf(phase_pdf, state->min_ray_pdf); - - /* update path state */ - path_state_next(kg, state, label); - - /* Russian roulette termination of volume ray scattering. */ - float probability = path_state_continuation_probability(kg, state, *throughput); - - if (probability == 0.0f) { - return false; - } - else if (probability != 1.0f) { - /* Use dimension from the previous bounce, has not been used yet. */ - float terminate = path_state_rng_1D(kg, state, PRNG_TERMINATE - PRNG_BOUNCE_NUM); - - if (terminate >= probability) { - return false; - } - - *throughput /= probability; - } - - /* setup ray */ - ray->P = sd->P; - ray->D = phase_omega_in; - ray->t = FLT_MAX; - -# ifdef __RAY_DIFFERENTIALS__ - ray->dP = sd->dP; - ray->dD = phase_domega_in; -# endif - - return true; -} - -# if !defined(__SPLIT_KERNEL__) && || defined(__VOLUME_DECOUPLED__) -ccl_device void kernel_branched_path_volume_connect_light(const KernelGlobals *kg, - ShaderData *sd, - ShaderData *emission_sd, - float3 throughput, - ccl_addr_space PathState *state, - PathRadiance *L, - bool sample_all_lights, - Ray *ray, - const VolumeSegment *segment) -{ -# ifdef __EMISSION__ - BsdfEval L_light ccl_optional_struct_init; - - int num_lights = 1; - if (sample_all_lights) { - num_lights = kernel_data.integrator.num_all_lights; - if (kernel_data.integrator.pdf_triangles != 0.0f) { - num_lights += 1; - } - } - - for (int i = 0; i < num_lights; ++i) { - /* sample one light at random */ - int num_samples = 1; - int num_all_lights = 1; - uint lamp_rng_hash = state->rng_hash; - bool double_pdf = false; - bool is_mesh_light = false; - bool is_lamp = false; - - if (sample_all_lights) { - /* lamp sampling */ - is_lamp = i < kernel_data.integrator.num_all_lights; - if (is_lamp) { - if (UNLIKELY(light_select_reached_max_bounces(kg, i, state->bounce))) { - continue; - } - num_samples = light_select_num_samples(kg, i); - num_all_lights = kernel_data.integrator.num_all_lights; - lamp_rng_hash = cmj_hash(state->rng_hash, i); - double_pdf = kernel_data.integrator.pdf_triangles != 0.0f; - } - /* mesh light sampling */ - else { - num_samples = kernel_data.integrator.mesh_light_samples; - double_pdf = kernel_data.integrator.num_all_lights != 0; - is_mesh_light = true; - } - } - - float num_samples_inv = 1.0f / (num_samples * num_all_lights); - - for (int j = 0; j < num_samples; j++) { - Ray light_ray ccl_optional_struct_init; - light_ray.t = 0.0f; /* reset ray */ -# ifdef __OBJECT_MOTION__ - light_ray.time = sd->time; -# endif - bool has_emission = false; - - float3 tp = throughput; - - if (kernel_data.integrator.use_direct_light) { - /* sample random position on random light/triangle */ - float light_u, light_v; - path_branched_rng_2D( - kg, lamp_rng_hash, state, j, num_samples, PRNG_LIGHT_U, &light_u, &light_v); - - /* only sample triangle lights */ - if (is_mesh_light && double_pdf) { - light_u = 0.5f * light_u; - } - - LightSample ls ccl_optional_struct_init; - const int lamp = is_lamp ? i : -1; - light_sample(kg, lamp, light_u, light_v, sd->time, ray->P, state->bounce, &ls); - - /* sample position on volume segment */ - float rphase = path_branched_rng_1D( - kg, state->rng_hash, state, j, num_samples, PRNG_PHASE_CHANNEL); - float rscatter = path_branched_rng_1D( - kg, state->rng_hash, state, j, num_samples, PRNG_SCATTER_DISTANCE); - - VolumeIntegrateResult result = kernel_volume_decoupled_scatter(kg, - state, - ray, - sd, - &tp, - rphase, - rscatter, - segment, - (ls.t != FLT_MAX) ? &ls.P : - NULL, - false); - - if (result == VOLUME_PATH_SCATTERED) { - /* todo: split up light_sample so we don't have to call it again with new position */ - if (light_sample(kg, lamp, light_u, light_v, sd->time, sd->P, state->bounce, &ls)) { - if (double_pdf) { - ls.pdf *= 2.0f; - } - - /* sample random light */ - float terminate = path_branched_rng_light_termination( - kg, state->rng_hash, state, j, num_samples); - has_emission = direct_emission( - kg, sd, emission_sd, &ls, state, &light_ray, &L_light, &is_lamp, terminate); - } - } - } - - /* trace shadow ray */ - float3 shadow; - - const bool blocked = shadow_blocked(kg, sd, emission_sd, state, &light_ray, &shadow); - - if (has_emission && !blocked) { - /* accumulate */ - path_radiance_accum_light( - kg, L, state, tp * num_samples_inv, &L_light, shadow, num_samples_inv, is_lamp); - } - } - } -# endif /* __EMISSION__ */ -} -# endif /* __SPLIT_KERNEL__ */ - -#endif /* __VOLUME_SCATTER__ */ - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/kernel_volume.h b/intern/cycles/kernel/kernel_volume.h deleted file mode 100644 index efdc44e..0000000 --- a/intern/cycles/kernel/kernel_volume.h +++ /dev/null @@ -1,989 +0,0 @@ -/* - * Copyright 2011-2013 Blender Foundation - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -CCL_NAMESPACE_BEGIN - -/* Ignore paths that have volume throughput below this value, to avoid unnecessary work - * and precision issues. - * todo: this value could be tweaked or turned into a probability to avoid unnecessary - * work in volumes and subsurface scattering. */ -#define VOLUME_THROUGHPUT_EPSILON 1e-6f - -ccl_device_inline void kernel_volume_step_init(const KernelGlobals *kg, - ccl_addr_space PathState *state, - const float object_step_size, - float t, - float *step_size, - float *step_shade_offset, - float *steps_offset) -{ - const int max_steps = kernel_data.integrator.volume_max_steps; - float step = min(object_step_size, t); - - /* compute exact steps in advance for malloc */ - if (t > max_steps * step) { - step = t / (float)max_steps; - } - - *step_size = step; - - /* Perform shading at this offset within a step, to integrate over - * over the entire step segment. */ - *step_shade_offset = path_state_rng_1D_hash(kg, state, 0x1e31d8a4); - - /* Shift starting point of all segment by this random amount to avoid - * banding artifacts from the volume bounding shape. */ - *steps_offset = path_state_rng_1D_hash(kg, state, 0x3d22c7b3); -} - -/* Volume Shadows - * - * These functions are used to attenuate shadow rays to lights. Both absorption - * and scattering will block light, represented by the extinction coefficient. */ - -/* homogeneous volume: assume shader evaluation at the starts gives - * the extinction coefficient for the entire line segment */ -ccl_device void kernel_volume_shadow_homogeneous(const KernelGlobals *kg, - ccl_addr_space PathState *state, - Ray *ray, - ShaderData *sd, - float3 *throughput) -{ - float3 sigma_t = zero_float3(); - - if (volume_shader_extinction_sample(kg, sd, state, ray->P, &sigma_t)) - *throughput *= volume_color_transmittance(sigma_t, ray->t); -} - -/* heterogeneous volume: integrate stepping through the volume until we - * reach the end, get absorbed entirely, or run out of iterations */ -ccl_device void kernel_volume_shadow_heterogeneous(const KernelGlobals *kg, - ccl_addr_space PathState *state, - Ray *ray, - ShaderData *sd, - float3 *throughput, - const float object_step_size) -{ - float3 tp = *throughput; - - /* Prepare for stepping. - * For shadows we do not offset all segments, since the starting point is - * already a random distance inside the volume. It also appears to create - * banding artifacts for unknown reasons. */ - int max_steps = kernel_data.integrator.volume_max_steps; - float step_size, step_shade_offset, unused; - kernel_volume_step_init( - kg, state, object_step_size, ray->t, &step_size, &step_shade_offset, &unused); - const float steps_offset = 1.0f; - - /* compute extinction at the start */ - float t = 0.0f; - - float3 sum = zero_float3(); - - for (int i = 0; i < max_steps; i++) { - /* advance to new position */ - float new_t = min(ray->t, (i + steps_offset) * step_size); - float dt = new_t - t; - - float3 new_P = ray->P + ray->D * (t + dt * step_shade_offset); - float3 sigma_t = zero_float3(); - - /* compute attenuation over segment */ - if (volume_shader_extinction_sample(kg, sd, state, new_P, &sigma_t)) { - /* Compute expf() only for every Nth step, to save some calculations - * because exp(a)*exp(b) = exp(a+b), also do a quick VOLUME_THROUGHPUT_EPSILON - * check then. */ - sum += (-sigma_t * dt); - if ((i & 0x07) == 0) { /* ToDo: Other interval? */ - tp = *throughput * exp3(sum); - - /* stop if nearly all light is blocked */ - if (tp.x < VOLUME_THROUGHPUT_EPSILON && tp.y < VOLUME_THROUGHPUT_EPSILON && - tp.z < VOLUME_THROUGHPUT_EPSILON) - break; - } - } - - /* stop if at the end of the volume */ - t = new_t; - if (t == ray->t) { - /* Update throughput in case we haven't done it above */ - tp = *throughput * exp3(sum); - break; - } - } - - *throughput = tp; -} - -/* get the volume attenuation over line segment defined by ray, with the - * assumption that there are no surfaces blocking light between the endpoints */ -#if defined(__KERNEL_OPTIX__) && defined(__SHADER_RAYTRACE__) -ccl_device_inline void kernel_volume_shadow(const KernelGlobals *kg, - ShaderData *shadow_sd, - ccl_addr_space PathState *state, - Ray *ray, - float3 *throughput) -{ - optixDirectCall(1, kg, shadow_sd, state, ray, throughput); -} -extern "C" __device__ void __direct_callable__kernel_volume_shadow( -#else -ccl_device void kernel_volume_shadow( -#endif - const KernelGlobals *kg, - ShaderData *shadow_sd, - ccl_addr_space PathState *state, - Ray *ray, - float3 *throughput) -{ - shader_setup_from_volume(kg, shadow_sd, ray); - - float step_size = volume_stack_step_size(kg, state->volume_stack); - if (step_size != FLT_MAX) - kernel_volume_shadow_heterogeneous(kg, state, ray, shadow_sd, throughput, step_size); - else - kernel_volume_shadow_homogeneous(kg, state, ray, shadow_sd, throughput); -} - -#endif /* __VOLUME__ */ - -/* Equi-angular sampling as in: - * "Importance Sampling Techniques for Path Tracing in Participating Media" */ - -ccl_device float kernel_volume_equiangular_sample(Ray *ray, float3 light_P, float xi, float *pdf) -{ - float t = ray->t; - - float delta = dot((light_P - ray->P), ray->D); - float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta); - if (UNLIKELY(D == 0.0f)) { - *pdf = 0.0f; - return 0.0f; - } - float theta_a = -atan2f(delta, D); - float theta_b = atan2f(t - delta, D); - float t_ = D * tanf((xi * theta_b) + (1 - xi) * theta_a); - if (UNLIKELY(theta_b == theta_a)) { - *pdf = 0.0f; - return 0.0f; - } - *pdf = D / ((theta_b - theta_a) * (D * D + t_ * t_)); - - return min(t, delta + t_); /* min is only for float precision errors */ -} - -ccl_device float kernel_volume_equiangular_pdf(Ray *ray, float3 light_P, float sample_t) -{ - float delta = dot((light_P - ray->P), ray->D); - float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta); - if (UNLIKELY(D == 0.0f)) { - return 0.0f; - } - - float t = ray->t; - float t_ = sample_t - delta; - - float theta_a = -atan2f(delta, D); - float theta_b = atan2f(t - delta, D); - if (UNLIKELY(theta_b == theta_a)) { - return 0.0f; - } - - float pdf = D / ((theta_b - theta_a) * (D * D + t_ * t_)); - - return pdf; -} - -/* Distance sampling */ - -ccl_device float kernel_volume_distance_sample( - float max_t, float3 sigma_t, int channel, float xi, float3 *transmittance, float3 *pdf) -{ - /* xi is [0, 1[ so log(0) should never happen, division by zero is - * avoided because sample_sigma_t > 0 when SD_SCATTER is set */ - float sample_sigma_t = volume_channel_get(sigma_t, channel); - float3 full_transmittance = volume_color_transmittance(sigma_t, max_t); - float sample_transmittance = volume_channel_get(full_transmittance, channel); - - float sample_t = min(max_t, -logf(1.0f - xi * (1.0f - sample_transmittance)) / sample_sigma_t); - - *transmittance = volume_color_transmittance(sigma_t, sample_t); - *pdf = safe_divide_color(sigma_t * *transmittance, one_float3() - full_transmittance); - - /* todo: optimization: when taken together with hit/miss decision, - * the full_transmittance cancels out drops out and xi does not - * need to be remapped */ - - return sample_t; -} - -ccl_device float3 kernel_volume_distance_pdf(float max_t, float3 sigma_t, float sample_t) -{ - float3 full_transmittance = volume_color_transmittance(sigma_t, max_t); - float3 transmittance = volume_color_transmittance(sigma_t, sample_t); - - return safe_divide_color(sigma_t * transmittance, one_float3() - full_transmittance); -} - -/* Emission */ - -ccl_device float3 kernel_volume_emission_integrate(VolumeShaderCoefficients *coeff, - int closure_flag, - float3 transmittance, - float t) -{ - /* integral E * exp(-sigma_t * t) from 0 to t = E * (1 - exp(-sigma_t * t))/sigma_t - * this goes to E * t as sigma_t goes to zero - * - * todo: we should use an epsilon to avoid precision issues near zero sigma_t */ - float3 emission = coeff->emission; - - if (closure_flag & SD_EXTINCTION) { - float3 sigma_t = coeff->sigma_t; - - emission.x *= (sigma_t.x > 0.0f) ? (1.0f - transmittance.x) / sigma_t.x : t; - emission.y *= (sigma_t.y > 0.0f) ? (1.0f - transmittance.y) / sigma_t.y : t; - emission.z *= (sigma_t.z > 0.0f) ? (1.0f - transmittance.z) / sigma_t.z : t; - } - else - emission *= t; - - return emission; -} - -/* Volume Path */ - -#ifdef __VOLUME__ - -/* homogeneous volume: assume shader evaluation at the start gives - * the volume shading coefficient for the entire line segment */ -ccl_device VolumeIntegrateResult -kernel_volume_integrate_homogeneous(const KernelGlobals *kg, - ccl_addr_space PathState *state, - Ray *ray, - ShaderData *sd, - PathRadiance *L, - ccl_addr_space float3 *throughput, - bool probalistic_scatter) -{ - VolumeShaderCoefficients coeff ccl_optional_struct_init; - - if (!volume_shader_sample(kg, sd, state, ray->P, &coeff)) - return VOLUME_PATH_MISSED; - - int closure_flag = sd->flag; - float t = ray->t; - float3 new_tp; - -# ifdef __VOLUME_SCATTER__ - /* randomly scatter, and if we do t is shortened */ - if (closure_flag & SD_SCATTER) { - /* Sample channel, use MIS with balance heuristic. */ - float rphase = path_state_rng_1D(kg, state, PRNG_PHASE_CHANNEL); - float3 albedo = safe_divide_color(coeff.sigma_s, coeff.sigma_t); - float3 channel_pdf; - int channel = volume_sample_channel(albedo, *throughput, rphase, &channel_pdf); - - /* decide if we will hit or miss */ - bool scatter = true; - float xi = path_state_rng_1D(kg, state, PRNG_SCATTER_DISTANCE); - - if (probalistic_scatter) { - float sample_sigma_t = volume_channel_get(coeff.sigma_t, channel); - float sample_transmittance = expf(-sample_sigma_t * t); - - if (1.0f - xi >= sample_transmittance) { - scatter = true; - - /* rescale random number so we can reuse it */ - xi = 1.0f - (1.0f - xi - sample_transmittance) / (1.0f - sample_transmittance); - } - else - scatter = false; - } - - if (scatter) { - /* scattering */ - float3 pdf; - float3 transmittance; - float sample_t; - - /* distance sampling */ - sample_t = kernel_volume_distance_sample( - ray->t, coeff.sigma_t, channel, xi, &transmittance, &pdf); - - /* modify pdf for hit/miss decision */ - if (probalistic_scatter) - pdf *= one_float3() - volume_color_transmittance(coeff.sigma_t, t); - - new_tp = *throughput * coeff.sigma_s * transmittance / dot(channel_pdf, pdf); - t = sample_t; - } - else { - /* no scattering */ - float3 transmittance = volume_color_transmittance(coeff.sigma_t, t); - float pdf = dot(channel_pdf, transmittance); - new_tp = *throughput * transmittance / pdf; - } - } - else -# endif - if (closure_flag & SD_EXTINCTION) { - /* absorption only, no sampling needed */ - float3 transmittance = volume_color_transmittance(coeff.sigma_t, t); - new_tp = *throughput * transmittance; - } - else { - new_tp = *throughput; - } - - /* integrate emission attenuated by extinction */ - if (L && (closure_flag & SD_EMISSION)) { - float3 transmittance = volume_color_transmittance(coeff.sigma_t, ray->t); - float3 emission = kernel_volume_emission_integrate( - &coeff, closure_flag, transmittance, ray->t); - path_radiance_accum_emission(kg, L, state, *throughput, emission); - } - - /* modify throughput */ - if (closure_flag & SD_EXTINCTION) { - *throughput = new_tp; - - /* prepare to scatter to new direction */ - if (t < ray->t) { - /* adjust throughput and move to new location */ - sd->P = ray->P + t * ray->D; - - return VOLUME_PATH_SCATTERED; - } - } - - return VOLUME_PATH_ATTENUATED; -} - -/* heterogeneous volume distance sampling: integrate stepping through the - * volume until we reach the end, get absorbed entirely, or run out of - * iterations. this does probabilistically scatter or get transmitted through - * for path tracing where we don't want to branch. */ -ccl_device VolumeIntegrateResult -kernel_volume_integrate_heterogeneous_distance(const KernelGlobals *kg, - ccl_addr_space PathState *state, - Ray *ray, - ShaderData *sd, - PathRadiance *L, - ccl_addr_space float3 *throughput, - const float object_step_size) -{ - float3 tp = *throughput; - - /* Prepare for stepping. - * Using a different step offset for the first step avoids banding artifacts. */ - int max_steps = kernel_data.integrator.volume_max_steps; - float step_size, step_shade_offset, steps_offset; - kernel_volume_step_init( - kg, state, object_step_size, ray->t, &step_size, &step_shade_offset, &steps_offset); - - /* compute coefficients at the start */ - float t = 0.0f; - float3 accum_transmittance = one_float3(); - - /* pick random color channel, we use the Veach one-sample - * model with balance heuristic for the channels */ - float xi = path_state_rng_1D(kg, state, PRNG_SCATTER_DISTANCE); - float rphase = path_state_rng_1D(kg, state, PRNG_PHASE_CHANNEL); - bool has_scatter = false; - - for (int i = 0; i < max_steps; i++) { - /* advance to new position */ - float new_t = min(ray->t, (i + steps_offset) * step_size); - float dt = new_t - t; - - float3 new_P = ray->P + ray->D * (t + dt * step_shade_offset); - VolumeShaderCoefficients coeff ccl_optional_struct_init; - - /* compute segment */ - if (volume_shader_sample(kg, sd, state, new_P, &coeff)) { - int closure_flag = sd->flag; - float3 new_tp; - float3 transmittance; - bool scatter = false; - - /* distance sampling */ -# ifdef __VOLUME_SCATTER__ - if ((closure_flag & SD_SCATTER) || (has_scatter && (closure_flag & SD_EXTINCTION))) { - has_scatter = true; - - /* Sample channel, use MIS with balance heuristic. */ - float3 albedo = safe_divide_color(coeff.sigma_s, coeff.sigma_t); - float3 channel_pdf; - int channel = volume_sample_channel(albedo, tp, rphase, &channel_pdf); - - /* compute transmittance over full step */ - transmittance = volume_color_transmittance(coeff.sigma_t, dt); - - /* decide if we will scatter or continue */ - float sample_transmittance = volume_channel_get(transmittance, channel); - - if (1.0f - xi >= sample_transmittance) { - /* compute sampling distance */ - float sample_sigma_t = volume_channel_get(coeff.sigma_t, channel); - float new_dt = -logf(1.0f - xi) / sample_sigma_t; - new_t = t + new_dt; - - /* transmittance and pdf */ - float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt); - float3 pdf = coeff.sigma_t * new_transmittance; - - /* throughput */ - new_tp = tp * coeff.sigma_s * new_transmittance / dot(channel_pdf, pdf); - scatter = true; - } - else { - /* throughput */ - float pdf = dot(channel_pdf, transmittance); - new_tp = tp * transmittance / pdf; - - /* remap xi so we can reuse it and keep thing stratified */ - xi = 1.0f - (1.0f - xi) / sample_transmittance; - } - } - else -# endif - if (closure_flag & SD_EXTINCTION) { - /* absorption only, no sampling needed */ - transmittance = volume_color_transmittance(coeff.sigma_t, dt); - new_tp = tp * transmittance; - } - else { - transmittance = zero_float3(); - new_tp = tp; - } - - /* integrate emission attenuated by absorption */ - if (L && (closure_flag & SD_EMISSION)) { - float3 emission = kernel_volume_emission_integrate( - &coeff, closure_flag, transmittance, dt); - path_radiance_accum_emission(kg, L, state, tp, emission); - } - - /* modify throughput */ - if (closure_flag & SD_EXTINCTION) { - tp = new_tp; - - /* stop if nearly all light blocked */ - if (tp.x < VOLUME_THROUGHPUT_EPSILON && tp.y < VOLUME_THROUGHPUT_EPSILON && - tp.z < VOLUME_THROUGHPUT_EPSILON) { - tp = zero_float3(); - break; - } - } - - /* prepare to scatter to new direction */ - if (scatter) { - /* adjust throughput and move to new location */ - sd->P = ray->P + new_t * ray->D; - *throughput = tp; - - return VOLUME_PATH_SCATTERED; - } - else { - /* accumulate transmittance */ - accum_transmittance *= transmittance; - } - } - - /* stop if at the end of the volume */ - t = new_t; - if (t == ray->t) - break; - } - - *throughput = tp; - - return VOLUME_PATH_ATTENUATED; -} - -/* get the volume attenuation and emission over line segment defined by - * ray, with the assumption that there are no surfaces blocking light - * between the endpoints. distance sampling is used to decide if we will - * scatter or not. */ -ccl_device_noinline_cpu VolumeIntegrateResult -kernel_volume_integrate(const KernelGlobals *kg, - ccl_addr_space PathState *state, - ShaderData *sd, - Ray *ray, - PathRadiance *L, - ccl_addr_space float3 *throughput, - float step_size) -{ - shader_setup_from_volume(kg, sd, ray); - - if (step_size != FLT_MAX) - return kernel_volume_integrate_heterogeneous_distance( - kg, state, ray, sd, L, throughput, step_size); - else - return kernel_volume_integrate_homogeneous(kg, state, ray, sd, L, throughput, true); -} - -# ifndef __SPLIT_KERNEL__ -/* Decoupled Volume Sampling - * - * VolumeSegment is list of coefficients and transmittance stored at all steps - * through a volume. This can then later be used for decoupled sampling as in: - * "Importance Sampling Techniques for Path Tracing in Participating Media" - * - * On the GPU this is only supported (but currently not enabled) - * for homogeneous volumes (1 step), due to - * no support for malloc/free and too much stack usage with a fix size array. */ - -typedef struct VolumeStep { - float3 sigma_s; /* scatter coefficient */ - float3 sigma_t; /* extinction coefficient */ - float3 accum_transmittance; /* accumulated transmittance including this step */ - float3 cdf_distance; /* cumulative density function for distance sampling */ - float t; /* distance at end of this step */ - float shade_t; /* jittered distance where shading was done in step */ - int closure_flag; /* shader evaluation closure flags */ -} VolumeStep; - -typedef struct VolumeSegment { - VolumeStep stack_step; /* stack storage for homogeneous step, to avoid malloc */ - VolumeStep *steps; /* recorded steps */ - int numsteps; /* number of steps */ - int closure_flag; /* accumulated closure flags from all steps */ - - float3 accum_emission; /* accumulated emission at end of segment */ - float3 accum_transmittance; /* accumulated transmittance at end of segment */ - float3 accum_albedo; /* accumulated average albedo over segment */ - - int sampling_method; /* volume sampling method */ -} VolumeSegment; - -/* record volume steps to the end of the volume. - * - * it would be nice if we could only record up to the point that we need to scatter, - * but the entire segment is needed to do always scattering, rather than probabilistically - * hitting or missing the volume. if we don't know the transmittance at the end of the - * volume we can't generate stratified distance samples up to that transmittance */ -# ifdef __VOLUME_DECOUPLED__ -ccl_device void kernel_volume_decoupled_record(const KernelGlobals *kg, - PathState *state, - Ray *ray, - ShaderData *sd, - VolumeSegment *segment, - const float object_step_size) -{ - /* prepare for volume stepping */ - int max_steps; - float step_size, step_shade_offset, steps_offset; - - if (object_step_size != FLT_MAX) { - max_steps = kernel_data.integrator.volume_max_steps; - kernel_volume_step_init( - kg, state, object_step_size, ray->t, &step_size, &step_shade_offset, &steps_offset); - -# ifdef __KERNEL_CPU__ - /* NOTE: For the branched path tracing it's possible to have direct - * and indirect light integration both having volume segments allocated. - * We detect this using index in the pre-allocated memory. Currently we - * only support two segments allocated at a time, if more needed some - * modifications to the const KernelGlobals will be needed. - * - * This gives us restrictions that decoupled record should only happen - * in the stack manner, meaning if there's subsequent call of decoupled - * record it'll need to free memory before its caller frees memory. - */ - const int index = kg->decoupled_volume_steps_index; - assert(index < sizeof(kg->decoupled_volume_steps) / sizeof(*kg->decoupled_volume_steps)); - if (kg->decoupled_volume_steps[index] == NULL) { - kg->decoupled_volume_steps[index] = (VolumeStep *)malloc(sizeof(VolumeStep) * max_steps); - } - segment->steps = kg->decoupled_volume_steps[index]; - ++kg->decoupled_volume_steps_index; -# else - segment->steps = (VolumeStep *)malloc(sizeof(VolumeStep) * max_steps); -# endif - } - else { - max_steps = 1; - step_size = ray->t; - step_shade_offset = 0.0f; - steps_offset = 1.0f; - segment->steps = &segment->stack_step; - } - - /* init accumulation variables */ - float3 accum_emission = zero_float3(); - float3 accum_transmittance = one_float3(); - float3 accum_albedo = zero_float3(); - float3 cdf_distance = zero_float3(); - float t = 0.0f; - - segment->numsteps = 0; - segment->closure_flag = 0; - bool is_last_step_empty = false; - - VolumeStep *step = segment->steps; - - for (int i = 0; i < max_steps; i++, step++) { - /* advance to new position */ - float new_t = min(ray->t, (i + steps_offset) * step_size); - float dt = new_t - t; - - float3 new_P = ray->P + ray->D * (t + dt * step_shade_offset); - VolumeShaderCoefficients coeff ccl_optional_struct_init; - - /* compute segment */ - if (volume_shader_sample(kg, sd, state, new_P, &coeff)) { - int closure_flag = sd->flag; - float3 sigma_t = coeff.sigma_t; - - /* compute average albedo for channel sampling */ - if (closure_flag & SD_SCATTER) { - accum_albedo += (dt / ray->t) * safe_divide_color(coeff.sigma_s, sigma_t); - } - - /* compute accumulated transmittance */ - float3 transmittance = volume_color_transmittance(sigma_t, dt); - - /* compute emission attenuated by absorption */ - if (closure_flag & SD_EMISSION) { - float3 emission = kernel_volume_emission_integrate( - &coeff, closure_flag, transmittance, dt); - accum_emission += accum_transmittance * emission; - } - - accum_transmittance *= transmittance; - - /* compute pdf for distance sampling */ - float3 pdf_distance = dt * accum_transmittance * coeff.sigma_s; - cdf_distance = cdf_distance + pdf_distance; - - /* write step data */ - step->sigma_t = sigma_t; - step->sigma_s = coeff.sigma_s; - step->closure_flag = closure_flag; - - segment->closure_flag |= closure_flag; - - is_last_step_empty = false; - segment->numsteps++; - } - else { - if (is_last_step_empty) { - /* consecutive empty step, merge */ - step--; - } - else { - /* store empty step */ - step->sigma_t = zero_float3(); - step->sigma_s = zero_float3(); - step->closure_flag = 0; - - segment->numsteps++; - is_last_step_empty = true; - } - } - - step->accum_transmittance = accum_transmittance; - step->cdf_distance = cdf_distance; - step->t = new_t; - step->shade_t = t + dt * step_shade_offset; - - /* stop if at the end of the volume */ - t = new_t; - if (t == ray->t) - break; - - /* stop if nearly all light blocked */ - if (accum_transmittance.x < VOLUME_THROUGHPUT_EPSILON && - accum_transmittance.y < VOLUME_THROUGHPUT_EPSILON && - accum_transmittance.z < VOLUME_THROUGHPUT_EPSILON) - break; - } - - /* store total emission and transmittance */ - segment->accum_emission = accum_emission; - segment->accum_transmittance = accum_transmittance; - segment->accum_albedo = accum_albedo; - - /* normalize cumulative density function for distance sampling */ - VolumeStep *last_step = segment->steps + segment->numsteps - 1; - - if (!is_zero(last_step->cdf_distance)) { - VolumeStep *step = &segment->steps[0]; - int numsteps = segment->numsteps; - float3 inv_cdf_distance_sum = safe_invert_color(last_step->cdf_distance); - - for (int i = 0; i < numsteps; i++, step++) - step->cdf_distance *= inv_cdf_distance_sum; - } -} - -ccl_device void kernel_volume_decoupled_free(const KernelGlobals *kg, VolumeSegment *segment) -{ - if (segment->steps != &segment->stack_step) { -# ifdef __KERNEL_CPU__ - /* NOTE: We only allow free last allocated segment. - * No random order of alloc/free is supported. - */ - assert(kg->decoupled_volume_steps_index > 0); - assert(segment->steps == kg->decoupled_volume_steps[kg->decoupled_volume_steps_index - 1]); - --kg->decoupled_volume_steps_index; -# else - free(segment->steps); -# endif - } -} -# endif /* __VOLUME_DECOUPLED__ */ - -/* scattering for homogeneous and heterogeneous volumes, using decoupled ray - * marching. - * - * function is expected to return VOLUME_PATH_SCATTERED when probalistic_scatter is false */ -ccl_device VolumeIntegrateResult kernel_volume_decoupled_scatter(const KernelGlobals *kg, - PathState *state, - Ray *ray, - ShaderData *sd, - float3 *throughput, - float rphase, - float rscatter, - const VolumeSegment *segment, - const float3 *light_P, - bool probalistic_scatter) -{ - kernel_assert(segment->closure_flag & SD_SCATTER); - - /* Sample color channel, use MIS with balance heuristic. */ - float3 channel_pdf; - int channel = volume_sample_channel(segment->accum_albedo, *throughput, rphase, &channel_pdf); - - float xi = rscatter; - - /* probabilistic scattering decision based on transmittance */ - if (probalistic_scatter) { - float sample_transmittance = volume_channel_get(segment->accum_transmittance, channel); - - if (1.0f - xi >= sample_transmittance) { - /* rescale random number so we can reuse it */ - xi = 1.0f - (1.0f - xi - sample_transmittance) / (1.0f - sample_transmittance); - } - else { - *throughput /= sample_transmittance; - return VOLUME_PATH_MISSED; - } - } - - VolumeStep *step; - float3 transmittance; - float pdf, sample_t; - float mis_weight = 1.0f; - bool distance_sample = true; - bool use_mis = false; - - if (segment->sampling_method && light_P) { - if (segment->sampling_method == SD_VOLUME_MIS) { - /* multiple importance sample: randomly pick between - * equiangular and distance sampling strategy */ - if (xi < 0.5f) { - xi *= 2.0f; - } - else { - xi = (xi - 0.5f) * 2.0f; - distance_sample = false; - } - - use_mis = true; - } - else { - /* only equiangular sampling */ - distance_sample = false; - } - } - - /* distance sampling */ - if (distance_sample) { - /* find step in cdf */ - step = segment->steps; - - float prev_t = 0.0f; - float3 step_pdf_distance = one_float3(); - - if (segment->numsteps > 1) { - float prev_cdf = 0.0f; - float step_cdf = 1.0f; - float3 prev_cdf_distance = zero_float3(); - - for (int i = 0;; i++, step++) { - /* todo: optimize using binary search */ - step_cdf = volume_channel_get(step->cdf_distance, channel); - - if (xi < step_cdf || i == segment->numsteps - 1) - break; - - prev_cdf = step_cdf; - prev_t = step->t; - prev_cdf_distance = step->cdf_distance; - } - - /* remap xi so we can reuse it */ - xi = (xi - prev_cdf) / (step_cdf - prev_cdf); - - /* pdf for picking step */ - step_pdf_distance = step->cdf_distance - prev_cdf_distance; - } - - /* determine range in which we will sample */ - float step_t = step->t - prev_t; - - /* sample distance and compute transmittance */ - float3 distance_pdf; - sample_t = prev_t + kernel_volume_distance_sample( - step_t, step->sigma_t, channel, xi, &transmittance, &distance_pdf); - - /* modify pdf for hit/miss decision */ - if (probalistic_scatter) - distance_pdf *= one_float3() - segment->accum_transmittance; - - pdf = dot(channel_pdf, distance_pdf * step_pdf_distance); - - /* multiple importance sampling */ - if (use_mis) { - float equi_pdf = kernel_volume_equiangular_pdf(ray, *light_P, sample_t); - mis_weight = 2.0f * power_heuristic(pdf, equi_pdf); - } - } - /* equi-angular sampling */ - else { - /* sample distance */ - sample_t = kernel_volume_equiangular_sample(ray, *light_P, xi, &pdf); - - /* find step in which sampled distance is located */ - step = segment->steps; - - float prev_t = 0.0f; - float3 step_pdf_distance = one_float3(); - - if (segment->numsteps > 1) { - float3 prev_cdf_distance = zero_float3(); - - int numsteps = segment->numsteps; - int high = numsteps - 1; - int low = 0; - int mid; - - while (low < high) { - mid = (low + high) >> 1; - - if (sample_t < step[mid].t) - high = mid; - else if (sample_t >= step[mid + 1].t) - low = mid + 1; - else { - /* found our interval in step[mid] .. step[mid+1] */ - prev_t = step[mid].t; - prev_cdf_distance = step[mid].cdf_distance; - step += mid + 1; - break; - } - } - - if (low >= numsteps - 1) { - prev_t = step[numsteps - 1].t; - prev_cdf_distance = step[numsteps - 1].cdf_distance; - step += numsteps - 1; - } - - /* pdf for picking step with distance sampling */ - step_pdf_distance = step->cdf_distance - prev_cdf_distance; - } - - /* determine range in which we will sample */ - float step_t = step->t - prev_t; - float step_sample_t = sample_t - prev_t; - - /* compute transmittance */ - transmittance = volume_color_transmittance(step->sigma_t, step_sample_t); - - /* multiple importance sampling */ - if (use_mis) { - float3 distance_pdf3 = kernel_volume_distance_pdf(step_t, step->sigma_t, step_sample_t); - float distance_pdf = dot(channel_pdf, distance_pdf3 * step_pdf_distance); - mis_weight = 2.0f * power_heuristic(pdf, distance_pdf); - } - } - if (sample_t < 0.0f || pdf == 0.0f) { - return VOLUME_PATH_MISSED; - } - - /* compute transmittance up to this step */ - if (step != segment->steps) - transmittance *= (step - 1)->accum_transmittance; - - /* modify throughput */ - *throughput *= step->sigma_s * transmittance * (mis_weight / pdf); - - /* evaluate shader to create closures at shading point */ - if (segment->numsteps > 1) { - sd->P = ray->P + step->shade_t * ray->D; - - VolumeShaderCoefficients coeff; - volume_shader_sample(kg, sd, state, sd->P, &coeff); - } - - /* move to new position */ - sd->P = ray->P + sample_t * ray->D; - - return VOLUME_PATH_SCATTERED; -} -# endif /* __SPLIT_KERNEL */ - -/* decide if we need to use decoupled or not */ -ccl_device bool kernel_volume_use_decoupled(const KernelGlobals *kg, - bool heterogeneous, - bool direct, - int sampling_method) -{ - /* decoupled ray marching for heterogeneous volumes not supported on the GPU, - * which also means equiangular and multiple importance sampling is not - * support for that case */ - if (!kernel_data.integrator.volume_decoupled) - return false; - -# ifdef __KERNEL_GPU__ - if (heterogeneous) - return false; -# endif - - /* equiangular and multiple importance sampling only implemented for decoupled */ - if (sampling_method != 0) - return true; - - /* for all light sampling use decoupled, reusing shader evaluations is - * typically faster in that case */ - if (direct) - return kernel_data.integrator.sample_all_lights_direct; - else - return kernel_data.integrator.sample_all_lights_indirect; -} - -#endif /* __VOLUME__ */ - -CCL_NAMESPACE_END -- 2.25.1 From 3e94e1c3c98dd74d55eeecec97248916dbd989b4 Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Thu, 17 Jun 2021 17:56:33 +0200 Subject: [PATCH 6/7] Cycles X: refactor path termination in preparation for volume bounces To support similar logic as transparent and emissive surfaces. --- .../integrator/integrator_intersect_closest.h | 40 ++++++++++++++----- .../integrator/integrator_intersect_shadow.h | 4 +- .../integrator/integrator_shade_shadow.h | 2 +- .../integrator/integrator_shade_surface.h | 21 ++++++---- .../kernel/integrator/integrator_state_util.h | 11 +++++ intern/cycles/kernel/kernel_path_state.h | 13 ++---- intern/cycles/kernel/kernel_types.h | 29 ++++++++------ 7 files changed, 79 insertions(+), 41 deletions(-) diff --git a/intern/cycles/kernel/integrator/integrator_intersect_closest.h b/intern/cycles/kernel/integrator/integrator_intersect_closest.h index a4538e8..39a0cad 100644 --- a/intern/cycles/kernel/integrator/integrator_intersect_closest.h +++ b/intern/cycles/kernel/integrator/integrator_intersect_closest.h @@ -34,11 +34,17 @@ ccl_device_forceinline bool integrator_intersect_terminate(INTEGRATOR_STATE_ARGS const int shader_flags) { - /* Optional AO bounce termination. */ + /* Optional AO bounce termination. + * We continue evaluating emissive/transparent surfaces and volumes, similar + * to direct lighting. Only if we know there are none can we terminate the + * path immediately. */ if (path_state_ao_bounce(INTEGRATOR_STATE_PASS)) { if (shader_flags & (SD_HAS_TRANSPARENT_SHADOW | SD_HAS_EMISSION)) { INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_TERMINATE_AFTER_TRANSPARENT; } + else if (!integrator_state_volume_stack_is_empty(INTEGRATOR_STATE_PASS)) { + INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_TERMINATE_AFTER_VOLUME; + } else { return true; } @@ -52,15 +58,20 @@ ccl_device_forceinline bool integrator_intersect_terminate(INTEGRATOR_STATE_ARGS * and evaluating the shader when not needed. Only for emission and transparent * surfaces in front of emission do we need to evaluate the shader, since we * perform MIS as part of indirect rays. */ - const float probability = path_state_continuation_probability(INTEGRATOR_STATE_PASS); + const int path_flag = INTEGRATOR_STATE(path, flag); + const float probability = path_state_continuation_probability(INTEGRATOR_STATE_PASS, path_flag); if (probability != 1.0f) { const float terminate = path_state_rng_1D(kg, &rng_state, PRNG_TERMINATE); if (probability == 0.0f || terminate >= probability) { - if (shader_flags & (SD_HAS_TRANSPARENT_SHADOW | SD_HAS_EMISSION)) { - /* Mark path to be terminated right after shader evaluation. */ - INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_TERMINATE_IMMEDIATE; + if (shader_flags & SD_HAS_EMISSION) { + /* Mark path to be terminated right after shader evaluation on the surface. */ + INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_TERMINATE_ON_NEXT_SURFACE; + } + else if (!integrator_state_volume_stack_is_empty(INTEGRATOR_STATE_PASS)) { + /* TODO: only do this for emissive volumes. */ + INTEGRATOR_STATE_WRITE(path, flag) |= PATH_RAY_TERMINATE_IN_NEXT_VOLUME; } else { return true; @@ -130,11 +141,20 @@ ccl_device void integrator_intersect_closest(INTEGRATOR_STATE_ARGS) integrator_state_write_isect(INTEGRATOR_STATE_PASS, &isect); #ifdef __VOLUME__ - if (INTEGRATOR_STATE_ARRAY(volume_stack, 0, shader) != SHADER_NONE) { - /* Continue with volume kernel if we are inside a volume, regardless - * if we hit anything. */ - INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST, - DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME); + if (!integrator_state_volume_stack_is_empty(INTEGRATOR_STATE_PASS)) { + const int shader = (hit) ? intersection_get_shader(kg, &isect) : SHADER_NONE; + const int flags = (hit) ? kernel_tex_fetch(__shaders, shader).flags : 0; + + if (!integrator_intersect_terminate( + INTEGRATOR_STATE_PASS, &isect, flags)) { + /* Continue with volume kernel if we are inside a volume, regardless + * if we hit anything. */ + INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST, + DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME); + } + else { + INTEGRATOR_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST); + } return; } #endif diff --git a/intern/cycles/kernel/integrator/integrator_intersect_shadow.h b/intern/cycles/kernel/integrator/integrator_intersect_shadow.h index ea25e26..0582670 100644 --- a/intern/cycles/kernel/integrator/integrator_intersect_shadow.h +++ b/intern/cycles/kernel/integrator/integrator_intersect_shadow.h @@ -126,8 +126,8 @@ ccl_device void integrator_intersect_shadow(INTEGRATOR_STATE_ARGS) INTEGRATOR_STATE_PASS, &ray, visibility); #endif - if (opaque_hit && INTEGRATOR_STATE_ARRAY(shadow_volume_stack, 0, shader) == SHADER_NONE) { - /* Hit an opaque surface and no volumes, shadow path ends here. */ + if (opaque_hit) { + /* Hit an opaque surface, shadow path ends here. */ INTEGRATOR_SHADOW_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW); return; } diff --git a/intern/cycles/kernel/integrator/integrator_shade_shadow.h b/intern/cycles/kernel/integrator/integrator_shade_shadow.h index 15d7229..dc56d9e 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_shadow.h +++ b/intern/cycles/kernel/integrator/integrator_shade_shadow.h @@ -105,7 +105,7 @@ ccl_device_inline bool integrate_transparent_shadow(INTEGRATOR_STATE_ARGS, const /* Volume shaders. */ if (hit < num_recorded_hits || !shadow_intersections_remaining(num_hits)) { # ifdef __VOLUME__ - if (INTEGRATOR_STATE_ARRAY(shadow_volume_stack, 0, shader) != SHADER_NONE) { + if (!integrator_state_shadow_volume_stack_is_empty(INTEGRATOR_STATE_PASS)) { const float3 shadow = integrate_transparent_volume_shadow( INTEGRATOR_STATE_PASS, hit, num_recorded_hits); const float3 throughput = INTEGRATOR_STATE(shadow_path, throughput) * shadow; diff --git a/intern/cycles/kernel/integrator/integrator_shade_surface.h b/intern/cycles/kernel/integrator/integrator_shade_surface.h index 2c3da8c..8fc9b78 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_surface.h +++ b/intern/cycles/kernel/integrator/integrator_shade_surface.h @@ -297,14 +297,17 @@ ccl_device bool integrate_surface(INTEGRATOR_STATE_ARGS, if (!(sd.flag & SD_HAS_ONLY_VOLUME)) { #endif - const int path_flag = INTEGRATOR_STATE(path, flag); + { + const int path_flag = INTEGRATOR_STATE(path, flag); #ifdef __SUBSURFACE__ - /* Can skip shader evaluation for BSSRDF exit point without bump mapping. */ - if (!(path_flag & PATH_RAY_SUBSURFACE) || ((sd.flag & SD_HAS_BSSRDF_BUMP))) + /* Can skip shader evaluation for BSSRDF exit point without bump mapping. */ + if (!(path_flag & PATH_RAY_SUBSURFACE) || ((sd.flag & SD_HAS_BSSRDF_BUMP))) #endif - { - /* Evaluate shader. */ - shader_eval_surface(INTEGRATOR_STATE_PASS, &sd, render_buffer, path_flag); + { + /* Evaluate shader. */ + shader_eval_surface( + INTEGRATOR_STATE_PASS, &sd, render_buffer, path_flag); + } } #ifdef __SUBSURFACE__ @@ -344,7 +347,11 @@ ccl_device bool integrate_surface(INTEGRATOR_STATE_ARGS, /* Perform path termination. Most paths have already been terminated in * the intersect_closest kernel, this is just for emission and for dividing * throughput by the probability at the right moment. */ - const float probability = path_state_continuation_probability(INTEGRATOR_STATE_PASS); + const int path_flag = INTEGRATOR_STATE(path, flag); + const float probability = (path_flag & PATH_RAY_TERMINATE_ON_NEXT_SURFACE) ? + 0.0f : + path_state_continuation_probability(INTEGRATOR_STATE_PASS, + path_flag); if (probability == 0.0f) { return false; } diff --git a/intern/cycles/kernel/integrator/integrator_state_util.h b/intern/cycles/kernel/integrator/integrator_state_util.h index 2d38522..a15c187 100644 --- a/intern/cycles/kernel/integrator/integrator_state_util.h +++ b/intern/cycles/kernel/integrator/integrator_state_util.h @@ -113,6 +113,11 @@ ccl_device_forceinline void integrator_state_write_volume_stack(INTEGRATOR_STATE INTEGRATOR_STATE_ARRAY_WRITE(volume_stack, i, shader) = entry.shader; } +ccl_device_forceinline bool integrator_state_volume_stack_is_empty(INTEGRATOR_STATE_CONST_ARGS) +{ + return INTEGRATOR_STATE_ARRAY(volume_stack, 0, shader) == SHADER_NONE; +} + /* Shadow Intersection */ ccl_device_forceinline void integrator_state_write_shadow_isect( @@ -162,6 +167,12 @@ integrator_state_read_shadow_volume_stack(INTEGRATOR_STATE_CONST_ARGS, int i) return entry; } +ccl_device_forceinline bool integrator_state_shadow_volume_stack_is_empty( + INTEGRATOR_STATE_CONST_ARGS) +{ + return INTEGRATOR_STATE_ARRAY(shadow_volume_stack, 0, shader) == SHADER_NONE; +} + ccl_device_forceinline void integrator_state_write_shadow_volume_stack(INTEGRATOR_STATE_ARGS, int i, VolumeStack entry) diff --git a/intern/cycles/kernel/kernel_path_state.h b/intern/cycles/kernel/kernel_path_state.h index c2a156c..f736285 100644 --- a/intern/cycles/kernel/kernel_path_state.h +++ b/intern/cycles/kernel/kernel_path_state.h @@ -105,7 +105,7 @@ ccl_device_inline void path_state_next(INTEGRATOR_STATE_ARGS, int label) flag |= PATH_RAY_TRANSPARENT; if (transparent_bounce >= kernel_data.integrator.transparent_max_bounce) { - flag |= PATH_RAY_TERMINATE_IMMEDIATE; + flag |= PATH_RAY_TERMINATE_ON_NEXT_SURFACE; } if (!kernel_data.integrator.transparent_shadows) @@ -245,15 +245,10 @@ ccl_device_inline uint path_state_ray_visibility(INTEGRATOR_STATE_CONST_ARGS) return visibility; } -ccl_device_inline float path_state_continuation_probability(INTEGRATOR_STATE_CONST_ARGS) +ccl_device_inline float path_state_continuation_probability(INTEGRATOR_STATE_CONST_ARGS, + const uint32_t path_flag) { - const uint32_t flag = INTEGRATOR_STATE(path, flag); - - if (flag & PATH_RAY_TERMINATE_IMMEDIATE) { - /* Ray is to be terminated immediately. */ - return 0.0f; - } - else if (flag & PATH_RAY_TRANSPARENT) { + if (path_flag & PATH_RAY_TRANSPARENT) { const uint32_t transparent_bounce = INTEGRATOR_STATE(path, transparent_bounce); /* Do at least specified number of bounces without RR. */ if (transparent_bounce <= kernel_data.integrator.transparent_min_bounce) { diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index f8a4702..38b0f83 100644 --- a/intern/cycles/kernel/kernel_types.h +++ b/intern/cycles/kernel/kernel_types.h @@ -263,44 +263,49 @@ enum PathRayFlag { PATH_RAY_TRANSPARENT_BACKGROUND = (1 << 14), /* Terminate ray immediately at next bounce. */ - PATH_RAY_TERMINATE_IMMEDIATE = (1 << 15), + PATH_RAY_TERMINATE_ON_NEXT_SURFACE = (1 << 15), + PATH_RAY_TERMINATE_IN_NEXT_VOLUME = (1 << 16), /* Ray is to be terminated, but continue with transparent bounces and * emission as long as we encounter them. This is required to make the * MIS between direct and indirect light rays match, as shadow rays go * through transparent surfaces to reach emission too. */ - PATH_RAY_TERMINATE_AFTER_TRANSPARENT = (1 << 16), + PATH_RAY_TERMINATE_AFTER_TRANSPARENT = (1 << 17), + + /* Terminate ray immediately after volume shading. */ + PATH_RAY_TERMINATE_AFTER_VOLUME = (1 << 18), /* Ray is to be terminated. */ - PATH_RAY_TERMINATE = (PATH_RAY_TERMINATE_IMMEDIATE | PATH_RAY_TERMINATE_AFTER_TRANSPARENT), + PATH_RAY_TERMINATE = (PATH_RAY_TERMINATE_ON_NEXT_SURFACE | PATH_RAY_TERMINATE_IN_NEXT_VOLUME | + PATH_RAY_TERMINATE_AFTER_TRANSPARENT | PATH_RAY_TERMINATE_AFTER_VOLUME), /* Path and shader is being evaluated for direct lighting emission. */ - PATH_RAY_EMISSION = (1 << 17), + PATH_RAY_EMISSION = (1 << 19), /* Perform subsurface scattering. */ - PATH_RAY_SUBSURFACE = (1 << 18), + PATH_RAY_SUBSURFACE = (1 << 20), /* Contribute to denoising features. */ - PATH_RAY_DENOISING_FEATURES = (1 << 19), + PATH_RAY_DENOISING_FEATURES = (1 << 21), /* Render pass categories. */ - PATH_RAY_REFLECT_PASS = (1 << 20), - PATH_RAY_TRANSMISSION_PASS = (1 << 21), - PATH_RAY_VOLUME_PASS = (1 << 22), + PATH_RAY_REFLECT_PASS = (1 << 22), + PATH_RAY_TRANSMISSION_PASS = (1 << 23), + PATH_RAY_VOLUME_PASS = (1 << 24), PATH_RAY_ANY_PASS = (PATH_RAY_REFLECT_PASS | PATH_RAY_TRANSMISSION_PASS | PATH_RAY_VOLUME_PASS), /* Shadow ray is for a light or surface. */ - PATH_RAY_SHADOW_FOR_LIGHT = (1 << 23), + PATH_RAY_SHADOW_FOR_LIGHT = (1 << 25), /* A shadow catcher object was hit and the path was split into two. */ - PATH_RAY_SHADOW_CATCHER_HIT = (1 << 24), + PATH_RAY_SHADOW_CATCHER_HIT = (1 << 26), /* A shadow catcher object was hit and this path traces only shadow catchers, writing them into * their dedicated pass for later division. * * NOTE: Is not covered with `PATH_RAY_ANY_PASS` because shadow catcher does special handling * which is separate from the light passes. */ - PATH_RAY_SHADOW_CATCHER_PASS = (1 << 25), + PATH_RAY_SHADOW_CATCHER_PASS = (1 << 27), }; /* Configure ray visibility bits for rays and objects respectively, -- 2.25.1 From 98e38c30946fe3e9e5b4305422ed2e3b399e174b Mon Sep 17 00:00:00 2001 From: Brecht Van Lommel Date: Wed, 16 Jun 2021 17:24:27 +0200 Subject: [PATCH 7/7] Cycles X: more complete support for volumes Brings back volume rendering support to what it was before on the GPU. That means scattering, emissions and heterogeneous volumes with stepping. However no decoupled or equiangular sampling yet. Rather than separate homogeneous and heterogeneous code paths, there is only heterogeneous now. This ensure shader evaluation is only compiled once. Some optimizations may be possible for the homogeneous case, but that can be looked at later. --- .../integrator/integrator_shade_shadow.h | 24 +- .../integrator/integrator_shade_volume.h | 727 +++++++++++++++++- intern/cycles/kernel/kernel_emission.h | 9 - intern/cycles/kernel/kernel_path_state.h | 12 + 4 files changed, 728 insertions(+), 44 deletions(-) diff --git a/intern/cycles/kernel/integrator/integrator_shade_shadow.h b/intern/cycles/kernel/integrator/integrator_shade_shadow.h index dc56d9e..93b490e 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_shadow.h +++ b/intern/cycles/kernel/integrator/integrator_shade_shadow.h @@ -59,9 +59,10 @@ ccl_device_inline float3 integrate_transparent_surface_shadow(INTEGRATOR_STATE_A } # ifdef __VOLUME__ -ccl_device_inline float3 integrate_transparent_volume_shadow(INTEGRATOR_STATE_ARGS, - const int hit, - const int num_recorded_hits) +ccl_device_inline void integrate_transparent_volume_shadow(INTEGRATOR_STATE_ARGS, + const int hit, + const int num_recorded_hits, + float3 *ccl_restrict throughput) { /* TODO: deduplicate with surface, or does it not matter for memory usage? */ ShaderDataTinyStorage shadow_sd_storage; @@ -80,14 +81,11 @@ ccl_device_inline float3 integrate_transparent_volume_shadow(INTEGRATOR_STATE_AR shader_setup_from_volume(kg, shadow_sd, &ray); - /* Evaluate shader. */ - float3 sigma_a = zero_float3(); - if (!shadow_volume_shader_sample(INTEGRATOR_STATE_PASS, shadow_sd, &sigma_a)) { - return one_float3(); - } + const float step_size = volume_stack_step_size(INTEGRATOR_STATE_PASS, [=](const int i) { + return integrator_state_read_shadow_volume_stack(INTEGRATOR_STATE_PASS, i); + }); - /* Integrate extinction over segment. */ - return volume_color_transmittance(sigma_a, ray.t); + volume_shadow_heterogeneous(INTEGRATOR_STATE_PASS, &ray, shadow_sd, throughput, step_size); } # endif @@ -106,9 +104,9 @@ ccl_device_inline bool integrate_transparent_shadow(INTEGRATOR_STATE_ARGS, const if (hit < num_recorded_hits || !shadow_intersections_remaining(num_hits)) { # ifdef __VOLUME__ if (!integrator_state_shadow_volume_stack_is_empty(INTEGRATOR_STATE_PASS)) { - const float3 shadow = integrate_transparent_volume_shadow( - INTEGRATOR_STATE_PASS, hit, num_recorded_hits); - const float3 throughput = INTEGRATOR_STATE(shadow_path, throughput) * shadow; + float3 throughput = INTEGRATOR_STATE(shadow_path, throughput); + integrate_transparent_volume_shadow( + INTEGRATOR_STATE_PASS, hit, num_recorded_hits, &throughput); if (is_zero(throughput)) { return true; } diff --git a/intern/cycles/kernel/integrator/integrator_shade_volume.h b/intern/cycles/kernel/integrator/integrator_shade_volume.h index e7b0c18..e6939b9 100644 --- a/intern/cycles/kernel/integrator/integrator_shade_volume.h +++ b/intern/cycles/kernel/integrator/integrator_shade_volume.h @@ -31,6 +31,12 @@ typedef enum VolumeIntegrateResult { VOLUME_PATH_MISSED = 2 } VolumeIntegrateResult; +/* Ignore paths that have volume throughput below this value, to avoid unnecessary work + * and precision issues. + * todo: this value could be tweaked or turned into a probability to avoid unnecessary + * work in volumes and subsurface scattering. */ +# define VOLUME_THROUGHPUT_EPSILON 1e-6f + /* Volume shader properties * * extinction coefficient = absorption coefficient + scattering coefficient @@ -82,8 +88,9 @@ ccl_device_inline bool volume_shader_sample(INTEGRATOR_STATE_ARGS, for (int i = 0; i < sd->num_closure; i++) { const ShaderClosure *sc = &sd->closure[i]; - if (CLOSURE_IS_VOLUME(sc->type)) + if (CLOSURE_IS_VOLUME(sc->type)) { coeff->sigma_s += sc->weight; + } } } @@ -94,14 +101,700 @@ ccl_device_inline bool volume_shader_sample(INTEGRATOR_STATE_ARGS, return true; } + +ccl_device_forceinline void volume_step_init(const KernelGlobals *kg, + const RNGState *rng_state, + const float object_step_size, + float t, + float *step_size, + float *step_shade_offset, + float *steps_offset, + int *max_steps) +{ + if (object_step_size == FLT_MAX) { + /* Homogeneous volume. */ + *step_size = t; + *step_shade_offset = 0.0f; + *steps_offset = 1.0f; + *max_steps = 1; + } + else { + /* Heterogeneous volume. */ + *max_steps = kernel_data.integrator.volume_max_steps; + float step = min(object_step_size, t); + + /* compute exact steps in advance for malloc */ + if (t > *max_steps * step) { + step = t / (float)*max_steps; + } + + *step_size = step; + + /* Perform shading at this offset within a step, to integrate over + * over the entire step segment. */ + *step_shade_offset = path_state_rng_1D_hash(kg, rng_state, 0x1e31d8a4); + + /* Shift starting point of all segment by this random amount to avoid + * banding artifacts from the volume bounding shape. */ + *steps_offset = path_state_rng_1D_hash(kg, rng_state, 0x3d22c7b3); + } +} + +/* Volume Shadows + * + * These functions are used to attenuate shadow rays to lights. Both absorption + * and scattering will block light, represented by the extinction coefficient. */ + +# if 0 +/* homogeneous volume: assume shader evaluation at the starts gives + * the extinction coefficient for the entire line segment */ +ccl_device void volume_shadow_homogeneous(INTEGRATOR_STATE_ARGS, + Ray *ccl_restrict ray, + ShaderData *ccl_restrict sd, + float3 *ccl_restrict throughput) +{ + float3 sigma_t = zero_float3(); + + if (shadow_volume_shader_sample(INTEGRATOR_STATE_PASS, sd, &sigma_t)) { + *throughput *= volume_color_transmittance(sigma_t, ray->t); + } +} +# endif + +/* heterogeneous volume: integrate stepping through the volume until we + * reach the end, get absorbed entirely, or run out of iterations */ +ccl_device void volume_shadow_heterogeneous(INTEGRATOR_STATE_ARGS, + Ray *ccl_restrict ray, + ShaderData *ccl_restrict sd, + float3 *ccl_restrict throughput, + const float object_step_size) +{ + /* Load random number state. */ + RNGState rng_state; + shadow_path_state_rng_load(INTEGRATOR_STATE_PASS, &rng_state); + + float3 tp = *throughput; + + /* Prepare for stepping. + * For shadows we do not offset all segments, since the starting point is + * already a random distance inside the volume. It also appears to create + * banding artifacts for unknown reasons. */ + int max_steps; + float step_size, step_shade_offset, unused; + volume_step_init(kg, + &rng_state, + object_step_size, + ray->t, + &step_size, + &step_shade_offset, + &unused, + &max_steps); + const float steps_offset = 1.0f; + + /* compute extinction at the start */ + float t = 0.0f; + + float3 sum = zero_float3(); + + for (int i = 0; i < max_steps; i++) { + /* advance to new position */ + float new_t = min(ray->t, (i + steps_offset) * step_size); + float dt = new_t - t; + + float3 new_P = ray->P + ray->D * (t + dt * step_shade_offset); + float3 sigma_t = zero_float3(); + + /* compute attenuation over segment */ + sd->P = new_P; + if (shadow_volume_shader_sample(INTEGRATOR_STATE_PASS, sd, &sigma_t)) { + /* Compute expf() only for every Nth step, to save some calculations + * because exp(a)*exp(b) = exp(a+b), also do a quick VOLUME_THROUGHPUT_EPSILON + * check then. */ + sum += (-sigma_t * dt); + if ((i & 0x07) == 0) { /* ToDo: Other interval? */ + tp = *throughput * exp3(sum); + + /* stop if nearly all light is blocked */ + if (tp.x < VOLUME_THROUGHPUT_EPSILON && tp.y < VOLUME_THROUGHPUT_EPSILON && + tp.z < VOLUME_THROUGHPUT_EPSILON) + break; + } + } + + /* stop if at the end of the volume */ + t = new_t; + if (t == ray->t) { + /* Update throughput in case we haven't done it above */ + tp = *throughput * exp3(sum); + break; + } + } + + *throughput = tp; +} + +/* Equi-angular sampling as in: + * "Importance Sampling Techniques for Path Tracing in Participating Media" */ + +ccl_device float volume_equiangular_sample(Ray *ray, float3 light_P, float xi, float *pdf) +{ + float t = ray->t; + + float delta = dot((light_P - ray->P), ray->D); + float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta); + if (UNLIKELY(D == 0.0f)) { + *pdf = 0.0f; + return 0.0f; + } + float theta_a = -atan2f(delta, D); + float theta_b = atan2f(t - delta, D); + float t_ = D * tanf((xi * theta_b) + (1 - xi) * theta_a); + if (UNLIKELY(theta_b == theta_a)) { + *pdf = 0.0f; + return 0.0f; + } + *pdf = D / ((theta_b - theta_a) * (D * D + t_ * t_)); + + return min(t, delta + t_); /* min is only for float precision errors */ +} + +ccl_device float volume_equiangular_pdf(Ray *ray, float3 light_P, float sample_t) +{ + float delta = dot((light_P - ray->P), ray->D); + float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta); + if (UNLIKELY(D == 0.0f)) { + return 0.0f; + } + + float t = ray->t; + float t_ = sample_t - delta; + + float theta_a = -atan2f(delta, D); + float theta_b = atan2f(t - delta, D); + if (UNLIKELY(theta_b == theta_a)) { + return 0.0f; + } + + float pdf = D / ((theta_b - theta_a) * (D * D + t_ * t_)); + + return pdf; +} + +/* Distance sampling */ + +ccl_device float volume_distance_sample( + float max_t, float3 sigma_t, int channel, float xi, float3 *transmittance, float3 *pdf) +{ + /* xi is [0, 1[ so log(0) should never happen, division by zero is + * avoided because sample_sigma_t > 0 when SD_SCATTER is set */ + float sample_sigma_t = volume_channel_get(sigma_t, channel); + float3 full_transmittance = volume_color_transmittance(sigma_t, max_t); + float sample_transmittance = volume_channel_get(full_transmittance, channel); + + float sample_t = min(max_t, -logf(1.0f - xi * (1.0f - sample_transmittance)) / sample_sigma_t); + + *transmittance = volume_color_transmittance(sigma_t, sample_t); + *pdf = safe_divide_color(sigma_t * *transmittance, one_float3() - full_transmittance); + + /* todo: optimization: when taken together with hit/miss decision, + * the full_transmittance cancels out drops out and xi does not + * need to be remapped */ + + return sample_t; +} + +ccl_device float3 volume_distance_pdf(float max_t, float3 sigma_t, float sample_t) +{ + float3 full_transmittance = volume_color_transmittance(sigma_t, max_t); + float3 transmittance = volume_color_transmittance(sigma_t, sample_t); + + return safe_divide_color(sigma_t * transmittance, one_float3() - full_transmittance); +} + +/* Emission */ + +ccl_device float3 volume_emission_integrate(VolumeShaderCoefficients *coeff, + int closure_flag, + float3 transmittance, + float t) +{ + /* integral E * exp(-sigma_t * t) from 0 to t = E * (1 - exp(-sigma_t * t))/sigma_t + * this goes to E * t as sigma_t goes to zero + * + * todo: we should use an epsilon to avoid precision issues near zero sigma_t */ + float3 emission = coeff->emission; + + if (closure_flag & SD_EXTINCTION) { + float3 sigma_t = coeff->sigma_t; + + emission.x *= (sigma_t.x > 0.0f) ? (1.0f - transmittance.x) / sigma_t.x : t; + emission.y *= (sigma_t.y > 0.0f) ? (1.0f - transmittance.y) / sigma_t.y : t; + emission.z *= (sigma_t.z > 0.0f) ? (1.0f - transmittance.z) / sigma_t.z : t; + } + else + emission *= t; + + return emission; +} + +/* Volume Path */ + +# if 0 +/* homogeneous volume: assume shader evaluation at the start gives + * the volume shading coefficient for the entire line segment */ +ccl_device VolumeIntegrateResult +volume_integrate_homogeneous(INTEGRATOR_STATE_ARGS, + Ray *ccl_restrict ray, + ShaderData *ccl_restrict sd, + ccl_addr_space float3 *ccl_restrict throughput, + const RNGState *rng_state, + const bool probalistic_scatter, + ccl_global float *ccl_restrict render_buffer) +{ + /* Evaluate shader. */ + VolumeShaderCoefficients coeff ccl_optional_struct_init; + + if (!volume_shader_sample(INTEGRATOR_STATE_PASS, sd, &coeff)) { + return VOLUME_PATH_MISSED; + } + + const int closure_flag = sd->flag; + float t = ray->t; + float3 new_tp; + +# ifdef __VOLUME_SCATTER__ + /* randomly scatter, and if we do t is shortened */ + if (closure_flag & SD_SCATTER) { + /* Sample channel, use MIS with balance heuristic. */ + const float rphase = path_state_rng_1D(kg, rng_state, PRNG_PHASE_CHANNEL); + const float3 albedo = safe_divide_color(coeff.sigma_s, coeff.sigma_t); + float3 channel_pdf; + const int channel = volume_sample_channel(albedo, *throughput, rphase, &channel_pdf); + + /* decide if we will hit or miss */ + bool scatter = true; + float xi = path_state_rng_1D(kg, rng_state, PRNG_SCATTER_DISTANCE); + + if (probalistic_scatter) { + float sample_sigma_t = volume_channel_get(coeff.sigma_t, channel); + float sample_transmittance = expf(-sample_sigma_t * t); + + if (1.0f - xi >= sample_transmittance) { + scatter = true; + + /* rescale random number so we can reuse it */ + xi = 1.0f - (1.0f - xi - sample_transmittance) / (1.0f - sample_transmittance); + } + else + scatter = false; + } + + if (scatter) { + /* scattering */ + float3 pdf; + float3 transmittance; + float sample_t; + + /* distance sampling */ + sample_t = volume_distance_sample(ray->t, coeff.sigma_t, channel, xi, &transmittance, &pdf); + + /* modify pdf for hit/miss decision */ + if (probalistic_scatter) + pdf *= one_float3() - volume_color_transmittance(coeff.sigma_t, t); + + new_tp = *throughput * coeff.sigma_s * transmittance / dot(channel_pdf, pdf); + t = sample_t; + } + else { + /* no scattering */ + float3 transmittance = volume_color_transmittance(coeff.sigma_t, t); + float pdf = dot(channel_pdf, transmittance); + new_tp = *throughput * transmittance / pdf; + } + } + else +# endif + if (closure_flag & SD_EXTINCTION) { + /* absorption only, no sampling needed */ + float3 transmittance = volume_color_transmittance(coeff.sigma_t, t); + new_tp = *throughput * transmittance; + } + else { + new_tp = *throughput; + } + + /* integrate emission attenuated by extinction */ + if (closure_flag & SD_EMISSION) { + float3 transmittance = volume_color_transmittance(coeff.sigma_t, ray->t); + float3 emission = volume_emission_integrate(&coeff, closure_flag, transmittance, ray->t); + + kernel_accum_emission(INTEGRATOR_STATE_PASS, *throughput, emission, render_buffer); + } + + /* modify throughput */ + if (closure_flag & SD_EXTINCTION) { + *throughput = new_tp; + + /* prepare to scatter to new direction */ + if (t < ray->t) { + /* adjust throughput and move to new location */ + sd->P = ray->P + t * ray->D; + + return VOLUME_PATH_SCATTERED; + } + } + + return VOLUME_PATH_ATTENUATED; +} +# endif + +/* heterogeneous volume distance sampling: integrate stepping through the + * volume until we reach the end, get absorbed entirely, or run out of + * iterations. this does probabilistically scatter or get transmitted through + * for path tracing where we don't want to branch. */ +ccl_device VolumeIntegrateResult +volume_integrate_heterogeneous(INTEGRATOR_STATE_ARGS, + Ray *ccl_restrict ray, + ShaderData *ccl_restrict sd, + ccl_addr_space float3 *ccl_restrict throughput, + const RNGState *rng_state, + ccl_global float *ccl_restrict render_buffer, + const float object_step_size) +{ + float3 tp = *throughput; + + /* Prepare for stepping. + * Using a different step offset for the first step avoids banding artifacts. */ + int max_steps; + float step_size, step_shade_offset, steps_offset; + volume_step_init(kg, + rng_state, + object_step_size, + ray->t, + &step_size, + &step_shade_offset, + &steps_offset, + &max_steps); + + /* compute coefficients at the start */ + float t = 0.0f; + float3 accum_transmittance = one_float3(); + + /* pick random color channel, we use the Veach one-sample + * model with balance heuristic for the channels */ + float xi = path_state_rng_1D(kg, rng_state, PRNG_SCATTER_DISTANCE); + float rphase = path_state_rng_1D(kg, rng_state, PRNG_PHASE_CHANNEL); + bool has_scatter = false; + + for (int i = 0; i < max_steps; i++) { + /* advance to new position */ + float new_t = min(ray->t, (i + steps_offset) * step_size); + float dt = new_t - t; + + float3 new_P = ray->P + ray->D * (t + dt * step_shade_offset); + VolumeShaderCoefficients coeff ccl_optional_struct_init; + + /* compute segment */ + sd->P = new_P; + if (volume_shader_sample(INTEGRATOR_STATE_PASS, sd, &coeff)) { + int closure_flag = sd->flag; + float3 new_tp; + float3 transmittance; + bool scatter = false; + + /* distance sampling */ +# ifdef __VOLUME_SCATTER__ + if ((closure_flag & SD_SCATTER) || (has_scatter && (closure_flag & SD_EXTINCTION))) { + has_scatter = true; + + /* Sample channel, use MIS with balance heuristic. */ + float3 albedo = safe_divide_color(coeff.sigma_s, coeff.sigma_t); + float3 channel_pdf; + int channel = volume_sample_channel(albedo, tp, rphase, &channel_pdf); + + /* compute transmittance over full step */ + transmittance = volume_color_transmittance(coeff.sigma_t, dt); + + /* decide if we will scatter or continue */ + float sample_transmittance = volume_channel_get(transmittance, channel); + + if (1.0f - xi >= sample_transmittance) { + /* compute sampling distance */ + float sample_sigma_t = volume_channel_get(coeff.sigma_t, channel); + float new_dt = -logf(1.0f - xi) / sample_sigma_t; + new_t = t + new_dt; + + /* transmittance and pdf */ + float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt); + float3 pdf = coeff.sigma_t * new_transmittance; + + /* throughput */ + new_tp = tp * coeff.sigma_s * new_transmittance / dot(channel_pdf, pdf); + scatter = true; + } + else { + /* throughput */ + float pdf = dot(channel_pdf, transmittance); + new_tp = tp * transmittance / pdf; + + /* remap xi so we can reuse it and keep thing stratified */ + xi = 1.0f - (1.0f - xi) / sample_transmittance; + } + } + else +# endif + if (closure_flag & SD_EXTINCTION) { + /* absorption only, no sampling needed */ + transmittance = volume_color_transmittance(coeff.sigma_t, dt); + new_tp = tp * transmittance; + } + else { + transmittance = zero_float3(); + new_tp = tp; + } + + /* integrate emission attenuated by absorption */ + if (closure_flag & SD_EMISSION) { + float3 emission = volume_emission_integrate(&coeff, closure_flag, transmittance, dt); + kernel_accum_emission(INTEGRATOR_STATE_PASS, tp, emission, render_buffer); + } + + /* modify throughput */ + if (closure_flag & SD_EXTINCTION) { + tp = new_tp; + + /* stop if nearly all light blocked */ + if (tp.x < VOLUME_THROUGHPUT_EPSILON && tp.y < VOLUME_THROUGHPUT_EPSILON && + tp.z < VOLUME_THROUGHPUT_EPSILON) { + tp = zero_float3(); + break; + } + } + + /* prepare to scatter to new direction */ + if (scatter) { + /* adjust throughput and move to new location */ + sd->P = ray->P + new_t * ray->D; + *throughput = tp; + + return VOLUME_PATH_SCATTERED; + } + else { + /* accumulate transmittance */ + accum_transmittance *= transmittance; + } + } + + /* stop if at the end of the volume */ + t = new_t; + if (t == ray->t) + break; + } + + *throughput = tp; + + return VOLUME_PATH_ATTENUATED; +} + +# ifdef __EMISSION__ +/* Path tracing: sample point on light and evaluate light shader, then + * queue shadow ray to be traced. */ +ccl_device_forceinline void integrate_volume_direct_light(INTEGRATOR_STATE_ARGS, + ShaderData *sd, + const RNGState *rng_state) +{ + /* Test if there is a light or BSDF that needs direct light. */ + if (!kernel_data.integrator.use_direct_light) { + return; + } + + /* Sample position on a light. */ + LightSample ls ccl_optional_struct_init; + { + const int path_flag = INTEGRATOR_STATE(path, flag); + const uint bounce = INTEGRATOR_STATE(path, bounce); + float light_u, light_v; + path_state_rng_2D(kg, rng_state, PRNG_LIGHT_U, &light_u, &light_v); + + if (!light_sample(kg, light_u, light_v, sd->time, sd->P, bounce, path_flag, &ls)) { + return; + } + } + + if (ls.shader & SHADER_EXCLUDE_SCATTER) { + return; + } + + kernel_assert(ls.pdf != 0.0f); + + /* Evaluate light shader. + * + * TODO: can we reuse sd memory? In theory we can move this after + * integrate_surface_bounce, evaluate the BSDF, and only then evaluate + * the light shader. This could also move to its own kernel, for + * non-constant light sources. */ + ShaderDataTinyStorage emission_sd_storage; + ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage); + const float3 light_eval = light_sample_shader_eval( + INTEGRATOR_STATE_PASS, emission_sd, &ls, sd->time); + if (is_zero(light_eval)) { + return; + } + + /* Evaluate BSDF. */ + BsdfEval phase_eval ccl_optional_struct_init; + shader_volume_phase_eval(kg, sd, ls.D, &phase_eval, ls.pdf, ls.shader); + + bsdf_eval_mul3(&phase_eval, light_eval / ls.pdf); + + /* Path termination. */ + const float terminate = path_state_rng_light_termination(kg, rng_state); + if (light_sample_terminate(kg, &ls, &phase_eval, terminate)) { + return; + } + + /* Create shadow ray. */ + Ray ray ccl_optional_struct_init; + light_sample_to_shadow_ray(sd, &ls, &ray); + const bool is_light = light_sample_is_light(&ls); + + /* Write shadow ray and associated state to global memory. */ + integrator_state_write_shadow_ray(INTEGRATOR_STATE_PASS, &ray); + + /* Copy state from main path to shadow path. */ + const uint16_t bounce = INTEGRATOR_STATE(path, bounce); + const uint16_t transparent_bounce = INTEGRATOR_STATE(path, transparent_bounce); + uint32_t shadow_flag = INTEGRATOR_STATE(path, flag); + shadow_flag |= (is_light) ? PATH_RAY_SHADOW_FOR_LIGHT : 0; + shadow_flag |= PATH_RAY_VOLUME_PASS; + const float3 diffuse_glossy_ratio = (bounce == 0) ? one_float3() : + INTEGRATOR_STATE(path, diffuse_glossy_ratio); + const float3 throughput = INTEGRATOR_STATE(path, throughput) * bsdf_eval_sum(&phase_eval); + + INTEGRATOR_STATE_WRITE(shadow_path, flag) = shadow_flag; + INTEGRATOR_STATE_WRITE(shadow_path, bounce) = bounce; + INTEGRATOR_STATE_WRITE(shadow_path, transparent_bounce) = transparent_bounce; + INTEGRATOR_STATE_WRITE(shadow_path, diffuse_glossy_ratio) = diffuse_glossy_ratio; + INTEGRATOR_STATE_WRITE(shadow_path, throughput) = throughput; + + integrator_state_copy_volume_stack_to_shadow(INTEGRATOR_STATE_PASS); + + /* Branch off shadow kernel. */ + INTEGRATOR_SHADOW_PATH_INIT(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW); +} +# endif + +/* Path tracing: scatter in new direction using phase function. */ +ccl_device_forceinline bool integrate_volume_phase_scatter(INTEGRATOR_STATE_ARGS, + ShaderData *sd, + const RNGState *rng_state) +{ + float phase_u, phase_v; + path_state_rng_2D(kg, rng_state, PRNG_BSDF_U, &phase_u, &phase_v); + + /* Phase closure, sample direction. */ + float phase_pdf; + BsdfEval phase_eval ccl_optional_struct_init; + float3 phase_omega_in ccl_optional_struct_init; + differential3 phase_domega_in ccl_optional_struct_init; + + const int label = shader_volume_phase_sample( + kg, sd, phase_u, phase_v, &phase_eval, &phase_omega_in, &phase_domega_in, &phase_pdf); + + if (phase_pdf == 0.0f || bsdf_eval_is_zero(&phase_eval)) { + return false; + } + + /* Setup ray. */ + INTEGRATOR_STATE_WRITE(ray, P) = sd->P; + INTEGRATOR_STATE_WRITE(ray, D) = normalize(phase_omega_in); + INTEGRATOR_STATE_WRITE(ray, t) = FLT_MAX; + +# ifdef __RAY_DIFFERENTIALS__ + INTEGRATOR_STATE_WRITE(ray, dP) = differential_make_compact(sd->dP); + INTEGRATOR_STATE_WRITE(ray, dD) = differential_make_compact(phase_domega_in); +# endif + + /* Update throughput. */ + float3 throughput = INTEGRATOR_STATE(path, throughput); + throughput *= bsdf_eval_sum(&phase_eval) / phase_pdf; + INTEGRATOR_STATE_WRITE(path, throughput) = throughput; + INTEGRATOR_STATE_WRITE(path, diffuse_glossy_ratio) = one_float3(); + + /* Update path state */ + INTEGRATOR_STATE_WRITE(path, mis_ray_pdf) = phase_pdf; + INTEGRATOR_STATE_WRITE(path, mis_ray_t) = 0.0f; + INTEGRATOR_STATE_WRITE(path, min_ray_pdf) = fminf(phase_pdf, + INTEGRATOR_STATE(path, min_ray_pdf)); + + path_state_next(INTEGRATOR_STATE_PASS, label); + return true; +} + +/* get the volume attenuation and emission over line segment defined by + * ray, with the assumption that there are no surfaces blocking light + * between the endpoints. distance sampling is used to decide if we will + * scatter or not. */ +ccl_device VolumeIntegrateResult volume_integrate(INTEGRATOR_STATE_ARGS, + Ray *ccl_restrict ray, + ccl_global float *ccl_restrict render_buffer) +{ + ShaderData sd; + shader_setup_from_volume(kg, &sd, ray); + + float3 throughput = INTEGRATOR_STATE(path, throughput); + + /* Load random number state. */ + RNGState rng_state; + path_state_rng_load(INTEGRATOR_STATE_PASS, &rng_state); + + const float step_size = volume_stack_step_size(INTEGRATOR_STATE_PASS, [=](const int i) { + return integrator_state_read_volume_stack(INTEGRATOR_STATE_PASS, i); + }); + + VolumeIntegrateResult result = volume_integrate_heterogeneous( + INTEGRATOR_STATE_PASS, ray, &sd, &throughput, &rng_state, render_buffer, step_size); + + /* Perform path termination. The intersect_closest will have already marked this path + * to be terminated. That will shading evaluating to leave out any scattering closures, + * but emission and absorption are still handled for multiple importance sampling. */ + const uint32_t path_flag = INTEGRATOR_STATE(path, flag); + const float probability = (path_flag & PATH_RAY_TERMINATE_IN_NEXT_VOLUME) ? + 0.0f : + path_state_continuation_probability(INTEGRATOR_STATE_PASS, + path_flag); + if (probability == 0.0f) { + return VOLUME_PATH_MISSED; + } + else if (result == VOLUME_PATH_SCATTERED) { + /* Only divide throughput by probability if we scatter. For the attenuation + * case the next surface will already do this division. */ + if (probability != 1.0f) { + throughput /= probability; + } + } + + INTEGRATOR_STATE_WRITE(path, throughput) = throughput; + + if (result == VOLUME_PATH_SCATTERED) { + /* Direct light. */ + integrate_volume_direct_light(INTEGRATOR_STATE_PASS, &sd, &rng_state); + + /* Scatter. */ + if (!integrate_volume_phase_scatter(INTEGRATOR_STATE_PASS, &sd, &rng_state)) { + return VOLUME_PATH_MISSED; + } + } + + return result; +} + #endif ccl_device void integrator_shade_volume(INTEGRATOR_STATE_ARGS, ccl_global float *ccl_restrict render_buffer) { #ifdef __VOLUME__ - VolumeIntegrateResult result = VOLUME_PATH_ATTENUATED; - /* Setup shader data. */ Ray ray ccl_optional_struct_init; integrator_state_read_ray(INTEGRATOR_STATE_PASS, &ray); @@ -109,37 +802,27 @@ ccl_device void integrator_shade_volume(INTEGRATOR_STATE_ARGS, Intersection isect ccl_optional_struct_init; integrator_state_read_isect(INTEGRATOR_STATE_PASS, &isect); - ShaderData sd; - shader_setup_from_volume(kg, &sd, &ray); + /* Set ray length to current segment. */ + ray.t = (isect.prim != PRIM_NONE) ? isect.t : FLT_MAX; /* Clean volume stack for background rays. */ if (isect.prim == PRIM_NONE) { volume_stack_clean(INTEGRATOR_STATE_PASS); } - /* Evaluate shader. */ - /* TODO: implement scattering and heterogeneous media. */ - VolumeShaderCoefficients coeff ccl_optional_struct_init; - if (volume_shader_sample(INTEGRATOR_STATE_PASS, &sd, &coeff)) { - /* Integrate extinction over segment. */ - float3 throughput = INTEGRATOR_STATE(path, throughput); - throughput *= exp3(-coeff.sigma_t * isect.t); - INTEGRATOR_STATE_WRITE(path, throughput) = throughput; - } - - if (result == VOLUME_PATH_MISSED) { - /* End path. */ - INTEGRATOR_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME); - return; - } - else if (result == VOLUME_PATH_SCATTERED) { - /* TODO: handle path termination like intersect closest. */ + VolumeIntegrateResult result = volume_integrate(INTEGRATOR_STATE_PASS, &ray, render_buffer); + if (result == VOLUME_PATH_SCATTERED) { /* Queue intersect_closest kernel. */ INTEGRATOR_PATH_NEXT(DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST); return; } + else if (result == VOLUME_PATH_MISSED) { + /* End path. */ + INTEGRATOR_PATH_TERMINATE(DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME); + return; + } else { /* Continue to background, light or surface. */ if (isect.prim == PRIM_NONE) { diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h index c5173cc..eea54f9 100644 --- a/intern/cycles/kernel/kernel_emission.h +++ b/intern/cycles/kernel/kernel_emission.h @@ -105,15 +105,6 @@ ccl_device_inline bool light_sample_terminate(const KernelGlobals *ccl_restrict BsdfEval *eval, const float rand_terminate) { - /* TODO: move in volume phase evaluation. */ -#if 0 -# ifdef __PASSES__ - /* use visibility flag to skip lights */ - if (ls->shader & SHADER_EXCLUDE_SCATTER) - eval->volume = zero_float3(); -# endif -#endif - if (bsdf_eval_is_zero(eval)) { return true; } diff --git a/intern/cycles/kernel/kernel_path_state.h b/intern/cycles/kernel/kernel_path_state.h index f736285..b2f1fa7 100644 --- a/intern/cycles/kernel/kernel_path_state.h +++ b/intern/cycles/kernel/kernel_path_state.h @@ -67,6 +67,8 @@ ccl_device_inline void path_state_init_integrator(INTEGRATOR_STATE_ARGS, INTEGRATOR_STATE_ARRAY_WRITE(volume_stack, 0, object) = OBJECT_NONE; INTEGRATOR_STATE_ARRAY_WRITE(volume_stack, 0, shader) = kernel_data.background.volume_shader; + INTEGRATOR_STATE_ARRAY_WRITE(volume_stack, 1, object) = OBJECT_NONE; + INTEGRATOR_STATE_ARRAY_WRITE(volume_stack, 1, shader) = SHADER_NONE; #ifdef __DENOISING_FEATURES__ if (kernel_data.film.have_denoising_passes) { @@ -302,6 +304,16 @@ ccl_device_inline void path_state_rng_load(INTEGRATOR_STATE_CONST_ARGS, RNGState rng_state->sample = INTEGRATOR_STATE(path, sample); } +ccl_device_inline void shadow_path_state_rng_load(INTEGRATOR_STATE_CONST_ARGS, RNGState *rng_state) +{ + const uint shadow_bounces = INTEGRATOR_STATE_WRITE(shadow_path, transparent_bounce) - + INTEGRATOR_STATE_WRITE(path, transparent_bounce); + + rng_state->rng_hash = INTEGRATOR_STATE(path, rng_hash); + rng_state->rng_offset = INTEGRATOR_STATE(path, rng_offset) + PRNG_BOUNCE_NUM * shadow_bounces; + rng_state->sample = INTEGRATOR_STATE(path, sample); +} + ccl_device_inline float path_state_rng_1D(const KernelGlobals *kg, const RNGState *rng_state, int dimension) -- 2.25.1