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