19 #ifdef __BRANCHED_PATH__
28 int num_samples =
kernel_data.integrator.ao_samples;
29 float num_samples_inv = 1.0f / num_samples;
35 for (
int j = 0; j < num_samples; j++) {
45 if (
dot(sd->Ng, ao_D) > 0.0f && ao_pdf != 0.0f) {
52 light_ray.
time = sd->time;
53 light_ray.
dP = sd->dP;
58 kg,
L,
state, throughput * num_samples_inv, ao_alpha, ao_bsdf, ao_shadow);
67 # ifndef __SPLIT_KERNEL__
83 kernel_volume_clean_stack(
kg,
state->volume_stack);
91 Ray volume_ray = *ray;
92 volume_ray.
t = (hit) ? isect->t : FLT_MAX;
94 float step_size = volume_stack_step_size(
kg,
state->volume_stack);
96 # ifdef __VOLUME_DECOUPLED__
100 VolumeSegment volume_segment;
102 shader_setup_from_volume(
kg, sd, &volume_ray);
103 kernel_volume_decoupled_record(
kg,
state, &volume_ray, sd, &volume_segment, step_size);
106 if (volume_segment.closure_flag &
SD_SCATTER) {
107 volume_segment.sampling_method = volume_stack_sampling_method(
kg,
state->volume_stack);
111 kernel_branched_path_volume_connect_light(
112 kg, sd, emission_sd, *throughput,
state,
L,
all, &volume_ray, &volume_segment);
115 int num_samples =
kernel_data.integrator.volume_samples;
116 float num_samples_inv = 1.0f / num_samples;
118 for (
int j = 0; j < num_samples; j++) {
133 kg, &ps, &pray, sd, &tp, rphase, rscatter, &volume_segment,
NULL,
false);
136 kernel_path_volume_bounce(
kg, sd, &tp, &ps, &
L->state, &pray)) {
137 kernel_path_indirect(
kg, indirect_sd, emission_sd, &pray, tp * num_samples_inv, &ps,
L);
150 *throughput *= volume_segment.accum_transmittance;
153 kernel_volume_decoupled_free(
kg, &volume_segment);
159 int num_samples =
kernel_data.integrator.volume_samples;
160 float num_samples_inv = 1.0f / num_samples;
165 for (
int j = 0; j < num_samples; j++) {
168 float3 tp = (*throughput) * num_samples_inv;
174 kg, &ps, sd, &volume_ray,
L, &tp, step_size);
176 # ifdef __VOLUME_SCATTER__
180 kernel_path_volume_connect_light(
kg, sd, emission_sd, tp,
state,
L);
182 if (kernel_path_volume_bounce(
kg, sd, &tp, &ps, &
L->state, &pray)) {
183 kernel_path_indirect(
kg, indirect_sd, emission_sd, &pray, tp, &ps,
L);
195 kernel_volume_shadow(
kg, emission_sd,
state, &volume_ray, throughput);
206 float num_samples_adjust,
210 float sum_sample_weight = 0.0f;
211 # ifdef __DENOISING_FEATURES__
212 if (
state->denoising_feature_weight > 0.0f) {
213 for (
int i = 0; i < sd->num_closure; i++) {
221 sum_sample_weight += sc->sample_weight;
225 sum_sample_weight = 1.0f;
229 for (
int i = 0; i < sd->num_closure; i++) {
240 num_samples =
kernel_data.integrator.diffuse_samples;
244 num_samples =
kernel_data.integrator.glossy_samples;
246 num_samples =
kernel_data.integrator.transmission_samples;
248 num_samples =
ceil_to_int(num_samples_adjust * num_samples);
250 float num_samples_inv = num_samples_adjust / num_samples;
252 for (
int j = 0; j < num_samples; j++) {
256 # ifdef __SHADOW_TRICKS__
257 float shadow_transparency =
L->shadow_transparency;
262 if (!kernel_branched_path_surface_bounce(
263 kg, sd, sc, j, num_samples, &tp, &ps, &
L->state, &bsdf_ray, sum_sample_weight)) {
269 kernel_path_indirect(
kg, indirect_sd, emission_sd, &bsdf_ray, tp * num_samples_inv, &ps,
L);
276 # ifdef __SHADOW_TRICKS__
277 L->shadow_transparency = shadow_transparency;
283 # ifdef __SUBSURFACE__
284 ccl_device void kernel_branched_path_subsurface_scatter(KernelGlobals *
kg,
293 for (
int i = 0; i < sd->num_closure; i++) {
301 int num_samples =
kernel_data.integrator.subsurface_samples * 3;
302 float num_samples_inv = 1.0f / num_samples;
307 for (
int j = 0; j < num_samples; j++) {
310 hit_state.
rng_hash = bssrdf_rng_hash;
313 float bssrdf_u, bssrdf_v;
316 kg, &ss_isect, sd, &hit_state, sc, &lcg_state, bssrdf_u, bssrdf_v,
true);
321 Ray volume_ray = *ray;
322 bool need_update_volume_stack =
kernel_data.integrator.use_volumes &&
327 for (
int hit = 0; hit < num_hits; hit++) {
331 float bssrdf_roughness =
bssrdf->roughness;
333 kg, &ss_isect, hit, &bssrdf_sd, &hit_state, bssrdf_type, bssrdf_roughness);
336 if (need_update_volume_stack) {
345 kernel_volume_stack_update_for_subsurface(
355 kernel_branched_path_surface_connect_light(
356 kg, &bssrdf_sd, emission_sd, &hit_state, throughput, num_samples_inv,
L,
all);
361 kernel_branched_path_surface_indirect_light(
362 kg, &bssrdf_sd, indirect_sd, emission_sd, throughput, num_samples_inv, &hit_state,
L);
369 ccl_device void kernel_branched_path_integrate(KernelGlobals *
kg,
402 kernel_branched_path_volume(
403 kg, &sd, &
state, &ray, &throughput, &isect, hit, &indirect_sd, emission_sd,
L);
421 shader_merge_closures(&sd);
435 if (probability == 0.0f) {
438 else if (probability != 1.0f) {
441 if (terminate >= probability)
444 throughput /= probability;
448 # ifdef __DENOISING_FEATURES__
449 kernel_update_denoising_features(
kg, &sd, &
state,
L);
454 if (
kernel_data.integrator.use_ambient_occlusion) {
455 kernel_branched_path_ao(
kg, &sd, emission_sd,
L, &
state, throughput);
459 # ifdef __SUBSURFACE__
462 kernel_branched_path_subsurface_scatter(
463 kg, &sd, &indirect_sd, emission_sd,
L, &
state, &ray, throughput);
474 kernel_branched_path_surface_connect_light(
475 kg, &sd, emission_sd, &hit_state, throughput, 1.0f,
L,
all);
480 kernel_branched_path_surface_indirect_light(
481 kg, &sd, &indirect_sd, emission_sd, throughput, 1.0f, &hit_state,
L);
495 if (!path_state_volume_next(
kg, &
state)) {
502 ray.
t -= sd.ray_length;
504 # ifdef __RAY_DIFFERENTIALS__
506 ray.
dD.
dx = -sd.dI.dx;
507 ray.
dD.
dy = -sd.dI.dy;
512 kernel_volume_stack_enter_exit(
kg, &sd,
state.volume_stack);
521 int index = offset +
x +
y *
stride;
524 buffer += index * pass_stride;
529 if ((*aux).w > 0.0f) {
_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_inline void path_radiance_sum_indirect(PathRadiance *L)
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_reset_indirect(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_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_device_noinline_cpu
#define ccl_device_inline
#define CCL_NAMESPACE_END
ccl_device differential3 differential3_zero()
ccl_device_inline uint cmj_hash(uint i, uint p)
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_result(KernelGlobals *kg, ccl_global float *buffer, int sample, PathRadiance *L)
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_NAMESPACE_BEGIN ccl_device_forceinline bool kernel_path_scene_intersect(KernelGlobals *kg, ccl_addr_space PathState *state, Ray *ray, Intersection *isect, PathRadiance *L)
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 void path_state_next(KernelGlobals *kg, ccl_addr_space PathState *state, int label)
ccl_device_inline float path_state_continuation_probability(KernelGlobals *kg, ccl_addr_space PathState *state, const float3 throughput)
ccl_device_inline void path_state_branch(ccl_addr_space PathState *state, int branch, int num_branches)
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 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_device_inline uint lcg_state_init(PathState *state, uint scramble)
ccl_device_inline void path_branched_rng_2D(KernelGlobals *kg, uint rng_hash, const ccl_addr_space PathState *state, int branch, int num_branches, int dimension, float *fx, float *fy)
CCL_NAMESPACE_BEGIN ccl_device_noinline void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray)
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_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
ccl_device_inline int subsurface_scatter_multi_intersect(KernelGlobals *kg, LocalIntersection *ss_isect, ShaderData *sd, ccl_addr_space PathState *state, const ShaderClosure *sc, uint *lcg_state, float bssrdf_u, float bssrdf_v, bool all)
ccl_device_noinline void subsurface_scatter_multi_setup(KernelGlobals *kg, LocalIntersection *ss_isect, int hit, ShaderData *sd, ccl_addr_space PathState *state, ClosureType type, float roughness)
#define AS_SHADER_DATA(shader_data_tiny_storage)
@ PATH_RAY_SHADOW_CATCHER
@ SD_OBJECT_INTERSECTS_VOLUME
#define VOLUME_STACK_SIZE
static void sample(SocketReader *reader, int x, int y, float color[4])
closure color bssrdf(string method, normal N, vector radius, color albedo) BUILTIN
VolumeStack volume_stack[VOLUME_STACK_SIZE]
#define CLOSURE_IS_BSDF_TRANSPARENT(type)
#define CLOSURE_IS_BSDF(type)
#define CLOSURE_IS_BSDF_GLOSSY(type)
#define CLOSURE_IS_BSDF_BSSRDF(type)
#define CLOSURE_IS_BSDF_DIFFUSE(type)
#define CLOSURE_IS_BSSRDF(type)
__forceinline bool all(const avxb &b)
ccl_device_inline int ceil_to_int(float f)
ccl_device_inline float dot(const float2 &a, const float2 &b)
ccl_device_inline float2 normalize_len(const float2 &a, float *t)
ccl_device_inline bool is_zero(const float2 &a)
ccl_device_inline float3 one_float3()