38 #if defined(__VOLUME__) || defined(__SUBSURFACE__)
74 #ifdef __KERNEL_DEBUG__
76 L->debug_data.num_bvh_traversed_nodes += isect->num_traversed_nodes;
77 L->debug_data.num_bvh_traversed_instances += isect->num_traversed_instances;
78 L->debug_data.num_bvh_intersections += isect->num_intersections;
80 L->debug_data.num_ray_bounces++;
101 light_ray.P = ray->
P -
state->ray_t * ray->
D;
102 state->ray_t += isect->t;
103 light_ray.D = ray->
D;
104 light_ray.t =
state->ray_t;
105 light_ray.time = ray->
time;
106 light_ray.dD = ray->
dD;
107 light_ray.dP = ray->
dP;
125 L->transparent +=
average(throughput);
136 throughput *=
kernel_data.background.ao_bounces_factor;
139 #ifdef __BACKGROUND__
146 #ifndef __SPLIT_KERNEL__
163 kernel_volume_clean_stack(
kg,
state->volume_stack);
171 Ray volume_ray = *ray;
172 volume_ray.
t = (hit) ? isect->t : FLT_MAX;
174 float step_size = volume_stack_step_size(
kg,
state->volume_stack);
176 # ifdef __VOLUME_DECOUPLED__
177 int sampling_method = volume_stack_sampling_method(
kg,
state->volume_stack);
179 bool decoupled = kernel_volume_use_decoupled(
kg, step_size, direct, sampling_method);
183 VolumeSegment volume_segment;
185 shader_setup_from_volume(
kg, sd, &volume_ray);
186 kernel_volume_decoupled_record(
kg,
state, &volume_ray, sd, &volume_segment, step_size);
188 volume_segment.sampling_method = sampling_method;
197 if (volume_segment.closure_flag &
SD_SCATTER) {
201 kernel_branched_path_volume_connect_light(
202 kg, sd, emission_sd, *throughput,
state,
L,
all, &volume_ray, &volume_segment);
210 result = kernel_volume_decoupled_scatter(
211 kg,
state, &volume_ray, sd, throughput, rphase, rscatter, &volume_segment,
NULL,
true);
215 kernel_volume_decoupled_free(
kg, &volume_segment);
218 if (kernel_path_volume_bounce(
kg, sd, throughput,
state, &
L->state, ray))
224 *throughput *= volume_segment.accum_transmittance;
232 kg,
state, sd, &volume_ray,
L, throughput, step_size);
234 # ifdef __VOLUME_SCATTER__
237 kernel_path_volume_connect_light(
kg, sd, emission_sd, *throughput,
state,
L);
240 if (kernel_path_volume_bounce(
kg, sd, throughput,
state, &
L->state, ray))
265 #ifdef __SHADOW_TRICKS__
274 path_radiance_accum_shadowcatcher(
L, throughput, bg);
289 L->transparent +=
average(holdout_weight * throughput);
302 if (
kernel_data.integrator.filter_glossy != FLT_MAX) {
305 if (blur_pdf < 1.0f) {
306 float blur_roughness =
sqrtf(1.0f - blur_pdf) * 0.5f;
323 #ifdef __KERNEL_OPTIX__
340 float bsdf_u, bsdf_v;
344 float ao_factor =
kernel_data.background.ao_factor;
352 if (
dot(sd->Ng, ao_D) > 0.0f && ao_pdf != 0.0f) {
359 light_ray.
time = sd->time;
360 light_ray.
dP = sd->dP;
372 #ifndef __SPLIT_KERNEL__
374 # if defined(__BRANCHED_PATH__) || defined(__BAKING__)
384 # ifdef __SUBSURFACE__
386 kernel_path_subsurface_init_indirect(&ss_indirect);
403 kg, sd,
state, ray, &throughput, &isect, hit, emission_sd,
L);
444 if (probability == 0.0f) {
447 else if (probability != 1.0f) {
450 if (terminate >= probability)
453 throughput /= probability;
456 # ifdef __DENOISING_FEATURES__
457 kernel_update_denoising_features(
kg, sd,
state,
L);
462 if (
kernel_data.integrator.use_ambient_occlusion) {
467 # ifdef __SUBSURFACE__
471 if (kernel_path_subsurface_scatter(
472 kg, sd, emission_sd,
L,
state, ray, &throughput, &ss_indirect)) {
478 # if defined(__EMISSION__)
481 kernel_branched_path_surface_connect_light(
482 kg, sd, emission_sd,
state, throughput, 1.0f,
L,
all);
493 # ifdef __SUBSURFACE__
498 kernel_path_subsurface_setup_indirect(
kg, &ss_indirect,
state, ray,
L, &throughput);
522 # ifdef __SUBSURFACE__
524 kernel_path_subsurface_init_indirect(&ss_indirect);
541 kg, &sd,
state, ray, &throughput, &isect, hit, emission_sd,
L);
582 if (probability == 0.0f) {
585 else if (probability != 1.0f) {
587 if (terminate >= probability)
590 throughput /= probability;
593 # ifdef __DENOISING_FEATURES__
594 kernel_update_denoising_features(
kg, &sd,
state,
L);
599 if (
kernel_data.integrator.use_ambient_occlusion) {
604 # ifdef __SUBSURFACE__
608 if (kernel_path_subsurface_scatter(
609 kg, &sd, emission_sd,
L,
state, ray, &throughput, &ss_indirect)) {
629 # ifdef __SUBSURFACE__
634 kernel_path_subsurface_setup_indirect(
kg, &ss_indirect,
state, ray,
L, &throughput);
649 int index = offset +
x +
y *
stride;
652 buffer += index * pass_stride;
657 if ((*aux).w > 0.0f) {
684 # ifdef __KERNEL_OPTIX__
687 for (
int i = 0; i <
sizeof(
L); ++i)
688 reinterpret_cast<unsigned char *
>(&
L)[-pass_stride + i] = 0;
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint y
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint GLsizei GLsizei GLenum type _GL_VOID_RET _GL_VOID GLsizei GLenum GLenum const void *pixels _GL_VOID_RET _GL_VOID const void *pointer _GL_VOID_RET _GL_VOID GLdouble v _GL_VOID_RET _GL_VOID GLfloat v _GL_VOID_RET _GL_VOID GLint GLint i2 _GL_VOID_RET _GL_VOID GLint j _GL_VOID_RET _GL_VOID GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble GLdouble GLdouble zFar _GL_VOID_RET _GL_UINT GLdouble *equation _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLenum GLfloat *v _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLfloat *values _GL_VOID_RET _GL_VOID GLushort *values _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLenum GLdouble *params _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLsizei stride
ccl_device_inline float3 ray_offset(float3 P, float3 Ng)
ccl_device_intersect bool scene_intersect(KernelGlobals *kg, const Ray *ray, const uint visibility, Intersection *isect)
CCL_NAMESPACE_BEGIN ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, const ShaderData *sd)
ccl_device_inline void path_radiance_init(KernelGlobals *kg, PathRadiance *L)
ccl_device_inline void path_radiance_accum_ao(KernelGlobals *kg, PathRadiance *L, ccl_addr_space PathState *state, float3 throughput, float3 alpha, float3 bsdf, float3 ao)
ccl_device_inline void path_radiance_accum_background(KernelGlobals *kg, PathRadiance *L, ccl_addr_space PathState *state, float3 throughput, float3 value)
ccl_device_inline void path_radiance_accum_emission(KernelGlobals *kg, PathRadiance *L, ccl_addr_space PathState *state, float3 throughput, float3 value)
ccl_device_inline void path_radiance_accum_total_ao(PathRadiance *L, ccl_addr_space PathState *state, float3 throughput, float3 bsdf)
#define ccl_device_forceinline
#define ccl_optional_struct_init
#define ccl_device_inline
#define ccl_device_noinline
#define CCL_NAMESPACE_END
ccl_device differential3 differential3_zero()
ccl_device_noinline_cpu float3 indirect_primitive_emission(KernelGlobals *kg, ShaderData *sd, float t, int path_flag, float bsdf_pdf)
ccl_device_noinline_cpu void indirect_lamp_emission(KernelGlobals *kg, ShaderData *emission_sd, ccl_addr_space PathState *state, PathRadiance *L, Ray *ray, float3 throughput)
ccl_device_noinline_cpu float3 indirect_background(KernelGlobals *kg, ShaderData *emission_sd, ccl_addr_space PathState *state, ccl_global float *buffer, ccl_addr_space Ray *ray)
ccl_device_inline void sample_cos_hemisphere(const float3 N, float randu, float randv, float3 *omega_in, float *pdf)
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 void kernel_path_trace(KernelGlobals *kg, ccl_global float *buffer, int sample, int x, int y, int offset, int stride)
ccl_device_forceinline bool kernel_path_shader_apply(KernelGlobals *kg, ShaderData *sd, ccl_addr_space PathState *state, ccl_addr_space Ray *ray, float3 throughput, ShaderData *emission_sd, PathRadiance *L, ccl_global float *buffer)
ccl_device_forceinline void kernel_path_background(KernelGlobals *kg, ccl_addr_space PathState *state, ccl_addr_space Ray *ray, float3 throughput, ShaderData *sd, ccl_global float *buffer, PathRadiance *L)
ccl_device_forceinline void kernel_path_integrate(KernelGlobals *kg, PathState *state, float3 throughput, Ray *ray, PathRadiance *L, ccl_global float *buffer, ShaderData *emission_sd)
CCL_NAMESPACE_BEGIN ccl_device_forceinline bool kernel_path_scene_intersect(KernelGlobals *kg, ccl_addr_space PathState *state, Ray *ray, Intersection *isect, PathRadiance *L)
ccl_device_forceinline void kernel_path_lamp_emission(KernelGlobals *kg, ccl_addr_space PathState *state, Ray *ray, float3 throughput, ccl_addr_space Intersection *isect, ShaderData *emission_sd, PathRadiance *L)
ccl_device_noinline void kernel_path_ao(KernelGlobals *kg, ShaderData *sd, ShaderData *emission_sd, PathRadiance *L, ccl_addr_space PathState *state, float3 throughput, float3 ao_alpha)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_path_trace_setup(KernelGlobals *kg, int sample, int x, int y, uint *rng_hash, ccl_addr_space Ray *ray)
ccl_device_inline bool path_state_ao_bounce(KernelGlobals *kg, ccl_addr_space PathState *state)
ccl_device_inline float path_state_continuation_probability(KernelGlobals *kg, ccl_addr_space PathState *state, const float3 throughput)
CCL_NAMESPACE_BEGIN ccl_device_inline void path_state_init(KernelGlobals *kg, ShaderData *stack_sd, ccl_addr_space PathState *state, uint rng_hash, int sample, ccl_addr_space Ray *ray)
ccl_device_inline uint path_state_ray_visibility(KernelGlobals *kg, ccl_addr_space PathState *state)
ccl_device bool kernel_path_surface_bounce(KernelGlobals *kg, ShaderData *sd, ccl_addr_space float3 *throughput, ccl_addr_space PathState *state, PathRadianceState *L_state, ccl_addr_space Ray *ray)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_path_surface_connect_light(KernelGlobals *kg, ShaderData *sd, ShaderData *emission_sd, float3 throughput, ccl_addr_space PathState *state, PathRadiance *L)
#define PROFILING_INIT(kg, event)
ccl_device_inline void path_state_rng_2D(KernelGlobals *kg, const ccl_addr_space PathState *state, int dimension, float *fx, float *fy)
ccl_device_inline float path_state_rng_1D(KernelGlobals *kg, const ccl_addr_space PathState *state, int dimension)
CCL_NAMESPACE_BEGIN ccl_device_noinline void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray)
ccl_device void shader_bsdf_blur(KernelGlobals *kg, ShaderData *sd, float roughness)
ccl_device float3 shader_bsdf_alpha(KernelGlobals *kg, ShaderData *sd)
ccl_device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd, ccl_addr_space PathState *state, ccl_global float *buffer, int path_flag)
ccl_device float3 shader_bsdf_ao(KernelGlobals *kg, ShaderData *sd, float ao_factor, float3 *N_)
ccl_device float3 shader_holdout_apply(KernelGlobals *kg, ShaderData *sd)
ccl_device_inline void shader_prepare_closures(ShaderData *sd, ccl_addr_space PathState *state)
ccl_device_inline bool shadow_blocked(KernelGlobals *kg, ShaderData *sd, ShaderData *shadow_sd, ccl_addr_space PathState *state, Ray *ray, float3 *shadow)
__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 AS_SHADER_DATA(shader_data_tiny_storage)
@ PATH_RAY_STORE_SHADOW_INFO
@ PATH_RAY_SHADOW_CATCHER
@ PATH_RAY_TRANSPARENT_BACKGROUND
@ SD_OBJECT_SHADOW_CATCHER
static void sample(SocketReader *reader, int x, int y, float color[4])
__forceinline bool all(const avxb &b)
ccl_device_inline float dot(const float2 &a, const float2 &b)
ccl_device_inline float average(const float2 &a)
ccl_device_inline float3 one_float3()
ccl_device_inline float3 zero_float3()
ccl_device_inline bool isequal_float3(const float3 a, const float3 b)
@ PROFILING_PATH_INTEGRATE
@ PROFILING_SCENE_INTERSECT
@ PROFILING_INDIRECT_EMISSION