21 #ifdef __DENOISING_FEATURES__
27 float path_total_shaded)
42 float value = path_total_shaded /
max(path_total, 1e-7f);
51 if (
state->denoising_feature_weight == 0.0f) {
65 float sum_weight = 0.0f, sum_nonspecular_weight = 0.0f;
67 for (
int i = 0; i < sd->num_closure; i++) {
74 normal += sc->N * sc->sample_weight;
75 sum_weight += sc->sample_weight;
77 float3 closure_albedo = sc->weight;
95 diffuse_albedo += closure_albedo;
96 sum_nonspecular_weight += sc->sample_weight;
99 specular_albedo += closure_albedo;
104 if ((sum_weight == 0.0f) || (sum_nonspecular_weight * 4.0f > sum_weight)) {
105 if (sum_weight != 0.0f) {
115 state->denoising_feature_throughput * diffuse_albedo);
117 state->denoising_feature_weight = 0.0f;
120 state->denoising_feature_throughput *= specular_albedo;
125 #ifdef __KERNEL_DEBUG__
131 if (flag &
PASSMASK(BVH_TRAVERSED_NODES)) {
133 L->debug_data.num_bvh_traversed_nodes);
135 if (flag &
PASSMASK(BVH_TRAVERSED_INSTANCES)) {
137 L->debug_data.num_bvh_traversed_instances);
139 if (flag &
PASSMASK(BVH_INTERSECTIONS)) {
141 L->debug_data.num_bvh_intersections);
145 L->debug_data.num_ray_bounces);
150 #ifdef __KERNEL_CPU__
151 # define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name) \
152 kernel_write_id_pass_cpu(buffer, depth * 2, id, matte_weight, kg->coverage_##name)
157 (*map)[
id] += matte_weight;
161 # define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name) \
162 kernel_write_id_slots_gpu(buffer, depth * 2, id, matte_weight)
181 int path_flag =
state->flag;
189 if (!((flag | light_flag) &
PASS_ANY))
195 if (
state->sample == 0) {
229 const float matte_weight =
average(throughput) *
231 if (matte_weight > 0.0f) {
236 cryptomatte_buffer,
kernel_data.film.cryptomatte_depth,
id, matte_weight,
object);
246 cryptomatte_buffer,
kernel_data.film.cryptomatte_depth,
id, matte_weight, asset);
261 float mist_inv_depth =
kernel_data.film.mist_inv_depth;
264 float mist =
saturate((depth - mist_start) * mist_inv_depth);
267 float mist_falloff =
kernel_data.film.mist_falloff;
269 if (mist_falloff == 1.0f)
271 else if (mist_falloff == 2.0f)
273 else if (mist_falloff == 0.5f)
276 mist =
powf(mist, mist_falloff);
295 if (light_flag &
PASSMASK(DIFFUSE_INDIRECT))
297 if (light_flag &
PASSMASK(GLOSSY_INDIRECT))
299 if (light_flag &
PASSMASK(TRANSMISSION_INDIRECT))
301 L->indirect_transmission);
302 if (light_flag &
PASSMASK(VOLUME_INDIRECT))
304 if (light_flag &
PASSMASK(DIFFUSE_DIRECT))
306 if (light_flag &
PASSMASK(GLOSSY_DIRECT))
308 if (light_flag &
PASSMASK(TRANSMISSION_DIRECT))
310 L->direct_transmission);
311 if (light_flag &
PASSMASK(VOLUME_DIRECT))
314 if (light_flag &
PASSMASK(EMISSION))
316 if (light_flag &
PASSMASK(BACKGROUND))
321 if (light_flag &
PASSMASK(DIFFUSE_COLOR))
323 if (light_flag &
PASSMASK(GLOSSY_COLOR))
325 if (light_flag &
PASSMASK(TRANSMISSION_COLOR))
327 L->color_transmission);
328 if (light_flag &
PASSMASK(SHADOW)) {
356 #ifdef __DENOISING_FEATURES__
358 # ifdef __SHADOW_TRICKS__
359 kernel_write_denoising_shadow(
kg,
365 kernel_write_denoising_shadow(
371 kernel_write_pass_float3_variance(
373 kernel_write_pass_float3_unaligned(
buffer +
kernel_data.film.pass_denoising_clean, clean);
376 kernel_write_pass_float3_variance(
buffer +
kernel_data.film.pass_denoising_data +
381 kernel_write_pass_float3_variance(
buffer +
kernel_data.film.pass_denoising_data +
383 L->denoising_normal);
384 kernel_write_pass_float3_variance(
buffer +
kernel_data.film.pass_denoising_data +
386 L->denoising_albedo);
387 kernel_write_pass_float_variance(
392 #ifdef __KERNEL_DEBUG__
393 kernel_write_debug_passes(
kg,
buffer,
L);
401 kernel_data.integrator.adaptive_threshold > 0.0f) {
404 make_float4(L_sum.
x * 2.0f, L_sum.
y * 2.0f, L_sum.
z * 2.0f, 0.0f));
406 #ifdef __KERNEL_CPU__
409 const int step =
kernel_data.integrator.adaptive_step;
411 if ((
sample & (step - 1)) == (step - 1)) {
424 #ifdef __ATOMIC_PASS_WRITE__
ATOMIC_INLINE uint32_t atomic_fetch_and_or_uint32(uint32_t *p, uint32_t x)
CCL_NAMESPACE_BEGIN ccl_device_inline float bsdf_get_specular_roughness_squared(const ShaderClosure *sc)
ccl_device float3 bsdf_principled_hair_albedo(ShaderClosure *sc)
unordered_map< float, float > CoverageMap
static CCL_NAMESPACE_BEGIN const double alpha
ccl_device_inline float object_cryptomatte_asset_id(KernelGlobals *kg, int object)
ccl_device_inline float object_pass_id(KernelGlobals *kg, int object)
ccl_device int shader_pass_id(KernelGlobals *kg, const ShaderData *sd)
ccl_device_inline float object_cryptomatte_id(KernelGlobals *kg, int object)
ccl_device_inline float3 primitive_uv(KernelGlobals *kg, ShaderData *sd)
ccl_device_inline float4 primitive_motion_vector(KernelGlobals *kg, ShaderData *sd)
IconTextureDrawCall normal
CCL_NAMESPACE_BEGIN ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, const ShaderData *sd)
ccl_device_inline void path_radiance_split_denoising(KernelGlobals *kg, PathRadiance *L, float3 *noisy, float3 *clean)
ccl_device_inline float3 path_radiance_clamp_and_sum(KernelGlobals *kg, PathRadiance *L, float *alpha)
CCL_NAMESPACE_BEGIN ccl_device void kernel_do_adaptive_stopping(KernelGlobals *kg, ccl_global float *buffer, int sample)
ccl_device_inline float camera_distance(KernelGlobals *kg, float3 P)
ccl_device_inline float camera_z_depth(KernelGlobals *kg, float3 P)
#define ccl_device_inline
#define CCL_NAMESPACE_END
#define make_float4(x, y, z, w)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_write_id_slots(ccl_global float *buffer, int num_slots, float id, float weight)
ccl_device_inline void kernel_write_light_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L)
#define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name)
ccl_device_inline void kernel_write_data_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L, ShaderData *sd, ccl_addr_space PathState *state, float3 throughput)
ccl_device_inline void kernel_write_result(KernelGlobals *kg, ccl_global float *buffer, int sample, PathRadiance *L)
ccl_device_inline size_t kernel_write_id_slots_gpu(ccl_global float *buffer, size_t depth, float id, float matte_weight)
#define PROFILING_OBJECT(object)
#define PROFILING_INIT(kg, event)
ccl_device_inline bool sample_is_even(int pattern, int sample)
ccl_device float3 shader_bsdf_alpha(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd)
ccl_device float shader_cryptomatte_id(KernelGlobals *kg, int shader)
ccl_device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_bsdf_average_normal(KernelGlobals *kg, ShaderData *sd)
__kernel void ccl_constant KernelData ccl_global void ccl_global char ccl_global int ccl_global char ccl_global unsigned int ccl_global float * buffer
#define PASSMASK_COMPONENT(comp)
@ PATH_RAY_SINGLE_PASS_DONE
@ DENOISING_PASS_SHADOW_B
@ DENOISING_PASS_SHADOW_A
ccl_device_inline void kernel_write_pass_float3(ccl_global float *buffer, float3 value)
ccl_device_inline void kernel_write_pass_float4(ccl_global float *buffer, float4 value)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_write_pass_float(ccl_global float *buffer, float value)
static void sample(SocketReader *reader, int x, int y, float color[4])
#define CLOSURE_IS_BSDF_MICROFACET_FRESNEL(type)
#define CLOSURE_IS_BSDF_OR_BSSRDF(type)
@ CLOSURE_BSDF_HAIR_PRINCIPLED_ID
@ CLOSURE_BSDF_PRINCIPLED_SHEEN_ID
ccl_device_inline float ensure_finite(float v)
ccl_device_inline float saturate(float a)
ccl_device_inline float sqr(float a)
ccl_device_inline float average(const float2 &a)
ccl_device_inline float3 zero_float3()
ccl_device_inline float3 ensure_finite3(float3 v)