Blender  V2.93
kernel_passes.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 
20 
21 #ifdef __DENOISING_FEATURES__
22 
23 ccl_device_inline void kernel_write_denoising_shadow(KernelGlobals *kg,
24  ccl_global float *buffer,
25  int sample,
26  float path_total,
27  float path_total_shaded)
28 {
29  if (kernel_data.film.pass_denoising_data == 0)
30  return;
31 
32  buffer += sample_is_even(kernel_data.integrator.sampling_pattern, sample) ?
35 
36  path_total = ensure_finite(path_total);
37  path_total_shaded = ensure_finite(path_total_shaded);
38 
39  kernel_write_pass_float(buffer, path_total);
40  kernel_write_pass_float(buffer + 1, path_total_shaded);
41 
42  float value = path_total_shaded / max(path_total, 1e-7f);
43  kernel_write_pass_float(buffer + 2, value * value);
44 }
45 
46 ccl_device_inline void kernel_update_denoising_features(KernelGlobals *kg,
47  ShaderData *sd,
49  PathRadiance *L)
50 {
51  if (state->denoising_feature_weight == 0.0f) {
52  return;
53  }
54 
55  L->denoising_depth += ensure_finite(state->denoising_feature_weight * sd->ray_length);
56 
57  /* Skip implicitly transparent surfaces. */
58  if (sd->flag & SD_HAS_ONLY_VOLUME) {
59  return;
60  }
61 
63  float3 diffuse_albedo = zero_float3();
64  float3 specular_albedo = zero_float3();
65  float sum_weight = 0.0f, sum_nonspecular_weight = 0.0f;
66 
67  for (int i = 0; i < sd->num_closure; i++) {
68  ShaderClosure *sc = &sd->closure[i];
69 
70  if (!CLOSURE_IS_BSDF_OR_BSSRDF(sc->type))
71  continue;
72 
73  /* All closures contribute to the normal feature, but only diffuse-like ones to the albedo. */
74  normal += sc->N * sc->sample_weight;
75  sum_weight += sc->sample_weight;
76 
77  float3 closure_albedo = sc->weight;
78  /* Closures that include a Fresnel term typically have weights close to 1 even though their
79  * actual contribution is significantly lower.
80  * To account for this, we scale their weight by the average fresnel factor (the same is also
81  * done for the sample weight in the BSDF setup, so we don't need to scale that here). */
82  if (CLOSURE_IS_BSDF_MICROFACET_FRESNEL(sc->type)) {
83  MicrofacetBsdf *bsdf = (MicrofacetBsdf *)sc;
84  closure_albedo *= bsdf->extra->fresnel_color;
85  }
86  else if (sc->type == CLOSURE_BSDF_PRINCIPLED_SHEEN_ID) {
88  closure_albedo *= bsdf->avg_value;
89  }
90  else if (sc->type == CLOSURE_BSDF_HAIR_PRINCIPLED_ID) {
91  closure_albedo *= bsdf_principled_hair_albedo(sc);
92  }
93 
94  if (bsdf_get_specular_roughness_squared(sc) > sqr(0.075f)) {
95  diffuse_albedo += closure_albedo;
96  sum_nonspecular_weight += sc->sample_weight;
97  }
98  else {
99  specular_albedo += closure_albedo;
100  }
101  }
102 
103  /* Wait for next bounce if 75% or more sample weight belongs to specular-like closures. */
104  if ((sum_weight == 0.0f) || (sum_nonspecular_weight * 4.0f > sum_weight)) {
105  if (sum_weight != 0.0f) {
106  normal /= sum_weight;
107  }
108 
109  /* Transform normal into camera space. */
110  const Transform worldtocamera = kernel_data.cam.worldtocamera;
111  normal = transform_direction(&worldtocamera, normal);
112 
113  L->denoising_normal += ensure_finite3(state->denoising_feature_weight * normal);
114  L->denoising_albedo += ensure_finite3(state->denoising_feature_weight *
115  state->denoising_feature_throughput * diffuse_albedo);
116 
117  state->denoising_feature_weight = 0.0f;
118  }
119  else {
120  state->denoising_feature_throughput *= specular_albedo;
121  }
122 }
123 #endif /* __DENOISING_FEATURES__ */
124 
125 #ifdef __KERNEL_DEBUG__
126 ccl_device_inline void kernel_write_debug_passes(KernelGlobals *kg,
127  ccl_global float *buffer,
128  PathRadiance *L)
129 {
130  int flag = kernel_data.film.pass_flag;
131  if (flag & PASSMASK(BVH_TRAVERSED_NODES)) {
132  kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_traversed_nodes,
133  L->debug_data.num_bvh_traversed_nodes);
134  }
135  if (flag & PASSMASK(BVH_TRAVERSED_INSTANCES)) {
136  kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_traversed_instances,
137  L->debug_data.num_bvh_traversed_instances);
138  }
139  if (flag & PASSMASK(BVH_INTERSECTIONS)) {
140  kernel_write_pass_float(buffer + kernel_data.film.pass_bvh_intersections,
141  L->debug_data.num_bvh_intersections);
142  }
143  if (flag & PASSMASK(RAY_BOUNCES)) {
144  kernel_write_pass_float(buffer + kernel_data.film.pass_ray_bounces,
145  L->debug_data.num_ray_bounces);
146  }
147 }
148 #endif /* __KERNEL_DEBUG__ */
149 
150 #ifdef __KERNEL_CPU__
151 # define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name) \
152  kernel_write_id_pass_cpu(buffer, depth * 2, id, matte_weight, kg->coverage_##name)
153 ccl_device_inline size_t kernel_write_id_pass_cpu(
154  float *buffer, size_t depth, float id, float matte_weight, CoverageMap *map)
155 {
156  if (map) {
157  (*map)[id] += matte_weight;
158  return 0;
159  }
160 #else /* __KERNEL_CPU__ */
161 # define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name) \
162  kernel_write_id_slots_gpu(buffer, depth * 2, id, matte_weight)
164  size_t depth,
165  float id,
166  float matte_weight)
167 {
168 #endif /* __KERNEL_CPU__ */
169  kernel_write_id_slots(buffer, depth, id, matte_weight);
170  return depth * 2;
171 }
172 
174  ccl_global float *buffer,
175  PathRadiance *L,
176  ShaderData *sd,
178  float3 throughput)
179 {
180 #ifdef __PASSES__
181  int path_flag = state->flag;
182 
183  if (!(path_flag & PATH_RAY_CAMERA))
184  return;
185 
186  int flag = kernel_data.film.pass_flag;
187  int light_flag = kernel_data.film.light_pass_flag;
188 
189  if (!((flag | light_flag) & PASS_ANY))
190  return;
191 
192  if (!(path_flag & PATH_RAY_SINGLE_PASS_DONE)) {
193  if (!(sd->flag & SD_TRANSPARENT) || kernel_data.film.pass_alpha_threshold == 0.0f ||
194  average(shader_bsdf_alpha(kg, sd)) >= kernel_data.film.pass_alpha_threshold) {
195  if (state->sample == 0) {
196  if (flag & PASSMASK(DEPTH)) {
197  float depth = camera_z_depth(kg, sd->P);
198  kernel_write_pass_float(buffer + kernel_data.film.pass_depth, depth);
199  }
200  if (flag & PASSMASK(OBJECT_ID)) {
201  float id = object_pass_id(kg, sd->object);
202  kernel_write_pass_float(buffer + kernel_data.film.pass_object_id, id);
203  }
204  if (flag & PASSMASK(MATERIAL_ID)) {
205  float id = shader_pass_id(kg, sd);
206  kernel_write_pass_float(buffer + kernel_data.film.pass_material_id, id);
207  }
208  }
209 
210  if (flag & PASSMASK(NORMAL)) {
212  kernel_write_pass_float3(buffer + kernel_data.film.pass_normal, normal);
213  }
214  if (flag & PASSMASK(UV)) {
215  float3 uv = primitive_uv(kg, sd);
216  kernel_write_pass_float3(buffer + kernel_data.film.pass_uv, uv);
217  }
218  if (flag & PASSMASK(MOTION)) {
219  float4 speed = primitive_motion_vector(kg, sd);
220  kernel_write_pass_float4(buffer + kernel_data.film.pass_motion, speed);
221  kernel_write_pass_float(buffer + kernel_data.film.pass_motion_weight, 1.0f);
222  }
223 
225  }
226  }
227 
228  if (kernel_data.film.cryptomatte_passes) {
229  const float matte_weight = average(throughput) *
230  (1.0f - average(shader_bsdf_transparency(kg, sd)));
231  if (matte_weight > 0.0f) {
232  ccl_global float *cryptomatte_buffer = buffer + kernel_data.film.pass_cryptomatte;
233  if (kernel_data.film.cryptomatte_passes & CRYPT_OBJECT) {
234  float id = object_cryptomatte_id(kg, sd->object);
235  cryptomatte_buffer += WRITE_ID_SLOT(
236  cryptomatte_buffer, kernel_data.film.cryptomatte_depth, id, matte_weight, object);
237  }
238  if (kernel_data.film.cryptomatte_passes & CRYPT_MATERIAL) {
239  float id = shader_cryptomatte_id(kg, sd->shader);
240  cryptomatte_buffer += WRITE_ID_SLOT(
241  cryptomatte_buffer, kernel_data.film.cryptomatte_depth, id, matte_weight, material);
242  }
243  if (kernel_data.film.cryptomatte_passes & CRYPT_ASSET) {
244  float id = object_cryptomatte_asset_id(kg, sd->object);
245  cryptomatte_buffer += WRITE_ID_SLOT(
246  cryptomatte_buffer, kernel_data.film.cryptomatte_depth, id, matte_weight, asset);
247  }
248  }
249  }
250 
251  if (light_flag & PASSMASK_COMPONENT(DIFFUSE))
252  L->color_diffuse += shader_bsdf_diffuse(kg, sd) * throughput;
253  if (light_flag & PASSMASK_COMPONENT(GLOSSY))
254  L->color_glossy += shader_bsdf_glossy(kg, sd) * throughput;
255  if (light_flag & PASSMASK_COMPONENT(TRANSMISSION))
256  L->color_transmission += shader_bsdf_transmission(kg, sd) * throughput;
257 
258  if (light_flag & PASSMASK(MIST)) {
259  /* bring depth into 0..1 range */
260  float mist_start = kernel_data.film.mist_start;
261  float mist_inv_depth = kernel_data.film.mist_inv_depth;
262 
263  float depth = camera_distance(kg, sd->P);
264  float mist = saturate((depth - mist_start) * mist_inv_depth);
265 
266  /* falloff */
267  float mist_falloff = kernel_data.film.mist_falloff;
268 
269  if (mist_falloff == 1.0f)
270  ;
271  else if (mist_falloff == 2.0f)
272  mist = mist * mist;
273  else if (mist_falloff == 0.5f)
274  mist = sqrtf(mist);
275  else
276  mist = powf(mist, mist_falloff);
277 
278  /* modulate by transparency */
280  L->mist += (1.0f - mist) * average(throughput * alpha);
281  }
282 #endif
283 }
284 
286  ccl_global float *buffer,
287  PathRadiance *L)
288 {
289 #ifdef __PASSES__
290  int light_flag = kernel_data.film.light_pass_flag;
291 
292  if (!kernel_data.film.use_light_pass)
293  return;
294 
295  if (light_flag & PASSMASK(DIFFUSE_INDIRECT))
296  kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_indirect, L->indirect_diffuse);
297  if (light_flag & PASSMASK(GLOSSY_INDIRECT))
298  kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_indirect, L->indirect_glossy);
299  if (light_flag & PASSMASK(TRANSMISSION_INDIRECT))
300  kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_indirect,
301  L->indirect_transmission);
302  if (light_flag & PASSMASK(VOLUME_INDIRECT))
303  kernel_write_pass_float3(buffer + kernel_data.film.pass_volume_indirect, L->indirect_volume);
304  if (light_flag & PASSMASK(DIFFUSE_DIRECT))
305  kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_direct, L->direct_diffuse);
306  if (light_flag & PASSMASK(GLOSSY_DIRECT))
307  kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_direct, L->direct_glossy);
308  if (light_flag & PASSMASK(TRANSMISSION_DIRECT))
309  kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_direct,
310  L->direct_transmission);
311  if (light_flag & PASSMASK(VOLUME_DIRECT))
312  kernel_write_pass_float3(buffer + kernel_data.film.pass_volume_direct, L->direct_volume);
313 
314  if (light_flag & PASSMASK(EMISSION))
315  kernel_write_pass_float3(buffer + kernel_data.film.pass_emission, L->emission);
316  if (light_flag & PASSMASK(BACKGROUND))
317  kernel_write_pass_float3(buffer + kernel_data.film.pass_background, L->background);
318  if (light_flag & PASSMASK(AO))
319  kernel_write_pass_float3(buffer + kernel_data.film.pass_ao, L->ao);
320 
321  if (light_flag & PASSMASK(DIFFUSE_COLOR))
322  kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_color, L->color_diffuse);
323  if (light_flag & PASSMASK(GLOSSY_COLOR))
324  kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_color, L->color_glossy);
325  if (light_flag & PASSMASK(TRANSMISSION_COLOR))
326  kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_color,
327  L->color_transmission);
328  if (light_flag & PASSMASK(SHADOW)) {
329  float3 shadow = L->shadow;
331  buffer + kernel_data.film.pass_shadow,
332  make_float4(shadow.x, shadow.y, shadow.z, kernel_data.film.pass_shadow_scale));
333  }
334  if (light_flag & PASSMASK(MIST))
335  kernel_write_pass_float(buffer + kernel_data.film.pass_mist, 1.0f - L->mist);
336 #endif
337 }
338 
340  ccl_global float *buffer,
341  int sample,
342  PathRadiance *L)
343 {
346 
347  float alpha;
349 
350  if (kernel_data.film.pass_flag & PASSMASK(COMBINED)) {
351  kernel_write_pass_float4(buffer, make_float4(L_sum.x, L_sum.y, L_sum.z, alpha));
352  }
353 
355 
356 #ifdef __DENOISING_FEATURES__
357  if (kernel_data.film.pass_denoising_data) {
358 # ifdef __SHADOW_TRICKS__
359  kernel_write_denoising_shadow(kg,
360  buffer + kernel_data.film.pass_denoising_data,
361  sample,
362  average(L->path_total),
363  average(L->path_total_shaded));
364 # else
365  kernel_write_denoising_shadow(
366  kg, buffer + kernel_data.film.pass_denoising_data, sample, 0.0f, 0.0f);
367 # endif
368  if (kernel_data.film.pass_denoising_clean) {
369  float3 noisy, clean;
370  path_radiance_split_denoising(kg, L, &noisy, &clean);
371  kernel_write_pass_float3_variance(
372  buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_COLOR, noisy);
373  kernel_write_pass_float3_unaligned(buffer + kernel_data.film.pass_denoising_clean, clean);
374  }
375  else {
376  kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data +
378  ensure_finite3(L_sum));
379  }
380 
381  kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data +
383  L->denoising_normal);
384  kernel_write_pass_float3_variance(buffer + kernel_data.film.pass_denoising_data +
386  L->denoising_albedo);
387  kernel_write_pass_float_variance(
388  buffer + kernel_data.film.pass_denoising_data + DENOISING_PASS_DEPTH, L->denoising_depth);
389  }
390 #endif /* __DENOISING_FEATURES__ */
391 
392 #ifdef __KERNEL_DEBUG__
393  kernel_write_debug_passes(kg, buffer, L);
394 #endif
395 
396  /* Adaptive Sampling. Fill the additional buffer with the odd samples and calculate our stopping
397  criteria. This is the heuristic from "A hierarchical automatic stopping condition for Monte
398  Carlo global illumination" except that here it is applied per pixel and not in hierarchical
399  tiles. */
400  if (kernel_data.film.pass_adaptive_aux_buffer &&
401  kernel_data.integrator.adaptive_threshold > 0.0f) {
402  if (sample_is_even(kernel_data.integrator.sampling_pattern, sample)) {
403  kernel_write_pass_float4(buffer + kernel_data.film.pass_adaptive_aux_buffer,
404  make_float4(L_sum.x * 2.0f, L_sum.y * 2.0f, L_sum.z * 2.0f, 0.0f));
405  }
406 #ifdef __KERNEL_CPU__
407  if ((sample > kernel_data.integrator.adaptive_min_samples) &&
408  kernel_data.integrator.adaptive_stop_per_sample) {
409  const int step = kernel_data.integrator.adaptive_step;
410 
411  if ((sample & (step - 1)) == (step - 1)) {
413  }
414  }
415 #endif
416  }
417 
418  /* Write the sample count as negative numbers initially to mark the samples as in progress.
419  * Once the tile has finished rendering, the sign gets flipped and all the pixel values
420  * are scaled as if they were taken at a uniform sample count. */
421  if (kernel_data.film.pass_sample_count) {
422  /* Make sure it's a negative number. In progressive refine mode, this bit gets flipped between
423  * passes. */
424 #ifdef __ATOMIC_PASS_WRITE__
425  atomic_fetch_and_or_uint32((ccl_global uint *)(buffer + kernel_data.film.pass_sample_count),
426  0x80000000);
427 #else
428  if (buffer[kernel_data.film.pass_sample_count] > 0) {
429  buffer[kernel_data.film.pass_sample_count] *= -1.0f;
430  }
431 #endif
432  kernel_write_pass_float(buffer + kernel_data.film.pass_sample_count, -1.0f);
433  }
434 }
435 
unsigned int uint
Definition: BLI_sys_types.h:83
Group RGB to NORMAL
ATOMIC_INLINE uint32_t atomic_fetch_and_or_uint32(uint32_t *p, uint32_t x)
CCL_NAMESPACE_BEGIN ccl_device_inline float bsdf_get_specular_roughness_squared(const ShaderClosure *sc)
Definition: bsdf.h:42
ccl_device float3 bsdf_principled_hair_albedo(ShaderClosure *sc)
unordered_map< float, float > CoverageMap
Definition: coverage.h:26
Material material
static CCL_NAMESPACE_BEGIN const double alpha
ccl_device_inline float object_cryptomatte_asset_id(KernelGlobals *kg, int object)
Definition: geom_object.h:353
ccl_device_inline float object_pass_id(KernelGlobals *kg, int object)
Definition: geom_object.h:230
ccl_device int shader_pass_id(KernelGlobals *kg, const ShaderData *sd)
Definition: geom_object.h:338
ccl_device_inline float object_cryptomatte_id(KernelGlobals *kg, int object)
Definition: geom_object.h:345
ccl_device_inline float3 primitive_uv(KernelGlobals *kg, ShaderData *sd)
ccl_device_inline float4 primitive_motion_vector(KernelGlobals *kg, ShaderData *sd)
IconTextureDrawCall normal
CCL_NAMESPACE_BEGIN ccl_device float3 shader_bsdf_transparency(KernelGlobals *kg, const ShaderData *sd)
ccl_device_inline void path_radiance_split_denoising(KernelGlobals *kg, PathRadiance *L, float3 *noisy, float3 *clean)
ccl_device_inline float3 path_radiance_clamp_and_sum(KernelGlobals *kg, PathRadiance *L, float *alpha)
CCL_NAMESPACE_BEGIN ccl_device void kernel_do_adaptive_stopping(KernelGlobals *kg, ccl_global float *buffer, int sample)
ccl_device_inline float camera_distance(KernelGlobals *kg, float3 P)
ccl_device_inline float camera_z_depth(KernelGlobals *kg, float3 P)
#define kernel_data
#define ccl_addr_space
#define ccl_device_inline
#define ccl_global
#define powf(x, y)
#define CCL_NAMESPACE_END
#define make_float4(x, y, z, w)
#define sqrtf(x)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_write_id_slots(ccl_global float *buffer, int num_slots, float id, float weight)
ccl_device_inline void kernel_write_light_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L)
#define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name)
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_inline size_t kernel_write_id_slots_gpu(ccl_global float *buffer, size_t depth, float id, float matte_weight)
#define PROFILING_OBJECT(object)
#define PROFILING_INIT(kg, event)
ccl_device_inline bool sample_is_even(int pattern, int sample)
ccl_device float3 shader_bsdf_alpha(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd)
ccl_device float shader_cryptomatte_id(KernelGlobals *kg, int shader)
ccl_device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd)
ccl_device float3 shader_bsdf_average_normal(KernelGlobals *kg, ShaderData *sd)
__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_ONLY_VOLUME
Definition: kernel_types.h:875
@ SD_TRANSPARENT
Definition: kernel_types.h:859
#define PASSMASK_COMPONENT(comp)
Definition: kernel_types.h:343
#define PRIM_NONE
Definition: kernel_types.h:60
@ PATH_RAY_SINGLE_PASS_DONE
Definition: kernel_types.h:304
@ PATH_RAY_CAMERA
Definition: kernel_types.h:266
@ CRYPT_ASSET
Definition: kernel_types.h:404
@ CRYPT_OBJECT
Definition: kernel_types.h:402
@ CRYPT_MATERIAL
Definition: kernel_types.h:403
ShaderData
#define PASSMASK(pass)
Definition: kernel_types.h:341
ShaderClosure
Definition: kernel_types.h:831
#define PASS_ANY
Definition: kernel_types.h:398
@ DENOISING_PASS_ALBEDO
Definition: kernel_types.h:411
@ DENOISING_PASS_DEPTH
Definition: kernel_types.h:413
@ DENOISING_PASS_SHADOW_B
Definition: kernel_types.h:416
@ DENOISING_PASS_COLOR
Definition: kernel_types.h:417
@ DENOISING_PASS_NORMAL
Definition: kernel_types.h:409
@ DENOISING_PASS_SHADOW_A
Definition: kernel_types.h:415
ccl_device_inline void kernel_write_pass_float3(ccl_global float *buffer, float3 value)
ccl_device_inline void kernel_write_pass_float4(ccl_global float *buffer, float4 value)
CCL_NAMESPACE_BEGIN ccl_device_inline void kernel_write_pass_float(ccl_global float *buffer, float value)
static ulong state[N]
#define L
static void sample(SocketReader *reader, int x, int y, float color[4])
MicrofacetExtra * extra
float z
Definition: sky_float3.h:35
float y
Definition: sky_float3.h:35
float x
Definition: sky_float3.h:35
#define CLOSURE_IS_BSDF_MICROFACET_FRESNEL(type)
Definition: svm_types.h:617
#define CLOSURE_IS_BSDF_OR_BSSRDF(type)
Definition: svm_types.h:622
@ CLOSURE_BSDF_HAIR_PRINCIPLED_ID
Definition: svm_types.h:565
@ CLOSURE_BSDF_PRINCIPLED_SHEEN_ID
Definition: svm_types.h:538
float max
ccl_device_inline float ensure_finite(float v)
Definition: util_math.h:277
ccl_device_inline float saturate(float a)
Definition: util_math.h:315
ccl_device_inline float sqr(float a)
Definition: util_math.h:651
ccl_device_inline float average(const float2 &a)
ccl_device_inline float3 zero_float3()
ccl_device_inline float3 ensure_finite3(float3 v)
@ PROFILING_WRITE_RESULT
ccl_device_inline float3 transform_direction(const Transform *t, const float3 a)