Blender V4.5
kernel/device/metal/bvh.h
Go to the documentation of this file.
1/* SPDX-FileCopyrightText: 2021-2022 Blender Foundation
2 *
3 * SPDX-License-Identifier: Apache-2.0 */
4
5/* MetalRT implementation of ray-scene intersection. */
6
7#pragma once
8
9#include "kernel/bvh/types.h"
10#include "kernel/bvh/util.h"
11
12CCL_NAMESPACE_BEGIN
13
14/* Payload types.
15 *
16 * Best practice is to minimize the size of MetalRT payloads to avoid heavy spilling during
17 * intersection tests.
18 */
19
24
25struct MetalRTIntersectionLocalPayload_single_hit {
26 int self_prim;
27#if defined(__METALRT_MOTION__)
28 int self_object;
29#endif
30};
31
32struct MetalRTIntersectionLocalPayload {
33 int self_prim;
34#if defined(__METALRT_MOTION__)
35 int self_object;
36#endif
37 uint lcg_state;
38 uint hit_prim[LOCAL_MAX_HITS];
39 float hit_t[LOCAL_MAX_HITS];
40 float hit_u[LOCAL_MAX_HITS];
41 float hit_v[LOCAL_MAX_HITS];
42 uint max_hits : 3;
43 uint num_hits : 3;
44 uint has_lcg_state : 1;
45};
46static_assert(LOCAL_MAX_HITS < 8,
47 "MetalRTIntersectionLocalPayload max_hits & num_hits bitfields are too small");
48
52
62
63#ifdef __HAIR__
64ccl_device_forceinline bool curve_ribbon_accept(KernelGlobals kg,
65 const float u,
66 float t,
67 const ccl_private Ray *ray,
68 const int object,
69 const int prim,
70 const int type)
71{
72 KernelCurve kcurve = kernel_data_fetch(curves, prim);
73
74 int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
75 int k1 = k0 + 1;
76 int ka = max(k0 - 1, kcurve.first_key);
77 int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
78
79 /* We can ignore motion blur here because we don't need the positions, and it doesn't affect the
80 * radius. */
81 float radius[4];
82 radius[0] = kernel_data_fetch(curve_keys, ka).w;
83 radius[1] = kernel_data_fetch(curve_keys, k0).w;
84 radius[2] = kernel_data_fetch(curve_keys, k1).w;
85 radius[3] = kernel_data_fetch(curve_keys, kb).w;
86 const float r = metal::catmull_rom(u, radius[0], radius[1], radius[2], radius[3]);
87
88 /* MPJ TODO: Can we ignore motion and/or object transforms here? Depends on scaling? */
89 float3 ray_P = ray->P;
90 float3 ray_D = ray->D;
91 if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
92 float3 idir;
93# if defined(__METALRT_MOTION__)
94 bvh_instance_motion_push(nullptr, object, ray, &ray_P, &ray_D, &idir);
95# else
96 bvh_instance_push(nullptr, object, ray, &ray_P, &ray_D, &idir);
97# endif
98 }
99
100 /* ignore self intersections */
101 const float avoidance_factor = 2.0f;
102 return t * len(ray_D) > avoidance_factor * r;
103}
104
105ccl_device_forceinline float curve_ribbon_v(KernelGlobals kg,
106 const float u,
107 float t,
108 const ccl_private Ray *ray,
109 const int object,
110 const int prim,
111 const int type)
112{
113# if defined(__METALRT_MOTION__)
114 float time = ray->time;
115# else
116 float time = 0.0f;
117# endif
118
119 const bool is_motion = (type & PRIMITIVE_MOTION);
120
121 KernelCurve kcurve = kernel_data_fetch(curves, prim);
122
123 int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
124 int k1 = k0 + 1;
125 int ka = max(k0 - 1, kcurve.first_key);
126 int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
127
128 float4 curve[4];
129 if (!is_motion) {
130 curve[0] = kernel_data_fetch(curve_keys, ka);
131 curve[1] = kernel_data_fetch(curve_keys, k0);
132 curve[2] = kernel_data_fetch(curve_keys, k1);
133 curve[3] = kernel_data_fetch(curve_keys, kb);
134 }
135 else {
136 motion_curve_keys(kg, object, time, ka, k0, k1, kb, curve);
137 }
138
139 float3 ray_P = ray->P;
140 float3 ray_D = ray->D;
141 if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
142 float3 idir;
143# if defined(__METALRT_MOTION__)
144 bvh_instance_motion_push(nullptr, object, ray, &ray_P, &ray_D, &idir);
145# else
146 bvh_instance_push(nullptr, object, ray, &ray_P, &ray_D, &idir);
147# endif
148 }
149
150 const float4 P_curve4 = metal::catmull_rom(u, curve[0], curve[1], curve[2], curve[3]);
151 const float r_curve = P_curve4.w;
152
153 float3 P = ray_P + ray_D * t;
154 const float3 P_curve = make_float3(P_curve4);
155
156 const float4 dPdu4 = metal::catmull_rom_derivative(u, curve[0], curve[1], curve[2], curve[3]);
157 const float3 dPdu = make_float3(dPdu4);
158
159 const float3 tangent = normalize(dPdu);
160 const float3 bitangent = normalize(cross(tangent, -ray_D));
161
162 float v = dot(P - P_curve, bitangent) / r_curve;
163 return clamp(v, -1.0, 1.0f);
164}
165#endif /* __HAIR__ */
166
167/* Scene intersection. */
168
169ccl_device_intersect bool scene_intersect(KernelGlobals kg,
170 const ccl_private Ray *ray,
171 const uint visibility,
172 ccl_private Intersection *isect)
173{
174 metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
175 metalrt_intersector_type metalrt_intersect;
176 metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
177 metalrt_intersect.assume_geometry_type(
178 metal::raytracing::geometry_type::triangle |
179 (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
180 metal::raytracing::geometry_type::none) |
181 (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
182 metal::raytracing::geometry_type::none));
183
184 typename metalrt_intersector_type::result_type intersection;
185
186 MetalRTIntersectionPayload payload;
187 payload.self_prim = ray->self.prim;
188 payload.self_object = ray->self.object;
189
190#if defined(__METALRT_MOTION__)
191 intersection = metalrt_intersect.intersect(r,
192 metal_ancillaries->accel_struct,
193 visibility,
194 ray->time,
195 metal_ancillaries->ift_default,
196 payload);
197#else
198 intersection = metalrt_intersect.intersect(
199 r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_default, payload);
200#endif
201
202 if (intersection.type == intersection_type::none) {
203 isect->t = ray->tmax;
204 isect->type = PRIMITIVE_NONE;
205
206 return false;
207 }
208
209 isect->object = intersection.instance_id;
210 isect->t = intersection.distance;
211 if (intersection.type == intersection_type::triangle) {
212 isect->prim = intersection.primitive_id + intersection.user_instance_id;
213 isect->type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
214 isect->u = intersection.triangle_barycentric_coord.x;
215 isect->v = intersection.triangle_barycentric_coord.y;
216 }
217#ifdef __HAIR__
218 else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) {
219 int prim = intersection.primitive_id + intersection.user_instance_id;
220 const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
221 isect->prim = segment.prim;
222 isect->type = segment.type;
223 isect->u = intersection.curve_parameter;
224
225 if (segment.type & PRIMITIVE_CURVE_RIBBON) {
226 isect->v = curve_ribbon_v(kg,
227 intersection.curve_parameter,
228 intersection.distance,
229 ray,
230 intersection.instance_id,
231 segment.prim,
232 segment.type);
233 }
234 else {
235 isect->v = 0.0f;
236 }
237 }
238#endif /* __HAIR__ */
239#ifdef __POINTCLOUD__
240 else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) {
241 const int object = intersection.instance_id;
242 const uint prim = intersection.primitive_id + intersection.user_instance_id;
243 const int prim_type = kernel_data_fetch(objects, object).primitive_type;
244
245 if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
246 float3 idir;
247# if defined(__METALRT_MOTION__)
248 bvh_instance_motion_push(nullptr, object, ray, &r.origin, &r.direction, &idir);
249# else
250 bvh_instance_push(nullptr, object, ray, &r.origin, &r.direction, &idir);
251# endif
252 }
253
254 if (prim_type & PRIMITIVE_POINT) {
255 if (!point_intersect(nullptr,
256 isect,
257 r.origin,
258 r.direction,
259 ray->tmin,
260 ray->tmax,
261 object,
262 prim,
263 ray->time,
264 prim_type))
265 {
266 /* Shouldn't get here */
267 kernel_assert(!"Intersection mismatch");
268 isect->t = ray->tmax;
269 isect->type = PRIMITIVE_NONE;
270 return false;
271 }
272 return true;
273 }
274 }
275#endif /* __POINTCLOUD__ */
276
277 return true;
278}
279
280ccl_device_intersect bool scene_intersect_shadow(KernelGlobals kg,
281 const ccl_private Ray *ray,
282 const uint visibility)
283{
284 metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
285 metalrt_intersector_type metalrt_intersect;
286 metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
287 metalrt_intersect.assume_geometry_type(
288 metal::raytracing::geometry_type::triangle |
289 (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
290 metal::raytracing::geometry_type::none) |
291 (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
292 metal::raytracing::geometry_type::none));
293
294 typename metalrt_intersector_type::result_type intersection;
295
296 metalrt_intersect.accept_any_intersection(true);
297
298 MetalRTIntersectionShadowPayload payload;
299 payload.self = ray->self;
300
301#if defined(__METALRT_MOTION__)
302 intersection = metalrt_intersect.intersect(r,
303 metal_ancillaries->accel_struct,
304 visibility,
305 ray->time,
306 metal_ancillaries->ift_shadow,
307 payload);
308#else
309 intersection = metalrt_intersect.intersect(
310 r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_shadow, payload);
311#endif
312 return (intersection.type != intersection_type::none);
313}
314
315#ifdef __BVH_LOCAL__
316template<bool single_hit = false>
317ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
318 const ccl_private Ray *ray,
319 ccl_private LocalIntersection *local_isect,
320 const int local_object,
321 ccl_private uint *lcg_state,
322 const int max_hits)
323{
324 uint primitive_id_offset = kernel_data_fetch(object_prim_offset, local_object);
325
326 metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
327
328# if defined(__METALRT_MOTION__)
329 metalrt_intersector_type metalrt_intersect;
330 typename metalrt_intersector_type::result_type intersection;
331# else
332 metalrt_blas_intersector_type metalrt_intersect;
333 typename metalrt_blas_intersector_type::result_type intersection;
334
335 if (!(kernel_data_fetch(object_flag, local_object) & SD_OBJECT_TRANSFORM_APPLIED)) {
336 /* Transform the ray into object's local space. */
337 Transform itfm = kernel_data_fetch(objects, local_object).itfm;
338 r.origin = transform_point(&itfm, r.origin);
339 r.direction = transform_direction(&itfm, r.direction);
340 }
341# endif
342
343 metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
344
345 if (single_hit) {
346 MetalRTIntersectionLocalPayload_single_hit payload;
347 payload.self_prim = ray->self.prim - primitive_id_offset;
348
349# if defined(__METALRT_MOTION__)
350 /* We can't skip over the top-level BVH in the motion blur case, so still need to do
351 * the self-object check. */
352 payload.self_object = local_object;
353 metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
354 intersection = metalrt_intersect.intersect(r,
355 metal_ancillaries->accel_struct,
356 ~0,
357 ray->time,
358 metal_ancillaries->ift_local_single_hit_mblur,
359 payload);
360# else
361 /* We only need custom intersection filtering (i.e. non_opaque) if we are performing a
362 * self-primitive intersection check. */
363 metalrt_intersect.force_opacity((ray->self.prim == PRIM_NONE) ?
364 metal::raytracing::forced_opacity::opaque :
365 metal::raytracing::forced_opacity::non_opaque);
366 intersection = metalrt_intersect.intersect(
367 r,
368 metal_ancillaries->blas_accel_structs[local_object].blas,
369 metal_ancillaries->ift_local_single_hit,
370 payload);
371# endif
372
373 if (intersection.type == intersection_type::none) {
374 local_isect->num_hits = 0;
375 return false;
376 }
377
378 uint prim = intersection.primitive_id + primitive_id_offset;
379 int prim_type = kernel_data_fetch(objects, local_object).primitive_type;
380
381 local_isect->num_hits = 1;
382 local_isect->hits[0].prim = prim;
383 local_isect->hits[0].type = prim_type;
384 local_isect->hits[0].object = local_object;
385 local_isect->hits[0].u = intersection.triangle_barycentric_coord.x;
386 local_isect->hits[0].v = intersection.triangle_barycentric_coord.y;
387 local_isect->hits[0].t = intersection.distance;
388
389 const packed_uint3 tri_vindex = kernel_data_fetch(tri_vindex, prim);
390 const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex.x));
391 const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex.y));
392 const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex.z));
393 local_isect->Ng[0] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
394 return true;
395 }
396 else {
397 MetalRTIntersectionLocalPayload payload;
398 payload.self_prim = ray->self.prim - primitive_id_offset;
399 payload.max_hits = max_hits;
400 payload.num_hits = 0;
401 if (lcg_state) {
402 payload.has_lcg_state = 1;
403 payload.lcg_state = *lcg_state;
404 }
405
406 metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
407
408# if defined(__METALRT_MOTION__)
409 /* We can't skip over the top-level BVH in the motion blur case, so still need to do
410 * the self-object check. */
411 payload.self_object = local_object;
412 intersection = metalrt_intersect.intersect(r,
413 metal_ancillaries->accel_struct,
414 ~0,
415 ray->time,
416 metal_ancillaries->ift_local_mblur,
417 payload);
418# else
419 intersection = metalrt_intersect.intersect(
420 r,
421 metal_ancillaries->blas_accel_structs[local_object].blas,
422 metal_ancillaries->ift_local,
423 payload);
424# endif
425
426 if (max_hits == 0) {
427 /* Special case for when no hit information is requested, just report that something was hit
428 */
429 return (intersection.type != intersection_type::none);
430 }
431
432 if (lcg_state) {
433 *lcg_state = payload.lcg_state;
434 }
435
436 const int num_hits = payload.num_hits;
437 if (local_isect) {
438
439 /* Record geometric normal */
440 int prim_type = kernel_data_fetch(objects, local_object).primitive_type;
441
442 local_isect->num_hits = num_hits;
443 for (int hit = 0; hit < num_hits; hit++) {
444 uint prim = payload.hit_prim[hit] + primitive_id_offset;
445 local_isect->hits[hit].prim = prim;
446 local_isect->hits[hit].t = payload.hit_t[hit];
447 local_isect->hits[hit].u = payload.hit_u[hit];
448 local_isect->hits[hit].v = payload.hit_v[hit];
449 local_isect->hits[hit].object = local_object;
450 local_isect->hits[hit].type = prim_type;
451
452 const packed_uint3 tri_vindex = kernel_data_fetch(tri_vindex, prim);
453 const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex.x));
454 const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex.y));
455 const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex.z));
456 local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
457 }
458 }
459 return num_hits > 0;
460 }
461}
462#endif
463
464#ifdef __SHADOW_RECORD_ALL__
465ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
466 IntegratorShadowState state,
467 const ccl_private Ray *ray,
468 const uint visibility,
469 const uint max_hits,
470 ccl_private uint *num_recorded_hits,
471 ccl_private float *throughput)
472{
473 metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
474 metalrt_intersector_type metalrt_intersect;
475 metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
476 metalrt_intersect.assume_geometry_type(
477 metal::raytracing::geometry_type::triangle |
478 (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
479 metal::raytracing::geometry_type::none) |
480 (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
481 metal::raytracing::geometry_type::none));
482
483 MetalRTIntersectionShadowAllPayload payload;
484 payload.self = ray->self;
485 payload.max_hits = max_hits;
486 payload.num_hits = 0;
487 payload.num_recorded_hits = 0;
488 payload.throughput = 1.0f;
489 payload.result = false;
490 payload.state = state;
491
492 typename metalrt_intersector_type::result_type intersection;
493
494# if defined(__METALRT_MOTION__)
495 intersection = metalrt_intersect.intersect(r,
496 metal_ancillaries->accel_struct,
497 visibility,
498 ray->time,
499 metal_ancillaries->ift_shadow_all,
500 payload);
501# else
502 intersection = metalrt_intersect.intersect(
503 r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_shadow_all, payload);
504# endif
505
506 *num_recorded_hits = payload.num_recorded_hits;
507 *throughput = payload.throughput;
508
509 return payload.result;
510}
511#endif
512
513#ifdef __VOLUME__
514ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
515 const ccl_private Ray *ray,
516 ccl_private Intersection *isect,
517 const uint visibility)
518{
519 metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
520 metalrt_intersector_type metalrt_intersect;
521 metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
522 metalrt_intersect.set_geometry_cull_mode(metal::raytracing::geometry_cull_mode::bounding_box |
523 metal::raytracing::geometry_cull_mode::curve);
524 metalrt_intersect.assume_geometry_type(
525 metal::raytracing::geometry_type::triangle |
526 (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
527 metal::raytracing::geometry_type::none) |
528 (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
529 metal::raytracing::geometry_type::none));
530
531 MetalRTIntersectionShadowPayload payload;
532 payload.self = ray->self;
533
534 typename metalrt_intersector_type::result_type intersection;
535
536# if defined(__METALRT_MOTION__)
537 intersection = metalrt_intersect.intersect(r,
538 metal_ancillaries->accel_struct,
539 visibility,
540 ray->time,
541 metal_ancillaries->ift_volume,
542 payload);
543# else
544 intersection = metalrt_intersect.intersect(
545 r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_volume, payload);
546# endif
547
548 if (intersection.type == intersection_type::triangle) {
549 isect->prim = intersection.primitive_id + intersection.user_instance_id;
550 isect->type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
551 isect->u = intersection.triangle_barycentric_coord.x;
552 isect->v = intersection.triangle_barycentric_coord.y;
553 isect->object = intersection.instance_id;
554 isect->t = intersection.distance;
555 return true;
556 }
557 return false;
558}
559#endif
560
561CCL_NAMESPACE_END
uint
unsigned int uint
Definition BLI_sys_types.h:64
v
ATTR_WARN_UNUSED_RESULT const BMVert * v
Definition bmesh_query_inline.hh:23
dot
dot(value.rgb, luminance_coefficients)") DEFINE_VALUE("REDUCE(lhs
kernel_assert
#define kernel_assert(cond)
Definition device/cpu/compat.h:20
kernel_data
#define kernel_data
Definition device/cpu/globals.h:97
ccl_device_forceinline
#define ccl_device_forceinline
kernel_data_fetch
#define kernel_data_fetch(name, index)
Definition device/cpu/globals.h:95
PRIM_NONE
#define PRIM_NONE
PRIMITIVE_UNPACK_SEGMENT
#define PRIMITIVE_UNPACK_SEGMENT(type)
LOCAL_MAX_HITS
#define LOCAL_MAX_HITS
ccl_private
#define ccl_private
KernelGlobals
const ThreadKernelGlobalsCPU * KernelGlobals
Definition device/cpu/globals.h:92
CCL_NAMESPACE_END
#define CCL_NAMESPACE_END
Definition device/cuda/compat.h:10
make_float3
ccl_device_forceinline float3 make_float3(const float x, const float y, const float z)
Definition device/metal/compat.h:229
float4
VecBase< float, 4 > float4
Definition gpu_glsl_cpp_stubs.hh:365
normalize
VecBase< float, D > normalize(VecOp< float, D >) RET
float3
VecBase< float, 3 > float3
Definition gpu_glsl_cpp_stubs.hh:364
cross
VecBase< float, 3 > cross(VecOp< float, 3 >, VecOp< float, 3 >) RET
clamp
constexpr T clamp(T, U, U) RET
ccl_device_intersect
#define ccl_device_intersect
Definition kernel/bvh/types.h:13
scene_intersect
ccl_device_intersect bool scene_intersect(KernelGlobals kg, const ccl_private Ray *ray, const uint visibility, ccl_private Intersection *isect)
Definition kernel/device/metal/bvh.h:169
scene_intersect_shadow
ccl_device_intersect bool scene_intersect_shadow(KernelGlobals kg, const ccl_private Ray *ray, const uint visibility)
Definition kernel/device/metal/bvh.h:280
bvh_instance_push
ccl_device_inline void bvh_instance_push(KernelGlobals kg, const int object, const ccl_private Ray *ray, ccl_private float3 *P, ccl_private float3 *dir, ccl_private float3 *idir)
Definition kernel/geom/object.h:460
PRIMITIVE_MOTION
@ PRIMITIVE_MOTION
Definition kernel/types.h:813
PRIMITIVE_NONE
@ PRIMITIVE_NONE
Definition kernel/types.h:805
PRIMITIVE_CURVE_RIBBON
@ PRIMITIVE_CURVE_RIBBON
Definition kernel/types.h:808
PRIMITIVE_POINT
@ PRIMITIVE_POINT
Definition kernel/types.h:809
SD_OBJECT_TRANSFORM_APPLIED
@ SD_OBJECT_TRANSFORM_APPLIED
Definition kernel/types.h:1080
state
static ulong state[N]
Definition mathutils_noise.cc:56
CCL_NAMESPACE_BEGIN
Definition python.cpp:37
Freestyle::intersection
Intersection< segment > intersection
Definition ViewMapBuilder.cpp:2307
min
#define min(a, b)
Definition sort.cc:36
IntegratorShadowState
IntegratorShadowStateCPU * IntegratorShadowState
Definition state.h:230
MetalRTIntersectionLocalPayload_single_hit
Definition kernel/device/metal/bvh.h:25
MetalRTIntersectionLocalPayload_single_hit::self_prim
int self_prim
Definition kernel/device/metal/bvh.h:26
packed_uint3
Definition types_uint3.h:44
packed_uint3::y
uint y
Definition types_uint3.h:45
packed_uint3::z
uint z
Definition types_uint3.h:45
packed_uint3::x
uint x
Definition types_uint3.h:45
max
max
Definition text_draw.cc:251
transform_direction
ccl_device_inline float3 transform_direction(const ccl_private Transform *t, const float3 a)
Definition transform.h:87
transform_point
ccl_device_inline float3 transform_point(const ccl_private Transform *t, const float3 a)
Definition transform.h:56
P
#define P
Definition uvedit_clipboard_graph_iso.cc:24
len
uint len
Definition uvedit_unwrap_ops.cc:2080
Generated on for Blender by  doxygen 1.15.0