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