Blender  V2.93
kernel_path.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 
17 #ifdef __OSL__
18 # include "kernel/osl/osl_shader.h"
19 #endif
20 
21 // clang-format off
22 #include "kernel/kernel_random.h"
26 #include "kernel/kernel_camera.h"
27 
28 #include "kernel/geom/geom.h"
29 #include "kernel/bvh/bvh.h"
30 
33 #include "kernel/kernel_shader.h"
34 #include "kernel/kernel_light.h"
36 #include "kernel/kernel_passes.h"
37 
38 #if defined(__VOLUME__) || defined(__SUBSURFACE__)
39 # include "kernel/kernel_volume.h"
40 #endif
41 
42 #ifdef __SUBSURFACE__
44 #endif
45 
47 #include "kernel/kernel_shadow.h"
48 #include "kernel/kernel_emission.h"
53 // clang-format on
54 
56 
59  Ray *ray,
60  Intersection *isect,
61  PathRadiance *L)
62 {
64 
65  uint visibility = path_state_ray_visibility(kg, state);
66 
68  visibility = PATH_RAY_SHADOW;
69  ray->t = kernel_data.background.ao_distance;
70  }
71 
72  bool hit = scene_intersect(kg, ray, visibility, isect);
73 
74 #ifdef __KERNEL_DEBUG__
75  if (state->flag & PATH_RAY_CAMERA) {
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;
79  }
80  L->debug_data.num_ray_bounces++;
81 #endif /* __KERNEL_DEBUG__ */
82 
83  return hit;
84 }
85 
88  Ray *ray,
89  float3 throughput,
91  ShaderData *emission_sd,
92  PathRadiance *L)
93 {
95 
96 #ifdef __LAMP_MIS__
97  if (kernel_data.integrator.use_lamp_mis && !(state->flag & PATH_RAY_CAMERA)) {
98  /* ray starting from previous non-transparent bounce */
99  Ray light_ray ccl_optional_struct_init;
100 
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;
108 
109  /* intersect with lamp */
110  indirect_lamp_emission(kg, emission_sd, state, L, &light_ray, throughput);
111  }
112 #endif /* __LAMP_MIS__ */
113 }
114 
117  ccl_addr_space Ray *ray,
118  float3 throughput,
119  ShaderData *sd,
120  ccl_global float *buffer,
121  PathRadiance *L)
122 {
123  /* eval background shader if nothing hit */
124  if (kernel_data.background.transparent && (state->flag & PATH_RAY_TRANSPARENT_BACKGROUND)) {
125  L->transparent += average(throughput);
126 
127 #ifdef __PASSES__
128  if (!(kernel_data.film.light_pass_flag & PASSMASK(BACKGROUND)))
129 #endif /* __PASSES__ */
130  return;
131  }
132 
133  /* When using the ao bounces approximation, adjust background
134  * shader intensity with ao factor. */
135  if (path_state_ao_bounce(kg, state)) {
136  throughput *= kernel_data.background.ao_bounces_factor;
137  }
138 
139 #ifdef __BACKGROUND__
140  /* sample background shader */
141  float3 L_background = indirect_background(kg, sd, state, buffer, ray);
142  path_radiance_accum_background(kg, L, state, throughput, L_background);
143 #endif /* __BACKGROUND__ */
144 }
145 
146 #ifndef __SPLIT_KERNEL__
147 
148 # ifdef __VOLUME__
149 ccl_device_forceinline VolumeIntegrateResult kernel_path_volume(KernelGlobals *kg,
150  ShaderData *sd,
151  PathState *state,
152  Ray *ray,
153  float3 *throughput,
155  bool hit,
156  ShaderData *emission_sd,
157  PathRadiance *L)
158 {
160 
161  /* Sanitize volume stack. */
162  if (!hit) {
163  kernel_volume_clean_stack(kg, state->volume_stack);
164  }
165 
166  if (state->volume_stack[0].shader == SHADER_NONE) {
167  return VOLUME_PATH_ATTENUATED;
168  }
169 
170  /* volume attenuation, emission, scatter */
171  Ray volume_ray = *ray;
172  volume_ray.t = (hit) ? isect->t : FLT_MAX;
173 
174  float step_size = volume_stack_step_size(kg, state->volume_stack);
175 
176 # ifdef __VOLUME_DECOUPLED__
177  int sampling_method = volume_stack_sampling_method(kg, state->volume_stack);
178  bool direct = (state->flag & PATH_RAY_CAMERA) != 0;
179  bool decoupled = kernel_volume_use_decoupled(kg, step_size, direct, sampling_method);
180 
181  if (decoupled) {
182  /* cache steps along volume for repeated sampling */
183  VolumeSegment volume_segment;
184 
185  shader_setup_from_volume(kg, sd, &volume_ray);
186  kernel_volume_decoupled_record(kg, state, &volume_ray, sd, &volume_segment, step_size);
187 
188  volume_segment.sampling_method = sampling_method;
189 
190  /* emission */
191  if (volume_segment.closure_flag & SD_EMISSION)
192  path_radiance_accum_emission(kg, L, state, *throughput, volume_segment.accum_emission);
193 
194  /* scattering */
196 
197  if (volume_segment.closure_flag & SD_SCATTER) {
198  int all = kernel_data.integrator.sample_all_lights_indirect;
199 
200  /* direct light sampling */
201  kernel_branched_path_volume_connect_light(
202  kg, sd, emission_sd, *throughput, state, L, all, &volume_ray, &volume_segment);
203 
204  /* indirect sample. if we use distance sampling and take just
205  * one sample for direct and indirect light, we could share
206  * this computation, but makes code a bit complex */
207  float rphase = path_state_rng_1D(kg, state, PRNG_PHASE_CHANNEL);
208  float rscatter = path_state_rng_1D(kg, state, PRNG_SCATTER_DISTANCE);
209 
210  result = kernel_volume_decoupled_scatter(
211  kg, state, &volume_ray, sd, throughput, rphase, rscatter, &volume_segment, NULL, true);
212  }
213 
214  /* free cached steps */
215  kernel_volume_decoupled_free(kg, &volume_segment);
216 
217  if (result == VOLUME_PATH_SCATTERED) {
218  if (kernel_path_volume_bounce(kg, sd, throughput, state, &L->state, ray))
219  return VOLUME_PATH_SCATTERED;
220  else
221  return VOLUME_PATH_MISSED;
222  }
223  else {
224  *throughput *= volume_segment.accum_transmittance;
225  }
226  }
227  else
228 # endif /* __VOLUME_DECOUPLED__ */
229  {
230  /* integrate along volume segment with distance sampling */
231  VolumeIntegrateResult result = kernel_volume_integrate(
232  kg, state, sd, &volume_ray, L, throughput, step_size);
233 
234 # ifdef __VOLUME_SCATTER__
235  if (result == VOLUME_PATH_SCATTERED) {
236  /* direct lighting */
237  kernel_path_volume_connect_light(kg, sd, emission_sd, *throughput, state, L);
238 
239  /* indirect light bounce */
240  if (kernel_path_volume_bounce(kg, sd, throughput, state, &L->state, ray))
241  return VOLUME_PATH_SCATTERED;
242  else
243  return VOLUME_PATH_MISSED;
244  }
245 # endif /* __VOLUME_SCATTER__ */
246  }
247 
248  return VOLUME_PATH_ATTENUATED;
249 }
250 # endif /* __VOLUME__ */
251 
252 #endif /* __SPLIT_KERNEL__ */
253 
255  ShaderData *sd,
257  ccl_addr_space Ray *ray,
258  float3 throughput,
259  ShaderData *emission_sd,
260  PathRadiance *L,
261  ccl_global float *buffer)
262 {
264 
265 #ifdef __SHADOW_TRICKS__
266  if ((sd->object_flag & SD_OBJECT_SHADOW_CATCHER)) {
269 
270  float3 bg = zero_float3();
271  if (!kernel_data.background.transparent) {
272  bg = indirect_background(kg, emission_sd, state, NULL, ray);
273  }
274  path_radiance_accum_shadowcatcher(L, throughput, bg);
275  }
276  }
277  else if (state->flag & PATH_RAY_SHADOW_CATCHER) {
278  /* Only update transparency after shadow catcher bounce. */
279  L->shadow_transparency *= average(shader_bsdf_transparency(kg, sd));
280  }
281 #endif /* __SHADOW_TRICKS__ */
282 
283  /* holdout */
284 #ifdef __HOLDOUT__
285  if (((sd->flag & SD_HOLDOUT) || (sd->object_flag & SD_OBJECT_HOLDOUT_MASK)) &&
287  const float3 holdout_weight = shader_holdout_apply(kg, sd);
288  if (kernel_data.background.transparent) {
289  L->transparent += average(holdout_weight * throughput);
290  }
291  if (isequal_float3(holdout_weight, one_float3())) {
292  return false;
293  }
294  }
295 #endif /* __HOLDOUT__ */
296 
297  /* holdout mask objects do not write data passes */
298  kernel_write_data_passes(kg, buffer, L, sd, state, throughput);
299 
300  /* blurring of bsdf after bounces, for rays that have a small likelihood
301  * of following this particular path (diffuse, rough glossy) */
302  if (kernel_data.integrator.filter_glossy != FLT_MAX) {
303  float blur_pdf = kernel_data.integrator.filter_glossy * state->min_ray_pdf;
304 
305  if (blur_pdf < 1.0f) {
306  float blur_roughness = sqrtf(1.0f - blur_pdf) * 0.5f;
307  shader_bsdf_blur(kg, sd, blur_roughness);
308  }
309  }
310 
311 #ifdef __EMISSION__
312  /* emission */
313  if (sd->flag & SD_EMISSION) {
315  kg, sd, sd->ray_length, state->flag, state->ray_pdf);
316  path_radiance_accum_emission(kg, L, state, throughput, emission);
317  }
318 #endif /* __EMISSION__ */
319 
320  return true;
321 }
322 
323 #ifdef __KERNEL_OPTIX__
324 ccl_device_inline /* inline trace calls */
325 #else
327 #endif
328  void
329  kernel_path_ao(KernelGlobals *kg,
330  ShaderData *sd,
331  ShaderData *emission_sd,
332  PathRadiance *L,
334  float3 throughput,
335  float3 ao_alpha)
336 {
338 
339  /* todo: solve correlation */
340  float bsdf_u, bsdf_v;
341 
342  path_state_rng_2D(kg, state, PRNG_BSDF_U, &bsdf_u, &bsdf_v);
343 
344  float ao_factor = kernel_data.background.ao_factor;
345  float3 ao_N;
346  float3 ao_bsdf = shader_bsdf_ao(kg, sd, ao_factor, &ao_N);
347  float3 ao_D;
348  float ao_pdf;
349 
350  sample_cos_hemisphere(ao_N, bsdf_u, bsdf_v, &ao_D, &ao_pdf);
351 
352  if (dot(sd->Ng, ao_D) > 0.0f && ao_pdf != 0.0f) {
353  Ray light_ray;
354  float3 ao_shadow;
355 
356  light_ray.P = ray_offset(sd->P, sd->Ng);
357  light_ray.D = ao_D;
358  light_ray.t = kernel_data.background.ao_distance;
359  light_ray.time = sd->time;
360  light_ray.dP = sd->dP;
361  light_ray.dD = differential3_zero();
362 
363  if (!shadow_blocked(kg, sd, emission_sd, state, &light_ray, &ao_shadow)) {
364  path_radiance_accum_ao(kg, L, state, throughput, ao_alpha, ao_bsdf, ao_shadow);
365  }
366  else {
367  path_radiance_accum_total_ao(L, state, throughput, ao_bsdf);
368  }
369  }
370 }
371 
372 #ifndef __SPLIT_KERNEL__
373 
374 # if defined(__BRANCHED_PATH__) || defined(__BAKING__)
375 
376 ccl_device void kernel_path_indirect(KernelGlobals *kg,
377  ShaderData *sd,
378  ShaderData *emission_sd,
379  Ray *ray,
380  float3 throughput,
381  PathState *state,
382  PathRadiance *L)
383 {
384 # ifdef __SUBSURFACE__
385  SubsurfaceIndirectRays ss_indirect;
386  kernel_path_subsurface_init_indirect(&ss_indirect);
387 
388  for (;;) {
389 # endif /* __SUBSURFACE__ */
390 
391  /* path iteration */
392  for (;;) {
393  /* Find intersection with objects in scene. */
394  Intersection isect;
395  bool hit = kernel_path_scene_intersect(kg, state, ray, &isect, L);
396 
397  /* Find intersection with lamps and compute emission for MIS. */
398  kernel_path_lamp_emission(kg, state, ray, throughput, &isect, sd, L);
399 
400 # ifdef __VOLUME__
401  /* Volume integration. */
402  VolumeIntegrateResult result = kernel_path_volume(
403  kg, sd, state, ray, &throughput, &isect, hit, emission_sd, L);
404 
405  if (result == VOLUME_PATH_SCATTERED) {
406  continue;
407  }
408  else if (result == VOLUME_PATH_MISSED) {
409  break;
410  }
411 # endif /* __VOLUME__*/
412 
413  /* Shade background. */
414  if (!hit) {
415  kernel_path_background(kg, state, ray, throughput, sd, NULL, L);
416  break;
417  }
418  else if (path_state_ao_bounce(kg, state)) {
419  break;
420  }
421 
422  /* Setup shader data. */
423  shader_setup_from_ray(kg, sd, &isect, ray);
424 
425  /* Skip most work for volume bounding surface. */
426 # ifdef __VOLUME__
427  if (!(sd->flag & SD_HAS_ONLY_VOLUME)) {
428 # endif
429 
430  /* Evaluate shader. */
431  shader_eval_surface(kg, sd, state, NULL, state->flag);
433 
434  /* Apply shadow catcher, holdout, emission. */
435  if (!kernel_path_shader_apply(kg, sd, state, ray, throughput, emission_sd, L, NULL)) {
436  break;
437  }
438 
439  /* path termination. this is a strange place to put the termination, it's
440  * mainly due to the mixed in MIS that we use. gives too many unneeded
441  * shader evaluations, only need emission if we are going to terminate */
442  float probability = path_state_continuation_probability(kg, state, throughput);
443 
444  if (probability == 0.0f) {
445  break;
446  }
447  else if (probability != 1.0f) {
448  float terminate = path_state_rng_1D(kg, state, PRNG_TERMINATE);
449 
450  if (terminate >= probability)
451  break;
452 
453  throughput /= probability;
454  }
455 
456 # ifdef __DENOISING_FEATURES__
457  kernel_update_denoising_features(kg, sd, state, L);
458 # endif
459 
460 # ifdef __AO__
461  /* ambient occlusion */
462  if (kernel_data.integrator.use_ambient_occlusion) {
463  kernel_path_ao(kg, sd, emission_sd, L, state, throughput, zero_float3());
464  }
465 # endif /* __AO__ */
466 
467 # ifdef __SUBSURFACE__
468  /* bssrdf scatter to a different location on the same object, replacing
469  * the closures with a diffuse BSDF */
470  if (sd->flag & SD_BSSRDF) {
471  if (kernel_path_subsurface_scatter(
472  kg, sd, emission_sd, L, state, ray, &throughput, &ss_indirect)) {
473  break;
474  }
475  }
476 # endif /* __SUBSURFACE__ */
477 
478 # if defined(__EMISSION__)
479  int all = (kernel_data.integrator.sample_all_lights_indirect) ||
480  (state->flag & PATH_RAY_SHADOW_CATCHER);
481  kernel_branched_path_surface_connect_light(
482  kg, sd, emission_sd, state, throughput, 1.0f, L, all);
483 # endif /* defined(__EMISSION__) */
484 
485 # ifdef __VOLUME__
486  }
487 # endif
488 
489  if (!kernel_path_surface_bounce(kg, sd, &throughput, state, &L->state, ray))
490  break;
491  }
492 
493 # ifdef __SUBSURFACE__
494  /* Trace indirect subsurface rays by restarting the loop. this uses less
495  * stack memory than invoking kernel_path_indirect.
496  */
497  if (ss_indirect.num_rays) {
498  kernel_path_subsurface_setup_indirect(kg, &ss_indirect, state, ray, L, &throughput);
499  }
500  else {
501  break;
502  }
503  }
504 # endif /* __SUBSURFACE__ */
505 }
506 
507 # endif /* defined(__BRANCHED_PATH__) || defined(__BAKING__) */
508 
510  PathState *state,
511  float3 throughput,
512  Ray *ray,
513  PathRadiance *L,
514  ccl_global float *buffer,
515  ShaderData *emission_sd)
516 {
518 
519  /* Shader data memory used for both volumes and surfaces, saves stack space. */
520  ShaderData sd;
521 
522 # ifdef __SUBSURFACE__
523  SubsurfaceIndirectRays ss_indirect;
524  kernel_path_subsurface_init_indirect(&ss_indirect);
525 
526  for (;;) {
527 # endif /* __SUBSURFACE__ */
528 
529  /* path iteration */
530  for (;;) {
531  /* Find intersection with objects in scene. */
532  Intersection isect;
533  bool hit = kernel_path_scene_intersect(kg, state, ray, &isect, L);
534 
535  /* Find intersection with lamps and compute emission for MIS. */
536  kernel_path_lamp_emission(kg, state, ray, throughput, &isect, &sd, L);
537 
538 # ifdef __VOLUME__
539  /* Volume integration. */
540  VolumeIntegrateResult result = kernel_path_volume(
541  kg, &sd, state, ray, &throughput, &isect, hit, emission_sd, L);
542 
543  if (result == VOLUME_PATH_SCATTERED) {
544  continue;
545  }
546  else if (result == VOLUME_PATH_MISSED) {
547  break;
548  }
549 # endif /* __VOLUME__*/
550 
551  /* Shade background. */
552  if (!hit) {
553  kernel_path_background(kg, state, ray, throughput, &sd, buffer, L);
554  break;
555  }
556  else if (path_state_ao_bounce(kg, state)) {
557  break;
558  }
559 
560  /* Setup shader data. */
561  shader_setup_from_ray(kg, &sd, &isect, ray);
562 
563  /* Skip most work for volume bounding surface. */
564 # ifdef __VOLUME__
565  if (!(sd.flag & SD_HAS_ONLY_VOLUME)) {
566 # endif
567 
568  /* Evaluate shader. */
569  shader_eval_surface(kg, &sd, state, buffer, state->flag);
571 
572  /* Apply shadow catcher, holdout, emission. */
573  if (!kernel_path_shader_apply(kg, &sd, state, ray, throughput, emission_sd, L, buffer)) {
574  break;
575  }
576 
577  /* path termination. this is a strange place to put the termination, it's
578  * mainly due to the mixed in MIS that we use. gives too many unneeded
579  * shader evaluations, only need emission if we are going to terminate */
580  float probability = path_state_continuation_probability(kg, state, throughput);
581 
582  if (probability == 0.0f) {
583  break;
584  }
585  else if (probability != 1.0f) {
586  float terminate = path_state_rng_1D(kg, state, PRNG_TERMINATE);
587  if (terminate >= probability)
588  break;
589 
590  throughput /= probability;
591  }
592 
593 # ifdef __DENOISING_FEATURES__
594  kernel_update_denoising_features(kg, &sd, state, L);
595 # endif
596 
597 # ifdef __AO__
598  /* ambient occlusion */
599  if (kernel_data.integrator.use_ambient_occlusion) {
600  kernel_path_ao(kg, &sd, emission_sd, L, state, throughput, shader_bsdf_alpha(kg, &sd));
601  }
602 # endif /* __AO__ */
603 
604 # ifdef __SUBSURFACE__
605  /* bssrdf scatter to a different location on the same object, replacing
606  * the closures with a diffuse BSDF */
607  if (sd.flag & SD_BSSRDF) {
608  if (kernel_path_subsurface_scatter(
609  kg, &sd, emission_sd, L, state, ray, &throughput, &ss_indirect)) {
610  break;
611  }
612  }
613 # endif /* __SUBSURFACE__ */
614 
615 # ifdef __EMISSION__
616  /* direct lighting */
617  kernel_path_surface_connect_light(kg, &sd, emission_sd, throughput, state, L);
618 # endif /* __EMISSION__ */
619 
620 # ifdef __VOLUME__
621  }
622 # endif
623 
624  /* compute direct lighting and next bounce */
625  if (!kernel_path_surface_bounce(kg, &sd, &throughput, state, &L->state, ray))
626  break;
627  }
628 
629 # ifdef __SUBSURFACE__
630  /* Trace indirect subsurface rays by restarting the loop. this uses less
631  * stack memory than invoking kernel_path_indirect.
632  */
633  if (ss_indirect.num_rays) {
634  kernel_path_subsurface_setup_indirect(kg, &ss_indirect, state, ray, L, &throughput);
635  }
636  else {
637  break;
638  }
639  }
640 # endif /* __SUBSURFACE__ */
641 }
642 
644  KernelGlobals *kg, ccl_global float *buffer, int sample, int x, int y, int offset, int stride)
645 {
647 
648  /* buffer offset */
649  int index = offset + x + y * stride;
650  int pass_stride = kernel_data.film.pass_stride;
651 
652  buffer += index * pass_stride;
653 
654  if (kernel_data.film.pass_adaptive_aux_buffer) {
655  ccl_global float4 *aux = (ccl_global float4 *)(buffer +
656  kernel_data.film.pass_adaptive_aux_buffer);
657  if ((*aux).w > 0.0f) {
658  return;
659  }
660  }
661 
662  /* Initialize random numbers and sample ray. */
663  uint rng_hash;
664  Ray ray;
665 
666  kernel_path_trace_setup(kg, sample, x, y, &rng_hash, &ray);
667 
668  if (ray.t == 0.0f) {
669  return;
670  }
671 
672  /* Initialize state. */
673  float3 throughput = one_float3();
674 
675  PathRadiance L;
677 
678  ShaderDataTinyStorage emission_sd_storage;
679  ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
680 
682  path_state_init(kg, emission_sd, &state, rng_hash, sample, &ray);
683 
684 # ifdef __KERNEL_OPTIX__
685  /* Force struct into local memory to avoid costly spilling on trace calls. */
686  if (pass_stride < 0) /* This is never executed and just prevents the compiler from doing SROA. */
687  for (int i = 0; i < sizeof(L); ++i)
688  reinterpret_cast<unsigned char *>(&L)[-pass_stride + i] = 0;
689 # endif
690 
691  /* Integrate. */
692  kernel_path_integrate(kg, &state, throughput, &ray, &L, buffer, emission_sd);
693 
695 }
696 
697 #endif /* __SPLIT_KERNEL__ */
698 
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_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 kernel_data
#define ccl_addr_space
#define ccl_device_forceinline
#define ccl_optional_struct_init
#define ccl_device
#define ccl_device_inline
#define ccl_global
#define ccl_device_noinline
#define CCL_NAMESPACE_END
#define sqrtf(x)
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)
Definition: kernel_path.h:643
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_device_forceinline void kernel_path_integrate(KernelGlobals *kg, PathState *state, float3 throughput, Ray *ray, PathRadiance *L, ccl_global float *buffer, ShaderData *emission_sd)
Definition: kernel_path.h:509
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_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)
Definition: kernel_path.h:86
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)
Definition: kernel_path.h:329
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)
Definition: kernel_shader.h:59
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
@ SD_BSSRDF
Definition: kernel_types.h:851
@ SD_HAS_ONLY_VOLUME
Definition: kernel_types.h:875
@ SD_SCATTER
Definition: kernel_types.h:857
@ SD_HOLDOUT
Definition: kernel_types.h:853
@ SD_EMISSION
Definition: kernel_types.h:845
#define AS_SHADER_DATA(shader_data_tiny_storage)
#define SHADER_NONE
Definition: kernel_types.h:58
@ 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_STORE_SHADOW_INFO
Definition: kernel_types.h:308
@ PATH_RAY_SHADOW
Definition: kernel_types.h:284
@ PATH_RAY_SHADOW_CATCHER
Definition: kernel_types.h:306
@ PATH_RAY_TRANSPARENT_BACKGROUND
Definition: kernel_types.h:310
@ PATH_RAY_CAMERA
Definition: kernel_types.h:266
ShaderDataTinyStorage
ShaderData
@ SD_OBJECT_HOLDOUT_MASK
Definition: kernel_types.h:904
@ SD_OBJECT_SHADOW_CATCHER
Definition: kernel_types.h:918
#define PASSMASK(pass)
Definition: kernel_types.h:341
VolumeIntegrateResult
Definition: kernel_volume.h:27
@ VOLUME_PATH_MISSED
Definition: kernel_volume.h:30
@ VOLUME_PATH_SCATTERED
Definition: kernel_volume.h:28
@ VOLUME_PATH_ATTENUATED
Definition: kernel_volume.h:29
static ulong state[N]
#define L
static void sample(SocketReader *reader, int x, int y, float color[4])
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
__forceinline bool all(const avxb &b)
Definition: util_avxb.h:214
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_AO
@ PROFILING_SHADER_APPLY
@ PROFILING_VOLUME
@ PROFILING_SCENE_INTERSECT
@ PROFILING_RAY_SETUP
@ PROFILING_INDIRECT_EMISSION