Blender  V2.93
kernel_path_branched.h
Go to the documentation of this file.
1 /*
2  * Copyright 2011-2013 Blender Foundation
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
18 
19 #ifdef __BRANCHED_PATH__
20 
21 ccl_device_inline void kernel_branched_path_ao(KernelGlobals *kg,
22  ShaderData *sd,
23  ShaderData *emission_sd,
24  PathRadiance *L,
26  float3 throughput)
27 {
28  int num_samples = kernel_data.integrator.ao_samples;
29  float num_samples_inv = 1.0f / num_samples;
30  float ao_factor = kernel_data.background.ao_factor;
31  float3 ao_N;
32  float3 ao_bsdf = shader_bsdf_ao(kg, sd, ao_factor, &ao_N);
33  float3 ao_alpha = shader_bsdf_alpha(kg, sd);
34 
35  for (int j = 0; j < num_samples; j++) {
36  float bsdf_u, bsdf_v;
38  kg, state->rng_hash, state, j, num_samples, PRNG_BSDF_U, &bsdf_u, &bsdf_v);
39 
40  float3 ao_D;
41  float ao_pdf;
42 
43  sample_cos_hemisphere(ao_N, bsdf_u, bsdf_v, &ao_D, &ao_pdf);
44 
45  if (dot(sd->Ng, ao_D) > 0.0f && ao_pdf != 0.0f) {
46  Ray light_ray;
47  float3 ao_shadow;
48 
49  light_ray.P = ray_offset(sd->P, sd->Ng);
50  light_ray.D = ao_D;
51  light_ray.t = kernel_data.background.ao_distance;
52  light_ray.time = sd->time;
53  light_ray.dP = sd->dP;
54  light_ray.dD = differential3_zero();
55 
56  if (!shadow_blocked(kg, sd, emission_sd, state, &light_ray, &ao_shadow)) {
58  kg, L, state, throughput * num_samples_inv, ao_alpha, ao_bsdf, ao_shadow);
59  }
60  else {
61  path_radiance_accum_total_ao(L, state, throughput * num_samples_inv, ao_bsdf);
62  }
63  }
64  }
65 }
66 
67 # ifndef __SPLIT_KERNEL__
68 
69 # ifdef __VOLUME__
70 ccl_device_forceinline void kernel_branched_path_volume(KernelGlobals *kg,
71  ShaderData *sd,
73  Ray *ray,
74  float3 *throughput,
76  bool hit,
77  ShaderData *indirect_sd,
78  ShaderData *emission_sd,
79  PathRadiance *L)
80 {
81  /* Sanitize volume stack. */
82  if (!hit) {
83  kernel_volume_clean_stack(kg, state->volume_stack);
84  }
85 
86  if (state->volume_stack[0].shader == SHADER_NONE) {
87  return;
88  }
89 
90  /* volume attenuation, emission, scatter */
91  Ray volume_ray = *ray;
92  volume_ray.t = (hit) ? isect->t : FLT_MAX;
93 
94  float step_size = volume_stack_step_size(kg, state->volume_stack);
95 
96 # ifdef __VOLUME_DECOUPLED__
97  /* decoupled ray marching only supported on CPU */
98  if (kernel_data.integrator.volume_decoupled) {
99  /* cache steps along volume for repeated sampling */
100  VolumeSegment volume_segment;
101 
102  shader_setup_from_volume(kg, sd, &volume_ray);
103  kernel_volume_decoupled_record(kg, state, &volume_ray, sd, &volume_segment, step_size);
104 
105  /* direct light sampling */
106  if (volume_segment.closure_flag & SD_SCATTER) {
107  volume_segment.sampling_method = volume_stack_sampling_method(kg, state->volume_stack);
108 
109  int all = kernel_data.integrator.sample_all_lights_direct;
110 
111  kernel_branched_path_volume_connect_light(
112  kg, sd, emission_sd, *throughput, state, L, all, &volume_ray, &volume_segment);
113 
114  /* indirect light sampling */
115  int num_samples = kernel_data.integrator.volume_samples;
116  float num_samples_inv = 1.0f / num_samples;
117 
118  for (int j = 0; j < num_samples; j++) {
119  PathState ps = *state;
120  Ray pray = *ray;
121  float3 tp = *throughput;
122 
123  /* branch RNG state */
124  path_state_branch(&ps, j, num_samples);
125 
126  /* scatter sample. if we use distance sampling and take just one
127  * sample for direct and indirect light, we could share this
128  * computation, but makes code a bit complex */
129  float rphase = path_state_rng_1D(kg, &ps, PRNG_PHASE_CHANNEL);
130  float rscatter = path_state_rng_1D(kg, &ps, PRNG_SCATTER_DISTANCE);
131 
132  VolumeIntegrateResult result = kernel_volume_decoupled_scatter(
133  kg, &ps, &pray, sd, &tp, rphase, rscatter, &volume_segment, NULL, false);
134 
135  if (result == VOLUME_PATH_SCATTERED &&
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);
138 
139  /* for render passes, sum and reset indirect light pass variables
140  * for the next samples */
143  }
144  }
145  }
146 
147  /* emission and transmittance */
148  if (volume_segment.closure_flag & SD_EMISSION)
149  path_radiance_accum_emission(kg, L, state, *throughput, volume_segment.accum_emission);
150  *throughput *= volume_segment.accum_transmittance;
151 
152  /* free cached steps */
153  kernel_volume_decoupled_free(kg, &volume_segment);
154  }
155  else
156 # endif /* __VOLUME_DECOUPLED__ */
157  {
158  /* GPU: no decoupled ray marching, scatter probabilistically. */
159  int num_samples = kernel_data.integrator.volume_samples;
160  float num_samples_inv = 1.0f / num_samples;
161 
162  /* todo: we should cache the shader evaluations from stepping
163  * through the volume, for now we redo them multiple times */
164 
165  for (int j = 0; j < num_samples; j++) {
166  PathState ps = *state;
167  Ray pray = *ray;
168  float3 tp = (*throughput) * num_samples_inv;
169 
170  /* branch RNG state */
171  path_state_branch(&ps, j, num_samples);
172 
173  VolumeIntegrateResult result = kernel_volume_integrate(
174  kg, &ps, sd, &volume_ray, L, &tp, step_size);
175 
176 # ifdef __VOLUME_SCATTER__
177  if (result == VOLUME_PATH_SCATTERED) {
178  /* todo: support equiangular, MIS and all light sampling.
179  * alternatively get decoupled ray marching working on the GPU */
180  kernel_path_volume_connect_light(kg, sd, emission_sd, tp, state, L);
181 
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);
184 
185  /* for render passes, sum and reset indirect light pass variables
186  * for the next samples */
189  }
190  }
191 # endif /* __VOLUME_SCATTER__ */
192  }
193 
194  /* todo: avoid this calculation using decoupled ray marching */
195  kernel_volume_shadow(kg, emission_sd, state, &volume_ray, throughput);
196  }
197 }
198 # endif /* __VOLUME__ */
199 
200 /* bounce off surface and integrate indirect light */
201 ccl_device_noinline_cpu void kernel_branched_path_surface_indirect_light(KernelGlobals *kg,
202  ShaderData *sd,
203  ShaderData *indirect_sd,
204  ShaderData *emission_sd,
205  float3 throughput,
206  float num_samples_adjust,
207  PathState *state,
208  PathRadiance *L)
209 {
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++) {
214  const ShaderClosure *sc = &sd->closure[i];
215 
216  /* transparency is not handled here, but in outer loop */
217  if (!CLOSURE_IS_BSDF(sc->type) || CLOSURE_IS_BSDF_TRANSPARENT(sc->type)) {
218  continue;
219  }
220 
221  sum_sample_weight += sc->sample_weight;
222  }
223  }
224  else {
225  sum_sample_weight = 1.0f;
226  }
227 # endif /* __DENOISING_FEATURES__ */
228 
229  for (int i = 0; i < sd->num_closure; i++) {
230  const ShaderClosure *sc = &sd->closure[i];
231 
232  /* transparency is not handled here, but in outer loop */
233  if (!CLOSURE_IS_BSDF(sc->type) || CLOSURE_IS_BSDF_TRANSPARENT(sc->type)) {
234  continue;
235  }
236 
237  int num_samples;
238 
239  if (CLOSURE_IS_BSDF_DIFFUSE(sc->type))
240  num_samples = kernel_data.integrator.diffuse_samples;
241  else if (CLOSURE_IS_BSDF_BSSRDF(sc->type))
242  num_samples = 1;
243  else if (CLOSURE_IS_BSDF_GLOSSY(sc->type))
244  num_samples = kernel_data.integrator.glossy_samples;
245  else
246  num_samples = kernel_data.integrator.transmission_samples;
247 
248  num_samples = ceil_to_int(num_samples_adjust * num_samples);
249 
250  float num_samples_inv = num_samples_adjust / num_samples;
251 
252  for (int j = 0; j < num_samples; j++) {
253  PathState ps = *state;
254  float3 tp = throughput;
255  Ray bsdf_ray;
256 # ifdef __SHADOW_TRICKS__
257  float shadow_transparency = L->shadow_transparency;
258 # endif
259 
260  ps.rng_hash = cmj_hash(state->rng_hash, i);
261 
262  if (!kernel_branched_path_surface_bounce(
263  kg, sd, sc, j, num_samples, &tp, &ps, &L->state, &bsdf_ray, sum_sample_weight)) {
264  continue;
265  }
266 
267  ps.rng_hash = state->rng_hash;
268 
269  kernel_path_indirect(kg, indirect_sd, emission_sd, &bsdf_ray, tp * num_samples_inv, &ps, L);
270 
271  /* for render passes, sum and reset indirect light pass variables
272  * for the next samples */
275 
276 # ifdef __SHADOW_TRICKS__
277  L->shadow_transparency = shadow_transparency;
278 # endif
279  }
280  }
281 }
282 
283 # ifdef __SUBSURFACE__
284 ccl_device void kernel_branched_path_subsurface_scatter(KernelGlobals *kg,
285  ShaderData *sd,
286  ShaderData *indirect_sd,
287  ShaderData *emission_sd,
288  PathRadiance *L,
289  PathState *state,
290  Ray *ray,
291  float3 throughput)
292 {
293  for (int i = 0; i < sd->num_closure; i++) {
294  ShaderClosure *sc = &sd->closure[i];
295 
296  if (!CLOSURE_IS_BSSRDF(sc->type))
297  continue;
298 
299  /* set up random number generator */
300  uint lcg_state = lcg_state_init(state, 0x68bc21eb);
301  int num_samples = kernel_data.integrator.subsurface_samples * 3;
302  float num_samples_inv = 1.0f / num_samples;
303  uint bssrdf_rng_hash = cmj_hash(state->rng_hash, i);
304 
305  /* do subsurface scatter step with copy of shader data, this will
306  * replace the BSSRDF with a diffuse BSDF closure */
307  for (int j = 0; j < num_samples; j++) {
308  PathState hit_state = *state;
309  path_state_branch(&hit_state, j, num_samples);
310  hit_state.rng_hash = bssrdf_rng_hash;
311 
312  LocalIntersection ss_isect;
313  float bssrdf_u, bssrdf_v;
314  path_state_rng_2D(kg, &hit_state, PRNG_BSDF_U, &bssrdf_u, &bssrdf_v);
315  int num_hits = subsurface_scatter_multi_intersect(
316  kg, &ss_isect, sd, &hit_state, sc, &lcg_state, bssrdf_u, bssrdf_v, true);
317 
318  hit_state.rng_offset += PRNG_BOUNCE_NUM;
319 
320 # ifdef __VOLUME__
321  Ray volume_ray = *ray;
322  bool need_update_volume_stack = kernel_data.integrator.use_volumes &&
323  sd->object_flag & SD_OBJECT_INTERSECTS_VOLUME;
324 # endif /* __VOLUME__ */
325 
326  /* compute lighting with the BSDF closure */
327  for (int hit = 0; hit < num_hits; hit++) {
328  ShaderData bssrdf_sd = *sd;
329  Bssrdf *bssrdf = (Bssrdf *)sc;
330  ClosureType bssrdf_type = sc->type;
331  float bssrdf_roughness = bssrdf->roughness;
333  kg, &ss_isect, hit, &bssrdf_sd, &hit_state, bssrdf_type, bssrdf_roughness);
334 
335 # ifdef __VOLUME__
336  if (need_update_volume_stack) {
337  /* Setup ray from previous surface point to the new one. */
338  float3 P = ray_offset(bssrdf_sd.P, -bssrdf_sd.Ng);
339  volume_ray.D = normalize_len(P - volume_ray.P, &volume_ray.t);
340 
341  for (int k = 0; k < VOLUME_STACK_SIZE; k++) {
342  hit_state.volume_stack[k] = state->volume_stack[k];
343  }
344 
345  kernel_volume_stack_update_for_subsurface(
346  kg, emission_sd, &volume_ray, hit_state.volume_stack);
347  }
348 # endif /* __VOLUME__ */
349 
350 # ifdef __EMISSION__
351  /* direct light */
352  if (kernel_data.integrator.use_direct_light) {
353  int all = (kernel_data.integrator.sample_all_lights_direct) ||
354  (hit_state.flag & PATH_RAY_SHADOW_CATCHER);
355  kernel_branched_path_surface_connect_light(
356  kg, &bssrdf_sd, emission_sd, &hit_state, throughput, num_samples_inv, L, all);
357  }
358 # endif /* __EMISSION__ */
359 
360  /* indirect light */
361  kernel_branched_path_surface_indirect_light(
362  kg, &bssrdf_sd, indirect_sd, emission_sd, throughput, num_samples_inv, &hit_state, L);
363  }
364  }
365  }
366 }
367 # endif /* __SUBSURFACE__ */
368 
369 ccl_device void kernel_branched_path_integrate(KernelGlobals *kg,
370  uint rng_hash,
371  int sample,
372  Ray ray,
373  ccl_global float *buffer,
374  PathRadiance *L)
375 {
376  /* initialize */
377  float3 throughput = one_float3();
378 
380 
381  /* shader data memory used for both volumes and surfaces, saves stack space */
382  ShaderData sd;
383  /* shader data used by emission, shadows, volume stacks, indirect path */
384  ShaderDataTinyStorage emission_sd_storage;
385  ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
386  ShaderData indirect_sd;
387 
389  path_state_init(kg, emission_sd, &state, rng_hash, sample, &ray);
390 
391  /* Main Loop
392  * Here we only handle transparency intersections from the camera ray.
393  * Indirect bounces are handled in kernel_branched_path_surface_indirect_light().
394  */
395  for (;;) {
396  /* Find intersection with objects in scene. */
397  Intersection isect;
398  bool hit = kernel_path_scene_intersect(kg, &state, &ray, &isect, L);
399 
400 # ifdef __VOLUME__
401  /* Volume integration. */
402  kernel_branched_path_volume(
403  kg, &sd, &state, &ray, &throughput, &isect, hit, &indirect_sd, emission_sd, L);
404 # endif /* __VOLUME__ */
405 
406  /* Shade background. */
407  if (!hit) {
408  kernel_path_background(kg, &state, &ray, throughput, &sd, buffer, L);
409  break;
410  }
411 
412  /* Setup and evaluate shader. */
413  shader_setup_from_ray(kg, &sd, &isect, &ray);
414 
415  /* Skip most work for volume bounding surface. */
416 # ifdef __VOLUME__
417  if (!(sd.flag & SD_HAS_ONLY_VOLUME)) {
418 # endif
419 
420  shader_eval_surface(kg, &sd, &state, buffer, state.flag);
421  shader_merge_closures(&sd);
422 
423  /* Apply shadow catcher, holdout, emission. */
424  if (!kernel_path_shader_apply(kg, &sd, &state, &ray, throughput, emission_sd, L, buffer)) {
425  break;
426  }
427 
428  /* transparency termination */
429  if (state.flag & PATH_RAY_TRANSPARENT) {
430  /* path termination. this is a strange place to put the termination, it's
431  * mainly due to the mixed in MIS that we use. gives too many unneeded
432  * shader evaluations, only need emission if we are going to terminate */
433  float probability = path_state_continuation_probability(kg, &state, throughput);
434 
435  if (probability == 0.0f) {
436  break;
437  }
438  else if (probability != 1.0f) {
439  float terminate = path_state_rng_1D(kg, &state, PRNG_TERMINATE);
440 
441  if (terminate >= probability)
442  break;
443 
444  throughput /= probability;
445  }
446  }
447 
448 # ifdef __DENOISING_FEATURES__
449  kernel_update_denoising_features(kg, &sd, &state, L);
450 # endif
451 
452 # ifdef __AO__
453  /* ambient occlusion */
454  if (kernel_data.integrator.use_ambient_occlusion) {
455  kernel_branched_path_ao(kg, &sd, emission_sd, L, &state, throughput);
456  }
457 # endif /* __AO__ */
458 
459 # ifdef __SUBSURFACE__
460  /* bssrdf scatter to a different location on the same object */
461  if (sd.flag & SD_BSSRDF) {
462  kernel_branched_path_subsurface_scatter(
463  kg, &sd, &indirect_sd, emission_sd, L, &state, &ray, throughput);
464  }
465 # endif /* __SUBSURFACE__ */
466 
467  PathState hit_state = state;
468 
469 # ifdef __EMISSION__
470  /* direct light */
471  if (kernel_data.integrator.use_direct_light) {
472  int all = (kernel_data.integrator.sample_all_lights_direct) ||
474  kernel_branched_path_surface_connect_light(
475  kg, &sd, emission_sd, &hit_state, throughput, 1.0f, L, all);
476  }
477 # endif /* __EMISSION__ */
478 
479  /* indirect light */
480  kernel_branched_path_surface_indirect_light(
481  kg, &sd, &indirect_sd, emission_sd, throughput, 1.0f, &hit_state, L);
482 
483  /* continue in case of transparency */
484  throughput *= shader_bsdf_transparency(kg, &sd);
485 
486  if (is_zero(throughput))
487  break;
488 
489  /* Update Path State */
491 
492 # ifdef __VOLUME__
493  }
494  else {
495  if (!path_state_volume_next(kg, &state)) {
496  break;
497  }
498  }
499 # endif
500 
501  ray.P = ray_offset(sd.P, -sd.Ng);
502  ray.t -= sd.ray_length; /* clipping works through transparent */
503 
504 # ifdef __RAY_DIFFERENTIALS__
505  ray.dP = sd.dP;
506  ray.dD.dx = -sd.dI.dx;
507  ray.dD.dy = -sd.dI.dy;
508 # endif /* __RAY_DIFFERENTIALS__ */
509 
510 # ifdef __VOLUME__
511  /* enter/exit volume */
512  kernel_volume_stack_enter_exit(kg, &sd, state.volume_stack);
513 # endif /* __VOLUME__ */
514  }
515 }
516 
517 ccl_device void kernel_branched_path_trace(
518  KernelGlobals *kg, ccl_global float *buffer, int sample, int x, int y, int offset, int stride)
519 {
520  /* buffer offset */
521  int index = offset + x + y * stride;
522  int pass_stride = kernel_data.film.pass_stride;
523 
524  buffer += index * pass_stride;
525 
526  if (kernel_data.film.pass_adaptive_aux_buffer) {
527  ccl_global float4 *aux = (ccl_global float4 *)(buffer +
528  kernel_data.film.pass_adaptive_aux_buffer);
529  if ((*aux).w > 0.0f) {
530  return;
531  }
532  }
533 
534  /* initialize random numbers and ray */
535  uint rng_hash;
536  Ray ray;
537 
538  kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &ray);
539 
540  /* integrate */
541  PathRadiance L;
542 
543  if (ray.t != 0.0f) {
544  kernel_branched_path_integrate(kg, rng_hash, sample, ray, buffer, &L);
546  }
547 }
548 
549 # endif /* __SPLIT_KERNEL__ */
550 
551 #endif /* __BRANCHED_PATH__ */
552 
unsigned int uint
Definition: BLI_sys_types.h:83
_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 kernel_data
#define ccl_addr_space
#define ccl_device_forceinline
#define ccl_device
#define ccl_device_noinline_cpu
#define ccl_device_inline
#define ccl_global
#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)
Definition: kernel_path.h:254
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)
Definition: kernel_path.h:115
CCL_NAMESPACE_BEGIN ccl_device_forceinline bool kernel_path_scene_intersect(KernelGlobals *kg, ccl_addr_space PathState *state, Ray *ray, Intersection *isect, PathRadiance *L)
Definition: kernel_path.h:57
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)
Definition: kernel_shader.h:59
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)
@ SD_BSSRDF
Definition: kernel_types.h:851
@ SD_HAS_ONLY_VOLUME
Definition: kernel_types.h:875
@ SD_SCATTER
Definition: kernel_types.h:857
@ SD_EMISSION
Definition: kernel_types.h:845
#define AS_SHADER_DATA(shader_data_tiny_storage)
#define SHADER_NONE
Definition: kernel_types.h:58
@ PRNG_BOUNCE_NUM
Definition: kernel_types.h:248
@ PRNG_BSDF_U
Definition: kernel_types.h:240
@ PRNG_SCATTER_DISTANCE
Definition: kernel_types.h:247
@ PRNG_TERMINATE
Definition: kernel_types.h:245
@ PRNG_PHASE_CHANNEL
Definition: kernel_types.h:246
@ PATH_RAY_SHADOW_CATCHER
Definition: kernel_types.h:306
@ PATH_RAY_TRANSPARENT
Definition: kernel_types.h:272
ShaderDataTinyStorage
ShaderData
@ SD_OBJECT_INTERSECTS_VOLUME
Definition: kernel_types.h:914
@ LABEL_TRANSPARENT
Definition: kernel_types.h:333
#define VOLUME_STACK_SIZE
Definition: kernel_types.h:64
ShaderClosure
Definition: kernel_types.h:831
VolumeIntegrateResult
Definition: kernel_volume.h:27
@ VOLUME_PATH_SCATTERED
Definition: kernel_volume.h:28
static float P(float k)
Definition: math_interp.c:41
static ulong state[N]
#define L
static void sample(SocketReader *reader, int x, int y, float color[4])
closure color bssrdf(string method, normal N, vector radius, color albedo) BUILTIN
Definition: bssrdf.h:22
VolumeStack volume_stack[VOLUME_STACK_SIZE]
float t
Definition: kernel_types.h:649
differential3 dD
Definition: kernel_types.h:660
float3 P
Definition: kernel_types.h:647
float time
Definition: kernel_types.h:650
differential3 dP
Definition: kernel_types.h:659
float3 D
Definition: kernel_types.h:648
#define CLOSURE_IS_BSDF_TRANSPARENT(type)
Definition: svm_types.h:608
#define CLOSURE_IS_BSDF(type)
Definition: svm_types.h:595
#define CLOSURE_IS_BSDF_GLOSSY(type)
Definition: svm_types.h:598
#define CLOSURE_IS_BSDF_BSSRDF(type)
Definition: svm_types.h:603
#define CLOSURE_IS_BSDF_DIFFUSE(type)
Definition: svm_types.h:596
ClosureType
Definition: svm_types.h:527
#define CLOSURE_IS_BSSRDF(type)
Definition: svm_types.h:623
__forceinline bool all(const avxb &b)
Definition: util_avxb.h:214
ccl_device_inline int ceil_to_int(float f)
Definition: util_math.h:342
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()