Blender  V2.93
kernel_bake.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 __BAKING__
20 
21 ccl_device_noinline void compute_light_pass(
22  KernelGlobals *kg, ShaderData *sd, PathRadiance *L, uint rng_hash, int pass_filter, int sample)
23 {
24  kernel_assert(kernel_data.film.use_light_pass);
25 
26  float3 throughput = one_float3();
27 
28  /* Emission and indirect shader data memory used by various functions. */
29  ShaderDataTinyStorage emission_sd_storage;
30  ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
31  ShaderData indirect_sd;
32 
33  /* Init radiance. */
35 
36  /* Init path state. */
38  path_state_init(kg, emission_sd, &state, rng_hash, sample, NULL);
39 
40  /* Evaluate surface shader. */
41  shader_eval_surface(kg, sd, &state, NULL, state.flag);
42 
43  /* TODO, disable more closures we don't need besides transparent */
45 
46  /* Init ray. */
47  Ray ray;
48  ray.P = sd->P + sd->Ng;
49  ray.D = -sd->Ng;
50  ray.t = FLT_MAX;
51 # ifdef __CAMERA_MOTION__
52  ray.time = 0.5f;
53 # endif
54 
55 # ifdef __BRANCHED_PATH__
56  if (!kernel_data.integrator.branched) {
57  /* regular path tracer */
58 # endif
59 
60  /* sample ambient occlusion */
61  if (pass_filter & BAKE_FILTER_AO) {
62  kernel_path_ao(kg, sd, emission_sd, L, &state, throughput, shader_bsdf_alpha(kg, sd));
63  }
64 
65  /* sample emission */
66  if ((pass_filter & BAKE_FILTER_EMISSION) && (sd->flag & SD_EMISSION)) {
67  float3 emission = indirect_primitive_emission(kg, sd, 0.0f, state.flag, state.ray_pdf);
68  path_radiance_accum_emission(kg, L, &state, throughput, emission);
69  }
70 
71  bool is_sss_sample = false;
72 
73 # ifdef __SUBSURFACE__
74  /* sample subsurface scattering */
75  if ((pass_filter & BAKE_FILTER_DIFFUSE) && (sd->flag & SD_BSSRDF)) {
76  /* When mixing BSSRDF and BSDF closures we should skip BSDF lighting
77  * if scattering was successful. */
78  SubsurfaceIndirectRays ss_indirect;
79  kernel_path_subsurface_init_indirect(&ss_indirect);
80  if (kernel_path_subsurface_scatter(
81  kg, sd, emission_sd, L, &state, &ray, &throughput, &ss_indirect)) {
82  while (ss_indirect.num_rays) {
83  kernel_path_subsurface_setup_indirect(kg, &ss_indirect, &state, &ray, L, &throughput);
84  kernel_path_indirect(kg, &indirect_sd, emission_sd, &ray, throughput, &state, L);
85  }
86  is_sss_sample = true;
87  }
88  }
89 # endif
90 
91  /* sample light and BSDF */
92  if (!is_sss_sample && (pass_filter & (BAKE_FILTER_DIRECT | BAKE_FILTER_INDIRECT))) {
93  kernel_path_surface_connect_light(kg, sd, emission_sd, throughput, &state, L);
94 
95  if (kernel_path_surface_bounce(kg, sd, &throughput, &state, &L->state, &ray)) {
96 # ifdef __LAMP_MIS__
97  state.ray_t = 0.0f;
98 # endif
99  /* compute indirect light */
100  kernel_path_indirect(kg, &indirect_sd, emission_sd, &ray, throughput, &state, L);
101 
102  /* sum and reset indirect light pass variables for the next samples */
105  }
106  }
107 # ifdef __BRANCHED_PATH__
108  }
109  else {
110  /* branched path tracer */
111 
112  /* sample ambient occlusion */
113  if (pass_filter & BAKE_FILTER_AO) {
114  kernel_branched_path_ao(kg, sd, emission_sd, L, &state, throughput);
115  }
116 
117  /* sample emission */
118  if ((pass_filter & BAKE_FILTER_EMISSION) && (sd->flag & SD_EMISSION)) {
119  float3 emission = indirect_primitive_emission(kg, sd, 0.0f, state.flag, state.ray_pdf);
120  path_radiance_accum_emission(kg, L, &state, throughput, emission);
121  }
122 
123 # ifdef __SUBSURFACE__
124  /* sample subsurface scattering */
125  if ((pass_filter & BAKE_FILTER_DIFFUSE) && (sd->flag & SD_BSSRDF)) {
126  /* When mixing BSSRDF and BSDF closures we should skip BSDF lighting
127  * if scattering was successful. */
128  kernel_branched_path_subsurface_scatter(
129  kg, sd, &indirect_sd, emission_sd, L, &state, &ray, throughput);
130  }
131 # endif
132 
133  /* sample light and BSDF */
134  if (pass_filter & (BAKE_FILTER_DIRECT | BAKE_FILTER_INDIRECT)) {
135 # if defined(__EMISSION__)
136  /* direct light */
137  if (kernel_data.integrator.use_direct_light) {
138  int all = kernel_data.integrator.sample_all_lights_direct;
139  kernel_branched_path_surface_connect_light(
140  kg, sd, emission_sd, &state, throughput, 1.0f, L, all);
141  }
142 # endif
143 
144  /* indirect light */
145  kernel_branched_path_surface_indirect_light(
146  kg, sd, &indirect_sd, emission_sd, throughput, 1.0f, &state, L);
147  }
148  }
149 # endif
150 }
151 
152 /* this helps with AA but it's not the real solution as it does not AA the geometry
153  * but it's better than nothing, thus committed */
154 ccl_device_inline float bake_clamp_mirror_repeat(float u, float max)
155 {
156  /* use mirror repeat (like opengl texture) so that if the barycentric
157  * coordinate goes past the end of the triangle it is not always clamped
158  * to the same value, gives ugly patterns */
159  u /= max;
160  float fu = floorf(u);
161  u = u - fu;
162 
163  return ((((int)fu) & 1) ? 1.0f - u : u) * max;
164 }
165 
166 ccl_device_inline float3 kernel_bake_shader_bsdf(KernelGlobals *kg,
167  ShaderData *sd,
168  const ShaderEvalType type)
169 {
170  switch (type) {
171  case SHADER_EVAL_DIFFUSE:
172  return shader_bsdf_diffuse(kg, sd);
173  case SHADER_EVAL_GLOSSY:
174  return shader_bsdf_glossy(kg, sd);
176  return shader_bsdf_transmission(kg, sd);
177  default:
178  kernel_assert(!"Unknown bake type passed to BSDF evaluate");
179  return zero_float3();
180  }
181 }
182 
183 ccl_device float3 kernel_bake_evaluate_direct_indirect(KernelGlobals *kg,
184  ShaderData *sd,
185  PathState *state,
186  float3 direct,
187  float3 indirect,
188  const ShaderEvalType type,
189  const int pass_filter)
190 {
191  float3 color;
192  const bool is_color = (pass_filter & BAKE_FILTER_COLOR) != 0;
193  const bool is_direct = (pass_filter & BAKE_FILTER_DIRECT) != 0;
194  const bool is_indirect = (pass_filter & BAKE_FILTER_INDIRECT) != 0;
195  float3 out = zero_float3();
196 
197  if (is_color) {
198  if (is_direct || is_indirect) {
199  /* Leave direct and diffuse channel colored. */
200  color = one_float3();
201  }
202  else {
203  /* surface color of the pass only */
204  shader_eval_surface(kg, sd, state, NULL, 0);
205  return kernel_bake_shader_bsdf(kg, sd, type);
206  }
207  }
208  else {
209  shader_eval_surface(kg, sd, state, NULL, 0);
210  color = kernel_bake_shader_bsdf(kg, sd, type);
211  }
212 
213  if (is_direct) {
214  out += safe_divide_even_color(direct, color);
215  }
216 
217  if (is_indirect) {
218  out += safe_divide_even_color(indirect, color);
219  }
220 
221  return out;
222 }
223 
224 ccl_device void kernel_bake_evaluate(
225  KernelGlobals *kg, ccl_global float *buffer, int sample, int x, int y, int offset, int stride)
226 {
227  /* Setup render buffers. */
228  const int index = offset + x + y * stride;
229  const int pass_stride = kernel_data.film.pass_stride;
230  buffer += index * pass_stride;
231 
232  ccl_global float *primitive = buffer + kernel_data.film.pass_bake_primitive;
233  ccl_global float *differential = buffer + kernel_data.film.pass_bake_differential;
234  ccl_global float *output = buffer + kernel_data.film.pass_combined;
235 
236  int seed = __float_as_uint(primitive[0]);
237  int prim = __float_as_uint(primitive[1]);
238  if (prim == -1)
239  return;
240 
241  prim += kernel_data.bake.tri_offset;
242 
243  /* Random number generator. */
244  uint rng_hash = hash_uint(seed) ^ kernel_data.integrator.seed;
245  int num_samples = kernel_data.integrator.aa_samples;
246 
247  float filter_x, filter_y;
248  if (sample == 0) {
249  filter_x = filter_y = 0.5f;
250  }
251  else {
252  path_rng_2D(kg, rng_hash, sample, num_samples, PRNG_FILTER_U, &filter_x, &filter_y);
253  }
254 
255  /* Barycentric UV with sub-pixel offset. */
256  float u = primitive[2];
257  float v = primitive[3];
258 
259  float dudx = differential[0];
260  float dudy = differential[1];
261  float dvdx = differential[2];
262  float dvdy = differential[3];
263 
264  if (sample > 0) {
265  u = bake_clamp_mirror_repeat(u + dudx * (filter_x - 0.5f) + dudy * (filter_y - 0.5f), 1.0f);
266  v = bake_clamp_mirror_repeat(v + dvdx * (filter_x - 0.5f) + dvdy * (filter_y - 0.5f),
267  1.0f - u);
268  }
269 
270  /* Shader data setup. */
271  int object = kernel_data.bake.object_index;
272  int shader;
273  float3 P, Ng;
274 
275  triangle_point_normal(kg, object, prim, u, v, &P, &Ng, &shader);
276 
277  ShaderData sd;
279  kg,
280  &sd,
281  P,
282  Ng,
283  Ng,
284  shader,
285  object,
286  prim,
287  u,
288  v,
289  1.0f,
290  0.5f,
291  !(kernel_tex_fetch(__object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED),
292  LAMP_NONE);
293  sd.I = sd.N;
294 
295  /* Setup differentials. */
296  sd.dP.dx = sd.dPdu * dudx + sd.dPdv * dvdx;
297  sd.dP.dy = sd.dPdu * dudy + sd.dPdv * dvdy;
298  sd.du.dx = dudx;
299  sd.du.dy = dudy;
300  sd.dv.dx = dvdx;
301  sd.dv.dy = dvdy;
302 
303  /* Set RNG state for shaders that use sampling. */
304  PathState state = {0};
305  state.rng_hash = rng_hash;
306  state.rng_offset = 0;
307  state.sample = sample;
308  state.num_samples = num_samples;
309  state.min_ray_pdf = FLT_MAX;
310 
311  /* Light passes if we need more than color. */
312  PathRadiance L;
313  int pass_filter = kernel_data.bake.pass_filter;
314 
315  if (kernel_data.bake.pass_filter & ~BAKE_FILTER_COLOR)
316  compute_light_pass(kg, &sd, &L, rng_hash, pass_filter, sample);
317 
318  float3 out = zero_float3();
319 
321  switch (type) {
322  /* data passes */
323  case SHADER_EVAL_NORMAL:
325  case SHADER_EVAL_EMISSION: {
326  if (type != SHADER_EVAL_NORMAL || (sd.flag & SD_HAS_BUMP)) {
327  int path_flag = (type == SHADER_EVAL_EMISSION) ? PATH_RAY_EMISSION : 0;
328  shader_eval_surface(kg, &sd, &state, NULL, path_flag);
329  }
330 
331  if (type == SHADER_EVAL_NORMAL) {
332  float3 N = sd.N;
333  if (sd.flag & SD_HAS_BUMP) {
335  }
336 
337  /* encoding: normal = (2 * color) - 1 */
338  out = N * 0.5f + make_float3(0.5f, 0.5f, 0.5f);
339  }
340  else if (type == SHADER_EVAL_ROUGHNESS) {
343  }
344  else {
345  out = shader_emissive_eval(&sd);
346  }
347  break;
348  }
349  case SHADER_EVAL_UV: {
350  out = primitive_uv(kg, &sd);
351  break;
352  }
353 # ifdef __PASSES__
354  /* light passes */
355  case SHADER_EVAL_AO: {
356  out = L.ao;
357  break;
358  }
359  case SHADER_EVAL_COMBINED: {
360  if ((pass_filter & BAKE_FILTER_COMBINED) == BAKE_FILTER_COMBINED) {
361  float alpha;
363  break;
364  }
365 
367  out += L.direct_diffuse;
369  out += L.indirect_diffuse;
370 
372  out += L.direct_glossy;
374  out += L.indirect_glossy;
375 
377  out += L.direct_transmission;
379  out += L.indirect_transmission;
380 
381  if ((pass_filter & BAKE_FILTER_EMISSION) != 0)
382  out += L.emission;
383 
384  break;
385  }
386  case SHADER_EVAL_SHADOW: {
387  out = L.shadow;
388  break;
389  }
390  case SHADER_EVAL_DIFFUSE: {
391  out = kernel_bake_evaluate_direct_indirect(
392  kg, &sd, &state, L.direct_diffuse, L.indirect_diffuse, type, pass_filter);
393  break;
394  }
395  case SHADER_EVAL_GLOSSY: {
396  out = kernel_bake_evaluate_direct_indirect(
397  kg, &sd, &state, L.direct_glossy, L.indirect_glossy, type, pass_filter);
398  break;
399  }
401  out = kernel_bake_evaluate_direct_indirect(
402  kg, &sd, &state, L.direct_transmission, L.indirect_transmission, type, pass_filter);
403  break;
404  }
405 # endif
406 
407  /* extra */
409  /* setup ray */
410  Ray ray;
411 
412  ray.P = zero_float3();
413  ray.D = normalize(P);
414  ray.t = 0.0f;
415 # ifdef __CAMERA_MOTION__
416  ray.time = 0.5f;
417 # endif
418 
419 # ifdef __RAY_DIFFERENTIALS__
420  ray.dD = differential3_zero();
421  ray.dP = differential3_zero();
422 # endif
423 
424  /* setup shader data */
425  shader_setup_from_background(kg, &sd, &ray);
426 
427  /* evaluate */
428  int path_flag = 0; /* we can't know which type of BSDF this is for */
429  shader_eval_surface(kg, &sd, &state, NULL, path_flag | PATH_RAY_EMISSION);
430  out = shader_background_eval(&sd);
431  break;
432  }
433  default: {
434  /* no real shader, returning the position of the verts for debugging */
435  out = normalize(P);
436  break;
437  }
438  }
439 
440  /* write output */
441  const float4 result = make_float4(out.x, out.y, out.z, 1.0f);
443 }
444 
445 #endif /* __BAKING__ */
446 
448  ccl_global uint4 *input,
449  ccl_global float4 *output,
450  int i)
451 {
452  ShaderData sd;
453  PathState state = {0};
454  uint4 in = input[i];
455 
456  /* setup shader data */
457  int object = in.x;
458  int prim = in.y;
459  float u = __uint_as_float(in.z);
460  float v = __uint_as_float(in.w);
461 
462  shader_setup_from_displace(kg, &sd, object, prim, u, v);
463 
464  /* evaluate */
465  float3 P = sd.P;
467  float3 D = sd.P - P;
468 
470 
471  /* write output */
472  output[i] += make_float4(D.x, D.y, D.z, 0.0f);
473 }
474 
476  ccl_global uint4 *input,
477  ccl_global float4 *output,
478  int i)
479 {
480  ShaderData sd;
481  PathState state = {0};
482  uint4 in = input[i];
483 
484  /* setup ray */
485  Ray ray;
486  float u = __uint_as_float(in.x);
487  float v = __uint_as_float(in.y);
488 
489  ray.P = zero_float3();
490  ray.D = equirectangular_to_direction(u, v);
491  ray.t = 0.0f;
492 #ifdef __CAMERA_MOTION__
493  ray.time = 0.5f;
494 #endif
495 
496 #ifdef __RAY_DIFFERENTIALS__
497  ray.dD = differential3_zero();
498  ray.dP = differential3_zero();
499 #endif
500 
501  /* setup shader data */
502  shader_setup_from_background(kg, &sd, &ray);
503 
504  /* evaluate */
505  int path_flag = 0; /* we can't know which type of BSDF this is for */
506  shader_eval_surface(kg, &sd, &state, NULL, path_flag | PATH_RAY_EMISSION);
507  float3 color = shader_background_eval(&sd);
508 
509  /* write output */
510  output[i] += make_float4(color.x, color.y, color.z, 0.0f);
511 }
512 
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 type
_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
ATTR_WARN_UNUSED_RESULT const BMVert * v
#define output
static unsigned long seed
Definition: btSoftBody.h:39
static CCL_NAMESPACE_BEGIN const double alpha
ccl_device_inline void object_inverse_dir_transform(KernelGlobals *kg, const ShaderData *sd, float3 *D)
Definition: geom_object.h:190
ccl_device_inline float3 primitive_uv(KernelGlobals *kg, ShaderData *sd)
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 void path_radiance_sum_indirect(PathRadiance *L)
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_emission(KernelGlobals *kg, PathRadiance *L, ccl_addr_space PathState *state, float3 throughput, float3 value)
ccl_device_inline float3 path_radiance_clamp_and_sum(KernelGlobals *kg, PathRadiance *L, float *alpha)
ccl_device void kernel_background_evaluate(KernelGlobals *kg, ccl_global uint4 *input, ccl_global float4 *output, int i)
Definition: kernel_bake.h:475
CCL_NAMESPACE_BEGIN ccl_device void kernel_displace_evaluate(KernelGlobals *kg, ccl_global uint4 *input, ccl_global float4 *output, int i)
Definition: kernel_bake.h:447
#define kernel_data
#define kernel_assert(cond)
#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 floorf(x)
#define make_float4(x, y, z, w)
#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 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 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 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 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)
ccl_device float3 equirectangular_to_direction(float u, float v)
ccl_device_forceinline void path_rng_2D(KernelGlobals *kg, uint rng_hash, int sample, int num_samples, int dimension, float *fx, float *fy)
Definition: kernel_random.h:91
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 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 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 void shader_setup_from_background(KernelGlobals *kg, ShaderData *sd, const Ray *ray)
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)
__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_HAS_BUMP
Definition: kernel_types.h:887
@ SD_BSSRDF
Definition: kernel_types.h:851
@ SD_EMISSION
Definition: kernel_types.h:845
#define AS_SHADER_DATA(shader_data_tiny_storage)
@ PRNG_FILTER_U
Definition: kernel_types.h:230
@ PATH_RAY_EMISSION
Definition: kernel_types.h:321
ShaderDataTinyStorage
ShaderData
@ SD_OBJECT_TRANSFORM_APPLIED
Definition: kernel_types.h:908
@ BAKE_FILTER_TRANSMISSION_DIRECT
Definition: kernel_types.h:452
@ BAKE_FILTER_DIFFUSE_INDIRECT
Definition: kernel_types.h:453
@ BAKE_FILTER_GLOSSY_INDIRECT
Definition: kernel_types.h:454
@ BAKE_FILTER_COMBINED
Definition: kernel_types.h:447
@ BAKE_FILTER_DIFFUSE_DIRECT
Definition: kernel_types.h:450
@ BAKE_FILTER_TRANSMISSION_INDIRECT
Definition: kernel_types.h:455
@ BAKE_FILTER_GLOSSY_DIRECT
Definition: kernel_types.h:451
#define LAMP_NONE
Definition: kernel_types.h:61
ShaderEvalType
Definition: kernel_types.h:196
@ SHADER_EVAL_DIFFUSE
Definition: kernel_types.h:218
@ SHADER_EVAL_TRANSMISSION
Definition: kernel_types.h:220
@ SHADER_EVAL_GLOSSY
Definition: kernel_types.h:219
@ SHADER_EVAL_EMISSION
Definition: kernel_types.h:210
@ SHADER_EVAL_UV
Definition: kernel_types.h:205
@ SHADER_EVAL_ROUGHNESS
Definition: kernel_types.h:206
@ SHADER_EVAL_SHADOW
Definition: kernel_types.h:217
@ SHADER_EVAL_AO
Definition: kernel_types.h:215
@ SHADER_EVAL_NORMAL
Definition: kernel_types.h:204
@ SHADER_EVAL_ENVIRONMENT
Definition: kernel_types.h:223
@ SHADER_EVAL_COMBINED
Definition: kernel_types.h:216
@ BAKE_FILTER_EMISSION
Definition: kernel_types.h:442
@ BAKE_FILTER_DIFFUSE
Definition: kernel_types.h:439
@ BAKE_FILTER_DIRECT
Definition: kernel_types.h:436
@ BAKE_FILTER_INDIRECT
Definition: kernel_types.h:437
@ BAKE_FILTER_AO
Definition: kernel_types.h:443
@ BAKE_FILTER_COLOR
Definition: kernel_types.h:438
ccl_device_inline void kernel_write_pass_float4(ccl_global float *buffer, float4 value)
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])
static const pxr::TfToken roughness("roughness", pxr::TfToken::Immortal)
params N
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
float z
Definition: sky_float3.h:35
float y
Definition: sky_float3.h:35
float x
Definition: sky_float3.h:35
float max
__forceinline bool all(const avxb &b)
Definition: util_avxb.h:214
ccl_device_inline uint hash_uint(uint kx)
Definition: util_hash.h:72
ccl_device_inline float __uint_as_float(uint i)
Definition: util_math.h:232
ccl_device_inline uint __float_as_uint(float f)
Definition: util_math.h:222
ccl_device_inline float3 safe_divide_even_color(float3 a, float3 b)
Definition: util_math.h:525
ccl_device_inline float2 normalize(const float2 &a)
ccl_device_inline float3 one_float3()
ccl_device_inline float3 zero_float3()
BLI_INLINE float D(const float *data, const int res[3], int x, int y, int z)
Definition: voxel.c:29