diff options
Diffstat (limited to 'intern/cycles/kernel/kernel_volume.h')
| -rw-r--r-- | intern/cycles/kernel/kernel_volume.h | 1440 |
1 files changed, 0 insertions, 1440 deletions
diff --git a/intern/cycles/kernel/kernel_volume.h b/intern/cycles/kernel/kernel_volume.h deleted file mode 100644 index f6b34be040e..00000000000 --- a/intern/cycles/kernel/kernel_volume.h +++ /dev/null @@ -1,1440 +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. - */ - -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 - -/* Events for probalistic scattering */ - -typedef enum VolumeIntegrateResult { - VOLUME_PATH_SCATTERED = 0, - VOLUME_PATH_ATTENUATED = 1, - VOLUME_PATH_MISSED = 2 -} VolumeIntegrateResult; - -/* Volume shader properties - * - * extinction coefficient = absorption coefficient + scattering coefficient - * sigma_t = sigma_a + sigma_s */ - -typedef struct VolumeShaderCoefficients { - float3 sigma_t; - float3 sigma_s; - float3 emission; -} VolumeShaderCoefficients; - -#ifdef __VOLUME__ - -/* evaluate shader to get extinction coefficient at P */ -ccl_device_inline bool volume_shader_extinction_sample(KernelGlobals *kg, - ShaderData *sd, - ccl_addr_space PathState *state, - float3 P, - float3 *extinction) -{ - sd->P = P; - shader_eval_volume(kg, sd, state, state->volume_stack, PATH_RAY_SHADOW); - - if (sd->flag & SD_EXTINCTION) { - const float density = object_volume_density(kg, sd->object); - *extinction = sd->closure_transparent_extinction * density; - return true; - } - else { - return false; - } -} - -/* evaluate shader to get absorption, scattering and emission at P */ -ccl_device_inline bool volume_shader_sample(KernelGlobals *kg, - ShaderData *sd, - ccl_addr_space PathState *state, - float3 P, - VolumeShaderCoefficients *coeff) -{ - sd->P = P; - shader_eval_volume(kg, sd, state, state->volume_stack, state->flag); - - if (!(sd->flag & (SD_EXTINCTION | SD_SCATTER | SD_EMISSION))) - return false; - - coeff->sigma_s = zero_float3(); - coeff->sigma_t = (sd->flag & SD_EXTINCTION) ? sd->closure_transparent_extinction : zero_float3(); - coeff->emission = (sd->flag & SD_EMISSION) ? sd->closure_emission_background : zero_float3(); - - if (sd->flag & SD_SCATTER) { - for (int i = 0; i < sd->num_closure; i++) { - const ShaderClosure *sc = &sd->closure[i]; - - if (CLOSURE_IS_VOLUME(sc->type)) - coeff->sigma_s += sc->weight; - } - } - - const float density = object_volume_density(kg, sd->object); - coeff->sigma_s *= density; - coeff->sigma_t *= density; - coeff->emission *= density; - - return true; -} - -#endif /* __VOLUME__ */ - -ccl_device float3 volume_color_transmittance(float3 sigma, float t) -{ - return exp3(-sigma * t); -} - -ccl_device float kernel_volume_channel_get(float3 value, int channel) -{ - return (channel == 0) ? value.x : ((channel == 1) ? value.y : value.z); -} - -#ifdef __VOLUME__ - -ccl_device float volume_stack_step_size(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(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(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(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(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(KernelGlobals *kg, - ShaderData *shadow_sd, - ccl_addr_space PathState *state, - Ray *ray, - float3 *throughput) -{ - optixDirectCall<void>(1, kg, shadow_sd, state, ray, throughput); -} -extern "C" __device__ void __direct_callable__kernel_volume_shadow( -# else -ccl_device_noinline void kernel_volume_shadow( -# endif - 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 = kernel_volume_channel_get(sigma_t, channel); - float3 full_transmittance = volume_color_transmittance(sigma_t, max_t); - float sample_transmittance = kernel_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 */ - -ccl_device int kernel_volume_sample_channel(float3 albedo, - float3 throughput, - float rand, - float3 *pdf) -{ - /* Sample color channel proportional to throughput and single scattering - * albedo, to significantly reduce noise with many bounce, following: - * - * "Practical and Controllable Subsurface Scattering for Production Path - * Tracing". Matt Jen-Yuan Chiang, Peter Kutz, Brent Burley. SIGGRAPH 2016. */ - float3 weights = fabs(throughput * albedo); - float sum_weights = weights.x + weights.y + weights.z; - float3 weights_pdf; - - if (sum_weights > 0.0f) { - weights_pdf = weights / sum_weights; - } - else { - weights_pdf = make_float3(1.0f / 3.0f, 1.0f / 3.0f, 1.0f / 3.0f); - } - - *pdf = weights_pdf; - - /* OpenCL does not support -> on float3, so don't use pdf->x. */ - if (rand < weights_pdf.x) { - return 0; - } - else if (rand < weights_pdf.x + weights_pdf.y) { - return 1; - } - else { - return 2; - } -} - -#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(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 = kernel_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 = kernel_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(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 = kernel_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 = kernel_volume_channel_get(transmittance, channel); - - if (1.0f - xi >= sample_transmittance) { - /* compute sampling distance */ - float sample_sigma_t = kernel_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(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(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 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(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(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 = kernel_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 = kernel_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 = kernel_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(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; -} - -/* Volume Stack - * - * This is an array of object/shared ID's that the current segment of the path - * is inside of. */ - -ccl_device void kernel_volume_stack_init(KernelGlobals *kg, - ShaderData *stack_sd, - ccl_addr_space const PathState *state, - ccl_addr_space const Ray *ray, - ccl_addr_space VolumeStack *stack) -{ - /* NULL ray happens in the baker, does it need proper initialization of - * camera in volume? - */ - if (!kernel_data.cam.is_inside_volume || ray == NULL) { - /* Camera is guaranteed to be in the air, only take background volume - * into account in this case. - */ - if (kernel_data.background.volume_shader != SHADER_NONE) { - stack[0].shader = kernel_data.background.volume_shader; - stack[0].object = PRIM_NONE; - stack[1].shader = SHADER_NONE; - } - else { - stack[0].shader = SHADER_NONE; - } - return; - } - - kernel_assert(state->flag & PATH_RAY_CAMERA); - - Ray volume_ray = *ray; - volume_ray.t = FLT_MAX; - - const uint visibility = (state->flag & PATH_RAY_ALL_VISIBILITY); - int stack_index = 0, enclosed_index = 0; - -# ifdef __VOLUME_RECORD_ALL__ - Intersection hits[2 * VOLUME_STACK_SIZE + 1]; - uint num_hits = scene_intersect_volume_all( - kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, visibility); - if (num_hits > 0) { - int enclosed_volumes[VOLUME_STACK_SIZE]; - Intersection *isect = hits; - - qsort(hits, num_hits, sizeof(Intersection), intersections_compare); - - for (uint hit = 0; hit < num_hits; ++hit, ++isect) { - shader_setup_from_ray(kg, stack_sd, isect, &volume_ray); - if (stack_sd->flag & SD_BACKFACING) { - bool need_add = true; - for (int i = 0; i < enclosed_index && need_add; ++i) { - /* If ray exited the volume and never entered to that volume - * it means that camera is inside such a volume. - */ - if (enclosed_volumes[i] == stack_sd->object) { - need_add = false; - } - } - for (int i = 0; i < stack_index && need_add; ++i) { - /* Don't add intersections twice. */ - if (stack[i].object == stack_sd->object) { - need_add = false; - break; - } - } - if (need_add && stack_index < VOLUME_STACK_SIZE - 1) { - stack[stack_index].object = stack_sd->object; - stack[stack_index].shader = stack_sd->shader; - ++stack_index; - } - } - else { - /* If ray from camera enters the volume, this volume shouldn't - * be added to the stack on exit. - */ - enclosed_volumes[enclosed_index++] = stack_sd->object; - } - } - } -# else - int enclosed_volumes[VOLUME_STACK_SIZE]; - int step = 0; - - while (stack_index < VOLUME_STACK_SIZE - 1 && enclosed_index < VOLUME_STACK_SIZE - 1 && - step < 2 * VOLUME_STACK_SIZE) { - Intersection isect; - if (!scene_intersect_volume(kg, &volume_ray, &isect, visibility)) { - break; - } - - shader_setup_from_ray(kg, stack_sd, &isect, &volume_ray); - if (stack_sd->flag & SD_BACKFACING) { - /* If ray exited the volume and never entered to that volume - * it means that camera is inside such a volume. - */ - bool need_add = true; - for (int i = 0; i < enclosed_index && need_add; ++i) { - /* If ray exited the volume and never entered to that volume - * it means that camera is inside such a volume. - */ - if (enclosed_volumes[i] == stack_sd->object) { - need_add = false; - } - } - for (int i = 0; i < stack_index && need_add; ++i) { - /* Don't add intersections twice. */ - if (stack[i].object == stack_sd->object) { - need_add = false; - break; - } - } - if (need_add) { - stack[stack_index].object = stack_sd->object; - stack[stack_index].shader = stack_sd->shader; - ++stack_index; - } - } - else { - /* If ray from camera enters the volume, this volume shouldn't - * be added to the stack on exit. - */ - enclosed_volumes[enclosed_index++] = stack_sd->object; - } - - /* Move ray forward. */ - volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng); - ++step; - } -# endif - /* stack_index of 0 means quick checks outside of the kernel gave false - * positive, nothing to worry about, just we've wasted quite a few of - * ticks just to come into conclusion that camera is in the air. - * - * In this case we're doing the same above -- check whether background has - * volume. - */ - if (stack_index == 0 && kernel_data.background.volume_shader == SHADER_NONE) { - stack[0].shader = kernel_data.background.volume_shader; - stack[0].object = OBJECT_NONE; - stack[1].shader = SHADER_NONE; - } - else { - stack[stack_index].shader = SHADER_NONE; - } -} - -ccl_device void kernel_volume_stack_enter_exit(KernelGlobals *kg, - ShaderData *sd, - ccl_addr_space VolumeStack *stack) -{ - /* todo: we should have some way for objects to indicate if they want the - * world shader to work inside them. excluding it by default is problematic - * because non-volume objects can't be assumed to be closed manifolds */ - - if (!(sd->flag & SD_HAS_VOLUME)) - return; - - if (sd->flag & SD_BACKFACING) { - /* exit volume object: remove from stack */ - for (int i = 0; stack[i].shader != SHADER_NONE; i++) { - if (stack[i].object == sd->object) { - /* shift back next stack entries */ - do { - stack[i] = stack[i + 1]; - i++; - } while (stack[i].shader != SHADER_NONE); - - return; - } - } - } - else { - /* enter volume object: add to stack */ - int i; - - for (i = 0; stack[i].shader != SHADER_NONE; i++) { - /* already in the stack? then we have nothing to do */ - if (stack[i].object == sd->object) - return; - } - - /* if we exceed the stack limit, ignore */ - if (i >= VOLUME_STACK_SIZE - 1) - return; - - /* add to the end of the stack */ - stack[i].shader = sd->shader; - stack[i].object = sd->object; - stack[i + 1].shader = SHADER_NONE; - } -} - -# ifdef __SUBSURFACE__ -ccl_device void kernel_volume_stack_update_for_subsurface(KernelGlobals *kg, - ShaderData *stack_sd, - Ray *ray, - ccl_addr_space VolumeStack *stack) -{ - kernel_assert(kernel_data.integrator.use_volumes); - - Ray volume_ray = *ray; - -# ifdef __VOLUME_RECORD_ALL__ - Intersection hits[2 * VOLUME_STACK_SIZE + 1]; - uint num_hits = scene_intersect_volume_all( - kg, &volume_ray, hits, 2 * VOLUME_STACK_SIZE, PATH_RAY_ALL_VISIBILITY); - if (num_hits > 0) { - Intersection *isect = hits; - - qsort(hits, num_hits, sizeof(Intersection), intersections_compare); - - for (uint hit = 0; hit < num_hits; ++hit, ++isect) { - shader_setup_from_ray(kg, stack_sd, isect, &volume_ray); - kernel_volume_stack_enter_exit(kg, stack_sd, stack); - } - } -# else - Intersection isect; - int step = 0; - float3 Pend = ray->P + ray->D * ray->t; - while (step < 2 * VOLUME_STACK_SIZE && - scene_intersect_volume(kg, &volume_ray, &isect, PATH_RAY_ALL_VISIBILITY)) { - shader_setup_from_ray(kg, stack_sd, &isect, &volume_ray); - kernel_volume_stack_enter_exit(kg, stack_sd, stack); - - /* Move ray forward. */ - volume_ray.P = ray_offset(stack_sd->P, -stack_sd->Ng); - if (volume_ray.t != FLT_MAX) { - volume_ray.D = normalize_len(Pend - volume_ray.P, &volume_ray.t); - } - ++step; - } -# endif -} -# endif - -/* Clean stack after the last bounce. - * - * It is expected that all volumes are closed manifolds, so at the time when ray - * hits nothing (for example, it is a last bounce which goes to environment) the - * only expected volume in the stack is the world's one. All the rest volume - * entries should have been exited already. - * - * This isn't always true because of ray intersection precision issues, which - * could lead us to an infinite non-world volume in the stack, causing render - * artifacts. - * - * Use this function after the last bounce to get rid of all volumes apart from - * the world's one after the last bounce to avoid render artifacts. - */ -ccl_device_inline void kernel_volume_clean_stack(KernelGlobals *kg, - ccl_addr_space VolumeStack *volume_stack) -{ - if (kernel_data.background.volume_shader != SHADER_NONE) { - /* Keep the world's volume in stack. */ - volume_stack[1].shader = SHADER_NONE; - } - else { - volume_stack[0].shader = SHADER_NONE; - } -} - -#endif /* __VOLUME__ */ - -CCL_NAMESPACE_END |