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