Blender  V3.3
kernel/device/optix/bvh.h
Go to the documentation of this file.
1 /* SPDX-License-Identifier: Apache-2.0
2  * Copyright 2021-2022 Blender Foundation */
3 
4 /* OptiX implementation of ray-scene intersection. */
5 
6 #pragma once
7 
8 #include "kernel/bvh/types.h"
9 #include "kernel/bvh/util.h"
10 
11 #define OPTIX_DEFINE_ABI_VERSION_ONLY
12 #include <optix_function_table.h>
13 
15 
16 /* Utilities. */
17 
19 {
20  return pointer_unpack_from_uint<T>(optixGetPayload_0(), optixGetPayload_1());
21 }
23 {
24  return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3());
25 }
26 
28 {
29  return (T *)(((uint64_t)optixGetPayload_7() << 32) | optixGetPayload_6());
30 }
31 
33 {
34 #ifdef __OBJECT_MOTION__
35  /* Always get the instance ID from the TLAS
36  * There might be a motion transform node between TLAS and BLAS which does not have one. */
37  return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
38 #else
39  return optixGetInstanceId();
40 #endif
41 }
42 
43 /* Hit/miss functions. */
44 
45 extern "C" __global__ void __miss__kernel_optix_miss()
46 {
47  /* 'kernel_path_lamp_emission' checks intersection distance, so need to set it even on a miss. */
48  optixSetPayload_0(__float_as_uint(optixGetRayTmax()));
49  optixSetPayload_5(PRIMITIVE_NONE);
50 }
51 
52 extern "C" __global__ void __anyhit__kernel_optix_local_hit()
53 {
54 #if defined(__HAIR__) || defined(__POINTCLOUD__)
55  if (!optixIsTriangleHit()) {
56  /* Ignore curves and points. */
57  return optixIgnoreIntersection();
58  }
59 #endif
60 
61 #ifdef __BVH_LOCAL__
62  const int object = get_object_id();
63  if (object != optixGetPayload_4() /* local_object */) {
64  /* Only intersect with matching object. */
65  return optixIgnoreIntersection();
66  }
67 
68  const int prim = optixGetPrimitiveIndex();
69  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
70  if (intersection_skip_self_local(ray->self, prim)) {
71  return optixIgnoreIntersection();
72  }
73 
74  const uint max_hits = optixGetPayload_5();
75  if (max_hits == 0) {
76  /* Special case for when no hit information is requested, just report that something was hit */
77  optixSetPayload_5(true);
78  return optixTerminateRay();
79  }
80 
81  int hit = 0;
82  uint *const lcg_state = get_payload_ptr_0<uint>();
83  LocalIntersection *const local_isect = get_payload_ptr_2<LocalIntersection>();
84 
85  if (lcg_state) {
86  for (int i = min(max_hits, local_isect->num_hits) - 1; i >= 0; --i) {
87  if (optixGetRayTmax() == local_isect->hits[i].t) {
88  return optixIgnoreIntersection();
89  }
90  }
91 
92  hit = local_isect->num_hits++;
93 
94  if (local_isect->num_hits > max_hits) {
95  hit = lcg_step_uint(lcg_state) % local_isect->num_hits;
96  if (hit >= max_hits) {
97  return optixIgnoreIntersection();
98  }
99  }
100  }
101  else {
102  if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) {
103  /* Record closest intersection only.
104  * Do not terminate ray here, since there is no guarantee about distance ordering in any-hit.
105  */
106  return optixIgnoreIntersection();
107  }
108 
109  local_isect->num_hits = 1;
110  }
111 
112  Intersection *isect = &local_isect->hits[hit];
113  isect->t = optixGetRayTmax();
114  isect->prim = prim;
115  isect->object = get_object_id();
116  isect->type = kernel_data_fetch(objects, isect->object).primitive_type;
117 
118  const float2 barycentrics = optixGetTriangleBarycentrics();
119  isect->u = barycentrics.x;
120  isect->v = barycentrics.y;
121 
122  /* Record geometric normal. */
123  const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w;
124  const float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0);
125  const float3 tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1);
126  const float3 tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
127  local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
128 
129  /* Continue tracing (without this the trace call would return after the first hit). */
130  optixIgnoreIntersection();
131 #endif
132 }
133 
134 extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
135 {
136 #ifdef __SHADOW_RECORD_ALL__
137  int prim = optixGetPrimitiveIndex();
138  const uint object = get_object_id();
139 # ifdef __VISIBILITY_FLAG__
140  const uint visibility = optixGetPayload_4();
141  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
142  return optixIgnoreIntersection();
143  }
144 # endif
145 
146  float u = 0.0f, v = 0.0f;
147  int type = 0;
148  if (optixIsTriangleHit()) {
149  /* Triangle. */
150  const float2 barycentrics = optixGetTriangleBarycentrics();
151  u = barycentrics.x;
152  v = barycentrics.y;
153  type = kernel_data_fetch(objects, object).primitive_type;
154  }
155 # ifdef __HAIR__
156  else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
157  /* Curve. */
158  u = __uint_as_float(optixGetAttribute_0());
159  v = __uint_as_float(optixGetAttribute_1());
160 
161  const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
162  type = segment.type;
163  prim = segment.prim;
164 
165 # if OPTIX_ABI_VERSION < 55
166  /* Filter out curve end-caps. */
167  if (u == 0.0f || u == 1.0f) {
168  return optixIgnoreIntersection();
169  }
170 # endif
171  }
172 # endif
173  else {
174  /* Point. */
175  type = kernel_data_fetch(objects, object).primitive_type;
176  u = 0.0f;
177  v = 0.0f;
178  }
179 
180  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
181  if (intersection_skip_self_shadow(ray->self, object, prim)) {
182  return optixIgnoreIntersection();
183  }
184 
185 # ifndef __TRANSPARENT_SHADOWS__
186  /* No transparent shadows support compiled in, make opaque. */
187  optixSetPayload_5(true);
188  return optixTerminateRay();
189 # else
190  const uint max_hits = optixGetPayload_3();
191  const uint num_hits_packed = optixGetPayload_2();
192  const uint num_recorded_hits = uint16_unpack_from_uint_0(num_hits_packed);
193  const uint num_hits = uint16_unpack_from_uint_1(num_hits_packed);
194 
195  /* If no transparent shadows, all light is blocked and we can stop immediately. */
196  if (num_hits >= max_hits ||
198  optixSetPayload_5(true);
199  return optixTerminateRay();
200  }
201 
202  /* Always use baked shadow transparency for curves. */
203  if (type & PRIMITIVE_CURVE) {
204  float throughput = __uint_as_float(optixGetPayload_1());
205  throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, type, u);
206  optixSetPayload_1(__float_as_uint(throughput));
207  optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits, num_hits + 1));
208 
209  if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) {
210  optixSetPayload_5(true);
211  return optixTerminateRay();
212  }
213  else {
214  /* Continue tracing. */
215  optixIgnoreIntersection();
216  return;
217  }
218  }
219 
220  /* Record transparent intersection. */
221  optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits + 1, num_hits + 1));
222 
223  uint record_index = num_recorded_hits;
224 
225  const IntegratorShadowState state = optixGetPayload_0();
226 
227  const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE);
228  if (record_index >= max_record_hits) {
229  /* If maximum number of hits reached, find a hit to replace. */
230  float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t);
231  uint max_recorded_hit = 0;
232 
233  for (int i = 1; i < max_record_hits; i++) {
234  const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t);
235  if (isect_t > max_recorded_t) {
236  max_recorded_t = isect_t;
237  max_recorded_hit = i;
238  }
239  }
240 
241  if (optixGetRayTmax() >= max_recorded_t) {
242  /* Accept hit, so that OptiX won't consider any more hits beyond the distance of the
243  * current hit anymore. */
244  return;
245  }
246 
247  record_index = max_recorded_hit;
248  }
249 
250  INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u;
251  INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v;
252  INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = optixGetRayTmax();
253  INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim;
254  INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object;
255  INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type;
256 
257  /* Continue tracing. */
258  optixIgnoreIntersection();
259 # endif /* __TRANSPARENT_SHADOWS__ */
260 #endif /* __SHADOW_RECORD_ALL__ */
261 }
262 
263 extern "C" __global__ void __anyhit__kernel_optix_volume_test()
264 {
265 #if defined(__HAIR__) || defined(__POINTCLOUD__)
266  if (!optixIsTriangleHit()) {
267  /* Ignore curves. */
268  return optixIgnoreIntersection();
269  }
270 #endif
271 
272  const uint object = get_object_id();
273 #ifdef __VISIBILITY_FLAG__
274  const uint visibility = optixGetPayload_4();
275  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
276  return optixIgnoreIntersection();
277  }
278 #endif
279 
280  if ((kernel_data_fetch(object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
281  return optixIgnoreIntersection();
282  }
283 
284  const int prim = optixGetPrimitiveIndex();
285  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
286  if (intersection_skip_self(ray->self, object, prim)) {
287  return optixIgnoreIntersection();
288  }
289 }
290 
291 extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
292 {
293 #ifdef __HAIR__
294 # if OPTIX_ABI_VERSION < 55
295  if (optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE) {
296  /* Filter out curve end-caps. */
297  const float u = __uint_as_float(optixGetAttribute_0());
298  if (u == 0.0f || u == 1.0f) {
299  return optixIgnoreIntersection();
300  }
301  }
302 # endif
303 #endif
304 
305  const uint object = get_object_id();
306  const uint visibility = optixGetPayload_4();
307 #ifdef __VISIBILITY_FLAG__
308  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
309  return optixIgnoreIntersection();
310  }
311 #endif
312 
313  int prim = optixGetPrimitiveIndex();
314  if (optixIsTriangleHit()) {
315  /* Triangle. */
316  }
317 #ifdef __HAIR__
318  else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
319  /* Curve. */
320  prim = kernel_data_fetch(curve_segments, prim).prim;
321  }
322 #endif
323 
324  ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
325 
326  if (visibility & PATH_RAY_SHADOW_OPAQUE) {
327  if (intersection_skip_self_shadow(ray->self, object, prim)) {
328  return optixIgnoreIntersection();
329  }
330  else {
331  /* Shadow ray early termination. */
332  return optixTerminateRay();
333  }
334  }
335  else {
336  if (intersection_skip_self(ray->self, object, prim)) {
337  return optixIgnoreIntersection();
338  }
339  }
340 }
341 
342 extern "C" __global__ void __closesthit__kernel_optix_hit()
343 {
344  const int object = get_object_id();
345  const int prim = optixGetPrimitiveIndex();
346 
347  optixSetPayload_0(__float_as_uint(optixGetRayTmax())); /* Intersection distance */
348  optixSetPayload_4(object);
349 
350  if (optixIsTriangleHit()) {
351  const float2 barycentrics = optixGetTriangleBarycentrics();
352  optixSetPayload_1(__float_as_uint(barycentrics.x));
353  optixSetPayload_2(__float_as_uint(barycentrics.y));
354  optixSetPayload_3(prim);
355  optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type);
356  }
357  else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) {
358  const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
359  optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
360  optixSetPayload_2(optixGetAttribute_1());
361  optixSetPayload_3(segment.prim);
362  optixSetPayload_5(segment.type);
363  }
364  else {
365  optixSetPayload_1(0);
366  optixSetPayload_2(0);
367  optixSetPayload_3(prim);
368  optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type);
369  }
370 }
371 
372 /* Custom primitive intersection functions. */
373 
374 #ifdef __HAIR__
375 ccl_device_inline void optix_intersection_curve(const int prim, const int type)
376 {
377  const int object = get_object_id();
378 
379 # ifdef __VISIBILITY_FLAG__
380  const uint visibility = optixGetPayload_4();
381  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
382  return;
383  }
384 # endif
385 
386  const float3 ray_P = optixGetObjectRayOrigin();
387  const float3 ray_D = optixGetObjectRayDirection();
388  const float ray_tmin = optixGetRayTmin();
389 
390 # ifdef __OBJECT_MOTION__
391  const float time = optixGetRayTime();
392 # else
393  const float time = 0.0f;
394 # endif
395 
396  Intersection isect;
397  isect.t = optixGetRayTmax();
398 
399  if (curve_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) {
400  static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
401  optixReportIntersection(isect.t,
402  type & PRIMITIVE_ALL,
403  __float_as_int(isect.u), /* Attribute_0 */
404  __float_as_int(isect.v)); /* Attribute_1 */
405  }
406 }
407 
408 extern "C" __global__ void __intersection__curve_ribbon()
409 {
410  const KernelCurveSegment segment = kernel_data_fetch(curve_segments, optixGetPrimitiveIndex());
411  const int prim = segment.prim;
412  const int type = segment.type;
413  if (type & PRIMITIVE_CURVE_RIBBON) {
414  optix_intersection_curve(prim, type);
415  }
416 }
417 
418 #endif
419 
420 #ifdef __POINTCLOUD__
421 extern "C" __global__ void __intersection__point()
422 {
423  const int prim = optixGetPrimitiveIndex();
424  const int object = get_object_id();
425  const int type = kernel_data_fetch(objects, object).primitive_type;
426 
427 # ifdef __VISIBILITY_FLAG__
428  const uint visibility = optixGetPayload_4();
429  if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
430  return;
431  }
432 # endif
433 
434  const float3 ray_P = optixGetObjectRayOrigin();
435  const float3 ray_D = optixGetObjectRayDirection();
436  const float ray_tmin = optixGetRayTmin();
437 
438 # ifdef __OBJECT_MOTION__
439  const float time = optixGetRayTime();
440 # else
441  const float time = 0.0f;
442 # endif
443 
444  Intersection isect;
445  isect.t = optixGetRayTmax();
446 
447  if (point_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) {
448  static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
449  optixReportIntersection(isect.t, type & PRIMITIVE_ALL);
450  }
451 }
452 #endif
453 
454 /* Scene intersection. */
455 
456 ccl_device_intersect bool scene_intersect(KernelGlobals kg,
457  ccl_private const Ray *ray,
458  const uint visibility,
459  ccl_private Intersection *isect)
460 {
461  uint p0 = 0;
462  uint p1 = 0;
463  uint p2 = 0;
464  uint p3 = 0;
465  uint p4 = visibility;
466  uint p5 = PRIMITIVE_NONE;
467  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
468  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
469 
470  uint ray_mask = visibility & 0xFF;
471  uint ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
472  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
473  ray_mask = 0xFF;
474  }
475  else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
476  ray_flags |= OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT;
477  }
478 
479  optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
480  ray->P,
481  ray->D,
482  ray->tmin,
483  ray->tmax,
484  ray->time,
485  ray_mask,
486  ray_flags,
487  0, /* SBT offset for PG_HITD */
488  0,
489  0,
490  p0,
491  p1,
492  p2,
493  p3,
494  p4,
495  p5,
496  p6,
497  p7);
498 
499  isect->t = __uint_as_float(p0);
500  isect->u = __uint_as_float(p1);
501  isect->v = __uint_as_float(p2);
502  isect->prim = p3;
503  isect->object = p4;
504  isect->type = p5;
505 
506  return p5 != PRIMITIVE_NONE;
507 }
508 
509 #ifdef __BVH_LOCAL__
510 ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
511  ccl_private const Ray *ray,
512  ccl_private LocalIntersection *local_isect,
513  int local_object,
514  ccl_private uint *lcg_state,
515  int max_hits)
516 {
517  uint p0 = pointer_pack_to_uint_0(lcg_state);
518  uint p1 = pointer_pack_to_uint_1(lcg_state);
519  uint p2 = pointer_pack_to_uint_0(local_isect);
520  uint p3 = pointer_pack_to_uint_1(local_isect);
521  uint p4 = local_object;
522  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
523  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
524 
525  /* Is set to zero on miss or if ray is aborted, so can be used as return value. */
526  uint p5 = max_hits;
527 
528  if (local_isect) {
529  local_isect->num_hits = 0; /* Initialize hit count to zero. */
530  }
531  optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
532  ray->P,
533  ray->D,
534  ray->tmin,
535  ray->tmax,
536  ray->time,
537  0xFF,
538  /* Need to always call into __anyhit__kernel_optix_local_hit. */
539  OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
540  2, /* SBT offset for PG_HITL */
541  0,
542  0,
543  p0,
544  p1,
545  p2,
546  p3,
547  p4,
548  p5,
549  p6,
550  p7);
551 
552  return p5;
553 }
554 #endif
555 
556 #ifdef __SHADOW_RECORD_ALL__
557 ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
558  IntegratorShadowState state,
559  ccl_private const Ray *ray,
560  uint visibility,
561  uint max_hits,
562  ccl_private uint *num_recorded_hits,
563  ccl_private float *throughput)
564 {
565  uint p0 = state;
566  uint p1 = __float_as_uint(1.0f); /* Throughput. */
567  uint p2 = 0; /* Number of hits. */
568  uint p3 = max_hits;
569  uint p4 = visibility;
570  uint p5 = false;
571  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
572  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
573 
574  uint ray_mask = visibility & 0xFF;
575  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
576  ray_mask = 0xFF;
577  }
578 
579  optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
580  ray->P,
581  ray->D,
582  ray->tmin,
583  ray->tmax,
584  ray->time,
585  ray_mask,
586  /* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */
587  OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
588  1, /* SBT offset for PG_HITS */
589  0,
590  0,
591  p0,
592  p1,
593  p2,
594  p3,
595  p4,
596  p5,
597  p6,
598  p7);
599 
600  *num_recorded_hits = uint16_unpack_from_uint_0(p2);
601  *throughput = __uint_as_float(p1);
602 
603  return p5;
604 }
605 #endif
606 
607 #ifdef __VOLUME__
608 ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
609  ccl_private const Ray *ray,
610  ccl_private Intersection *isect,
611  const uint visibility)
612 {
613  uint p0 = 0;
614  uint p1 = 0;
615  uint p2 = 0;
616  uint p3 = 0;
617  uint p4 = visibility;
618  uint p5 = PRIMITIVE_NONE;
619  uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
620  uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
621 
622  uint ray_mask = visibility & 0xFF;
623  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
624  ray_mask = 0xFF;
625  }
626 
627  optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
628  ray->P,
629  ray->D,
630  ray->tmin,
631  ray->tmax,
632  ray->time,
633  ray_mask,
634  /* Need to always call into __anyhit__kernel_optix_volume_test. */
635  OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
636  3, /* SBT offset for PG_HITV */
637  0,
638  0,
639  p0,
640  p1,
641  p2,
642  p3,
643  p4,
644  p5,
645  p6,
646  p7);
647 
648  isect->t = __uint_as_float(p0);
649  isect->u = __uint_as_float(p1);
650  isect->v = __uint_as_float(p2);
651  isect->prim = p3;
652  isect->object = p4;
653  isect->type = p5;
654 
655  return p5 != PRIMITIVE_NONE;
656 }
657 #endif
658 
659 CCL_NAMESPACE_END
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 type
_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 t
#define C
Definition: RandGen.cpp:25
ATTR_WARN_UNUSED_RESULT const BMVert * v
#define ccl_device_forceinline
Definition: cuda/compat.h:35
#define ccl_private
Definition: cuda/compat.h:48
#define ccl_device_inline
Definition: cuda/compat.h:34
ccl_device_forceinline int intersection_get_shader_flags(KernelGlobals kg, const int prim, const int type)
ccl_device_inline bool intersection_skip_self_local(ccl_private const RaySelfPrimitives &self, const int prim)
#define CURVE_SHADOW_TRANSPARENCY_CUTOFF
ccl_device_inline float intersection_curve_shadow_transparency(KernelGlobals kg, const int object, const int prim, const int type, const float u)
ccl_device_inline bool intersection_skip_self_shadow(ccl_private const RaySelfPrimitives &self, const int object, const int prim)
ccl_device_inline bool intersection_skip_self(ccl_private const RaySelfPrimitives &self, const int object, const int prim)
double time
#define kernel_data_fetch(name, index)
const int state
CCL_NAMESPACE_BEGIN ccl_device_forceinline T * get_payload_ptr_0()
__global__ void __anyhit__kernel_optix_volume_test()
__global__ void __miss__kernel_optix_miss()
__global__ void __anyhit__kernel_optix_visibility_test()
__global__ void __anyhit__kernel_optix_local_hit()
ccl_device_forceinline T * get_payload_ptr_6()
ccl_device_forceinline T * get_payload_ptr_2()
__global__ void __anyhit__kernel_optix_shadow_all_hit()
__global__ void __closesthit__kernel_optix_hit()
ccl_device_forceinline int get_object_id()
@ SD_HAS_TRANSPARENT_SHADOW
Definition: kernel/types.h:766
@ PRIMITIVE_MOTION
Definition: kernel/types.h:558
@ PRIMITIVE_NONE
Definition: kernel/types.h:550
@ PRIMITIVE_CURVE
Definition: kernel/types.h:564
@ PRIMITIVE_POINT
Definition: kernel/types.h:554
@ PATH_RAY_SHADOW_OPAQUE
Definition: kernel/types.h:204
#define INTEGRATOR_SHADOW_ISECT_SIZE
Definition: kernel/types.h:53
@ SD_OBJECT_HAS_VOLUME
Definition: kernel/types.h:812
CCL_NAMESPACE_BEGIN ccl_device uint lcg_step_uint(T rng)
Definition: lcg.h:11
#define T
Segment< FEdge *, Vec3r > segment
vec_base< T, 3 > cross(const vec_base< T, 3 > &a, const vec_base< T, 3 > &b)
vec_base< T, Size > normalize(const vec_base< T, Size > &v)
#define min(a, b)
Definition: sort.c:35
#define INTEGRATOR_STATE_ARRAY_WRITE(state, nested_struct, array_index, member)
Definition: state.h:159
IntegratorShadowStateCPU *ccl_restrict IntegratorShadowState
Definition: state.h:149
#define INTEGRATOR_STATE_ARRAY(state, nested_struct, array_index, member)
Definition: state.h:157
unsigned __int64 uint64_t
Definition: stdint.h:90
struct Intersection hits[LOCAL_MAX_HITS]
Definition: kernel/types.h:973
float3 Ng[LOCAL_MAX_HITS]
Definition: kernel/types.h:974
float x
Definition: types_float2.h:15
float y
Definition: types_float2.h:15
ccl_device_inline float __uint_as_float(uint i)
Definition: util/math.h:273
ccl_device_inline uint __float_as_uint(float f)
Definition: util/math.h:263
ccl_device_inline uint uint16_unpack_from_uint_1(const uint i)
Definition: util/math.h:341
ccl_device_inline uint uint16_pack_to_uint(const uint a, const uint b)
Definition: util/math.h:331
ccl_device_inline uint uint16_unpack_from_uint_0(const uint i)
Definition: util/math.h:336