11 #define OPTIX_DEFINE_ABI_VERSION_ONLY
12 #include <optix_function_table.h>
20 return pointer_unpack_from_uint<T>(optixGetPayload_0(), optixGetPayload_1());
24 return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3());
29 return (
T *)(((
uint64_t)optixGetPayload_7() << 32) | optixGetPayload_6());
34 #ifdef __OBJECT_MOTION__
37 return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
39 return optixGetInstanceId();
54 #if defined(__HAIR__) || defined(__POINTCLOUD__)
55 if (!optixIsTriangleHit()) {
57 return optixIgnoreIntersection();
63 if (
object != optixGetPayload_4() ) {
65 return optixIgnoreIntersection();
68 const int prim = optixGetPrimitiveIndex();
71 return optixIgnoreIntersection();
74 const uint max_hits = optixGetPayload_5();
77 optixSetPayload_5(
true);
78 return optixTerminateRay();
82 uint *
const lcg_state = get_payload_ptr_0<uint>();
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();
94 if (local_isect->
num_hits > max_hits) {
96 if (hit >= max_hits) {
97 return optixIgnoreIntersection();
102 if (local_isect->
num_hits && optixGetRayTmax() > local_isect->
hits[0].
t) {
106 return optixIgnoreIntersection();
113 isect->
t = optixGetRayTmax();
118 const float2 barycentrics = optixGetTriangleBarycentrics();
119 isect->
u = barycentrics.
x;
120 isect->
v = barycentrics.
y;
130 optixIgnoreIntersection();
136 #ifdef __SHADOW_RECORD_ALL__
137 int prim = optixGetPrimitiveIndex();
139 # ifdef __VISIBILITY_FLAG__
140 const uint visibility = optixGetPayload_4();
142 return optixIgnoreIntersection();
146 float u = 0.0f,
v = 0.0f;
148 if (optixIsTriangleHit()) {
150 const float2 barycentrics = optixGetTriangleBarycentrics();
165 # if OPTIX_ABI_VERSION < 55
167 if (u == 0.0f || u == 1.0f) {
168 return optixIgnoreIntersection();
182 return optixIgnoreIntersection();
185 # ifndef __TRANSPARENT_SHADOWS__
187 optixSetPayload_5(
true);
188 return optixTerminateRay();
190 const uint max_hits = optixGetPayload_3();
191 const uint num_hits_packed = optixGetPayload_2();
196 if (num_hits >= max_hits ||
198 optixSetPayload_5(
true);
199 return optixTerminateRay();
210 optixSetPayload_5(
true);
211 return optixTerminateRay();
215 optixIgnoreIntersection();
223 uint record_index = num_recorded_hits;
228 if (record_index >= max_record_hits) {
231 uint max_recorded_hit = 0;
233 for (
int i = 1; i < max_record_hits; i++) {
235 if (isect_t > max_recorded_t) {
236 max_recorded_t = isect_t;
237 max_recorded_hit = i;
241 if (optixGetRayTmax() >= max_recorded_t) {
247 record_index = max_recorded_hit;
258 optixIgnoreIntersection();
265 #if defined(__HAIR__) || defined(__POINTCLOUD__)
266 if (!optixIsTriangleHit()) {
268 return optixIgnoreIntersection();
273 #ifdef __VISIBILITY_FLAG__
274 const uint visibility = optixGetPayload_4();
276 return optixIgnoreIntersection();
281 return optixIgnoreIntersection();
284 const int prim = optixGetPrimitiveIndex();
287 return optixIgnoreIntersection();
294 # if OPTIX_ABI_VERSION < 55
295 if (optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE) {
298 if (u == 0.0f || u == 1.0f) {
299 return optixIgnoreIntersection();
306 const uint visibility = optixGetPayload_4();
307 #ifdef __VISIBILITY_FLAG__
309 return optixIgnoreIntersection();
313 int prim = optixGetPrimitiveIndex();
314 if (optixIsTriangleHit()) {
328 return optixIgnoreIntersection();
332 return optixTerminateRay();
337 return optixIgnoreIntersection();
345 const int prim = optixGetPrimitiveIndex();
348 optixSetPayload_4(
object);
350 if (optixIsTriangleHit()) {
351 const float2 barycentrics = optixGetTriangleBarycentrics();
354 optixSetPayload_3(prim);
359 optixSetPayload_1(optixGetAttribute_0());
360 optixSetPayload_2(optixGetAttribute_1());
361 optixSetPayload_3(
segment.prim);
362 optixSetPayload_5(
segment.type);
365 optixSetPayload_1(0);
366 optixSetPayload_2(0);
367 optixSetPayload_3(prim);
379 # ifdef __VISIBILITY_FLAG__
380 const uint visibility = optixGetPayload_4();
386 const float3 ray_P = optixGetObjectRayOrigin();
387 const float3 ray_D = optixGetObjectRayDirection();
388 const float ray_tmin = optixGetRayTmin();
390 # ifdef __OBJECT_MOTION__
391 const float time = optixGetRayTime();
393 const float time = 0.0f;
397 isect.
t = optixGetRayTmax();
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 */
408 extern "C" __global__ void __intersection__curve_ribbon()
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);
420 #ifdef __POINTCLOUD__
421 extern "C" __global__ void __intersection__point()
423 const int prim = optixGetPrimitiveIndex();
424 const int object = get_object_id();
425 const int type = kernel_data_fetch(objects, object).primitive_type;
427 # ifdef __VISIBILITY_FLAG__
428 const uint visibility = optixGetPayload_4();
429 if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
434 const float3 ray_P = optixGetObjectRayOrigin();
435 const float3 ray_D = optixGetObjectRayDirection();
436 const float ray_tmin = optixGetRayTmin();
438 # ifdef __OBJECT_MOTION__
439 const float time = optixGetRayTime();
441 const float time = 0.0f;
445 isect.t = optixGetRayTmax();
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);
454 /* Scene intersection. */
456 ccl_device_intersect bool scene_intersect(KernelGlobals kg,
457 ccl_private const Ray *ray,
458 const uint visibility,
459 ccl_private Intersection *isect)
465 uint p4 = visibility;
466 uint p5 = PRIMITIVE_NONE;
467 uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
468 uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
470 uint ray_mask = visibility & 0xFF;
471 uint ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
472 if (0 == ray_mask && (visibility & ~0xFF) != 0) {
475 else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
476 ray_flags |= OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT;
479 optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
487 0, /* SBT offset for PG_HITD */
499 isect->t = __uint_as_float(p0);
500 isect->u = __uint_as_float(p1);
501 isect->v = __uint_as_float(p2);
506 return p5 != PRIMITIVE_NONE;
510 ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
511 ccl_private const Ray *ray,
512 ccl_private LocalIntersection *local_isect,
514 ccl_private uint *lcg_state,
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;
525 /* Is set to zero on miss or if ray is aborted, so can be used as return value. */
529 local_isect->num_hits = 0; /* Initialize hit count to zero. */
531 optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
538 /* Need to always call into __anyhit__kernel_optix_local_hit. */
539 OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
540 2, /* SBT offset for PG_HITL */
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,
562 ccl_private uint *num_recorded_hits,
563 ccl_private float *throughput)
566 uint p1 = __float_as_uint(1.0f); /* Throughput. */
567 uint p2 = 0; /* Number of hits. */
569 uint p4 = visibility;
571 uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
572 uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
574 uint ray_mask = visibility & 0xFF;
575 if (0 == ray_mask && (visibility & ~0xFF) != 0) {
579 optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
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 */
600 *num_recorded_hits = uint16_unpack_from_uint_0(p2);
601 *throughput = __uint_as_float(p1);
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)
617 uint p4 = visibility;
618 uint p5 = PRIMITIVE_NONE;
619 uint p6 = ((uint64_t)ray) & 0xFFFFFFFF;
620 uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF;
622 uint ray_mask = visibility & 0xFF;
623 if (0 == ray_mask && (visibility & ~0xFF) != 0) {
627 optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0,
634 /* Need to always call into __anyhit__kernel_optix_volume_test. */
635 OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
636 3, /* SBT offset for PG_HITV */
648 isect->t = __uint_as_float(p0);
649 isect->u = __uint_as_float(p1);
650 isect->v = __uint_as_float(p2);
655 return p5 != PRIMITIVE_NONE;
_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
ATTR_WARN_UNUSED_RESULT const BMVert * v
#define ccl_device_forceinline
#define ccl_device_inline
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)
#define kernel_data_fetch(name, index)
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
#define INTEGRATOR_SHADOW_ISECT_SIZE
CCL_NAMESPACE_BEGIN ccl_device uint lcg_step_uint(T rng)
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 INTEGRATOR_STATE_ARRAY_WRITE(state, nested_struct, array_index, member)
IntegratorShadowStateCPU *ccl_restrict IntegratorShadowState
#define INTEGRATOR_STATE_ARRAY(state, nested_struct, array_index, member)
unsigned __int64 uint64_t
struct Intersection hits[LOCAL_MAX_HITS]
float3 Ng[LOCAL_MAX_HITS]
ccl_device_inline float __uint_as_float(uint i)
ccl_device_inline uint __float_as_uint(float f)
ccl_device_inline uint uint16_unpack_from_uint_1(const uint i)
ccl_device_inline uint uint16_pack_to_uint(const uint a, const uint b)
ccl_device_inline uint uint16_unpack_from_uint_0(const uint i)