Blender  V2.93
kernel_shader.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 /*
18  * ShaderData, used in four steps:
19  *
20  * Setup from incoming ray, sampled position and background.
21  * Execute for surface, volume or displacement.
22  * Evaluate one or more closures.
23  * Release.
24  */
25 
26 // clang-format off
27 #include "kernel/closure/alloc.h"
29 #include "kernel/closure/bsdf.h"
31 // clang-format on
32 
33 #include "kernel/svm/svm.h"
34 
36 
37 /* ShaderData setup from incoming ray */
38 
39 #ifdef __OBJECT_MOTION__
40 ccl_device void shader_setup_object_transforms(KernelGlobals *kg, ShaderData *sd, float time)
41 {
42  if (sd->object_flag & SD_OBJECT_MOTION) {
43  sd->ob_tfm = object_fetch_transform_motion(kg, sd->object, time);
44  sd->ob_itfm = transform_quick_inverse(sd->ob_tfm);
45  }
46  else {
47  sd->ob_tfm = object_fetch_transform(kg, sd->object, OBJECT_TRANSFORM);
48  sd->ob_itfm = object_fetch_transform(kg, sd->object, OBJECT_INVERSE_TRANSFORM);
49  }
50 }
51 #endif
52 
53 #ifdef __KERNEL_OPTIX__
55 #else
57 #endif
58  void
59  shader_setup_from_ray(KernelGlobals *kg,
60  ShaderData *sd,
61  const Intersection *isect,
62  const Ray *ray)
63 {
65 
66  sd->object = (isect->object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, isect->prim) :
67  isect->object;
68  sd->lamp = LAMP_NONE;
69 
70  sd->type = isect->type;
71  sd->flag = 0;
72  sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
73 
74  /* matrices and time */
75 #ifdef __OBJECT_MOTION__
76  shader_setup_object_transforms(kg, sd, ray->time);
77 #endif
78  sd->time = ray->time;
79 
80  sd->prim = kernel_tex_fetch(__prim_index, isect->prim);
81  sd->ray_length = isect->t;
82 
83  sd->u = isect->u;
84  sd->v = isect->v;
85 
86 #ifdef __HAIR__
87  if (sd->type & PRIMITIVE_ALL_CURVE) {
88  /* curve */
89  curve_shader_setup(kg, sd, isect, ray);
90  }
91  else
92 #endif
93  if (sd->type & PRIMITIVE_TRIANGLE) {
94  /* static triangle */
95  float3 Ng = triangle_normal(kg, sd);
96  sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
97 
98  /* vectors */
99  sd->P = triangle_refine(kg, sd, isect, ray);
100  sd->Ng = Ng;
101  sd->N = Ng;
102 
103  /* smooth normal */
104  if (sd->shader & SHADER_SMOOTH_NORMAL)
105  sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);
106 
107 #ifdef __DPDU__
108  /* dPdu/dPdv */
109  triangle_dPdudv(kg, sd->prim, &sd->dPdu, &sd->dPdv);
110 #endif
111  }
112  else {
113  /* motion triangle */
114  motion_triangle_shader_setup(kg, sd, isect, ray, false);
115  }
116 
117  sd->I = -ray->D;
118 
119  sd->flag |= kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
120 
121  if (isect->object != OBJECT_NONE) {
122  /* instance transform */
123  object_normal_transform_auto(kg, sd, &sd->N);
124  object_normal_transform_auto(kg, sd, &sd->Ng);
125 #ifdef __DPDU__
126  object_dir_transform_auto(kg, sd, &sd->dPdu);
127  object_dir_transform_auto(kg, sd, &sd->dPdv);
128 #endif
129  }
130 
131  /* backfacing test */
132  bool backfacing = (dot(sd->Ng, sd->I) < 0.0f);
133 
134  if (backfacing) {
135  sd->flag |= SD_BACKFACING;
136  sd->Ng = -sd->Ng;
137  sd->N = -sd->N;
138 #ifdef __DPDU__
139  sd->dPdu = -sd->dPdu;
140  sd->dPdv = -sd->dPdv;
141 #endif
142  }
143 
144 #ifdef __RAY_DIFFERENTIALS__
145  /* differentials */
146  differential_transfer(&sd->dP, ray->dP, ray->D, ray->dD, sd->Ng, isect->t);
147  differential_incoming(&sd->dI, ray->dD);
148  differential_dudv(&sd->du, &sd->dv, sd->dPdu, sd->dPdv, sd->dP, sd->Ng);
149 #endif
150 
151  PROFILING_SHADER(sd->shader);
152  PROFILING_OBJECT(sd->object);
153 }
154 
155 /* ShaderData setup from BSSRDF scatter */
156 
157 #ifdef __SUBSURFACE__
158 # ifndef __KERNEL_CUDA__
160 # else
162 # endif
163  void
164  shader_setup_from_subsurface(KernelGlobals *kg,
165  ShaderData *sd,
166  const Intersection *isect,
167  const Ray *ray)
168 {
170 
171  const bool backfacing = sd->flag & SD_BACKFACING;
172 
173  /* object, matrices, time, ray_length stay the same */
174  sd->flag = 0;
175  sd->object_flag = kernel_tex_fetch(__object_flag, sd->object);
176  sd->prim = kernel_tex_fetch(__prim_index, isect->prim);
177  sd->type = isect->type;
178 
179  sd->u = isect->u;
180  sd->v = isect->v;
181 
182  /* fetch triangle data */
183  if (sd->type == PRIMITIVE_TRIANGLE) {
184  float3 Ng = triangle_normal(kg, sd);
185  sd->shader = kernel_tex_fetch(__tri_shader, sd->prim);
186 
187  /* static triangle */
188  sd->P = triangle_refine_local(kg, sd, isect, ray);
189  sd->Ng = Ng;
190  sd->N = Ng;
191 
192  if (sd->shader & SHADER_SMOOTH_NORMAL)
193  sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);
194 
195 # ifdef __DPDU__
196  /* dPdu/dPdv */
197  triangle_dPdudv(kg, sd->prim, &sd->dPdu, &sd->dPdv);
198 # endif
199  }
200  else {
201  /* motion triangle */
202  motion_triangle_shader_setup(kg, sd, isect, ray, true);
203  }
204 
205  sd->flag |= kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
206 
207  if (isect->object != OBJECT_NONE) {
208  /* instance transform */
209  object_normal_transform_auto(kg, sd, &sd->N);
210  object_normal_transform_auto(kg, sd, &sd->Ng);
211 # ifdef __DPDU__
212  object_dir_transform_auto(kg, sd, &sd->dPdu);
213  object_dir_transform_auto(kg, sd, &sd->dPdv);
214 # endif
215  }
216 
217  /* backfacing test */
218  if (backfacing) {
219  sd->flag |= SD_BACKFACING;
220  sd->Ng = -sd->Ng;
221  sd->N = -sd->N;
222 # ifdef __DPDU__
223  sd->dPdu = -sd->dPdu;
224  sd->dPdv = -sd->dPdv;
225 # endif
226  }
227 
228  /* should not get used in principle as the shading will only use a diffuse
229  * BSDF, but the shader might still access it */
230  sd->I = sd->N;
231 
232 # ifdef __RAY_DIFFERENTIALS__
233  /* differentials */
234  differential_dudv(&sd->du, &sd->dv, sd->dPdu, sd->dPdv, sd->dP, sd->Ng);
235  /* don't modify dP and dI */
236 # endif
237 
238  PROFILING_SHADER(sd->shader);
239 }
240 #endif
241 
242 /* ShaderData setup from position sampled on mesh */
243 
245  ShaderData *sd,
246  const float3 P,
247  const float3 Ng,
248  const float3 I,
249  int shader,
250  int object,
251  int prim,
252  float u,
253  float v,
254  float t,
255  float time,
256  bool object_space,
257  int lamp)
258 {
260 
261  /* vectors */
262  sd->P = P;
263  sd->N = Ng;
264  sd->Ng = Ng;
265  sd->I = I;
266  sd->shader = shader;
267  if (prim != PRIM_NONE)
268  sd->type = PRIMITIVE_TRIANGLE;
269  else if (lamp != LAMP_NONE)
270  sd->type = PRIMITIVE_LAMP;
271  else
272  sd->type = PRIMITIVE_NONE;
273 
274  /* primitive */
275  sd->object = object;
276  sd->lamp = LAMP_NONE;
277  /* currently no access to bvh prim index for strand sd->prim*/
278  sd->prim = prim;
279  sd->u = u;
280  sd->v = v;
281  sd->time = time;
282  sd->ray_length = t;
283 
284  sd->flag = kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
285  sd->object_flag = 0;
286  if (sd->object != OBJECT_NONE) {
287  sd->object_flag |= kernel_tex_fetch(__object_flag, sd->object);
288 
289 #ifdef __OBJECT_MOTION__
290  shader_setup_object_transforms(kg, sd, time);
291  }
292  else if (lamp != LAMP_NONE) {
293  sd->ob_tfm = lamp_fetch_transform(kg, lamp, false);
294  sd->ob_itfm = lamp_fetch_transform(kg, lamp, true);
295  sd->lamp = lamp;
296 #else
297  }
298  else if (lamp != LAMP_NONE) {
299  sd->lamp = lamp;
300 #endif
301  }
302 
303  /* transform into world space */
304  if (object_space) {
305  object_position_transform_auto(kg, sd, &sd->P);
306  object_normal_transform_auto(kg, sd, &sd->Ng);
307  sd->N = sd->Ng;
308  object_dir_transform_auto(kg, sd, &sd->I);
309  }
310 
311  if (sd->type & PRIMITIVE_TRIANGLE) {
312  /* smooth normal */
313  if (sd->shader & SHADER_SMOOTH_NORMAL) {
314  sd->N = triangle_smooth_normal(kg, Ng, sd->prim, sd->u, sd->v);
315 
316  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
317  object_normal_transform_auto(kg, sd, &sd->N);
318  }
319  }
320 
321  /* dPdu/dPdv */
322 #ifdef __DPDU__
323  triangle_dPdudv(kg, sd->prim, &sd->dPdu, &sd->dPdv);
324 
325  if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
326  object_dir_transform_auto(kg, sd, &sd->dPdu);
327  object_dir_transform_auto(kg, sd, &sd->dPdv);
328  }
329 #endif
330  }
331  else {
332 #ifdef __DPDU__
333  sd->dPdu = zero_float3();
334  sd->dPdv = zero_float3();
335 #endif
336  }
337 
338  /* backfacing test */
339  if (sd->prim != PRIM_NONE) {
340  bool backfacing = (dot(sd->Ng, sd->I) < 0.0f);
341 
342  if (backfacing) {
343  sd->flag |= SD_BACKFACING;
344  sd->Ng = -sd->Ng;
345  sd->N = -sd->N;
346 #ifdef __DPDU__
347  sd->dPdu = -sd->dPdu;
348  sd->dPdv = -sd->dPdv;
349 #endif
350  }
351  }
352 
353 #ifdef __RAY_DIFFERENTIALS__
354  /* no ray differentials here yet */
355  sd->dP = differential3_zero();
356  sd->dI = differential3_zero();
357  sd->du = differential_zero();
358  sd->dv = differential_zero();
359 #endif
360 
361  PROFILING_SHADER(sd->shader);
362  PROFILING_OBJECT(sd->object);
363 }
364 
365 /* ShaderData setup for displacement */
366 
368  KernelGlobals *kg, ShaderData *sd, int object, int prim, float u, float v)
369 {
370  float3 P, Ng, I = zero_float3();
371  int shader;
372 
373  triangle_point_normal(kg, object, prim, u, v, &P, &Ng, &shader);
374 
375  /* force smooth shading for displacement */
377 
379  kg,
380  sd,
381  P,
382  Ng,
383  I,
384  shader,
385  object,
386  prim,
387  u,
388  v,
389  0.0f,
390  0.5f,
391  !(kernel_tex_fetch(__object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED),
392  LAMP_NONE);
393 }
394 
395 /* ShaderData setup from ray into background */
396 
398  ShaderData *sd,
399  const Ray *ray)
400 {
402 
403  /* vectors */
404  sd->P = ray->D;
405  sd->N = -ray->D;
406  sd->Ng = -ray->D;
407  sd->I = -ray->D;
408  sd->shader = kernel_data.background.surface_shader;
409  sd->flag = kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
410  sd->object_flag = 0;
411  sd->time = ray->time;
412  sd->ray_length = 0.0f;
413 
414  sd->object = OBJECT_NONE;
415  sd->lamp = LAMP_NONE;
416  sd->prim = PRIM_NONE;
417  sd->u = 0.0f;
418  sd->v = 0.0f;
419 
420 #ifdef __DPDU__
421  /* dPdu/dPdv */
422  sd->dPdu = zero_float3();
423  sd->dPdv = zero_float3();
424 #endif
425 
426 #ifdef __RAY_DIFFERENTIALS__
427  /* differentials */
428  sd->dP = ray->dD;
429  differential_incoming(&sd->dI, sd->dP);
430  sd->du = differential_zero();
431  sd->dv = differential_zero();
432 #endif
433 
434  /* for NDC coordinates */
435  sd->ray_P = ray->P;
436 
437  PROFILING_SHADER(sd->shader);
438  PROFILING_OBJECT(sd->object);
439 }
440 
441 /* ShaderData setup from point inside volume */
442 
443 #ifdef __VOLUME__
444 ccl_device_inline void shader_setup_from_volume(KernelGlobals *kg, ShaderData *sd, const Ray *ray)
445 {
447 
448  /* vectors */
449  sd->P = ray->P;
450  sd->N = -ray->D;
451  sd->Ng = -ray->D;
452  sd->I = -ray->D;
453  sd->shader = SHADER_NONE;
454  sd->flag = 0;
455  sd->object_flag = 0;
456  sd->time = ray->time;
457  sd->ray_length = 0.0f; /* todo: can we set this to some useful value? */
458 
459  sd->object = OBJECT_NONE; /* todo: fill this for texture coordinates */
460  sd->lamp = LAMP_NONE;
461  sd->prim = PRIM_NONE;
462  sd->type = PRIMITIVE_NONE;
463 
464  sd->u = 0.0f;
465  sd->v = 0.0f;
466 
467 # ifdef __DPDU__
468  /* dPdu/dPdv */
469  sd->dPdu = zero_float3();
470  sd->dPdv = zero_float3();
471 # endif
472 
473 # ifdef __RAY_DIFFERENTIALS__
474  /* differentials */
475  sd->dP = ray->dD;
476  differential_incoming(&sd->dI, sd->dP);
477  sd->du = differential_zero();
478  sd->dv = differential_zero();
479 # endif
480 
481  /* for NDC coordinates */
482  sd->ray_P = ray->P;
483  sd->ray_dP = ray->dP;
484 
485  PROFILING_SHADER(sd->shader);
486  PROFILING_OBJECT(sd->object);
487 }
488 #endif /* __VOLUME__ */
489 
490 /* Merging */
491 
492 #if defined(__BRANCHED_PATH__) || defined(__VOLUME__)
493 ccl_device_inline void shader_merge_closures(ShaderData *sd)
494 {
495  /* merge identical closures, better when we sample a single closure at a time */
496  for (int i = 0; i < sd->num_closure; i++) {
497  ShaderClosure *sci = &sd->closure[i];
498 
499  for (int j = i + 1; j < sd->num_closure; j++) {
500  ShaderClosure *scj = &sd->closure[j];
501 
502  if (sci->type != scj->type)
503  continue;
504  if (!bsdf_merge(sci, scj))
505  continue;
506 
507  sci->weight += scj->weight;
508  sci->sample_weight += scj->sample_weight;
509 
510  int size = sd->num_closure - (j + 1);
511  if (size > 0) {
512  for (int k = 0; k < size; k++) {
513  scj[k] = scj[k + 1];
514  }
515  }
516 
517  sd->num_closure--;
518  kernel_assert(sd->num_closure >= 0);
519  j--;
520  }
521  }
522 }
523 #endif /* __BRANCHED_PATH__ || __VOLUME__ */
524 
525 /* Defensive sampling. */
526 
528 {
529  /* We can likely also do defensive sampling at deeper bounces, particularly
530  * for cases like a perfect mirror but possibly also others. This will need
531  * a good heuristic. */
532  if (state->bounce + state->transparent_bounce == 0 && sd->num_closure > 1) {
533  float sum = 0.0f;
534 
535  for (int i = 0; i < sd->num_closure; i++) {
536  ShaderClosure *sc = &sd->closure[i];
537  if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type)) {
538  sum += sc->sample_weight;
539  }
540  }
541 
542  for (int i = 0; i < sd->num_closure; i++) {
543  ShaderClosure *sc = &sd->closure[i];
544  if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type)) {
545  sc->sample_weight = max(sc->sample_weight, 0.125f * sum);
546  }
547  }
548  }
549 }
550 
551 /* BSDF */
552 
554  ShaderData *sd,
555  const float3 omega_in,
556  float *pdf,
557  const ShaderClosure *skip_sc,
558  BsdfEval *result_eval,
559  float sum_pdf,
560  float sum_sample_weight)
561 {
562  /* this is the veach one-sample model with balance heuristic, some pdf
563  * factors drop out when using balance heuristic weighting */
564  for (int i = 0; i < sd->num_closure; i++) {
565  const ShaderClosure *sc = &sd->closure[i];
566 
567  if (sc != skip_sc && CLOSURE_IS_BSDF(sc->type)) {
568  float bsdf_pdf = 0.0f;
569  float3 eval = bsdf_eval(kg, sd, sc, omega_in, &bsdf_pdf);
570 
571  if (bsdf_pdf != 0.0f) {
572  bsdf_eval_accum(result_eval, sc->type, eval * sc->weight, 1.0f);
573  sum_pdf += bsdf_pdf * sc->sample_weight;
574  }
575 
576  sum_sample_weight += sc->sample_weight;
577  }
578  }
579 
580  *pdf = (sum_sample_weight > 0.0f) ? sum_pdf / sum_sample_weight : 0.0f;
581 }
582 
583 #ifdef __BRANCHED_PATH__
584 ccl_device_inline void _shader_bsdf_multi_eval_branched(KernelGlobals *kg,
585  ShaderData *sd,
586  const float3 omega_in,
587  BsdfEval *result_eval,
588  float light_pdf,
589  bool use_mis)
590 {
591  for (int i = 0; i < sd->num_closure; i++) {
592  const ShaderClosure *sc = &sd->closure[i];
593  if (CLOSURE_IS_BSDF(sc->type)) {
594  float bsdf_pdf = 0.0f;
595  float3 eval = bsdf_eval(kg, sd, sc, omega_in, &bsdf_pdf);
596  if (bsdf_pdf != 0.0f) {
597  float mis_weight = use_mis ? power_heuristic(light_pdf, bsdf_pdf) : 1.0f;
598  bsdf_eval_accum(result_eval, sc->type, eval * sc->weight, mis_weight);
599  }
600  }
601  }
602 }
603 #endif /* __BRANCHED_PATH__ */
604 
605 #ifndef __KERNEL_CUDA__
607 #else
609 #endif
610  void
611  shader_bsdf_eval(KernelGlobals *kg,
612  ShaderData *sd,
613  const float3 omega_in,
614  BsdfEval *eval,
615  float light_pdf,
616  bool use_mis)
617 {
619 
620  bsdf_eval_init(eval, NBUILTIN_CLOSURES, zero_float3(), kernel_data.film.use_light_pass);
621 
622 #ifdef __BRANCHED_PATH__
623  if (kernel_data.integrator.branched)
624  _shader_bsdf_multi_eval_branched(kg, sd, omega_in, eval, light_pdf, use_mis);
625  else
626 #endif
627  {
628  float pdf;
629  _shader_bsdf_multi_eval(kg, sd, omega_in, &pdf, NULL, eval, 0.0f, 0.0f);
630  if (use_mis) {
631  float weight = power_heuristic(light_pdf, pdf);
632  bsdf_eval_mis(eval, weight);
633  }
634  }
635 }
636 
638 {
639  /* Note the sampling here must match shader_bssrdf_pick,
640  * since we reuse the same random number. */
641  int sampled = 0;
642 
643  if (sd->num_closure > 1) {
644  /* Pick a BSDF or based on sample weights. */
645  float sum = 0.0f;
646 
647  for (int i = 0; i < sd->num_closure; i++) {
648  const ShaderClosure *sc = &sd->closure[i];
649 
650  if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type)) {
651  sum += sc->sample_weight;
652  }
653  }
654 
655  float r = (*randu) * sum;
656  float partial_sum = 0.0f;
657 
658  for (int i = 0; i < sd->num_closure; i++) {
659  const ShaderClosure *sc = &sd->closure[i];
660 
661  if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type)) {
662  float next_sum = partial_sum + sc->sample_weight;
663 
664  if (r < next_sum) {
665  sampled = i;
666 
667  /* Rescale to reuse for direction sample, to better preserve stratification. */
668  *randu = (r - partial_sum) / sc->sample_weight;
669  break;
670  }
671 
672  partial_sum = next_sum;
673  }
674  }
675  }
676 
677  const ShaderClosure *sc = &sd->closure[sampled];
678  return CLOSURE_IS_BSDF(sc->type) ? sc : NULL;
679 }
680 
682  ccl_addr_space float3 *throughput,
683  float *randu)
684 {
685  /* Note the sampling here must match shader_bsdf_pick,
686  * since we reuse the same random number. */
687  int sampled = 0;
688 
689  if (sd->num_closure > 1) {
690  /* Pick a BSDF or BSSRDF or based on sample weights. */
691  float sum_bsdf = 0.0f;
692  float sum_bssrdf = 0.0f;
693 
694  for (int i = 0; i < sd->num_closure; i++) {
695  const ShaderClosure *sc = &sd->closure[i];
696 
697  if (CLOSURE_IS_BSDF(sc->type)) {
698  sum_bsdf += sc->sample_weight;
699  }
700  else if (CLOSURE_IS_BSSRDF(sc->type)) {
701  sum_bssrdf += sc->sample_weight;
702  }
703  }
704 
705  float r = (*randu) * (sum_bsdf + sum_bssrdf);
706  float partial_sum = 0.0f;
707 
708  for (int i = 0; i < sd->num_closure; i++) {
709  const ShaderClosure *sc = &sd->closure[i];
710 
711  if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type)) {
712  float next_sum = partial_sum + sc->sample_weight;
713 
714  if (r < next_sum) {
715  if (CLOSURE_IS_BSDF(sc->type)) {
716  *throughput *= (sum_bsdf + sum_bssrdf) / sum_bsdf;
717  return NULL;
718  }
719  else {
720  *throughput *= (sum_bsdf + sum_bssrdf) / sum_bssrdf;
721  sampled = i;
722 
723  /* Rescale to reuse for direction sample, to better preserve stratification. */
724  *randu = (r - partial_sum) / sc->sample_weight;
725  break;
726  }
727  }
728 
729  partial_sum = next_sum;
730  }
731  }
732  }
733 
734  const ShaderClosure *sc = &sd->closure[sampled];
735  return CLOSURE_IS_BSSRDF(sc->type) ? sc : NULL;
736 }
737 
739  ShaderData *sd,
740  float randu,
741  float randv,
743  float3 *omega_in,
744  differential3 *domega_in,
745  float *pdf)
746 {
748 
749  const ShaderClosure *sc = shader_bsdf_pick(sd, &randu);
750  if (sc == NULL) {
751  *pdf = 0.0f;
752  return LABEL_NONE;
753  }
754 
755  /* BSSRDF should already have been handled elsewhere. */
756  kernel_assert(CLOSURE_IS_BSDF(sc->type));
757 
758  int label;
759  float3 eval = zero_float3();
760 
761  *pdf = 0.0f;
762  label = bsdf_sample(kg, sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
763 
764  if (*pdf != 0.0f) {
765  bsdf_eval_init(bsdf_eval, sc->type, eval * sc->weight, kernel_data.film.use_light_pass);
766 
767  if (sd->num_closure > 1) {
768  float sweight = sc->sample_weight;
769  _shader_bsdf_multi_eval(kg, sd, *omega_in, pdf, sc, bsdf_eval, *pdf * sweight, sweight);
770  }
771  }
772 
773  return label;
774 }
775 
777  ShaderData *sd,
778  const ShaderClosure *sc,
779  float randu,
780  float randv,
782  float3 *omega_in,
783  differential3 *domega_in,
784  float *pdf)
785 {
787 
788  int label;
789  float3 eval = zero_float3();
790 
791  *pdf = 0.0f;
792  label = bsdf_sample(kg, sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
793 
794  if (*pdf != 0.0f)
795  bsdf_eval_init(bsdf_eval, sc->type, eval * sc->weight, kernel_data.film.use_light_pass);
796 
797  return label;
798 }
799 
801 {
802  float roughness = 0.0f;
803  float sum_weight = 0.0f;
804 
805  for (int i = 0; i < sd->num_closure; i++) {
806  ShaderClosure *sc = &sd->closure[i];
807 
808  if (CLOSURE_IS_BSDF(sc->type)) {
809  /* sqrt once to undo the squaring from multiplying roughness on the
810  * two axes, and once for the squared roughness convention. */
811  float weight = fabsf(average(sc->weight));
813  sum_weight += weight;
814  }
815  }
816 
817  return (sum_weight > 0.0f) ? roughness / sum_weight : 0.0f;
818 }
819 
820 ccl_device void shader_bsdf_blur(KernelGlobals *kg, ShaderData *sd, float roughness)
821 {
822  for (int i = 0; i < sd->num_closure; i++) {
823  ShaderClosure *sc = &sd->closure[i];
824 
825  if (CLOSURE_IS_BSDF(sc->type))
826  bsdf_blur(kg, sc, roughness);
827  }
828 }
829 
831 {
832  if (sd->flag & SD_HAS_ONLY_VOLUME) {
833  return one_float3();
834  }
835  else if (sd->flag & SD_TRANSPARENT) {
836  return sd->closure_transparent_extinction;
837  }
838  else {
839  return zero_float3();
840  }
841 }
842 
844 {
845  if (sd->flag & SD_TRANSPARENT) {
846  for (int i = 0; i < sd->num_closure; i++) {
847  ShaderClosure *sc = &sd->closure[i];
848 
849  if (sc->type == CLOSURE_BSDF_TRANSPARENT_ID) {
850  sc->sample_weight = 0.0f;
851  sc->weight = zero_float3();
852  }
853  }
854 
855  sd->flag &= ~SD_TRANSPARENT;
856  }
857 }
858 
860 {
862 
863  alpha = max(alpha, zero_float3());
864  alpha = min(alpha, one_float3());
865 
866  return alpha;
867 }
868 
870 {
871  float3 eval = zero_float3();
872 
873  for (int i = 0; i < sd->num_closure; i++) {
874  ShaderClosure *sc = &sd->closure[i];
875 
876  if (CLOSURE_IS_BSDF_DIFFUSE(sc->type) || CLOSURE_IS_BSSRDF(sc->type) ||
877  CLOSURE_IS_BSDF_BSSRDF(sc->type))
878  eval += sc->weight;
879  }
880 
881  return eval;
882 }
883 
885 {
886  float3 eval = zero_float3();
887 
888  for (int i = 0; i < sd->num_closure; i++) {
889  ShaderClosure *sc = &sd->closure[i];
890 
891  if (CLOSURE_IS_BSDF_GLOSSY(sc->type))
892  eval += sc->weight;
893  }
894 
895  return eval;
896 }
897 
899 {
900  float3 eval = zero_float3();
901 
902  for (int i = 0; i < sd->num_closure; i++) {
903  ShaderClosure *sc = &sd->closure[i];
904 
905  if (CLOSURE_IS_BSDF_TRANSMISSION(sc->type))
906  eval += sc->weight;
907  }
908 
909  return eval;
910 }
911 
913 {
914  float3 N = zero_float3();
915 
916  for (int i = 0; i < sd->num_closure; i++) {
917  ShaderClosure *sc = &sd->closure[i];
918  if (CLOSURE_IS_BSDF_OR_BSSRDF(sc->type))
919  N += sc->N * fabsf(average(sc->weight));
920  }
921 
922  return (is_zero(N)) ? sd->N : normalize(N);
923 }
924 
925 ccl_device float3 shader_bsdf_ao(KernelGlobals *kg, ShaderData *sd, float ao_factor, float3 *N_)
926 {
927  float3 eval = zero_float3();
928  float3 N = zero_float3();
929 
930  for (int i = 0; i < sd->num_closure; i++) {
931  ShaderClosure *sc = &sd->closure[i];
932 
933  if (CLOSURE_IS_BSDF_DIFFUSE(sc->type)) {
934  const DiffuseBsdf *bsdf = (const DiffuseBsdf *)sc;
935  eval += sc->weight * ao_factor;
936  N += bsdf->N * fabsf(average(sc->weight));
937  }
938  }
939 
940  *N_ = (is_zero(N)) ? sd->N : normalize(N);
941  return eval;
942 }
943 
944 #ifdef __SUBSURFACE__
945 ccl_device float3 shader_bssrdf_sum(ShaderData *sd, float3 *N_, float *texture_blur_)
946 {
947  float3 eval = zero_float3();
948  float3 N = zero_float3();
949  float texture_blur = 0.0f, weight_sum = 0.0f;
950 
951  for (int i = 0; i < sd->num_closure; i++) {
952  ShaderClosure *sc = &sd->closure[i];
953 
954  if (CLOSURE_IS_BSSRDF(sc->type)) {
955  const Bssrdf *bssrdf = (const Bssrdf *)sc;
956  float avg_weight = fabsf(average(sc->weight));
957 
958  N += bssrdf->N * avg_weight;
959  eval += sc->weight;
960  texture_blur += bssrdf->texture_blur * avg_weight;
961  weight_sum += avg_weight;
962  }
963  }
964 
965  if (N_)
966  *N_ = (is_zero(N)) ? sd->N : normalize(N);
967 
968  if (texture_blur_)
969  *texture_blur_ = safe_divide(texture_blur, weight_sum);
970 
971  return eval;
972 }
973 #endif /* __SUBSURFACE__ */
974 
975 /* Constant emission optimization */
976 
977 ccl_device bool shader_constant_emission_eval(KernelGlobals *kg, int shader, float3 *eval)
978 {
979  int shader_index = shader & SHADER_MASK;
980  int shader_flag = kernel_tex_fetch(__shaders, shader_index).flags;
981 
982  if (shader_flag & SD_HAS_CONSTANT_EMISSION) {
983  *eval = make_float3(kernel_tex_fetch(__shaders, shader_index).constant_emission[0],
984  kernel_tex_fetch(__shaders, shader_index).constant_emission[1],
985  kernel_tex_fetch(__shaders, shader_index).constant_emission[2]);
986 
987  return true;
988  }
989 
990  return false;
991 }
992 
993 /* Background */
994 
996 {
997  if (sd->flag & SD_EMISSION) {
998  return sd->closure_emission_background;
999  }
1000  else {
1001  return zero_float3();
1002  }
1003 }
1004 
1005 /* Emission */
1006 
1008 {
1009  if (sd->flag & SD_EMISSION) {
1010  return emissive_simple_eval(sd->Ng, sd->I) * sd->closure_emission_background;
1011  }
1012  else {
1013  return zero_float3();
1014  }
1015 }
1016 
1017 /* Holdout */
1018 
1020 {
1021  float3 weight = zero_float3();
1022 
1023  /* For objects marked as holdout, preserve transparency and remove all other
1024  * closures, replacing them with a holdout weight. */
1025  if (sd->object_flag & SD_OBJECT_HOLDOUT_MASK) {
1026  if ((sd->flag & SD_TRANSPARENT) && !(sd->flag & SD_HAS_ONLY_VOLUME)) {
1027  weight = one_float3() - sd->closure_transparent_extinction;
1028 
1029  for (int i = 0; i < sd->num_closure; i++) {
1030  ShaderClosure *sc = &sd->closure[i];
1031  if (!CLOSURE_IS_BSDF_TRANSPARENT(sc->type)) {
1032  sc->type = NBUILTIN_CLOSURES;
1033  }
1034  }
1035 
1036  sd->flag &= ~(SD_CLOSURE_FLAGS - (SD_TRANSPARENT | SD_BSDF));
1037  }
1038  else {
1039  weight = one_float3();
1040  }
1041  }
1042  else {
1043  for (int i = 0; i < sd->num_closure; i++) {
1044  ShaderClosure *sc = &sd->closure[i];
1045  if (CLOSURE_IS_HOLDOUT(sc->type)) {
1046  weight += sc->weight;
1047  }
1048  }
1049  }
1050 
1051  return weight;
1052 }
1053 
1054 /* Surface Evaluation */
1055 
1056 ccl_device void shader_eval_surface(KernelGlobals *kg,
1057  ShaderData *sd,
1059  ccl_global float *buffer,
1060  int path_flag)
1061 {
1063 
1064  /* If path is being terminated, we are tracing a shadow ray or evaluating
1065  * emission, then we don't need to store closures. The emission and shadow
1066  * shader data also do not have a closure array to save GPU memory. */
1067  int max_closures;
1068  if (path_flag & (PATH_RAY_TERMINATE | PATH_RAY_SHADOW | PATH_RAY_EMISSION)) {
1069  max_closures = 0;
1070  }
1071  else {
1072  max_closures = kernel_data.integrator.max_closures;
1073  }
1074 
1075  sd->num_closure = 0;
1076  sd->num_closure_left = max_closures;
1077 
1078 #ifdef __OSL__
1079  if (kg->osl) {
1080  if (sd->object == OBJECT_NONE && sd->lamp == LAMP_NONE) {
1081  OSLShader::eval_background(kg, sd, state, path_flag);
1082  }
1083  else {
1084  OSLShader::eval_surface(kg, sd, state, path_flag);
1085  }
1086  }
1087  else
1088 #endif
1089  {
1090 #ifdef __SVM__
1091  svm_eval_nodes(kg, sd, state, buffer, SHADER_TYPE_SURFACE, path_flag);
1092 #else
1093  if (sd->object == OBJECT_NONE) {
1094  sd->closure_emission_background = make_float3(0.8f, 0.8f, 0.8f);
1095  sd->flag |= SD_EMISSION;
1096  }
1097  else {
1098  DiffuseBsdf *bsdf = (DiffuseBsdf *)bsdf_alloc(
1099  sd, sizeof(DiffuseBsdf), make_float3(0.8f, 0.8f, 0.8f));
1100  if (bsdf != NULL) {
1101  bsdf->N = sd->N;
1102  sd->flag |= bsdf_diffuse_setup(bsdf);
1103  }
1104  }
1105 #endif
1106  }
1107 
1108  if (sd->flag & SD_BSDF_NEEDS_LCG) {
1109  sd->lcg_state = lcg_state_init_addrspace(state, 0xb4bc3953);
1110  }
1111 }
1112 
1113 /* Volume */
1114 
1115 #ifdef __VOLUME__
1116 
1117 ccl_device_inline void _shader_volume_phase_multi_eval(const ShaderData *sd,
1118  const float3 omega_in,
1119  float *pdf,
1120  int skip_phase,
1121  BsdfEval *result_eval,
1122  float sum_pdf,
1123  float sum_sample_weight)
1124 {
1125  for (int i = 0; i < sd->num_closure; i++) {
1126  if (i == skip_phase)
1127  continue;
1128 
1129  const ShaderClosure *sc = &sd->closure[i];
1130 
1131  if (CLOSURE_IS_PHASE(sc->type)) {
1132  float phase_pdf = 0.0f;
1133  float3 eval = volume_phase_eval(sd, sc, omega_in, &phase_pdf);
1134 
1135  if (phase_pdf != 0.0f) {
1136  bsdf_eval_accum(result_eval, sc->type, eval, 1.0f);
1137  sum_pdf += phase_pdf * sc->sample_weight;
1138  }
1139 
1140  sum_sample_weight += sc->sample_weight;
1141  }
1142  }
1143 
1144  *pdf = (sum_sample_weight > 0.0f) ? sum_pdf / sum_sample_weight : 0.0f;
1145 }
1146 
1147 ccl_device void shader_volume_phase_eval(
1148  KernelGlobals *kg, const ShaderData *sd, const float3 omega_in, BsdfEval *eval, float *pdf)
1149 {
1151 
1152  bsdf_eval_init(eval, NBUILTIN_CLOSURES, zero_float3(), kernel_data.film.use_light_pass);
1153 
1154  _shader_volume_phase_multi_eval(sd, omega_in, pdf, -1, eval, 0.0f, 0.0f);
1155 }
1156 
1157 ccl_device int shader_volume_phase_sample(KernelGlobals *kg,
1158  const ShaderData *sd,
1159  float randu,
1160  float randv,
1161  BsdfEval *phase_eval,
1162  float3 *omega_in,
1163  differential3 *domega_in,
1164  float *pdf)
1165 {
1167 
1168  int sampled = 0;
1169 
1170  if (sd->num_closure > 1) {
1171  /* pick a phase closure based on sample weights */
1172  float sum = 0.0f;
1173 
1174  for (sampled = 0; sampled < sd->num_closure; sampled++) {
1175  const ShaderClosure *sc = &sd->closure[sampled];
1176 
1177  if (CLOSURE_IS_PHASE(sc->type))
1178  sum += sc->sample_weight;
1179  }
1180 
1181  float r = randu * sum;
1182  float partial_sum = 0.0f;
1183 
1184  for (sampled = 0; sampled < sd->num_closure; sampled++) {
1185  const ShaderClosure *sc = &sd->closure[sampled];
1186 
1187  if (CLOSURE_IS_PHASE(sc->type)) {
1188  float next_sum = partial_sum + sc->sample_weight;
1189 
1190  if (r <= next_sum) {
1191  /* Rescale to reuse for BSDF direction sample. */
1192  randu = (r - partial_sum) / sc->sample_weight;
1193  break;
1194  }
1195 
1196  partial_sum = next_sum;
1197  }
1198  }
1199 
1200  if (sampled == sd->num_closure) {
1201  *pdf = 0.0f;
1202  return LABEL_NONE;
1203  }
1204  }
1205 
1206  /* todo: this isn't quite correct, we don't weight anisotropy properly
1207  * depending on color channels, even if this is perhaps not a common case */
1208  const ShaderClosure *sc = &sd->closure[sampled];
1209  int label;
1210  float3 eval = zero_float3();
1211 
1212  *pdf = 0.0f;
1213  label = volume_phase_sample(sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
1214 
1215  if (*pdf != 0.0f) {
1216  bsdf_eval_init(phase_eval, sc->type, eval, kernel_data.film.use_light_pass);
1217  }
1218 
1219  return label;
1220 }
1221 
1222 ccl_device int shader_phase_sample_closure(KernelGlobals *kg,
1223  const ShaderData *sd,
1224  const ShaderClosure *sc,
1225  float randu,
1226  float randv,
1227  BsdfEval *phase_eval,
1228  float3 *omega_in,
1229  differential3 *domega_in,
1230  float *pdf)
1231 {
1233 
1234  int label;
1235  float3 eval = zero_float3();
1236 
1237  *pdf = 0.0f;
1238  label = volume_phase_sample(sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
1239 
1240  if (*pdf != 0.0f)
1241  bsdf_eval_init(phase_eval, sc->type, eval, kernel_data.film.use_light_pass);
1242 
1243  return label;
1244 }
1245 
1246 /* Volume Evaluation */
1247 
1248 ccl_device_inline void shader_eval_volume(KernelGlobals *kg,
1249  ShaderData *sd,
1251  ccl_addr_space VolumeStack *stack,
1252  int path_flag)
1253 {
1254  /* If path is being terminated, we are tracing a shadow ray or evaluating
1255  * emission, then we don't need to store closures. The emission and shadow
1256  * shader data also do not have a closure array to save GPU memory. */
1257  int max_closures;
1258  if (path_flag & (PATH_RAY_TERMINATE | PATH_RAY_SHADOW | PATH_RAY_EMISSION)) {
1259  max_closures = 0;
1260  }
1261  else {
1262  max_closures = kernel_data.integrator.max_closures;
1263  }
1264 
1265  /* reset closures once at the start, we will be accumulating the closures
1266  * for all volumes in the stack into a single array of closures */
1267  sd->num_closure = 0;
1268  sd->num_closure_left = max_closures;
1269  sd->flag = 0;
1270  sd->object_flag = 0;
1271 
1272  for (int i = 0; stack[i].shader != SHADER_NONE; i++) {
1273  /* setup shaderdata from stack. it's mostly setup already in
1274  * shader_setup_from_volume, this switching should be quick */
1275  sd->object = stack[i].object;
1276  sd->lamp = LAMP_NONE;
1277  sd->shader = stack[i].shader;
1278 
1279  sd->flag &= ~SD_SHADER_FLAGS;
1280  sd->flag |= kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).flags;
1281  sd->object_flag &= ~SD_OBJECT_FLAGS;
1282 
1283  if (sd->object != OBJECT_NONE) {
1284  sd->object_flag |= kernel_tex_fetch(__object_flag, sd->object);
1285 
1286 # ifdef __OBJECT_MOTION__
1287  /* todo: this is inefficient for motion blur, we should be
1288  * caching matrices instead of recomputing them each step */
1289  shader_setup_object_transforms(kg, sd, sd->time);
1290 # endif
1291  }
1292 
1293  /* evaluate shader */
1294 # ifdef __SVM__
1295 # ifdef __OSL__
1296  if (kg->osl) {
1297  OSLShader::eval_volume(kg, sd, state, path_flag);
1298  }
1299  else
1300 # endif
1301  {
1302  svm_eval_nodes(kg, sd, state, NULL, SHADER_TYPE_VOLUME, path_flag);
1303  }
1304 # endif
1305 
1306  /* merge closures to avoid exceeding number of closures limit */
1307  if (i > 0)
1308  shader_merge_closures(sd);
1309  }
1310 }
1311 
1312 #endif /* __VOLUME__ */
1313 
1314 /* Displacement Evaluation */
1315 
1317  ShaderData *sd,
1319 {
1320  sd->num_closure = 0;
1321  sd->num_closure_left = 0;
1322 
1323  /* this will modify sd->P */
1324 #ifdef __SVM__
1325 # ifdef __OSL__
1326  if (kg->osl)
1328  else
1329 # endif
1330  {
1332  }
1333 #endif
1334 }
1335 
1336 /* Transparent Shadows */
1337 
1338 #ifdef __TRANSPARENT_SHADOWS__
1339 ccl_device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect)
1340 {
1341  int prim = kernel_tex_fetch(__prim_index, isect->prim);
1342  int shader = 0;
1343 
1344 # ifdef __HAIR__
1345  if (isect->type & PRIMITIVE_ALL_TRIANGLE) {
1346 # endif
1347  shader = kernel_tex_fetch(__tri_shader, prim);
1348 # ifdef __HAIR__
1349  }
1350  else {
1351  float4 str = kernel_tex_fetch(__curves, prim);
1352  shader = __float_as_int(str.z);
1353  }
1354 # endif
1355  int flag = kernel_tex_fetch(__shaders, (shader & SHADER_MASK)).flags;
1356 
1357  return (flag & SD_HAS_TRANSPARENT_SHADOW) != 0;
1358 }
1359 #endif /* __TRANSPARENT_SHADOWS__ */
1360 
1361 ccl_device float shader_cryptomatte_id(KernelGlobals *kg, int shader)
1362 {
1363  return kernel_tex_fetch(__shaders, (shader & SHADER_MASK)).cryptomatte_id;
1364 }
1365 
MINLINE float safe_sqrtf(float a)
MINLINE float safe_divide(float a, float b)
#define N_(msgid)
_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 const void *pointer _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_BOOL GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLushort pattern _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble u2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLdouble GLdouble v2 _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLdouble GLdouble nz _GL_VOID_RET _GL_VOID GLfloat GLfloat nz _GL_VOID_RET _GL_VOID GLint GLint nz _GL_VOID_RET _GL_VOID GLshort GLshort nz _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const GLfloat *values _GL_VOID_RET _GL_VOID GLsizei const GLushort *values _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID const GLuint const GLclampf *priorities _GL_VOID_RET _GL_VOID GLdouble y _GL_VOID_RET _GL_VOID GLfloat y _GL_VOID_RET _GL_VOID GLint y _GL_VOID_RET _GL_VOID GLshort y _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLfloat GLfloat z _GL_VOID_RET _GL_VOID GLint GLint z _GL_VOID_RET _GL_VOID GLshort GLshort z _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble w _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat w _GL_VOID_RET _GL_VOID GLint GLint GLint w _GL_VOID_RET _GL_VOID GLshort GLshort GLshort w _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble y2 _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat y2 _GL_VOID_RET _GL_VOID GLint GLint GLint y2 _GL_VOID_RET _GL_VOID GLshort GLshort GLshort y2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLuint *buffer _GL_VOID_RET _GL_VOID GLdouble t _GL_VOID_RET _GL_VOID GLfloat t _GL_VOID_RET _GL_VOID GLint t _GL_VOID_RET _GL_VOID GLshort t _GL_VOID_RET _GL_VOID GLdouble GLdouble r _GL_VOID_RET _GL_VOID GLfloat GLfloat r _GL_VOID_RET _GL_VOID GLint GLint r _GL_VOID_RET _GL_VOID GLshort GLshort r _GL_VOID_RET _GL_VOID GLdouble GLdouble r
_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 const void *pointer _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_BOOL GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLushort pattern _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble u2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLdouble GLdouble v2 _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLdouble GLdouble nz _GL_VOID_RET _GL_VOID GLfloat GLfloat nz _GL_VOID_RET _GL_VOID GLint GLint nz _GL_VOID_RET _GL_VOID GLshort GLshort nz _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const GLfloat *values _GL_VOID_RET _GL_VOID GLsizei const GLushort *values _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID const GLuint const GLclampf *priorities _GL_VOID_RET _GL_VOID GLdouble y _GL_VOID_RET _GL_VOID GLfloat y _GL_VOID_RET _GL_VOID GLint y _GL_VOID_RET _GL_VOID GLshort y _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLfloat GLfloat z _GL_VOID_RET _GL_VOID GLint GLint z _GL_VOID_RET _GL_VOID GLshort GLshort z _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble w _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat w _GL_VOID_RET _GL_VOID GLint GLint GLint w _GL_VOID_RET _GL_VOID GLshort GLshort GLshort w _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble y2 _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat y2 _GL_VOID_RET _GL_VOID GLint GLint GLint y2 _GL_VOID_RET _GL_VOID GLshort GLshort GLshort y2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLuint *buffer _GL_VOID_RET _GL_VOID GLdouble t _GL_VOID_RET _GL_VOID GLfloat t _GL_VOID_RET _GL_VOID GLint t _GL_VOID_RET _GL_VOID GLshort t _GL_VOID_RET _GL_VOID GLdouble t
ccl_device_inline ShaderClosure * bsdf_alloc(ShaderData *sd, int size, float3 weight)
Definition: alloc.h:58
ATTR_WARN_UNUSED_RESULT const BMVert * v
ccl_device_inline int bsdf_sample(KernelGlobals *kg, ShaderData *sd, const ShaderClosure *sc, float randu, float randv, float3 *eval, float3 *omega_in, differential3 *domega_in, float *pdf)
Definition: bsdf.h:112
ccl_device bool bsdf_merge(ShaderClosure *a, ShaderClosure *b)
Definition: bsdf.h:721
ccl_device_inline float bsdf_get_roughness_squared(const ShaderClosure *sc)
Definition: bsdf.h:56
ccl_device void bsdf_blur(KernelGlobals *kg, ShaderClosure *sc, float roughness)
Definition: bsdf.h:688
ccl_device float3 bsdf_eval(KernelGlobals *kg, ShaderData *sd, const ShaderClosure *sc, const float3 omega_in, float *pdf)
Definition: bsdf.h:485
ccl_device int bsdf_diffuse_setup(DiffuseBsdf *bsdf)
Definition: bsdf_diffuse.h:46
static DBVT_INLINE btScalar size(const btDbvtVolume &a)
Definition: btDbvt.cpp:52
static T sum(const btAlignedObjectArray< T > &items)
double time
const char * label
Light lamp
static CCL_NAMESPACE_BEGIN const double alpha
ccl_device float3 emissive_simple_eval(const float3 Ng, const float3 I)
Definition: emissive.h:76
#define str(s)
CCL_NAMESPACE_BEGIN ccl_device_noinline void motion_triangle_shader_setup(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray, bool is_local)
@ OBJECT_INVERSE_TRANSFORM
Definition: geom_object.h:31
@ OBJECT_TRANSFORM
Definition: geom_object.h:30
#define object_normal_transform_auto
Definition: geom_object.h:578
#define object_position_transform_auto
Definition: geom_object.h:576
#define object_dir_transform_auto
Definition: geom_object.h:577
ccl_device_inline Transform object_fetch_transform(KernelGlobals *kg, int object, enum ObjectTransform type)
Definition: geom_object.h:38
ccl_device_inline Transform lamp_fetch_transform(KernelGlobals *kg, int lamp, bool inverse)
Definition: geom_object.h:52
ccl_device_inline void triangle_dPdudv(KernelGlobals *kg, int prim, ccl_addr_space float3 *dPdu, ccl_addr_space float3 *dPdv)
Definition: geom_triangle.h:96
CCL_NAMESPACE_BEGIN ccl_device_inline float3 triangle_normal(KernelGlobals *kg, ShaderData *sd)
Definition: geom_triangle.h:26
ccl_device_inline void triangle_point_normal(KernelGlobals *kg, int object, int prim, float u, float v, float3 *P, float3 *Ng, int *shader)
Definition: geom_triangle.h:44
ccl_device_inline float3 triangle_smooth_normal(KernelGlobals *kg, float3 Ng, int prim, float u, float v)
Definition: geom_triangle.h:81
ccl_device_inline float3 triangle_refine_local(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray)
ccl_device_inline float3 triangle_refine(KernelGlobals *kg, ShaderData *sd, const Intersection *isect, const Ray *ray)
ccl_device int volume_phase_sample(const ShaderData *sd, const ShaderClosure *sc, float randu, float randv, float3 *eval, float3 *omega_in, differential3 *domega_in, float *pdf)
ccl_device float3 volume_phase_eval(const ShaderData *sd, const ShaderClosure *sc, float3 omega_in, float *pdf)
CCL_NAMESPACE_END CCL_NAMESPACE_BEGIN ccl_device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ccl_addr_space PathState *state, ccl_global float *buffer, ShaderType type, int path_flag)
ccl_device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 value, int use_light_pass)
ccl_device_inline void bsdf_eval_mis(BsdfEval *eval, float value)
ccl_device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 value, float mis_weight)
#define kernel_data
#define kernel_assert(cond)
#define ccl_addr_space
#define kernel_tex_fetch(tex, index)
#define ccl_device
#define ccl_device_inline
#define ccl_global
#define ccl_device_noinline
#define CCL_NAMESPACE_END
#define fabsf(x)
#define sqrtf(x)
#define make_float3(x, y, z)
void KERNEL_FUNCTION_FULL_NAME() shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int filter, int i, int offset, int sample)
ccl_device void differential_incoming(ccl_addr_space differential3 *dI, const differential3 dD)
ccl_device void differential_dudv(ccl_addr_space differential *du, ccl_addr_space differential *dv, float3 dPdu, float3 dPdv, differential3 dP, float3 Ng)
ccl_device differential differential_zero()
ccl_device differential3 differential3_zero()
CCL_NAMESPACE_BEGIN ccl_device void differential_transfer(ccl_addr_space differential3 *dP_, const differential3 dP, float3 D, const differential3 dD, float3 Ng, float t)
ccl_device float power_heuristic(float a, float b)
#define PROFILING_OBJECT(object)
#define PROFILING_INIT(kg, event)
#define PROFILING_SHADER(shader)
ccl_device_inline uint lcg_state_init_addrspace(ccl_addr_space PathState *state, uint scramble)
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_inline const ShaderClosure * shader_bssrdf_pick(ShaderData *sd, ccl_addr_space float3 *throughput, float *randu)
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 void shader_setup_from_displace(KernelGlobals *kg, ShaderData *sd, int object, int prim, float u, float v)
ccl_device_inline void _shader_bsdf_multi_eval(KernelGlobals *kg, ShaderData *sd, const float3 omega_in, float *pdf, const ShaderClosure *skip_sc, BsdfEval *result_eval, float sum_pdf, float sum_sample_weight)
ccl_device float3 shader_bsdf_ao(KernelGlobals *kg, ShaderData *sd, float ao_factor, float3 *N_)
ccl_device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd)
ccl_device_inline void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd, const float3 P, const float3 Ng, const float3 I, int shader, int object, int prim, float u, float v, float t, float time, bool object_space, int lamp)
ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, const ShaderData *sd)
ccl_device float3 shader_holdout_apply(KernelGlobals *kg, ShaderData *sd)
ccl_device void shader_bsdf_eval(KernelGlobals *kg, ShaderData *sd, const float3 omega_in, BsdfEval *eval, float light_pdf, bool use_mis)
ccl_device float shader_bsdf_average_roughness(ShaderData *sd)
ccl_device float3 shader_background_eval(ShaderData *sd)
ccl_device void shader_bsdf_disable_transparency(KernelGlobals *kg, ShaderData *sd)
ccl_device_inline const ShaderClosure * shader_bsdf_pick(ShaderData *sd, float *randu)
ccl_device_inline void shader_prepare_closures(ShaderData *sd, ccl_addr_space PathState *state)
ccl_device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData *sd, const Ray *ray)
ccl_device float shader_cryptomatte_id(KernelGlobals *kg, int shader)
ccl_device int shader_bsdf_sample_closure(KernelGlobals *kg, ShaderData *sd, const ShaderClosure *sc, float randu, float randv, BsdfEval *bsdf_eval, float3 *omega_in, differential3 *domega_in, float *pdf)
ccl_device bool shader_constant_emission_eval(KernelGlobals *kg, int shader, float3 *eval)
ccl_device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_emissive_eval(ShaderData *sd)
ccl_device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_bsdf_average_normal(KernelGlobals *kg, ShaderData *sd)
ccl_device void shader_eval_displacement(KernelGlobals *kg, ShaderData *sd, ccl_addr_space PathState *state)
ccl_device_inline int shader_bsdf_sample(KernelGlobals *kg, ShaderData *sd, float randu, float randv, BsdfEval *bsdf_eval, float3 *omega_in, differential3 *domega_in, float *pdf)
__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_CLOSURE_FLAGS
Definition: kernel_types.h:863
@ SD_HAS_TRANSPARENT_SHADOW
Definition: kernel_types.h:871
@ SD_SHADER_FLAGS
Definition: kernel_types.h:895
@ SD_HAS_CONSTANT_EMISSION
Definition: kernel_types.h:891
@ SD_BACKFACING
Definition: kernel_types.h:843
@ SD_BSDF_NEEDS_LCG
Definition: kernel_types.h:861
@ SD_HAS_ONLY_VOLUME
Definition: kernel_types.h:875
@ SD_BSDF
Definition: kernel_types.h:847
@ SD_TRANSPARENT
Definition: kernel_types.h:859
@ SD_EMISSION
Definition: kernel_types.h:845
@ PRIMITIVE_LAMP
Definition: kernel_types.h:695
@ PRIMITIVE_NONE
Definition: kernel_types.h:685
@ PRIMITIVE_ALL_TRIANGLE
Definition: kernel_types.h:697
@ PRIMITIVE_ALL_CURVE
Definition: kernel_types.h:698
@ PRIMITIVE_TRIANGLE
Definition: kernel_types.h:686
#define SHADER_NONE
Definition: kernel_types.h:58
#define PRIM_NONE
Definition: kernel_types.h:60
@ PATH_RAY_SHADOW
Definition: kernel_types.h:284
@ PATH_RAY_TERMINATE
Definition: kernel_types.h:319
@ PATH_RAY_EMISSION
Definition: kernel_types.h:321
#define OBJECT_NONE
Definition: kernel_types.h:59
ShaderData
@ SHADER_SMOOTH_NORMAL
Definition: kernel_types.h:581
@ SHADER_MASK
Definition: kernel_types.h:593
@ SD_OBJECT_MOTION
Definition: kernel_types.h:906
@ SD_OBJECT_HOLDOUT_MASK
Definition: kernel_types.h:904
@ SD_OBJECT_FLAGS
Definition: kernel_types.h:922
@ SD_OBJECT_TRANSFORM_APPLIED
Definition: kernel_types.h:908
@ LABEL_NONE
Definition: kernel_types.h:327
ShaderClosure
Definition: kernel_types.h:831
#define LAMP_NONE
Definition: kernel_types.h:61
static float P(float k)
Definition: math_interp.c:41
static ulong state[N]
static const pxr::TfToken roughness("roughness", pxr::TfToken::Immortal)
#define I
params N
#define min(a, b)
Definition: sort.c:51
closure color bssrdf(string method, normal N, vector radius, color albedo) BUILTIN
Definition: bssrdf.h:22
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
static void eval_displacement(SubdivDisplacement *displacement, const int ptex_face_index, const float u, const float v, const float dPdu[3], const float dPdv[3], float r_D[3])
#define CLOSURE_IS_HOLDOUT(type)
Definition: svm_types.h:631
#define CLOSURE_IS_BSDF_TRANSPARENT(type)
Definition: svm_types.h:608
#define CLOSURE_IS_BSDF_TRANSMISSION(type)
Definition: svm_types.h:601
#define CLOSURE_IS_BSDF(type)
Definition: svm_types.h:595
@ SHADER_TYPE_SURFACE
Definition: svm_types.h:512
@ SHADER_TYPE_VOLUME
Definition: svm_types.h:513
@ SHADER_TYPE_DISPLACEMENT
Definition: svm_types.h:514
#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
#define CLOSURE_IS_BSDF_OR_BSSRDF(type)
Definition: svm_types.h:622
@ CLOSURE_BSDF_TRANSPARENT_ID
Definition: svm_types.h:571
@ NBUILTIN_CLOSURES
Definition: svm_types.h:591
#define CLOSURE_IS_BSSRDF(type)
Definition: svm_types.h:623
#define CLOSURE_IS_PHASE(type)
Definition: svm_types.h:632
float max
ccl_device_inline int __float_as_int(float f)
Definition: util_math.h:202
ccl_device_inline float2 normalize(const float2 &a)
ccl_device_inline float dot(const float2 &a, const float2 &b)
ccl_device_inline bool is_zero(const float2 &a)
ccl_device_inline float average(const float2 &a)
ccl_device_inline float3 one_float3()
ccl_device_inline float3 zero_float3()
@ PROFILING_CLOSURE_EVAL
@ PROFILING_SHADER_SETUP
@ PROFILING_SHADER_EVAL
@ PROFILING_CLOSURE_SAMPLE
@ PROFILING_CLOSURE_VOLUME_SAMPLE
@ PROFILING_CLOSURE_VOLUME_EVAL
ccl_device_inline Transform transform_quick_inverse(Transform M)