Blender  V3.3
bevel.h
Go to the documentation of this file.
1 /* SPDX-License-Identifier: Apache-2.0
2  * Copyright 2011-2022 Blender Foundation */
3 
4 #pragma once
5 
6 #include "kernel/bvh/bvh.h"
9 
11 
12 #ifdef __SHADER_RAYTRACE__
13 
14 /* Planar Cubic BSSRDF falloff, reused for bevel.
15  *
16  * This is basically (Rm - x)^3, with some factors to normalize it. For sampling
17  * we integrate 2*pi*x * (Rm - x)^3, which gives us a quintic equation that as
18  * far as I can tell has no closed form solution. So we get an iterative solution
19  * instead with newton-raphson. */
20 
21 ccl_device float svm_bevel_cubic_eval(const float radius, float r)
22 {
23  const float Rm = radius;
24 
25  if (r >= Rm)
26  return 0.0f;
27 
28  /* integrate (2*pi*r * 10*(R - r)^3)/(pi * R^5) from 0 to R = 1 */
29  const float Rm5 = (Rm * Rm) * (Rm * Rm) * Rm;
30  const float f = Rm - r;
31  const float num = f * f * f;
32 
33  return (10.0f * num) / (Rm5 * M_PI_F);
34 }
35 
36 ccl_device float svm_bevel_cubic_pdf(const float radius, float r)
37 {
38  return svm_bevel_cubic_eval(radius, r);
39 }
40 
41 /* solve 10x^2 - 20x^3 + 15x^4 - 4x^5 - xi == 0 */
42 ccl_device_forceinline float svm_bevel_cubic_quintic_root_find(float xi)
43 {
44  /* newton-raphson iteration, usually succeeds in 2-4 iterations, except
45  * outside 0.02 ... 0.98 where it can go up to 10, so overall performance
46  * should not be too bad */
47  const float tolerance = 1e-6f;
48  const int max_iteration_count = 10;
49  float x = 0.25f;
50  int i;
51 
52  for (i = 0; i < max_iteration_count; i++) {
53  float x2 = x * x;
54  float x3 = x2 * x;
55  float nx = (1.0f - x);
56 
57  float f = 10.0f * x2 - 20.0f * x3 + 15.0f * x2 * x2 - 4.0f * x2 * x3 - xi;
58  float f_ = 20.0f * (x * nx) * (nx * nx);
59 
60  if (fabsf(f) < tolerance || f_ == 0.0f)
61  break;
62 
63  x = saturatef(x - f / f_);
64  }
65 
66  return x;
67 }
68 
69 ccl_device void svm_bevel_cubic_sample(const float radius,
70  float xi,
71  ccl_private float *r,
72  ccl_private float *h)
73 {
74  float Rm = radius;
75  float r_ = svm_bevel_cubic_quintic_root_find(xi);
76 
77  r_ *= Rm;
78  *r = r_;
79 
80  /* h^2 + r^2 = Rm^2 */
81  *h = safe_sqrtf(Rm * Rm - r_ * r_);
82 }
83 
84 /* Bevel shader averaging normals from nearby surfaces.
85  *
86  * Sampling strategy from: BSSRDF Importance Sampling, SIGGRAPH 2013
87  * http://library.imageworks.com/pdfs/imageworks-library-BSSRDF-sampling.pdf
88  */
89 
90 # ifdef __KERNEL_OPTIX__
91 extern "C" __device__ float3 __direct_callable__svm_node_bevel(
92 # else
93 ccl_device float3 svm_bevel(
94 # endif
95  KernelGlobals kg,
98  float radius,
99  int num_samples)
100 {
101  /* Early out if no sampling needed. */
102  if (radius <= 0.0f || num_samples < 1 || sd->object == OBJECT_NONE) {
103  return sd->N;
104  }
105 
106  /* Can't ray-trace from shaders like displacement, before BVH exists. */
107  if (kernel_data.bvh.bvh_layout == BVH_LAYOUT_NONE) {
108  return sd->N;
109  }
110 
111  /* Don't bevel for blurry indirect rays. */
112  if (INTEGRATOR_STATE(state, path, min_ray_pdf) < 8.0f) {
113  return sd->N;
114  }
115 
116  /* Setup for multi intersection. */
117  LocalIntersection isect;
118  uint lcg_state = lcg_state_init(INTEGRATOR_STATE(state, path, rng_hash),
119  INTEGRATOR_STATE(state, path, rng_offset),
120  INTEGRATOR_STATE(state, path, sample),
121  0x64c6a40e);
122 
123  /* Sample normals from surrounding points on surface. */
124  float3 sum_N = make_float3(0.0f, 0.0f, 0.0f);
125 
126  /* TODO: support ray-tracing in shadow shader evaluation? */
127  RNGState rng_state;
128  path_state_rng_load(state, &rng_state);
129 
130  for (int sample = 0; sample < num_samples; sample++) {
131  float disk_u, disk_v;
132  path_branched_rng_2D(kg, &rng_state, sample, num_samples, PRNG_BEVEL_U, &disk_u, &disk_v);
133 
134  /* Pick random axis in local frame and point on disk. */
135  float3 disk_N, disk_T, disk_B;
136  float pick_pdf_N, pick_pdf_T, pick_pdf_B;
137 
138  disk_N = sd->Ng;
139  make_orthonormals(disk_N, &disk_T, &disk_B);
140 
141  float axisu = disk_u;
142 
143  if (axisu < 0.5f) {
144  pick_pdf_N = 0.5f;
145  pick_pdf_T = 0.25f;
146  pick_pdf_B = 0.25f;
147  disk_u *= 2.0f;
148  }
149  else if (axisu < 0.75f) {
150  float3 tmp = disk_N;
151  disk_N = disk_T;
152  disk_T = tmp;
153  pick_pdf_N = 0.25f;
154  pick_pdf_T = 0.5f;
155  pick_pdf_B = 0.25f;
156  disk_u = (disk_u - 0.5f) * 4.0f;
157  }
158  else {
159  float3 tmp = disk_N;
160  disk_N = disk_B;
161  disk_B = tmp;
162  pick_pdf_N = 0.25f;
163  pick_pdf_T = 0.25f;
164  pick_pdf_B = 0.5f;
165  disk_u = (disk_u - 0.75f) * 4.0f;
166  }
167 
168  /* Sample point on disk. */
169  float phi = M_2PI_F * disk_u;
170  float disk_r = disk_v;
171  float disk_height;
172 
173  /* Perhaps find something better than Cubic BSSRDF, but happens to work well. */
174  svm_bevel_cubic_sample(radius, disk_r, &disk_r, &disk_height);
175 
176  float3 disk_P = (disk_r * cosf(phi)) * disk_T + (disk_r * sinf(phi)) * disk_B;
177 
178  /* Create ray. */
180  ray.P = sd->P + disk_N * disk_height + disk_P;
181  ray.D = -disk_N;
182  ray.tmin = 0.0f;
183  ray.tmax = 2.0f * disk_height;
184  ray.dP = differential_zero_compact();
185  ray.dD = differential_zero_compact();
186  ray.time = sd->time;
187  ray.self.object = OBJECT_NONE;
188  ray.self.prim = PRIM_NONE;
189  ray.self.light_object = OBJECT_NONE;
190  ray.self.light_prim = PRIM_NONE;
191 
192  /* Intersect with the same object. if multiple intersections are found it
193  * will use at most LOCAL_MAX_HITS hits, a random subset of all hits. */
194  scene_intersect_local(kg, &ray, &isect, sd->object, &lcg_state, LOCAL_MAX_HITS);
195 
196  int num_eval_hits = min(isect.num_hits, LOCAL_MAX_HITS);
197 
198  for (int hit = 0; hit < num_eval_hits; hit++) {
199  /* Quickly retrieve P and Ng without setting up ShaderData. */
200  float3 hit_P;
201  if (sd->type == PRIMITIVE_TRIANGLE) {
202  hit_P = triangle_point_from_uv(kg,
203  sd,
204  isect.hits[hit].object,
205  isect.hits[hit].prim,
206  isect.hits[hit].u,
207  isect.hits[hit].v);
208  }
209 # ifdef __OBJECT_MOTION__
210  else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
211  float3 verts[3];
212  motion_triangle_vertices(kg, sd->object, isect.hits[hit].prim, sd->time, verts);
214  sd,
215  isect.hits[hit].object,
216  isect.hits[hit].prim,
217  isect.hits[hit].u,
218  isect.hits[hit].v,
219  verts);
220  }
221 # endif /* __OBJECT_MOTION__ */
222 
223  /* Get geometric normal. */
224  float3 hit_Ng = isect.Ng[hit];
225  int object = isect.hits[hit].object;
226  int object_flag = kernel_data_fetch(object_flag, object);
227  if (object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
228  hit_Ng = -hit_Ng;
229  }
230 
231  /* Compute smooth normal. */
232  float3 N = hit_Ng;
233  int prim = isect.hits[hit].prim;
234  int shader = kernel_data_fetch(tri_shader, prim);
235 
236  if (shader & SHADER_SMOOTH_NORMAL) {
237  float u = isect.hits[hit].u;
238  float v = isect.hits[hit].v;
239 
240  if (sd->type == PRIMITIVE_TRIANGLE) {
241  N = triangle_smooth_normal(kg, N, prim, u, v);
242  }
243 # ifdef __OBJECT_MOTION__
244  else if (sd->type == PRIMITIVE_MOTION_TRIANGLE) {
245  N = motion_triangle_smooth_normal(kg, N, sd->object, prim, u, v, sd->time);
246  }
247 # endif /* __OBJECT_MOTION__ */
248  }
249 
250  /* Transform normals to world space. */
251  if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
252  object_normal_transform(kg, sd, &N);
253  object_normal_transform(kg, sd, &hit_Ng);
254  }
255 
256  /* Probability densities for local frame axes. */
257  float pdf_N = pick_pdf_N * fabsf(dot(disk_N, hit_Ng));
258  float pdf_T = pick_pdf_T * fabsf(dot(disk_T, hit_Ng));
259  float pdf_B = pick_pdf_B * fabsf(dot(disk_B, hit_Ng));
260 
261  /* Multiple importance sample between 3 axes, power heuristic
262  * found to be slightly better than balance heuristic. pdf_N
263  * in the MIS weight and denominator canceled out. */
264  float w = pdf_N / (sqr(pdf_N) + sqr(pdf_T) + sqr(pdf_B));
265  if (isect.num_hits > LOCAL_MAX_HITS) {
266  w *= isect.num_hits / (float)LOCAL_MAX_HITS;
267  }
268 
269  /* Real distance to sampled point. */
270  float r = len(hit_P - sd->P);
271 
272  /* Compute weight. */
273  float pdf = svm_bevel_cubic_pdf(radius, r);
274  float disk_pdf = svm_bevel_cubic_pdf(radius, disk_r);
275 
276  w *= pdf / disk_pdf;
277 
278  /* Sum normal and weight. */
279  sum_N += w * N;
280  }
281  }
282 
283  /* Normalize. */
284  float3 N = safe_normalize(sum_N);
285  return is_zero(N) ? sd->N : (sd->flag & SD_BACKFACING) ? -N : N;
286 }
287 
288 template<uint node_feature_mask, typename ConstIntegratorGenericState>
289 # if defined(__KERNEL_OPTIX__)
291 # else
293 # endif
294  void
295  svm_node_bevel(KernelGlobals kg,
296  ConstIntegratorGenericState state,
298  ccl_private float *stack,
299  uint4 node)
300 {
301  uint num_samples, radius_offset, normal_offset, out_offset;
302  svm_unpack_node_uchar4(node.y, &num_samples, &radius_offset, &normal_offset, &out_offset);
303 
304  float3 bevel_N = sd->N;
305 
306  IF_KERNEL_NODES_FEATURE(RAYTRACE)
307  {
308  float radius = stack_load_float(stack, radius_offset);
309 
310 # ifdef __KERNEL_OPTIX__
311  bevel_N = optixDirectCall<float3>(1, kg, state, sd, radius, num_samples);
312 # else
313  bevel_N = svm_bevel(kg, state, sd, radius, num_samples);
314 # endif
315 
316  if (stack_valid(normal_offset)) {
317  /* Preserve input normal. */
318  float3 ref_N = stack_load_float3(stack, normal_offset);
319  bevel_N = normalize(ref_N + (bevel_N - sd->N));
320  }
321  }
322 
323  stack_store_float3(stack, out_offset, bevel_N);
324 }
325 
326 #endif /* __SHADER_RAYTRACE__ */
327 
typedef float(TangentPoint)[2]
MINLINE float safe_sqrtf(float a)
unsigned int uint
Definition: BLI_sys_types.h:67
_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 x2
ATTR_WARN_UNUSED_RESULT const BMVert * v
SIMD_FORCE_INLINE const btScalar & w() const
Return the w value.
Definition: btQuadWord.h:119
#define ccl_device_forceinline
Definition: cuda/compat.h:35
#define ccl_optional_struct_init
Definition: cuda/compat.h:53
#define sinf(x)
Definition: cuda/compat.h:102
#define cosf(x)
Definition: cuda/compat.h:101
#define ccl_device
Definition: cuda/compat.h:32
#define ccl_private
Definition: cuda/compat.h:48
#define ccl_device_inline
Definition: cuda/compat.h:34
#define ccl_device_noinline
Definition: cuda/compat.h:40
#define CCL_NAMESPACE_END
Definition: cuda/compat.h:9
OperationNode * node
#define kernel_data
const KernelGlobalsCPU *ccl_restrict KernelGlobals
#define kernel_data_fetch(name, index)
ccl_device_forceinline float differential_zero_compact()
Definition: differential.h:112
int len
Definition: draw_manager.c:108
static float verts[][3]
const int state
ccl_gpu_kernel_postfix ccl_global float int int int int ccl_global const float int int int int int int int int int int int int num_samples
ccl_device_inline void object_normal_transform(KernelGlobals kg, ccl_private const ShaderData *sd, ccl_private float3 *N)
ccl_device_inline void stack_store_float3(ccl_private float *stack, uint a, float3 f)
CCL_NAMESPACE_BEGIN ccl_device_inline float3 stack_load_float3(ccl_private float *stack, uint a)
ccl_device_inline float stack_load_float(ccl_private float *stack, uint a)
ccl_device_forceinline void svm_unpack_node_uchar4(uint i, ccl_private uint *x, ccl_private uint *y, ccl_private uint *z, ccl_private uint *w)
ccl_device_inline bool stack_valid(uint a)
@ SD_BACKFACING
Definition: kernel/types.h:738
@ PRIMITIVE_MOTION_TRIANGLE
Definition: kernel/types.h:559
@ PRIMITIVE_TRIANGLE
Definition: kernel/types.h:551
#define IF_KERNEL_NODES_FEATURE(feature)
@ PRNG_BEVEL_U
Definition: kernel/types.h:174
#define PRIM_NONE
Definition: kernel/types.h:41
#define OBJECT_NONE
Definition: kernel/types.h:40
#define LOCAL_MAX_HITS
Definition: kernel/types.h:33
ShaderData
Definition: kernel/types.h:925
@ SHADER_SMOOTH_NORMAL
Definition: kernel/types.h:435
@ SD_OBJECT_TRANSFORM_APPLIED
Definition: kernel/types.h:808
@ SD_OBJECT_NEGATIVE_SCALE_APPLIED
Definition: kernel/types.h:810
@ BVH_LAYOUT_NONE
ccl_device_inline uint lcg_state_init(const uint rng_hash, const uint rng_offset, const uint sample, const uint scramble)
Definition: lcg.h:33
ccl_device_inline float2 safe_normalize(const float2 &a)
Definition: math_float2.h:201
#define N
#define fabsf(x)
Definition: metal/compat.h:219
#define __device__
Definition: metal/compat.h:248
#define make_float3(x, y, z)
Definition: metal/compat.h:204
ccl_device_inline float3 motion_triangle_smooth_normal(KernelGlobals kg, float3 Ng, int object, int prim, float u, float v, float time)
ccl_device_inline void motion_triangle_vertices(KernelGlobals kg, int object, int prim, float time, float3 verts[3])
CCL_NAMESPACE_BEGIN ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg, ccl_private ShaderData *sd, const int isect_object, const int isect_prim, const float u, const float v, float3 verts[3])
T dot(const vec_base< T, Size > &a, const vec_base< T, Size > &b)
vec_base< T, Size > normalize(const vec_base< T, Size > &v)
bool is_zero(const T &a)
ccl_device_inline void path_branched_rng_2D(KernelGlobals kg, ccl_private const RNGState *rng_state, int branch, int num_branches, int dimension, ccl_private float *fx, ccl_private float *fy)
Definition: path_state.h:340
ccl_device_inline void path_state_rng_load(ConstIntegratorState state, ccl_private RNGState *rng_state)
Definition: path_state.h:283
#define min(a, b)
Definition: sort.c:35
const IntegratorStateCPU *ccl_restrict ConstIntegratorState
Definition: state.h:148
#define INTEGRATOR_STATE(state, nested_struct, member)
Definition: state.h:154
struct Intersection hits[LOCAL_MAX_HITS]
Definition: kernel/types.h:973
float3 Ng[LOCAL_MAX_HITS]
Definition: kernel/types.h:974
ccl_device_inline float3 triangle_smooth_normal(KernelGlobals kg, float3 Ng, int prim, float u, float v)
Definition: triangle.h:92
ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg, ccl_private ShaderData *sd, const int isect_object, const int isect_prim, const float u, const float v)
ccl_device_inline float sqr(float a)
Definition: util/math.h:746
#define M_2PI_F
Definition: util/math.h:60
ccl_device_inline float saturatef(float a)
Definition: util/math.h:404
#define M_PI_F
Definition: util/math.h:34
ccl_device_inline void make_orthonormals(const float3 N, ccl_private float3 *a, ccl_private float3 *b)
Definition: util/math.h:566