diff options
author | Brecht Van Lommel <brecht@blender.org> | 2022-07-25 14:53:48 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2022-07-25 17:34:22 +0300 |
commit | 7a74d91e323c4d695b908ca4178837cee756eeaf (patch) | |
tree | 7e28e29d53b7905c1d464c49f68481a6380c3303 /intern/cycles/kernel | |
parent | c6ce70855a13c42a724755f2989dee756519bef0 (diff) |
Cleanup: move device BVH code to kernel/device/*/bvh.h
Having the OptiX/MetalRT/Embree/MetalRT implementations all in one file with
many #ifdefs became too confusing. Instead split it up per device, and also
move it together with device specific hit/filter/intersect functions and
associated data types.
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r-- | intern/cycles/kernel/CMakeLists.txt | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/bvh/bvh.h | 814 | ||||
-rw-r--r-- | intern/cycles/kernel/bvh/embree.h | 176 | ||||
-rw-r--r-- | intern/cycles/kernel/bvh/metal.h | 37 | ||||
-rw-r--r-- | intern/cycles/kernel/bvh/util.h | 15 | ||||
-rw-r--r-- | intern/cycles/kernel/device/cpu/bvh.h | 609 | ||||
-rw-r--r-- | intern/cycles/kernel/device/metal/bvh.h | 1123 | ||||
-rw-r--r-- | intern/cycles/kernel/device/metal/compat.h | 2 | ||||
-rw-r--r-- | intern/cycles/kernel/device/metal/kernel.metal | 708 | ||||
-rw-r--r-- | intern/cycles/kernel/device/optix/bvh.h | 646 | ||||
-rw-r--r-- | intern/cycles/kernel/device/optix/compat.h | 1 | ||||
-rw-r--r-- | intern/cycles/kernel/device/optix/kernel.cu | 421 | ||||
-rw-r--r-- | intern/cycles/kernel/integrator/intersect_volume_stack.h | 6 | ||||
-rw-r--r-- | intern/cycles/kernel/integrator/subsurface_random_walk.h | 1 | ||||
-rw-r--r-- | intern/cycles/kernel/types.h | 8 |
15 files changed, 2518 insertions, 2054 deletions
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index 21a78722c0d..94632dff200 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -42,6 +42,7 @@ set(SRC_KERNEL_DEVICE_ONEAPI ) set(SRC_KERNEL_DEVICE_CPU_HEADERS + device/cpu/bvh.h device/cpu/compat.h device/cpu/image.h device/cpu/globals.h @@ -71,11 +72,13 @@ set(SRC_KERNEL_DEVICE_HIP_HEADERS ) set(SRC_KERNEL_DEVICE_OPTIX_HEADERS + device/optix/bvh.h device/optix/compat.h device/optix/globals.h ) set(SRC_KERNEL_DEVICE_METAL_HEADERS + device/metal/bvh.h device/metal/compat.h device/metal/context_begin.h device/metal/context_end.h @@ -214,8 +217,6 @@ set(SRC_KERNEL_BVH_HEADERS bvh/util.h bvh/volume.h bvh/volume_all.h - bvh/embree.h - bvh/metal.h ) set(SRC_KERNEL_CAMERA_HEADERS diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index 387e74b9885..bcefe5d970c 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -1,40 +1,46 @@ /* SPDX-License-Identifier: Apache-2.0 * Copyright 2011-2022 Blender Foundation */ -/* BVH - * - * Bounding volume hierarchy for ray tracing. We compile different variations - * of the same BVH traversal function for faster rendering when some types of - * primitives are not needed, using #includes to work around the lack of - * C++ templates in OpenCL. - * - * Originally based on "Understanding the Efficiency of Ray Traversal on GPUs", - * the code has been extended and modified to support more primitives and work - * with CPU/CUDA/OpenCL. */ - #pragma once -#ifdef __EMBREE__ -# include "kernel/bvh/embree.h" -#endif - -#ifdef __METALRT__ -# include "kernel/bvh/metal.h" -#endif - #include "kernel/bvh/types.h" #include "kernel/bvh/util.h" #include "kernel/integrator/state_util.h" +/* Device specific accleration structures for ray tracing. */ + +#if defined(__EMBREE__) +# include "kernel/device/cpu/bvh.h" +#elif defined(__METALRT__) +# include "kernel/device/metal/bvh.h" +#elif defined(__KERNEL_OPTIX__) +# include "kernel/device/optix/bvh.h" +#else +# define __BVH2__ +#endif + CCL_NAMESPACE_BEGIN -#if !defined(__KERNEL_GPU_RAYTRACING__) +#ifdef __BVH2__ -/* Regular BVH traversal */ +/* BVH2 + * + * Bounding volume hierarchy for ray tracing, when no native acceleration + * structure is available for the device. + + * We compile different variations of the same BVH traversal function for + * faster rendering when some types of primitives are not needed, using #includes + * to work around the lack of C++ templates in OpenCL. + * + * Originally based on "Understanding the Efficiency of Ray Traversal on GPUs", + * the code has been extended and modified to support more primitives and work + * with CPU and various GPU kernel languages. */ # include "kernel/bvh/nodes.h" +/* Regular BVH traversal */ + # define BVH_FUNCTION_NAME bvh_intersect # define BVH_FUNCTION_FEATURES BVH_POINTCLOUD # include "kernel/bvh/traversal.h" @@ -57,261 +63,15 @@ CCL_NAMESPACE_BEGIN # include "kernel/bvh/traversal.h" # endif -/* Subsurface scattering BVH traversal */ - -# if defined(__BVH_LOCAL__) -# define BVH_FUNCTION_NAME bvh_intersect_local -# define BVH_FUNCTION_FEATURES BVH_HAIR -# include "kernel/bvh/local.h" - -# if defined(__OBJECT_MOTION__) -# define BVH_FUNCTION_NAME bvh_intersect_local_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_HAIR -# include "kernel/bvh/local.h" -# endif -# endif /* __BVH_LOCAL__ */ - -/* Volume BVH traversal */ - -# if defined(__VOLUME__) -# define BVH_FUNCTION_NAME bvh_intersect_volume -# define BVH_FUNCTION_FEATURES BVH_HAIR -# include "kernel/bvh/volume.h" - -# if defined(__OBJECT_MOTION__) -# define BVH_FUNCTION_NAME bvh_intersect_volume_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_HAIR -# include "kernel/bvh/volume.h" -# endif -# endif /* __VOLUME__ */ - -/* Record all intersections - Shadow BVH traversal */ - -# if defined(__SHADOW_RECORD_ALL__) -# define BVH_FUNCTION_NAME bvh_intersect_shadow_all -# define BVH_FUNCTION_FEATURES BVH_POINTCLOUD -# include "kernel/bvh/shadow_all.h" - -# if defined(__HAIR__) -# define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair -# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD -# include "kernel/bvh/shadow_all.h" -# endif - -# if defined(__OBJECT_MOTION__) -# define BVH_FUNCTION_NAME bvh_intersect_shadow_all_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD -# include "kernel/bvh/shadow_all.h" -# endif - -# if defined(__HAIR__) && defined(__OBJECT_MOTION__) -# define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair_motion -# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD -# include "kernel/bvh/shadow_all.h" -# endif - -# endif /* __SHADOW_RECORD_ALL__ */ - -/* Record all intersections - Volume BVH traversal. */ - -# if defined(__VOLUME_RECORD_ALL__) -# define BVH_FUNCTION_NAME bvh_intersect_volume_all -# define BVH_FUNCTION_FEATURES BVH_HAIR -# include "kernel/bvh/volume_all.h" - -# if defined(__OBJECT_MOTION__) -# define BVH_FUNCTION_NAME bvh_intersect_volume_all_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_HAIR -# include "kernel/bvh/volume_all.h" -# endif -# endif /* __VOLUME_RECORD_ALL__ */ - -# undef BVH_FEATURE -# undef BVH_NAME_JOIN -# undef BVH_NAME_EVAL -# undef BVH_FUNCTION_FULL_NAME - -#endif /* !defined(__KERNEL_GPU_RAYTRACING__) */ - -ccl_device_inline bool scene_intersect_valid(ccl_private const Ray *ray) -{ - /* NOTE: Due to some vectorization code non-finite origin point might - * cause lots of false-positive intersections which will overflow traversal - * stack. - * This code is a quick way to perform early output, to avoid crashes in - * such cases. - * From production scenes so far it seems it's enough to test first element - * only. - * Scene intersection may also called with empty rays for conditional trace - * calls that evaluate to false, so filter those out. - */ - return isfinite_safe(ray->P.x) && isfinite_safe(ray->D.x) && len_squared(ray->D) != 0.0f; -} - ccl_device_intersect bool scene_intersect(KernelGlobals kg, ccl_private const Ray *ray, const uint visibility, ccl_private Intersection *isect) { -#ifdef __KERNEL_OPTIX__ - uint p0 = 0; - uint p1 = 0; - uint p2 = 0; - uint p3 = 0; - uint p4 = visibility; - uint p5 = PRIMITIVE_NONE; - uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; - uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; - - uint ray_mask = visibility & 0xFF; - uint ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT; - if (0 == ray_mask && (visibility & ~0xFF) != 0) { - ray_mask = 0xFF; - } - else if (visibility & PATH_RAY_SHADOW_OPAQUE) { - ray_flags |= OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT; - } - - optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, - ray->P, - ray->D, - ray->tmin, - ray->tmax, - ray->time, - ray_mask, - ray_flags, - 0, /* SBT offset for PG_HITD */ - 0, - 0, - p0, - p1, - p2, - p3, - p4, - p5, - p6, - p7); - - isect->t = __uint_as_float(p0); - isect->u = __uint_as_float(p1); - isect->v = __uint_as_float(p2); - isect->prim = p3; - isect->object = p4; - isect->type = p5; - - return p5 != PRIMITIVE_NONE; -#elif defined(__METALRT__) - - if (!scene_intersect_valid(ray)) { - isect->t = ray->tmax; - isect->type = PRIMITIVE_NONE; + if (!intersection_ray_valid(ray)) { return false; } -# if defined(__KERNEL_DEBUG__) - if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { - isect->t = ray->tmax; - isect->type = PRIMITIVE_NONE; - kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); - return false; - } - - if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { - isect->t = ray->tmax; - isect->type = PRIMITIVE_NONE; - kernel_assert(!"Invalid ift_default"); - return false; - } -# endif - - metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); - metalrt_intersector_type metalrt_intersect; - - if (!kernel_data.bvh.have_curves) { - metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); - } - - MetalRTIntersectionPayload payload; - payload.self = ray->self; - payload.u = 0.0f; - payload.v = 0.0f; - payload.visibility = visibility; - - typename metalrt_intersector_type::result_type intersection; - - uint ray_mask = visibility & 0xFF; - if (0 == ray_mask && (visibility & ~0xFF) != 0) { - ray_mask = 0xFF; - /* No further intersector setup required: Default MetalRT behavior is any-hit. */ - } - else if (visibility & PATH_RAY_SHADOW_OPAQUE) { - /* No further intersector setup required: Shadow ray early termination is controlled by the - * intersection handler */ - } - -# if defined(__METALRT_MOTION__) - payload.time = ray->time; - intersection = metalrt_intersect.intersect(r, - metal_ancillaries->accel_struct, - ray_mask, - ray->time, - metal_ancillaries->ift_default, - payload); -# else - intersection = metalrt_intersect.intersect( - r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload); -# endif - - if (intersection.type == intersection_type::none) { - isect->t = ray->tmax; - isect->type = PRIMITIVE_NONE; - - return false; - } - - isect->t = intersection.distance; - - isect->prim = payload.prim; - isect->type = payload.type; - isect->object = intersection.user_instance_id; - - isect->t = intersection.distance; - if (intersection.type == intersection_type::triangle) { - isect->u = 1.0f - intersection.triangle_barycentric_coord.y - - intersection.triangle_barycentric_coord.x; - isect->v = intersection.triangle_barycentric_coord.x; - } - else { - isect->u = payload.u; - isect->v = payload.v; - } - - return isect->type != PRIMITIVE_NONE; - -#else - - if (!scene_intersect_valid(ray)) { - return false; - } - -# ifdef __EMBREE__ - if (kernel_data.device_bvh) { - isect->t = ray->tmax; - CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR); - IntersectContext rtc_ctx(&ctx); - RTCRayHit ray_hit; - ctx.ray = ray; - kernel_embree_setup_rayhit(*ray, ray_hit, visibility); - rtcIntersect1(kernel_data.device_bvh, &rtc_ctx.context, &ray_hit); - if (ray_hit.hit.geomID != RTC_INVALID_GEOMETRY_ID && - ray_hit.hit.primID != RTC_INVALID_GEOMETRY_ID) { - kernel_embree_convert_hit(kg, &ray_hit.ray, &ray_hit.hit, isect); - return true; - } - return false; - } -# endif /* __EMBREE__ */ - # ifdef __OBJECT_MOTION__ if (kernel_data.bvh.have_motion) { # ifdef __HAIR__ @@ -322,7 +82,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, return bvh_intersect_motion(kg, ray, isect, visibility); } -# endif /* __OBJECT_MOTION__ */ +# endif /* __OBJECT_MOTION__ */ # ifdef __HAIR__ if (kernel_data.bvh.have_curves) { @@ -331,10 +91,22 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, # endif /* __HAIR__ */ return bvh_intersect(kg, ray, isect, visibility); -#endif /* __KERNEL_OPTIX__ */ } -#ifdef __BVH_LOCAL__ +/* Single object BVH traversal, for SSS/AO/bevel. */ + +# ifdef __BVH_LOCAL__ + +# define BVH_FUNCTION_NAME bvh_intersect_local +# define BVH_FUNCTION_FEATURES BVH_HAIR +# include "kernel/bvh/local.h" + +# if defined(__OBJECT_MOTION__) +# define BVH_FUNCTION_NAME bvh_intersect_local_motion +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_HAIR +# include "kernel/bvh/local.h" +# endif + ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, ccl_private const Ray *ray, ccl_private LocalIntersection *local_isect, @@ -342,177 +114,48 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, ccl_private uint *lcg_state, int max_hits) { -# ifdef __KERNEL_OPTIX__ - uint p0 = pointer_pack_to_uint_0(lcg_state); - uint p1 = pointer_pack_to_uint_1(lcg_state); - uint p2 = pointer_pack_to_uint_0(local_isect); - uint p3 = pointer_pack_to_uint_1(local_isect); - uint p4 = local_object; - uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; - uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; - - /* Is set to zero on miss or if ray is aborted, so can be used as return value. */ - uint p5 = max_hits; - - if (local_isect) { - local_isect->num_hits = 0; /* Initialize hit count to zero. */ - } - optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, - ray->P, - ray->D, - ray->tmin, - ray->tmax, - ray->time, - 0xFF, - /* Need to always call into __anyhit__kernel_optix_local_hit. */ - OPTIX_RAY_FLAG_ENFORCE_ANYHIT, - 2, /* SBT offset for PG_HITL */ - 0, - 0, - p0, - p1, - p2, - p3, - p4, - p5, - p6, - p7); - - return p5; -# elif defined(__METALRT__) - if (!scene_intersect_valid(ray)) { - if (local_isect) { - local_isect->num_hits = 0; - } - return false; - } - -# if defined(__KERNEL_DEBUG__) - if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + if (!intersection_ray_valid(ray)) { if (local_isect) { local_isect->num_hits = 0; } - kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); return false; } - if (is_null_intersection_function_table(metal_ancillaries->ift_local)) { - if (local_isect) { - local_isect->num_hits = 0; - } - kernel_assert(!"Invalid ift_local"); - return false; +# ifdef __OBJECT_MOTION__ + if (kernel_data.bvh.have_motion) { + return bvh_intersect_local_motion(kg, ray, local_isect, local_object, lcg_state, max_hits); } -# endif - - metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); - metalrt_intersector_type metalrt_intersect; +# endif /* __OBJECT_MOTION__ */ + return bvh_intersect_local(kg, ray, local_isect, local_object, lcg_state, max_hits); +} +# endif - metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); - if (!kernel_data.bvh.have_curves) { - metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); - } +/* Transparent shadow BVH traversal, recording multiple intersections. */ - MetalRTIntersectionLocalPayload payload; - payload.self = ray->self; - payload.local_object = local_object; - payload.max_hits = max_hits; - payload.local_isect.num_hits = 0; - if (lcg_state) { - payload.has_lcg_state = true; - payload.lcg_state = *lcg_state; - } - payload.result = false; +# ifdef __SHADOW_RECORD_ALL__ - typename metalrt_intersector_type::result_type intersection; +# define BVH_FUNCTION_NAME bvh_intersect_shadow_all +# define BVH_FUNCTION_FEATURES BVH_POINTCLOUD +# include "kernel/bvh/shadow_all.h" -# if defined(__METALRT_MOTION__) - intersection = metalrt_intersect.intersect( - r, metal_ancillaries->accel_struct, 0xFF, ray->time, metal_ancillaries->ift_local, payload); -# else - intersection = metalrt_intersect.intersect( - r, metal_ancillaries->accel_struct, 0xFF, metal_ancillaries->ift_local, payload); +# if defined(__HAIR__) +# define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD +# include "kernel/bvh/shadow_all.h" # endif - if (lcg_state) { - *lcg_state = payload.lcg_state; - } - *local_isect = payload.local_isect; - - return payload.result; - -# else - - if (!scene_intersect_valid(ray)) { - if (local_isect) { - local_isect->num_hits = 0; - } - return false; - } - -# ifdef __EMBREE__ - if (kernel_data.device_bvh) { - const bool has_bvh = !(kernel_data_fetch(object_flag, local_object) & - SD_OBJECT_TRANSFORM_APPLIED); - CCLIntersectContext ctx( - kg, has_bvh ? CCLIntersectContext::RAY_SSS : CCLIntersectContext::RAY_LOCAL); - ctx.lcg_state = lcg_state; - ctx.max_hits = max_hits; - ctx.ray = ray; - ctx.local_isect = local_isect; - if (local_isect) { - local_isect->num_hits = 0; - } - ctx.local_object_id = local_object; - IntersectContext rtc_ctx(&ctx); - RTCRay rtc_ray; - kernel_embree_setup_ray(*ray, rtc_ray, PATH_RAY_ALL_VISIBILITY); - - /* If this object has its own BVH, use it. */ - if (has_bvh) { - RTCGeometry geom = rtcGetGeometry(kernel_data.device_bvh, local_object * 2); - if (geom) { - float3 P = ray->P; - float3 dir = ray->D; - float3 idir = ray->D; - bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir); - - rtc_ray.org_x = P.x; - rtc_ray.org_y = P.y; - rtc_ray.org_z = P.z; - rtc_ray.dir_x = dir.x; - rtc_ray.dir_y = dir.y; - rtc_ray.dir_z = dir.z; - rtc_ray.tnear = ray->tmin; - rtc_ray.tfar = ray->tmax; - RTCScene scene = (RTCScene)rtcGetGeometryUserData(geom); - kernel_assert(scene); - if (scene) { - rtcOccluded1(scene, &rtc_ctx.context, &rtc_ray); - } - } - } - else { - rtcOccluded1(kernel_data.device_bvh, &rtc_ctx.context, &rtc_ray); - } - - /* rtcOccluded1 sets tfar to -inf if a hit was found. */ - return (local_isect && local_isect->num_hits > 0) || (rtc_ray.tfar < 0); - ; - } -# endif /* __EMBREE__ */ +# if defined(__OBJECT_MOTION__) +# define BVH_FUNCTION_NAME bvh_intersect_shadow_all_motion +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD +# include "kernel/bvh/shadow_all.h" +# endif -# ifdef __OBJECT_MOTION__ - if (kernel_data.bvh.have_motion) { - return bvh_intersect_local_motion(kg, ray, local_isect, local_object, lcg_state, max_hits); - } -# endif /* __OBJECT_MOTION__ */ - return bvh_intersect_local(kg, ray, local_isect, local_object, lcg_state, max_hits); -# endif /* __KERNEL_OPTIX__ */ -} -#endif +# if defined(__HAIR__) && defined(__OBJECT_MOTION__) +# define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair_motion +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD +# include "kernel/bvh/shadow_all.h" +# endif -#ifdef __SHADOW_RECORD_ALL__ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, IntegratorShadowState state, ccl_private const Ray *ray, @@ -521,132 +164,12 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { -# ifdef __KERNEL_OPTIX__ - uint p0 = state; - uint p1 = __float_as_uint(1.0f); /* Throughput. */ - uint p2 = 0; /* Number of hits. */ - uint p3 = max_hits; - uint p4 = visibility; - uint p5 = false; - uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; - uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; - - uint ray_mask = visibility & 0xFF; - if (0 == ray_mask && (visibility & ~0xFF) != 0) { - ray_mask = 0xFF; - } - - optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, - ray->P, - ray->D, - ray->tmin, - ray->tmax, - ray->time, - ray_mask, - /* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */ - OPTIX_RAY_FLAG_ENFORCE_ANYHIT, - 1, /* SBT offset for PG_HITS */ - 0, - 0, - p0, - p1, - p2, - p3, - p4, - p5, - p6, - p7); - - *num_recorded_hits = uint16_unpack_from_uint_0(p2); - *throughput = __uint_as_float(p1); - - return p5; -# elif defined(__METALRT__) - - if (!scene_intersect_valid(ray)) { - return false; - } - -# if defined(__KERNEL_DEBUG__) - if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { - kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); - return false; - } - - if (is_null_intersection_function_table(metal_ancillaries->ift_shadow)) { - kernel_assert(!"Invalid ift_shadow"); - return false; - } -# endif - - metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); - metalrt_intersector_type metalrt_intersect; - - metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); - if (!kernel_data.bvh.have_curves) { - metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); - } - - MetalRTIntersectionShadowPayload payload; - payload.self = ray->self; - payload.visibility = visibility; - payload.max_hits = max_hits; - payload.num_hits = 0; - payload.num_recorded_hits = 0; - payload.throughput = 1.0f; - payload.result = false; - payload.state = state; - - uint ray_mask = visibility & 0xFF; - if (0 == ray_mask && (visibility & ~0xFF) != 0) { - ray_mask = 0xFF; - } - - typename metalrt_intersector_type::result_type intersection; - -# if defined(__METALRT_MOTION__) - payload.time = ray->time; - intersection = metalrt_intersect.intersect(r, - metal_ancillaries->accel_struct, - ray_mask, - ray->time, - metal_ancillaries->ift_shadow, - payload); -# else - intersection = metalrt_intersect.intersect( - r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_shadow, payload); -# endif - - *num_recorded_hits = payload.num_recorded_hits; - *throughput = payload.throughput; - - return payload.result; - -# else - if (!scene_intersect_valid(ray)) { + if (!intersection_ray_valid(ray)) { *num_recorded_hits = 0; *throughput = 1.0f; return false; } -# ifdef __EMBREE__ - if (kernel_data.device_bvh) { - CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_SHADOW_ALL); - Intersection *isect_array = (Intersection *)state->shadow_isect; - ctx.isect_s = isect_array; - ctx.max_hits = max_hits; - ctx.ray = ray; - IntersectContext rtc_ctx(&ctx); - RTCRay rtc_ray; - kernel_embree_setup_ray(*ray, rtc_ray, visibility); - rtcOccluded1(kernel_data.device_bvh, &rtc_ctx.context, &rtc_ray); - - *num_recorded_hits = ctx.num_recorded_hits; - *throughput = ctx.throughput; - return ctx.opaque_hit; - } -# endif /* __EMBREE__ */ - # ifdef __OBJECT_MOTION__ if (kernel_data.bvh.have_motion) { # ifdef __HAIR__ @@ -659,7 +182,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, return bvh_intersect_shadow_all_motion( kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); } -# endif /* __OBJECT_MOTION__ */ +# endif /* __OBJECT_MOTION__ */ # ifdef __HAIR__ if (kernel_data.bvh.have_curves) { @@ -670,180 +193,83 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, return bvh_intersect_shadow_all( kg, ray, state, visibility, max_hits, num_recorded_hits, throughput); -# endif /* __KERNEL_OPTIX__ */ } -#endif /* __SHADOW_RECORD_ALL__ */ +# endif /* __SHADOW_RECORD_ALL__ */ + +/* Volume BVH traversal, for initializing or updating the volume stack. */ + +# if defined(__VOLUME__) && !defined(__VOLUME_RECORD_ALL__) + +# define BVH_FUNCTION_NAME bvh_intersect_volume +# define BVH_FUNCTION_FEATURES BVH_HAIR +# include "kernel/bvh/volume.h" + +# if defined(__OBJECT_MOTION__) +# define BVH_FUNCTION_NAME bvh_intersect_volume_motion +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_HAIR +# include "kernel/bvh/volume.h" +# endif -#ifdef __VOLUME__ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, ccl_private const Ray *ray, ccl_private Intersection *isect, const uint visibility) { -# ifdef __KERNEL_OPTIX__ - uint p0 = 0; - uint p1 = 0; - uint p2 = 0; - uint p3 = 0; - uint p4 = visibility; - uint p5 = PRIMITIVE_NONE; - uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; - uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; - - uint ray_mask = visibility & 0xFF; - if (0 == ray_mask && (visibility & ~0xFF) != 0) { - ray_mask = 0xFF; - } - - optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, - ray->P, - ray->D, - ray->tmin, - ray->tmax, - ray->time, - ray_mask, - /* Need to always call into __anyhit__kernel_optix_volume_test. */ - OPTIX_RAY_FLAG_ENFORCE_ANYHIT, - 3, /* SBT offset for PG_HITV */ - 0, - 0, - p0, - p1, - p2, - p3, - p4, - p5, - p6, - p7); - - isect->t = __uint_as_float(p0); - isect->u = __uint_as_float(p1); - isect->v = __uint_as_float(p2); - isect->prim = p3; - isect->object = p4; - isect->type = p5; - - return p5 != PRIMITIVE_NONE; -# elif defined(__METALRT__) - - if (!scene_intersect_valid(ray)) { - return false; - } -# if defined(__KERNEL_DEBUG__) - if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { - kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + if (!intersection_ray_valid(ray)) { return false; } - if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { - kernel_assert(!"Invalid ift_default"); - return false; +# ifdef __OBJECT_MOTION__ + if (kernel_data.bvh.have_motion) { + return bvh_intersect_volume_motion(kg, ray, isect, visibility); } -# endif - - metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); - metalrt_intersector_type metalrt_intersect; +# endif /* __OBJECT_MOTION__ */ - metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); - if (!kernel_data.bvh.have_curves) { - metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); - } + return bvh_intersect_volume(kg, ray, isect, visibility); +} +# endif /* defined(__VOLUME__) && !defined(__VOLUME_RECORD_ALL__) */ - MetalRTIntersectionPayload payload; - payload.self = ray->self; - payload.visibility = visibility; +/* Volume BVH traversal, for initializing or updating the volume stack. + * Variation that records multiple intersections at once. */ - typename metalrt_intersector_type::result_type intersection; +# if defined(__VOLUME__) && defined(__VOLUME_RECORD_ALL__) - uint ray_mask = visibility & 0xFF; - if (0 == ray_mask && (visibility & ~0xFF) != 0) { - ray_mask = 0xFF; - } +# define BVH_FUNCTION_NAME bvh_intersect_volume_all +# define BVH_FUNCTION_FEATURES BVH_HAIR +# include "kernel/bvh/volume_all.h" -# if defined(__METALRT_MOTION__) - payload.time = ray->time; - intersection = metalrt_intersect.intersect(r, - metal_ancillaries->accel_struct, - ray_mask, - ray->time, - metal_ancillaries->ift_default, - payload); -# else - intersection = metalrt_intersect.intersect( - r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload); +# if defined(__OBJECT_MOTION__) +# define BVH_FUNCTION_NAME bvh_intersect_volume_all_motion +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_HAIR +# include "kernel/bvh/volume_all.h" # endif - if (intersection.type == intersection_type::none) { - return false; - } - - isect->prim = payload.prim; - isect->type = payload.type; - isect->object = intersection.user_instance_id; - - isect->t = intersection.distance; - if (intersection.type == intersection_type::triangle) { - isect->u = 1.0f - intersection.triangle_barycentric_coord.y - - intersection.triangle_barycentric_coord.x; - isect->v = intersection.triangle_barycentric_coord.x; - } - else { - isect->u = payload.u; - isect->v = payload.v; - } - - return isect->type != PRIMITIVE_NONE; - -# else - if (!scene_intersect_valid(ray)) { +ccl_device_intersect uint scene_intersect_volume(KernelGlobals kg, + ccl_private const Ray *ray, + ccl_private Intersection *isect, + const uint max_hits, + const uint visibility) +{ + if (!intersection_ray_valid(ray)) { return false; } # ifdef __OBJECT_MOTION__ if (kernel_data.bvh.have_motion) { - return bvh_intersect_volume_motion(kg, ray, isect, visibility); + return bvh_intersect_volume_all_motion(kg, ray, isect, max_hits, visibility); } # endif /* __OBJECT_MOTION__ */ - return bvh_intersect_volume(kg, ray, isect, visibility); -# endif /* __KERNEL_OPTIX__ */ + return bvh_intersect_volume_all(kg, ray, isect, max_hits, visibility); } -#endif /* __VOLUME__ */ -#ifdef __VOLUME_RECORD_ALL__ -ccl_device_intersect uint scene_intersect_volume_all(KernelGlobals kg, - ccl_private const Ray *ray, - ccl_private Intersection *isect, - const uint max_hits, - const uint visibility) -{ - if (!scene_intersect_valid(ray)) { - return false; - } +# endif /* defined(__VOLUME__) && defined(__VOLUME_RECORD_ALL__) */ -# ifdef __EMBREE__ - if (kernel_data.device_bvh) { - CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_VOLUME_ALL); - ctx.isect_s = isect; - ctx.max_hits = max_hits; - ctx.num_hits = 0; - ctx.ray = ray; - IntersectContext rtc_ctx(&ctx); - RTCRay rtc_ray; - kernel_embree_setup_ray(*ray, rtc_ray, visibility); - rtcOccluded1(kernel_data.device_bvh, &rtc_ctx.context, &rtc_ray); - return ctx.num_hits; - } -# endif /* __EMBREE__ */ - -# ifdef __OBJECT_MOTION__ - if (kernel_data.bvh.have_motion) { - return bvh_intersect_volume_all_motion(kg, ray, isect, max_hits, visibility); - } -# endif /* __OBJECT_MOTION__ */ +# undef BVH_FEATURE +# undef BVH_NAME_JOIN +# undef BVH_NAME_EVAL +# undef BVH_FUNCTION_FULL_NAME - return bvh_intersect_volume_all(kg, ray, isect, max_hits, visibility); -} -#endif /* __VOLUME_RECORD_ALL__ */ +#endif /* __BVH2__ */ CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/bvh/embree.h b/intern/cycles/kernel/bvh/embree.h deleted file mode 100644 index fecbccac2f8..00000000000 --- a/intern/cycles/kernel/bvh/embree.h +++ /dev/null @@ -1,176 +0,0 @@ -/* SPDX-License-Identifier: Apache-2.0 - * Copyright 2018-2022 Blender Foundation. */ - -#pragma once - -#include <embree3/rtcore_ray.h> -#include <embree3/rtcore_scene.h> - -#include "kernel/device/cpu/compat.h" -#include "kernel/device/cpu/globals.h" - -#include "kernel/bvh/util.h" - -#include "util/vector.h" - -CCL_NAMESPACE_BEGIN - -struct CCLIntersectContext { - typedef enum { - RAY_REGULAR = 0, - RAY_SHADOW_ALL = 1, - RAY_LOCAL = 2, - RAY_SSS = 3, - RAY_VOLUME_ALL = 4, - } RayType; - - KernelGlobals kg; - RayType type; - - /* For avoiding self intersections */ - const Ray *ray; - - /* for shadow rays */ - Intersection *isect_s; - uint max_hits; - uint num_hits; - uint num_recorded_hits; - float throughput; - float max_t; - bool opaque_hit; - - /* for SSS Rays: */ - LocalIntersection *local_isect; - int local_object_id; - uint *lcg_state; - - CCLIntersectContext(KernelGlobals kg_, RayType type_) - { - kg = kg_; - type = type_; - ray = NULL; - max_hits = 1; - num_hits = 0; - num_recorded_hits = 0; - throughput = 1.0f; - max_t = FLT_MAX; - opaque_hit = false; - isect_s = NULL; - local_isect = NULL; - local_object_id = -1; - lcg_state = NULL; - } -}; - -class IntersectContext { - public: - IntersectContext(CCLIntersectContext *ctx) - { - rtcInitIntersectContext(&context); - userRayExt = ctx; - } - RTCIntersectContext context; - CCLIntersectContext *userRayExt; -}; - -ccl_device_inline void kernel_embree_setup_ray(const Ray &ray, - RTCRay &rtc_ray, - const uint visibility) -{ - rtc_ray.org_x = ray.P.x; - rtc_ray.org_y = ray.P.y; - rtc_ray.org_z = ray.P.z; - rtc_ray.dir_x = ray.D.x; - rtc_ray.dir_y = ray.D.y; - rtc_ray.dir_z = ray.D.z; - rtc_ray.tnear = ray.tmin; - rtc_ray.tfar = ray.tmax; - rtc_ray.time = ray.time; - rtc_ray.mask = visibility; -} - -ccl_device_inline void kernel_embree_setup_rayhit(const Ray &ray, - RTCRayHit &rayhit, - const uint visibility) -{ - kernel_embree_setup_ray(ray, rayhit.ray, visibility); - rayhit.hit.geomID = RTC_INVALID_GEOMETRY_ID; - rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; -} - -ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg, - const RTCHit *hit, - const Ray *ray) -{ - bool status = false; - if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) { - const int oID = hit->instID[0] / 2; - if ((ray->self.object == oID) || (ray->self.light_object == oID)) { - RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( - rtcGetGeometry(kernel_data.device_bvh, hit->instID[0])); - const int pID = hit->primID + - (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); - status = intersection_skip_self_shadow(ray->self, oID, pID); - } - } - else { - const int oID = hit->geomID / 2; - if ((ray->self.object == oID) || (ray->self.light_object == oID)) { - const int pID = hit->primID + (intptr_t)rtcGetGeometryUserData( - rtcGetGeometry(kernel_data.device_bvh, hit->geomID)); - status = intersection_skip_self_shadow(ray->self, oID, pID); - } - } - - return status; -} - -ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg, - const RTCRay *ray, - const RTCHit *hit, - Intersection *isect) -{ - isect->t = ray->tfar; - if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) { - RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( - rtcGetGeometry(kernel_data.device_bvh, hit->instID[0])); - isect->prim = hit->primID + - (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); - isect->object = hit->instID[0] / 2; - } - else { - isect->prim = hit->primID + (intptr_t)rtcGetGeometryUserData( - rtcGetGeometry(kernel_data.device_bvh, hit->geomID)); - isect->object = hit->geomID / 2; - } - - const bool is_hair = hit->geomID & 1; - if (is_hair) { - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, isect->prim); - isect->type = segment.type; - isect->prim = segment.prim; - isect->u = hit->u; - isect->v = hit->v; - } - else { - isect->type = kernel_data_fetch(objects, isect->object).primitive_type; - isect->u = 1.0f - hit->v - hit->u; - isect->v = hit->u; - } -} - -ccl_device_inline void kernel_embree_convert_sss_hit( - KernelGlobals kg, const RTCRay *ray, const RTCHit *hit, Intersection *isect, int object) -{ - isect->u = 1.0f - hit->v - hit->u; - isect->v = hit->u; - isect->t = ray->tfar; - RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( - rtcGetGeometry(kernel_data.device_bvh, object * 2)); - isect->prim = hit->primID + - (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); - isect->object = object; - isect->type = kernel_data_fetch(objects, object).primitive_type; -} - -CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/bvh/metal.h b/intern/cycles/kernel/bvh/metal.h deleted file mode 100644 index 04289e259a7..00000000000 --- a/intern/cycles/kernel/bvh/metal.h +++ /dev/null @@ -1,37 +0,0 @@ -/* SPDX-License-Identifier: Apache-2.0 - * Copyright 2021-2022 Blender Foundation */ - -struct MetalRTIntersectionPayload { - RaySelfPrimitives self; - uint visibility; - float u, v; - int prim; - int type; -#if defined(__METALRT_MOTION__) - float time; -#endif -}; - -struct MetalRTIntersectionLocalPayload { - RaySelfPrimitives self; - uint local_object; - uint lcg_state; - short max_hits; - bool has_lcg_state; - bool result; - LocalIntersection local_isect; -}; - -struct MetalRTIntersectionShadowPayload { - RaySelfPrimitives self; - uint visibility; -#if defined(__METALRT_MOTION__) - float time; -#endif - int state; - float throughput; - short max_hits; - short num_hits; - short num_recorded_hits; - bool result; -}; diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h index 385e904d20f..02e927decd4 100644 --- a/intern/cycles/kernel/bvh/util.h +++ b/intern/cycles/kernel/bvh/util.h @@ -5,6 +5,21 @@ CCL_NAMESPACE_BEGIN +ccl_device_inline bool intersection_ray_valid(ccl_private const Ray *ray) +{ + /* NOTE: Due to some vectorization code non-finite origin point might + * cause lots of false-positive intersections which will overflow traversal + * stack. + * This code is a quick way to perform early output, to avoid crashes in + * such cases. + * From production scenes so far it seems it's enough to test first element + * only. + * Scene intersection may also called with empty rays for conditional trace + * calls that evaluate to false, so filter those out. + */ + return isfinite_safe(ray->P.x) && isfinite_safe(ray->D.x) && len_squared(ray->D) != 0.0f; +} + /* Offset intersection distance by the smallest possible amount, to skip * intersections at this distance. This works in cases where the ray start * position is unchanged and only tmin is updated, since for self diff --git a/intern/cycles/kernel/device/cpu/bvh.h b/intern/cycles/kernel/device/cpu/bvh.h new file mode 100644 index 00000000000..b5ea3d831f4 --- /dev/null +++ b/intern/cycles/kernel/device/cpu/bvh.h @@ -0,0 +1,609 @@ +/* SPDX-License-Identifier: Apache-2.0 + * Copyright 2021-2022 Blender Foundation */ + +/* CPU Embree implementation of ray-scene intersection. */ + +#pragma once + +#include <embree3/rtcore_ray.h> +#include <embree3/rtcore_scene.h> + +#include "kernel/device/cpu/compat.h" +#include "kernel/device/cpu/globals.h" + +#include "kernel/bvh/types.h" +#include "kernel/bvh/util.h" +#include "kernel/geom/object.h" +#include "kernel/integrator/state.h" +#include "kernel/sample/lcg.h" + +#include "util/vector.h" + +CCL_NAMESPACE_BEGIN + +#define EMBREE_IS_HAIR(x) (x & 1) + +/* Intersection context. */ + +struct CCLIntersectContext { + typedef enum { + RAY_REGULAR = 0, + RAY_SHADOW_ALL = 1, + RAY_LOCAL = 2, + RAY_SSS = 3, + RAY_VOLUME_ALL = 4, + } RayType; + + KernelGlobals kg; + RayType type; + + /* For avoiding self intersections */ + const Ray *ray; + + /* for shadow rays */ + Intersection *isect_s; + uint max_hits; + uint num_hits; + uint num_recorded_hits; + float throughput; + float max_t; + bool opaque_hit; + + /* for SSS Rays: */ + LocalIntersection *local_isect; + int local_object_id; + uint *lcg_state; + + CCLIntersectContext(KernelGlobals kg_, RayType type_) + { + kg = kg_; + type = type_; + ray = NULL; + max_hits = 1; + num_hits = 0; + num_recorded_hits = 0; + throughput = 1.0f; + max_t = FLT_MAX; + opaque_hit = false; + isect_s = NULL; + local_isect = NULL; + local_object_id = -1; + lcg_state = NULL; + } +}; + +class IntersectContext { + public: + IntersectContext(CCLIntersectContext *ctx) + { + rtcInitIntersectContext(&context); + userRayExt = ctx; + } + RTCIntersectContext context; + CCLIntersectContext *userRayExt; +}; + +/* Utilities. */ + +ccl_device_inline void kernel_embree_setup_ray(const Ray &ray, + RTCRay &rtc_ray, + const uint visibility) +{ + rtc_ray.org_x = ray.P.x; + rtc_ray.org_y = ray.P.y; + rtc_ray.org_z = ray.P.z; + rtc_ray.dir_x = ray.D.x; + rtc_ray.dir_y = ray.D.y; + rtc_ray.dir_z = ray.D.z; + rtc_ray.tnear = ray.tmin; + rtc_ray.tfar = ray.tmax; + rtc_ray.time = ray.time; + rtc_ray.mask = visibility; +} + +ccl_device_inline void kernel_embree_setup_rayhit(const Ray &ray, + RTCRayHit &rayhit, + const uint visibility) +{ + kernel_embree_setup_ray(ray, rayhit.ray, visibility); + rayhit.hit.geomID = RTC_INVALID_GEOMETRY_ID; + rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; +} + +ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg, + const RTCHit *hit, + const Ray *ray) +{ + bool status = false; + if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) { + const int oID = hit->instID[0] / 2; + if ((ray->self.object == oID) || (ray->self.light_object == oID)) { + RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( + rtcGetGeometry(kernel_data.device_bvh, hit->instID[0])); + const int pID = hit->primID + + (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); + status = intersection_skip_self_shadow(ray->self, oID, pID); + } + } + else { + const int oID = hit->geomID / 2; + if ((ray->self.object == oID) || (ray->self.light_object == oID)) { + const int pID = hit->primID + (intptr_t)rtcGetGeometryUserData( + rtcGetGeometry(kernel_data.device_bvh, hit->geomID)); + status = intersection_skip_self_shadow(ray->self, oID, pID); + } + } + + return status; +} + +ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg, + const RTCRay *ray, + const RTCHit *hit, + Intersection *isect) +{ + isect->t = ray->tfar; + if (hit->instID[0] != RTC_INVALID_GEOMETRY_ID) { + RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( + rtcGetGeometry(kernel_data.device_bvh, hit->instID[0])); + isect->prim = hit->primID + + (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); + isect->object = hit->instID[0] / 2; + } + else { + isect->prim = hit->primID + (intptr_t)rtcGetGeometryUserData( + rtcGetGeometry(kernel_data.device_bvh, hit->geomID)); + isect->object = hit->geomID / 2; + } + + const bool is_hair = hit->geomID & 1; + if (is_hair) { + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, isect->prim); + isect->type = segment.type; + isect->prim = segment.prim; + isect->u = hit->u; + isect->v = hit->v; + } + else { + isect->type = kernel_data_fetch(objects, isect->object).primitive_type; + isect->u = 1.0f - hit->v - hit->u; + isect->v = hit->u; + } +} + +ccl_device_inline void kernel_embree_convert_sss_hit( + KernelGlobals kg, const RTCRay *ray, const RTCHit *hit, Intersection *isect, int object) +{ + isect->u = 1.0f - hit->v - hit->u; + isect->v = hit->u; + isect->t = ray->tfar; + RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( + rtcGetGeometry(kernel_data.device_bvh, object * 2)); + isect->prim = hit->primID + + (intptr_t)rtcGetGeometryUserData(rtcGetGeometry(inst_scene, hit->geomID)); + isect->object = object; + isect->type = kernel_data_fetch(objects, object).primitive_type; +} + +/* Ray filter functions. */ + +/* This gets called by Embree at every valid ray/object intersection. + * Things like recording subsurface or shadow hits for later evaluation + * as well as filtering for volume objects happen here. + * Cycles' own BVH does that directly inside the traversal calls. */ +ccl_device void kernel_embree_filter_intersection_func(const RTCFilterFunctionNArguments *args) +{ + /* Current implementation in Cycles assumes only single-ray intersection queries. */ + assert(args->N == 1); + + RTCHit *hit = (RTCHit *)args->hit; + CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt; + const KernelGlobalsCPU *kg = ctx->kg; + const Ray *cray = ctx->ray; + + if (kernel_embree_is_self_intersection(kg, hit, cray)) { + *args->valid = 0; + } +} + +/* This gets called by Embree at every valid ray/object intersection. + * Things like recording subsurface or shadow hits for later evaluation + * as well as filtering for volume objects happen here. + * Cycles' own BVH does that directly inside the traversal calls. + */ +ccl_device void kernel_embree_filter_occluded_func(const RTCFilterFunctionNArguments *args) +{ + /* Current implementation in Cycles assumes only single-ray intersection queries. */ + assert(args->N == 1); + + const RTCRay *ray = (RTCRay *)args->ray; + RTCHit *hit = (RTCHit *)args->hit; + CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt; + const KernelGlobalsCPU *kg = ctx->kg; + const Ray *cray = ctx->ray; + + switch (ctx->type) { + case CCLIntersectContext::RAY_SHADOW_ALL: { + Intersection current_isect; + kernel_embree_convert_hit(kg, ray, hit, ¤t_isect); + if (intersection_skip_self_shadow(cray->self, current_isect.object, current_isect.prim)) { + *args->valid = 0; + return; + } + /* If no transparent shadows or max number of hits exceeded, all light is blocked. */ + const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type); + if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->num_hits >= ctx->max_hits) { + ctx->opaque_hit = true; + return; + } + + ++ctx->num_hits; + + /* Always use baked shadow transparency for curves. */ + if (current_isect.type & PRIMITIVE_CURVE) { + ctx->throughput *= intersection_curve_shadow_transparency( + kg, current_isect.object, current_isect.prim, current_isect.u); + + if (ctx->throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { + ctx->opaque_hit = true; + return; + } + else { + *args->valid = 0; + return; + } + } + + /* Test if we need to record this transparent intersection. */ + const uint max_record_hits = min(ctx->max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + if (ctx->num_recorded_hits < max_record_hits || ray->tfar < ctx->max_t) { + /* If maximum number of hits was reached, replace the intersection with the + * highest distance. We want to find the N closest intersections. */ + const uint num_recorded_hits = min(ctx->num_recorded_hits, max_record_hits); + uint isect_index = num_recorded_hits; + if (num_recorded_hits + 1 >= max_record_hits) { + float max_t = ctx->isect_s[0].t; + uint max_recorded_hit = 0; + + for (uint i = 1; i < num_recorded_hits; ++i) { + if (ctx->isect_s[i].t > max_t) { + max_recorded_hit = i; + max_t = ctx->isect_s[i].t; + } + } + + if (num_recorded_hits >= max_record_hits) { + isect_index = max_recorded_hit; + } + + /* Limit the ray distance and stop counting hits beyond this. + * TODO: is there some way we can tell Embree to stop intersecting beyond + * this distance when max number of hits is reached?. Or maybe it will + * become irrelevant if we make max_hits a very high number on the CPU. */ + ctx->max_t = max(current_isect.t, max_t); + } + + ctx->isect_s[isect_index] = current_isect; + } + + /* Always increase the number of recorded hits, even beyond the maximum, + * so that we can detect this and trace another ray if needed. */ + ++ctx->num_recorded_hits; + + /* This tells Embree to continue tracing. */ + *args->valid = 0; + break; + } + case CCLIntersectContext::RAY_LOCAL: + case CCLIntersectContext::RAY_SSS: { + /* Check if it's hitting the correct object. */ + Intersection current_isect; + if (ctx->type == CCLIntersectContext::RAY_SSS) { + kernel_embree_convert_sss_hit(kg, ray, hit, ¤t_isect, ctx->local_object_id); + } + else { + kernel_embree_convert_hit(kg, ray, hit, ¤t_isect); + if (ctx->local_object_id != current_isect.object) { + /* This tells Embree to continue tracing. */ + *args->valid = 0; + break; + } + } + if (intersection_skip_self_local(cray->self, current_isect.prim)) { + *args->valid = 0; + return; + } + + /* No intersection information requested, just return a hit. */ + if (ctx->max_hits == 0) { + break; + } + + /* Ignore curves. */ + if (EMBREE_IS_HAIR(hit->geomID)) { + /* This tells Embree to continue tracing. */ + *args->valid = 0; + break; + } + + LocalIntersection *local_isect = ctx->local_isect; + int hit_idx = 0; + + if (ctx->lcg_state) { + /* See triangle_intersect_subsurface() for the native equivalent. */ + for (int i = min((int)ctx->max_hits, local_isect->num_hits) - 1; i >= 0; --i) { + if (local_isect->hits[i].t == ray->tfar) { + /* This tells Embree to continue tracing. */ + *args->valid = 0; + return; + } + } + + local_isect->num_hits++; + + if (local_isect->num_hits <= ctx->max_hits) { + hit_idx = local_isect->num_hits - 1; + } + else { + /* reservoir sampling: if we are at the maximum number of + * hits, randomly replace element or skip it */ + hit_idx = lcg_step_uint(ctx->lcg_state) % local_isect->num_hits; + + if (hit_idx >= ctx->max_hits) { + /* This tells Embree to continue tracing. */ + *args->valid = 0; + return; + } + } + } + else { + /* Record closest intersection only. */ + if (local_isect->num_hits && current_isect.t > local_isect->hits[0].t) { + *args->valid = 0; + return; + } + + local_isect->num_hits = 1; + } + + /* record intersection */ + local_isect->hits[hit_idx] = current_isect; + local_isect->Ng[hit_idx] = normalize(make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)); + /* This tells Embree to continue tracing. */ + *args->valid = 0; + break; + } + case CCLIntersectContext::RAY_VOLUME_ALL: { + /* Append the intersection to the end of the array. */ + if (ctx->num_hits < ctx->max_hits) { + Intersection current_isect; + kernel_embree_convert_hit(kg, ray, hit, ¤t_isect); + if (intersection_skip_self(cray->self, current_isect.object, current_isect.prim)) { + *args->valid = 0; + return; + } + + Intersection *isect = &ctx->isect_s[ctx->num_hits]; + ++ctx->num_hits; + *isect = current_isect; + /* Only primitives from volume object. */ + uint tri_object = isect->object; + int object_flag = kernel_data_fetch(object_flag, tri_object); + if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) { + --ctx->num_hits; + } + /* This tells Embree to continue tracing. */ + *args->valid = 0; + } + break; + } + case CCLIntersectContext::RAY_REGULAR: + default: + if (kernel_embree_is_self_intersection(kg, hit, cray)) { + *args->valid = 0; + return; + } + break; + } +} + +ccl_device void kernel_embree_filter_func_backface_cull(const RTCFilterFunctionNArguments *args) +{ + const RTCRay *ray = (RTCRay *)args->ray; + RTCHit *hit = (RTCHit *)args->hit; + + /* Always ignore back-facing intersections. */ + if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z), + make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) { + *args->valid = 0; + return; + } + + CCLIntersectContext *ctx = ((IntersectContext *)args->context)->userRayExt; + const KernelGlobalsCPU *kg = ctx->kg; + const Ray *cray = ctx->ray; + + if (kernel_embree_is_self_intersection(kg, hit, cray)) { + *args->valid = 0; + } +} + +ccl_device void kernel_embree_filter_occluded_func_backface_cull( + const RTCFilterFunctionNArguments *args) +{ + const RTCRay *ray = (RTCRay *)args->ray; + RTCHit *hit = (RTCHit *)args->hit; + + /* Always ignore back-facing intersections. */ + if (dot(make_float3(ray->dir_x, ray->dir_y, ray->dir_z), + make_float3(hit->Ng_x, hit->Ng_y, hit->Ng_z)) > 0.0f) { + *args->valid = 0; + return; + } + + kernel_embree_filter_occluded_func(args); +} + +/* Scene intersection. */ + +ccl_device_intersect bool scene_intersect(KernelGlobals kg, + ccl_private const Ray *ray, + const uint visibility, + ccl_private Intersection *isect) +{ + if (!intersection_ray_valid(ray)) { + return false; + } + + if (!kernel_data.device_bvh) { + return false; + } + + isect->t = ray->tmax; + CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR); + IntersectContext rtc_ctx(&ctx); + RTCRayHit ray_hit; + ctx.ray = ray; + kernel_embree_setup_rayhit(*ray, ray_hit, visibility); + rtcIntersect1(kernel_data.device_bvh, &rtc_ctx.context, &ray_hit); + if (ray_hit.hit.geomID == RTC_INVALID_GEOMETRY_ID || + ray_hit.hit.primID == RTC_INVALID_GEOMETRY_ID) { + return false; + } + + kernel_embree_convert_hit(kg, &ray_hit.ray, &ray_hit.hit, isect); + return true; +} + +#ifdef __BVH_LOCAL__ +ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, + ccl_private const Ray *ray, + ccl_private LocalIntersection *local_isect, + int local_object, + ccl_private uint *lcg_state, + int max_hits) +{ + if (!intersection_ray_valid(ray)) { + if (local_isect) { + local_isect->num_hits = 0; + } + return false; + } + + if (!kernel_data.device_bvh) { + return false; + } + + const bool has_bvh = !(kernel_data_fetch(object_flag, local_object) & + SD_OBJECT_TRANSFORM_APPLIED); + CCLIntersectContext ctx(kg, + has_bvh ? CCLIntersectContext::RAY_SSS : CCLIntersectContext::RAY_LOCAL); + ctx.lcg_state = lcg_state; + ctx.max_hits = max_hits; + ctx.ray = ray; + ctx.local_isect = local_isect; + if (local_isect) { + local_isect->num_hits = 0; + } + ctx.local_object_id = local_object; + IntersectContext rtc_ctx(&ctx); + RTCRay rtc_ray; + kernel_embree_setup_ray(*ray, rtc_ray, PATH_RAY_ALL_VISIBILITY); + + /* If this object has its own BVH, use it. */ + if (has_bvh) { + RTCGeometry geom = rtcGetGeometry(kernel_data.device_bvh, local_object * 2); + if (geom) { + float3 P = ray->P; + float3 dir = ray->D; + float3 idir = ray->D; + bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir); + + rtc_ray.org_x = P.x; + rtc_ray.org_y = P.y; + rtc_ray.org_z = P.z; + rtc_ray.dir_x = dir.x; + rtc_ray.dir_y = dir.y; + rtc_ray.dir_z = dir.z; + rtc_ray.tnear = ray->tmin; + rtc_ray.tfar = ray->tmax; + RTCScene scene = (RTCScene)rtcGetGeometryUserData(geom); + kernel_assert(scene); + if (scene) { + rtcOccluded1(scene, &rtc_ctx.context, &rtc_ray); + } + } + } + else { + rtcOccluded1(kernel_data.device_bvh, &rtc_ctx.context, &rtc_ray); + } + + /* rtcOccluded1 sets tfar to -inf if a hit was found. */ + return (local_isect && local_isect->num_hits > 0) || (rtc_ray.tfar < 0); +} +#endif + +#ifdef __SHADOW_RECORD_ALL__ +ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, + IntegratorShadowStateCPU *state, + ccl_private const Ray *ray, + uint visibility, + uint max_hits, + ccl_private uint *num_recorded_hits, + ccl_private float *throughput) +{ + if (!intersection_ray_valid(ray)) { + *num_recorded_hits = 0; + *throughput = 1.0f; + return false; + } + + if (!kernel_data.device_bvh) { + return false; + } + + CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_SHADOW_ALL); + Intersection *isect_array = (Intersection *)state->shadow_isect; + ctx.isect_s = isect_array; + ctx.max_hits = max_hits; + ctx.ray = ray; + IntersectContext rtc_ctx(&ctx); + RTCRay rtc_ray; + kernel_embree_setup_ray(*ray, rtc_ray, visibility); + rtcOccluded1(kernel_data.device_bvh, &rtc_ctx.context, &rtc_ray); + + *num_recorded_hits = ctx.num_recorded_hits; + *throughput = ctx.throughput; + return ctx.opaque_hit; +} +#endif + +#ifdef __VOLUME__ +ccl_device_intersect uint scene_intersect_volume(KernelGlobals kg, + ccl_private const Ray *ray, + ccl_private Intersection *isect, + const uint max_hits, + const uint visibility) +{ + if (!intersection_ray_valid(ray)) { + return false; + } + + if (!kernel_data.device_bvh) { + return false; + } + + CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_VOLUME_ALL); + ctx.isect_s = isect; + ctx.max_hits = max_hits; + ctx.num_hits = 0; + ctx.ray = ray; + IntersectContext rtc_ctx(&ctx); + RTCRay rtc_ray; + kernel_embree_setup_ray(*ray, rtc_ray, visibility); + rtcOccluded1(kernel_data.device_bvh, &rtc_ctx.context, &rtc_ray); + return ctx.num_hits; +} +#endif + +CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/device/metal/bvh.h b/intern/cycles/kernel/device/metal/bvh.h new file mode 100644 index 00000000000..d3a0ab1b519 --- /dev/null +++ b/intern/cycles/kernel/device/metal/bvh.h @@ -0,0 +1,1123 @@ +/* SPDX-License-Identifier: Apache-2.0 + * Copyright 2021-2022 Blender Foundation */ + +/* MetalRT implementation of ray-scene intersection. */ + +#pragma once + +#include "kernel/bvh/types.h" +#include "kernel/bvh/util.h" + +CCL_NAMESPACE_BEGIN + +/* Payload types. */ + +struct MetalRTIntersectionPayload { + RaySelfPrimitives self; + uint visibility; + float u, v; + int prim; + int type; +#if defined(__METALRT_MOTION__) + float time; +#endif +}; + +struct MetalRTIntersectionLocalPayload { + RaySelfPrimitives self; + uint local_object; + uint lcg_state; + short max_hits; + bool has_lcg_state; + bool result; + LocalIntersection local_isect; +}; + +struct MetalRTIntersectionShadowPayload { + RaySelfPrimitives self; + uint visibility; +#if defined(__METALRT_MOTION__) + float time; +#endif + int state; + float throughput; + short max_hits; + short num_hits; + short num_recorded_hits; + bool result; +}; + +/* Intersection return types. */ + +/* For a bounding box intersection function. */ +struct BoundingBoxIntersectionResult { + bool accept [[accept_intersection]]; + bool continue_search [[continue_search]]; + float distance [[distance]]; +}; + +/* For a triangle intersection function. */ +struct TriangleIntersectionResult { + bool accept [[accept_intersection]]; + bool continue_search [[continue_search]]; +}; + +enum { METALRT_HIT_TRIANGLE, METALRT_HIT_BOUNDING_BOX }; + +/* Utilities. */ + +ccl_device_inline bool intersection_skip_self(ray_data const RaySelfPrimitives &self, + const int object, + const int prim) +{ + return (self.prim == prim) && (self.object == object); +} + +ccl_device_inline bool intersection_skip_self_shadow(ray_data const RaySelfPrimitives &self, + const int object, + const int prim) +{ + return ((self.prim == prim) && (self.object == object)) || + ((self.light_prim == prim) && (self.light_object == object)); +} + +ccl_device_inline bool intersection_skip_self_local(ray_data const RaySelfPrimitives &self, + const int prim) +{ + return (self.prim == prim); +} + +/* Hit functions. */ + +template<typename TReturn, uint intersection_type> +TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload, + const uint object, + const uint primitive_id, + const float2 barycentrics, + const float ray_tmax) +{ + TReturn result; + +#ifdef __BVH_LOCAL__ + uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + + if ((object != payload.local_object) || intersection_skip_self_local(payload.self, prim)) { + /* Only intersect with matching object and skip self-intersecton. */ + result.accept = false; + result.continue_search = true; + return result; + } + + const short max_hits = payload.max_hits; + if (max_hits == 0) { + /* Special case for when no hit information is requested, just report that something was hit */ + payload.result = true; + result.accept = true; + result.continue_search = false; + return result; + } + + int hit = 0; + if (payload.has_lcg_state) { + for (short i = min(max_hits, short(payload.local_isect.num_hits)) - 1; i >= 0; --i) { + if (ray_tmax == payload.local_isect.hits[i].t) { + result.accept = false; + result.continue_search = true; + return result; + } + } + + hit = payload.local_isect.num_hits++; + + if (payload.local_isect.num_hits > max_hits) { + hit = lcg_step_uint(&payload.lcg_state) % payload.local_isect.num_hits; + if (hit >= max_hits) { + result.accept = false; + result.continue_search = true; + return result; + } + } + } + else { + if (payload.local_isect.num_hits && ray_tmax > payload.local_isect.hits[0].t) { + /* Record closest intersection only. Do not terminate ray here, since there is no guarantee + * about distance ordering in any-hit */ + result.accept = false; + result.continue_search = true; + return result; + } + + payload.local_isect.num_hits = 1; + } + + ray_data Intersection *isect = &payload.local_isect.hits[hit]; + isect->t = ray_tmax; + isect->prim = prim; + isect->object = object; + isect->type = kernel_data_fetch(objects, object).primitive_type; + + isect->u = 1.0f - barycentrics.y - barycentrics.x; + isect->v = barycentrics.x; + + /* Record geometric normal */ + const uint tri_vindex = kernel_data_fetch(tri_vindex, isect->prim).w; + const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex + 0)); + const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex + 1)); + const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex + 2)); + payload.local_isect.Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); + + /* Continue tracing (without this the trace call would return after the first hit) */ + result.accept = false; + result.continue_search = true; + return result; +#endif +} + +[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult +__anyhit__cycles_metalrt_local_hit_tri( + constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]], + uint instance_id [[user_instance_id]], + uint primitive_id [[primitive_id]], + float2 barycentrics [[barycentric_coord]], + float ray_tmax [[distance]]) +{ + return metalrt_local_hit<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>( + launch_params_metal, payload, instance_id, primitive_id, barycentrics, ray_tmax); +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__anyhit__cycles_metalrt_local_hit_box(const float ray_tmax [[max_distance]]) +{ + /* unused function */ + BoundingBoxIntersectionResult result; + result.distance = ray_tmax; + result.accept = false; + result.continue_search = false; + return result; +} + +template<uint intersection_type> +bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, + uint object, + uint prim, + const float2 barycentrics, + const float ray_tmax) +{ +#ifdef __SHADOW_RECORD_ALL__ +# ifdef __VISIBILITY_FLAG__ + const uint visibility = payload.visibility; + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + /* continue search */ + return true; + } +# endif + + if (intersection_skip_self_shadow(payload.self, object, prim)) { + /* continue search */ + return true; + } + + float u = 0.0f, v = 0.0f; + int type = 0; + if (intersection_type == METALRT_HIT_TRIANGLE) { + u = 1.0f - barycentrics.y - barycentrics.x; + v = barycentrics.x; + type = kernel_data_fetch(objects, object).primitive_type; + } +# ifdef __HAIR__ + else { + u = barycentrics.x; + v = barycentrics.y; + + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + type = segment.type; + prim = segment.prim; + + /* Filter out curve endcaps */ + if (u == 0.0f || u == 1.0f) { + /* continue search */ + return true; + } + } +# endif + +# ifndef __TRANSPARENT_SHADOWS__ + /* No transparent shadows support compiled in, make opaque. */ + payload.result = true; + /* terminate ray */ + return false; +# else + short max_hits = payload.max_hits; + short num_hits = payload.num_hits; + short num_recorded_hits = payload.num_recorded_hits; + + MetalKernelContext context(launch_params_metal); + + /* If no transparent shadows, all light is blocked and we can stop immediately. */ + if (num_hits >= max_hits || + !(context.intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) { + payload.result = true; + /* terminate ray */ + return false; + } + + /* Always use baked shadow transparency for curves. */ + if (type & PRIMITIVE_CURVE) { + float throughput = payload.throughput; + throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, u); + payload.throughput = throughput; + payload.num_hits += 1; + + if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { + /* Accept result and terminate if throughput is sufficiently low */ + payload.result = true; + return false; + } + else { + return true; + } + } + + payload.num_hits += 1; + payload.num_recorded_hits += 1; + + uint record_index = num_recorded_hits; + + const IntegratorShadowState state = payload.state; + + const uint max_record_hits = min(uint(max_hits), INTEGRATOR_SHADOW_ISECT_SIZE); + if (record_index >= max_record_hits) { + /* If maximum number of hits reached, find a hit to replace. */ + float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); + uint max_recorded_hit = 0; + + for (int i = 1; i < max_record_hits; i++) { + const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); + if (isect_t > max_recorded_t) { + max_recorded_t = isect_t; + max_recorded_hit = i; + } + } + + if (ray_tmax >= max_recorded_t) { + /* Accept hit, so that we don't consider any more hits beyond the distance of the + * current hit anymore. */ + payload.result = true; + return true; + } + + record_index = max_recorded_hit; + } + + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = ray_tmax; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type; + + /* Continue tracing. */ +# endif /* __TRANSPARENT_SHADOWS__ */ +#endif /* __SHADOW_RECORD_ALL__ */ + + return true; +} + +[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult +__anyhit__cycles_metalrt_shadow_all_hit_tri( + constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], + unsigned int object [[user_instance_id]], + unsigned int primitive_id [[primitive_id]], + float2 barycentrics [[barycentric_coord]], + float ray_tmax [[distance]]) +{ + uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + + TriangleIntersectionResult result; + result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_TRIANGLE>( + launch_params_metal, payload, object, prim, barycentrics, ray_tmax); + result.accept = !result.continue_search; + return result; +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__anyhit__cycles_metalrt_shadow_all_hit_box(const float ray_tmax [[max_distance]]) +{ + /* unused function */ + BoundingBoxIntersectionResult result; + result.distance = ray_tmax; + result.accept = false; + result.continue_search = false; + return result; +} + +template<typename TReturnType, uint intersection_type> +inline TReturnType metalrt_visibility_test( + constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload, + const uint object, + const uint prim, + const float u) +{ + TReturnType result; + +#ifdef __HAIR__ + if (intersection_type == METALRT_HIT_BOUNDING_BOX) { + /* Filter out curve endcaps. */ + if (u == 0.0f || u == 1.0f) { + result.accept = false; + result.continue_search = true; + return result; + } + } +#endif + + uint visibility = payload.visibility; +#ifdef __VISIBILITY_FLAG__ + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + result.accept = false; + result.continue_search = true; + return result; + } +#endif + + /* Shadow ray early termination. */ + if (visibility & PATH_RAY_SHADOW_OPAQUE) { + if (intersection_skip_self_shadow(payload.self, object, prim)) { + result.accept = false; + result.continue_search = true; + return result; + } + else { + result.accept = true; + result.continue_search = false; + return result; + } + } + else { + if (intersection_skip_self(payload.self, object, prim)) { + result.accept = false; + result.continue_search = true; + return result; + } + } + + result.accept = true; + result.continue_search = true; + return result; +} + +[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult +__anyhit__cycles_metalrt_visibility_test_tri( + constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], + unsigned int object [[user_instance_id]], + unsigned int primitive_id [[primitive_id]]) +{ + uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + TriangleIntersectionResult result = + metalrt_visibility_test<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>( + launch_params_metal, payload, object, prim, 0.0f); + if (result.accept) { + payload.prim = prim; + payload.type = kernel_data_fetch(objects, object).primitive_type; + } + return result; +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance]]) +{ + /* Unused function */ + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + return result; +} + +/* Primitive intersection functions. */ + +#ifdef __HAIR__ +ccl_device_inline void metalrt_intersection_curve( + constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload, + const uint object, + const uint prim, + const uint type, + const float3 ray_P, + const float3 ray_D, + float time, + const float ray_tmin, + const float ray_tmax, + thread BoundingBoxIntersectionResult &result) +{ +# ifdef __VISIBILITY_FLAG__ + const uint visibility = payload.visibility; + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + return; + } +# endif + + Intersection isect; + isect.t = ray_tmax; + + MetalKernelContext context(launch_params_metal); + if (context.curve_intersect( + NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { + result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>( + launch_params_metal, payload, object, prim, isect.u); + if (result.accept) { + result.distance = isect.t; + payload.u = isect.u; + payload.v = isect.v; + payload.prim = prim; + payload.type = type; + } + } +} + +ccl_device_inline void metalrt_intersection_curve_shadow( + constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, + const uint object, + const uint prim, + const uint type, + float time, + const float ray_tmin, + const float ray_tmax, + thread BoundingBoxIntersectionResult &result) +{ + const uint visibility = payload.visibility; + + Intersection isect; + isect.t = ray_tmax; + + MetalKernelContext context(launch_params_metal); + if (context.curve_intersect( + NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { + result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>( + launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax); + result.accept = !result.continue_search; + } +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload + [[payload]], + const uint object [[user_instance_id]], + const uint primitive_id [[primitive_id]], + const float3 ray_P [[origin]], + const float3 ray_D [[direction]], + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) +{ + uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + + if (segment.type & PRIMITIVE_CURVE_RIBBON) { + metalrt_intersection_curve(launch_params_metal, + payload, + object, + segment.prim, + segment.type, + ray_P, + ray_D, +# if defined(__METALRT_MOTION__) + payload.time, +# else + 0.0f, +# endif + ray_tmin, + ray_tmax, + result); + } + + return result; +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__intersection__curve_ribbon_shadow( + constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], + const uint object [[user_instance_id]], + const uint primitive_id [[primitive_id]], + const float3 ray_P [[origin]], + const float3 ray_D [[direction]], + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) +{ + uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + + if (segment.type & PRIMITIVE_CURVE_RIBBON) { + metalrt_intersection_curve_shadow(launch_params_metal, + payload, + object, + segment.prim, + segment.type, + ray_P, + ray_D, +# if defined(__METALRT_MOTION__) + payload.time, +# else + 0.0f, +# endif + ray_tmin, + ray_tmax, + result); + } + + return result; +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload + [[payload]], + const uint object [[user_instance_id]], + const uint primitive_id [[primitive_id]], + const float3 ray_P [[origin]], + const float3 ray_D [[direction]], + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) +{ + uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + metalrt_intersection_curve(launch_params_metal, + payload, + object, + segment.prim, + segment.type, + ray_P, + ray_D, +# if defined(__METALRT_MOTION__) + payload.time, +# else + 0.0f, +# endif + ray_tmin, + ray_tmax, + result); + + return result; +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__intersection__curve_all_shadow( + constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], + const uint object [[user_instance_id]], + const uint primitive_id [[primitive_id]], + const float3 ray_P [[origin]], + const float3 ray_D [[direction]], + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) +{ + uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + + metalrt_intersection_curve_shadow(launch_params_metal, + payload, + object, + segment.prim, + segment.type, + ray_P, + ray_D, +# if defined(__METALRT_MOTION__) + payload.time, +# else + 0.0f, +# endif + ray_tmin, + ray_tmax, + result); + + return result; +} +#endif /* __HAIR__ */ + +#ifdef __POINTCLOUD__ +ccl_device_inline void metalrt_intersection_point( + constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload, + const uint object, + const uint prim, + const uint type, + const float3 ray_P, + const float3 ray_D, + float time, + const float ray_tmin, + const float ray_tmax, + thread BoundingBoxIntersectionResult &result) +{ +# ifdef __VISIBILITY_FLAG__ + const uint visibility = payload.visibility; + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + return; + } +# endif + + Intersection isect; + isect.t = ray_tmax; + + MetalKernelContext context(launch_params_metal); + if (context.point_intersect( + NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { + result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>( + launch_params_metal, payload, object, prim, isect.u); + if (result.accept) { + result.distance = isect.t; + payload.u = isect.u; + payload.v = isect.v; + payload.prim = prim; + payload.type = type; + } + } +} + +ccl_device_inline void metalrt_intersection_point_shadow( + constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, + const uint object, + const uint prim, + const uint type, + const float3 ray_P, + const float3 ray_D, + float time, + const float ray_tmin, + const float ray_tmax, + thread BoundingBoxIntersectionResult &result) +{ + const uint visibility = payload.visibility; + + Intersection isect; + isect.t = ray_tmax; + + MetalKernelContext context(launch_params_metal); + if (context.point_intersect( + NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { + result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>( + launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax); + result.accept = !result.continue_search; + + if (result.accept) { + result.distance = isect.t; + } + } +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], + const uint object [[user_instance_id]], + const uint primitive_id [[primitive_id]], + const float3 ray_origin [[origin]], + const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) +{ + const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + const int type = kernel_data_fetch(objects, object).primitive_type; + + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + + metalrt_intersection_point(launch_params_metal, + payload, + object, + prim, + type, + ray_origin, + ray_direction, +# if defined(__METALRT_MOTION__) + payload.time, +# else + 0.0f, +# endif + ray_tmin, + ray_tmax, + result); + + return result; +} + +[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult +__intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload + [[payload]], + const uint object [[user_instance_id]], + const uint primitive_id [[primitive_id]], + const float3 ray_origin [[origin]], + const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) +{ + const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); + const int type = kernel_data_fetch(objects, object).primitive_type; + + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + + metalrt_intersection_point_shadow(launch_params_metal, + payload, + object, + prim, + type, + ray_origin, + ray_direction, +# if defined(__METALRT_MOTION__) + payload.time, +# else + 0.0f, +# endif + ray_tmin, + ray_tmax, + result); + + return result; +} +#endif /* __POINTCLOUD__ */ + +/* Scene intersection. */ + +ccl_device_intersect bool scene_intersect(KernelGlobals kg, + ccl_private const Ray *ray, + const uint visibility, + ccl_private Intersection *isect) +{ + if (!scene_intersect_valid(ray)) { + isect->t = ray->tmax; + isect->type = PRIMITIVE_NONE; + return false; + } + +#if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + isect->t = ray->tmax; + isect->type = PRIMITIVE_NONE; + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { + isect->t = ray->tmax; + isect->type = PRIMITIVE_NONE; + kernel_assert(!"Invalid ift_default"); + return false; + } +#endif + + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); + metalrt_intersector_type metalrt_intersect; + + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionPayload payload; + payload.self = ray->self; + payload.u = 0.0f; + payload.v = 0.0f; + payload.visibility = visibility; + + typename metalrt_intersector_type::result_type intersection; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + /* No further intersector setup required: Default MetalRT behavior is any-hit. */ + } + else if (visibility & PATH_RAY_SHADOW_OPAQUE) { + /* No further intersector setup required: Shadow ray early termination is controlled by the + * intersection handler */ + } + +#if defined(__METALRT_MOTION__) + payload.time = ray->time; + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ray_mask, + ray->time, + metal_ancillaries->ift_default, + payload); +#else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload); +#endif + + if (intersection.type == intersection_type::none) { + isect->t = ray->tmax; + isect->type = PRIMITIVE_NONE; + + return false; + } + + isect->t = intersection.distance; + + isect->prim = payload.prim; + isect->type = payload.type; + isect->object = intersection.user_instance_id; + + isect->t = intersection.distance; + if (intersection.type == intersection_type::triangle) { + isect->u = 1.0f - intersection.triangle_barycentric_coord.y - + intersection.triangle_barycentric_coord.x; + isect->v = intersection.triangle_barycentric_coord.x; + } + else { + isect->u = payload.u; + isect->v = payload.v; + } + + return isect->type != PRIMITIVE_NONE; +} + +#ifdef __BVH_LOCAL__ +ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, + ccl_private const Ray *ray, + ccl_private LocalIntersection *local_isect, + int local_object, + ccl_private uint *lcg_state, + int max_hits) +{ + if (!intersection_ray_valid(ray)) { + if (local_isect) { + local_isect->num_hits = 0; + } + return false; + } + +# if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + if (local_isect) { + local_isect->num_hits = 0; + } + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_local)) { + if (local_isect) { + local_isect->num_hits = 0; + } + kernel_assert(!"Invalid ift_local"); + return false; + } +# endif + + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); + metalrt_intersector_type metalrt_intersect; + + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionLocalPayload payload; + payload.self = ray->self; + payload.local_object = local_object; + payload.max_hits = max_hits; + payload.local_isect.num_hits = 0; + if (lcg_state) { + payload.has_lcg_state = true; + payload.lcg_state = *lcg_state; + } + payload.result = false; + + typename metalrt_intersector_type::result_type intersection; + +# if defined(__METALRT_MOTION__) + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, 0xFF, ray->time, metal_ancillaries->ift_local, payload); +# else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, 0xFF, metal_ancillaries->ift_local, payload); +# endif + + if (lcg_state) { + *lcg_state = payload.lcg_state; + } + *local_isect = payload.local_isect; + + return payload.result; +} +#endif + +#ifdef __SHADOW_RECORD_ALL__ +ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, + IntegratorShadowState state, + ccl_private const Ray *ray, + uint visibility, + uint max_hits, + ccl_private uint *num_recorded_hits, + ccl_private float *throughput) +{ + if (!intersection_ray_valid(ray)) { + return false; + } + +# if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_shadow)) { + kernel_assert(!"Invalid ift_shadow"); + return false; + } +# endif + + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); + metalrt_intersector_type metalrt_intersect; + + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionShadowPayload payload; + payload.self = ray->self; + payload.visibility = visibility; + payload.max_hits = max_hits; + payload.num_hits = 0; + payload.num_recorded_hits = 0; + payload.throughput = 1.0f; + payload.result = false; + payload.state = state; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + } + + typename metalrt_intersector_type::result_type intersection; + +# if defined(__METALRT_MOTION__) + payload.time = ray->time; + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ray_mask, + ray->time, + metal_ancillaries->ift_shadow, + payload); +# else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_shadow, payload); +# endif + + *num_recorded_hits = payload.num_recorded_hits; + *throughput = payload.throughput; + + return payload.result; +} +#endif + +#ifdef __VOLUME__ +ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, + ccl_private const Ray *ray, + ccl_private Intersection *isect, + const uint visibility) +{ + if (!intersection_ray_valid(ray)) { + return false; + } + +# if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { + kernel_assert(!"Invalid ift_default"); + return false; + } +# endif + + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); + metalrt_intersector_type metalrt_intersect; + + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionPayload payload; + payload.self = ray->self; + payload.visibility = visibility; + + typename metalrt_intersector_type::result_type intersection; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + } + +# if defined(__METALRT_MOTION__) + payload.time = ray->time; + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ray_mask, + ray->time, + metal_ancillaries->ift_default, + payload); +# else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload); +# endif + + if (intersection.type == intersection_type::none) { + return false; + } + + isect->prim = payload.prim; + isect->type = payload.type; + isect->object = intersection.user_instance_id; + + isect->t = intersection.distance; + if (intersection.type == intersection_type::triangle) { + isect->u = 1.0f - intersection.triangle_barycentric_coord.y - + intersection.triangle_barycentric_coord.x; + isect->v = intersection.triangle_barycentric_coord.x; + } + else { + isect->u = payload.u; + isect->v = payload.v; + } + + return isect->type != PRIMITIVE_NONE; +} +#endif + +CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 0ed52074a90..80ee8ef5b57 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -260,8 +260,6 @@ void kernel_gpu_##name::run(thread MetalKernelContext& context, \ #ifdef __METALRT__ -# define __KERNEL_GPU_RAYTRACING__ - # if defined(__METALRT_MOTION__) # define METALRT_TAGS instancing, instance_motion, primitive_motion # else diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 8c6f2e1df5e..3df81fcf369 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -7,711 +7,3 @@ #include "kernel/device/metal/globals.h" #include "kernel/device/metal/function_constants.h" #include "kernel/device/gpu/kernel.h" - -/* MetalRT intersection handlers */ -#ifdef __METALRT__ - -/* Return type for a bounding box intersection function. */ -struct BoundingBoxIntersectionResult -{ - bool accept [[accept_intersection]]; - bool continue_search [[continue_search]]; - float distance [[distance]]; -}; - -/* Return type for a triangle intersection function. */ -struct TriangleIntersectionResult -{ - bool accept [[accept_intersection]]; - bool continue_search [[continue_search]]; -}; - -enum { METALRT_HIT_TRIANGLE, METALRT_HIT_BOUNDING_BOX }; - -ccl_device_inline bool intersection_skip_self(ray_data const RaySelfPrimitives& self, - const int object, - const int prim) -{ - return (self.prim == prim) && (self.object == object); -} - -ccl_device_inline bool intersection_skip_self_shadow(ray_data const RaySelfPrimitives& self, - const int object, - const int prim) -{ - return ((self.prim == prim) && (self.object == object)) || - ((self.light_prim == prim) && (self.light_object == object)); -} - -ccl_device_inline bool intersection_skip_self_local(ray_data const RaySelfPrimitives& self, - const int prim) -{ - return (self.prim == prim); -} - -template<typename TReturn, uint intersection_type> -TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload, - const uint object, - const uint primitive_id, - const float2 barycentrics, - const float ray_tmax) -{ - TReturn result; - -#ifdef __BVH_LOCAL__ - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - - if ((object != payload.local_object) || intersection_skip_self_local(payload.self, prim)) { - /* Only intersect with matching object and skip self-intersecton. */ - result.accept = false; - result.continue_search = true; - return result; - } - - const short max_hits = payload.max_hits; - if (max_hits == 0) { - /* Special case for when no hit information is requested, just report that something was hit */ - payload.result = true; - result.accept = true; - result.continue_search = false; - return result; - } - - int hit = 0; - if (payload.has_lcg_state) { - for (short i = min(max_hits, short(payload.local_isect.num_hits)) - 1; i >= 0; --i) { - if (ray_tmax == payload.local_isect.hits[i].t) { - result.accept = false; - result.continue_search = true; - return result; - } - } - - hit = payload.local_isect.num_hits++; - - if (payload.local_isect.num_hits > max_hits) { - hit = lcg_step_uint(&payload.lcg_state) % payload.local_isect.num_hits; - if (hit >= max_hits) { - result.accept = false; - result.continue_search = true; - return result; - } - } - } - else { - if (payload.local_isect.num_hits && ray_tmax > payload.local_isect.hits[0].t) { - /* Record closest intersection only. Do not terminate ray here, since there is no guarantee about distance ordering in any-hit */ - result.accept = false; - result.continue_search = true; - return result; - } - - payload.local_isect.num_hits = 1; - } - - ray_data Intersection *isect = &payload.local_isect.hits[hit]; - isect->t = ray_tmax; - isect->prim = prim; - isect->object = object; - isect->type = kernel_data_fetch(objects, object).primitive_type; - - isect->u = 1.0f - barycentrics.y - barycentrics.x; - isect->v = barycentrics.x; - - /* Record geometric normal */ - const uint tri_vindex = kernel_data_fetch(tri_vindex, isect->prim).w; - const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex + 0)); - const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex + 1)); - const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex + 2)); - payload.local_isect.Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); - - /* Continue tracing (without this the trace call would return after the first hit) */ - result.accept = false; - result.continue_search = true; - return result; -#endif -} - -[[intersection(triangle, triangle_data, METALRT_TAGS)]] -TriangleIntersectionResult -__anyhit__cycles_metalrt_local_hit_tri(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]], - uint instance_id [[user_instance_id]], - uint primitive_id [[primitive_id]], - float2 barycentrics [[barycentric_coord]], - float ray_tmax [[distance]]) -{ - return metalrt_local_hit<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>( - launch_params_metal, payload, instance_id, primitive_id, barycentrics, ray_tmax); -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_local_hit_box(const float ray_tmax [[max_distance]]) -{ - /* unused function */ - BoundingBoxIntersectionResult result; - result.distance = ray_tmax; - result.accept = false; - result.continue_search = false; - return result; -} - -template<uint intersection_type> -bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, - uint object, - uint prim, - const float2 barycentrics, - const float ray_tmax) -{ -#ifdef __SHADOW_RECORD_ALL__ -# ifdef __VISIBILITY_FLAG__ - const uint visibility = payload.visibility; - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - /* continue search */ - return true; - } -# endif - - if (intersection_skip_self_shadow(payload.self, object, prim)) { - /* continue search */ - return true; - } - - float u = 0.0f, v = 0.0f; - int type = 0; - if (intersection_type == METALRT_HIT_TRIANGLE) { - u = 1.0f - barycentrics.y - barycentrics.x; - v = barycentrics.x; - type = kernel_data_fetch(objects, object).primitive_type; - } -# ifdef __HAIR__ - else { - u = barycentrics.x; - v = barycentrics.y; - - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - type = segment.type; - prim = segment.prim; - - /* Filter out curve endcaps */ - if (u == 0.0f || u == 1.0f) { - /* continue search */ - return true; - } - } -# endif - -# ifndef __TRANSPARENT_SHADOWS__ - /* No transparent shadows support compiled in, make opaque. */ - payload.result = true; - /* terminate ray */ - return false; -# else - short max_hits = payload.max_hits; - short num_hits = payload.num_hits; - short num_recorded_hits = payload.num_recorded_hits; - - MetalKernelContext context(launch_params_metal); - - /* If no transparent shadows, all light is blocked and we can stop immediately. */ - if (num_hits >= max_hits || - !(context.intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) { - payload.result = true; - /* terminate ray */ - return false; - } - - /* Always use baked shadow transparency for curves. */ - if (type & PRIMITIVE_CURVE) { - float throughput = payload.throughput; - throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, u); - payload.throughput = throughput; - payload.num_hits += 1; - - if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { - /* Accept result and terminate if throughput is sufficiently low */ - payload.result = true; - return false; - } - else { - return true; - } - } - - payload.num_hits += 1; - payload.num_recorded_hits += 1; - - uint record_index = num_recorded_hits; - - const IntegratorShadowState state = payload.state; - - const uint max_record_hits = min(uint(max_hits), INTEGRATOR_SHADOW_ISECT_SIZE); - if (record_index >= max_record_hits) { - /* If maximum number of hits reached, find a hit to replace. */ - float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); - uint max_recorded_hit = 0; - - for (int i = 1; i < max_record_hits; i++) { - const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); - if (isect_t > max_recorded_t) { - max_recorded_t = isect_t; - max_recorded_hit = i; - } - } - - if (ray_tmax >= max_recorded_t) { - /* Accept hit, so that we don't consider any more hits beyond the distance of the - * current hit anymore. */ - payload.result = true; - return true; - } - - record_index = max_recorded_hit; - } - - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = ray_tmax; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type; - - /* Continue tracing. */ -# endif /* __TRANSPARENT_SHADOWS__ */ -#endif /* __SHADOW_RECORD_ALL__ */ - - return true; -} - -[[intersection(triangle, triangle_data, METALRT_TAGS)]] -TriangleIntersectionResult -__anyhit__cycles_metalrt_shadow_all_hit_tri(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], - unsigned int object [[user_instance_id]], - unsigned int primitive_id [[primitive_id]], - float2 barycentrics [[barycentric_coord]], - float ray_tmax [[distance]]) -{ - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - - TriangleIntersectionResult result; - result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_TRIANGLE>( - launch_params_metal, payload, object, prim, barycentrics, ray_tmax); - result.accept = !result.continue_search; - return result; -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_shadow_all_hit_box(const float ray_tmax [[max_distance]]) -{ - /* unused function */ - BoundingBoxIntersectionResult result; - result.distance = ray_tmax; - result.accept = false; - result.continue_search = false; - return result; -} - -template<typename TReturnType, uint intersection_type> -inline TReturnType metalrt_visibility_test(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload, - const uint object, - const uint prim, - const float u) -{ - TReturnType result; - -# ifdef __HAIR__ - if (intersection_type == METALRT_HIT_BOUNDING_BOX) { - /* Filter out curve endcaps. */ - if (u == 0.0f || u == 1.0f) { - result.accept = false; - result.continue_search = true; - return result; - } - } -# endif - - uint visibility = payload.visibility; -# ifdef __VISIBILITY_FLAG__ - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - result.accept = false; - result.continue_search = true; - return result; - } -# endif - - /* Shadow ray early termination. */ - if (visibility & PATH_RAY_SHADOW_OPAQUE) { - if (intersection_skip_self_shadow(payload.self, object, prim)) { - result.accept = false; - result.continue_search = true; - return result; - } - else { - result.accept = true; - result.continue_search = false; - return result; - } - } - else { - if (intersection_skip_self(payload.self, object, prim)) { - result.accept = false; - result.continue_search = true; - return result; - } - } - - result.accept = true; - result.continue_search = true; - return result; -} - -[[intersection(triangle, triangle_data, METALRT_TAGS)]] -TriangleIntersectionResult -__anyhit__cycles_metalrt_visibility_test_tri(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], - unsigned int object [[user_instance_id]], - unsigned int primitive_id [[primitive_id]]) -{ - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - TriangleIntersectionResult result = metalrt_visibility_test<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>( - launch_params_metal, payload, object, prim, 0.0f); - if (result.accept) { - payload.prim = prim; - payload.type = kernel_data_fetch(objects, object).primitive_type; - } - return result; -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance]]) -{ - /* Unused function */ - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; - return result; -} - -#ifdef __HAIR__ -ccl_device_inline -void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload, - const uint object, - const uint prim, - const uint type, - const float3 ray_P, - const float3 ray_D, - float time, - const float ray_tmin, - const float ray_tmax, - thread BoundingBoxIntersectionResult &result) -{ -# ifdef __VISIBILITY_FLAG__ - const uint visibility = payload.visibility; - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - return; - } -# endif - - Intersection isect; - isect.t = ray_tmax; - - MetalKernelContext context(launch_params_metal); - if (context.curve_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { - result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>( - launch_params_metal, payload, object, prim, isect.u); - if (result.accept) { - result.distance = isect.t; - payload.u = isect.u; - payload.v = isect.v; - payload.prim = prim; - payload.type = type; - } - } -} - -ccl_device_inline -void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, - const uint object, - const uint prim, - const uint type, - float time, - const float ray_tmin, - const float ray_tmax, - thread BoundingBoxIntersectionResult &result) -{ - const uint visibility = payload.visibility; - - Intersection isect; - isect.t = ray_tmax; - - MetalKernelContext context(launch_params_metal); - if (context.curve_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { - result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>( - launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax); - result.accept = !result.continue_search; - } -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], - const uint object [[user_instance_id]], - const uint primitive_id [[primitive_id]], - const float3 ray_P [[origin]], - const float3 ray_D [[direction]], - const float ray_tmin [[min_distance]], - const float ray_tmax [[max_distance]]) -{ - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; - - if (segment.type & PRIMITIVE_CURVE_RIBBON) { - metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D, -# if defined(__METALRT_MOTION__) - payload.time, -# else - 0.0f, -# endif - ray_tmin, ray_tmax, result); - } - - return result; -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], - const uint object [[user_instance_id]], - const uint primitive_id [[primitive_id]], - const float3 ray_P [[origin]], - const float3 ray_D [[direction]], - const float ray_tmin [[min_distance]], - const float ray_tmax [[max_distance]]) -{ - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; - - if (segment.type & PRIMITIVE_CURVE_RIBBON) { - metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D, -# if defined(__METALRT_MOTION__) - payload.time, -# else - 0.0f, -# endif - ray_tmin, ray_tmax, result); - } - - return result; -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], - const uint object [[user_instance_id]], - const uint primitive_id [[primitive_id]], - const float3 ray_P [[origin]], - const float3 ray_D [[direction]], - const float ray_tmin [[min_distance]], - const float ray_tmax [[max_distance]]) -{ - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; - metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D, -# if defined(__METALRT_MOTION__) - payload.time, -# else - 0.0f, -# endif - ray_tmin, ray_tmax, result); - - return result; -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], - const uint object [[user_instance_id]], - const uint primitive_id [[primitive_id]], - const float3 ray_P [[origin]], - const float3 ray_D [[direction]], - const float ray_tmin [[min_distance]], - const float ray_tmax [[max_distance]]) -{ - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; - - metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D, -# if defined(__METALRT_MOTION__) - payload.time, -# else - 0.0f, -# endif - ray_tmin, ray_tmax, result); - - return result; -} -#endif /* __HAIR__ */ - -#ifdef __POINTCLOUD__ -ccl_device_inline -void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload, - const uint object, - const uint prim, - const uint type, - const float3 ray_P, - const float3 ray_D, - float time, - const float ray_tmin, - const float ray_tmax, - thread BoundingBoxIntersectionResult &result) -{ -# ifdef __VISIBILITY_FLAG__ - const uint visibility = payload.visibility; - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - return; - } -# endif - - Intersection isect; - isect.t = ray_tmax; - - MetalKernelContext context(launch_params_metal); - if (context.point_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { - result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>( - launch_params_metal, payload, object, prim, isect.u); - if (result.accept) { - result.distance = isect.t; - payload.u = isect.u; - payload.v = isect.v; - payload.prim = prim; - payload.type = type; - } - } -} - -ccl_device_inline -void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, - const uint object, - const uint prim, - const uint type, - const float3 ray_P, - const float3 ray_D, - float time, - const float ray_tmin, - const float ray_tmax, - thread BoundingBoxIntersectionResult &result) -{ - const uint visibility = payload.visibility; - - Intersection isect; - isect.t = ray_tmax; - - MetalKernelContext context(launch_params_metal); - if (context.point_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { - result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>( - launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax); - result.accept = !result.continue_search; - - if (result.accept) { - result.distance = isect.t; - } - } -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], - const uint object [[user_instance_id]], - const uint primitive_id [[primitive_id]], - const float3 ray_origin [[origin]], - const float3 ray_direction [[direction]], - const float ray_tmin [[min_distance]], - const float ray_tmax [[max_distance]]) -{ - const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - const int type = kernel_data_fetch(objects, object).primitive_type; - - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; - - metalrt_intersection_point(launch_params_metal, payload, object, prim, type, ray_origin, ray_direction, -# if defined(__METALRT_MOTION__) - payload.time, -# else - 0.0f, -# endif - ray_tmin, ray_tmax, result); - - return result; -} - -[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] -BoundingBoxIntersectionResult -__intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], - const uint object [[user_instance_id]], - const uint primitive_id [[primitive_id]], - const float3 ray_origin [[origin]], - const float3 ray_direction [[direction]], - const float ray_tmin [[min_distance]], - const float ray_tmax [[max_distance]]) -{ - const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); - const int type = kernel_data_fetch(objects, object).primitive_type; - - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; - - metalrt_intersection_point_shadow(launch_params_metal, payload, object, prim, type, ray_origin, ray_direction, -# if defined(__METALRT_MOTION__) - payload.time, -# else - 0.0f, -# endif - ray_tmin, ray_tmax, result); - - return result; -} -#endif /* __POINTCLOUD__ */ -#endif /* __METALRT__ */ diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h new file mode 100644 index 00000000000..a1621277ec7 --- /dev/null +++ b/intern/cycles/kernel/device/optix/bvh.h @@ -0,0 +1,646 @@ +/* SPDX-License-Identifier: Apache-2.0 + * Copyright 2021-2022 Blender Foundation */ + +/* OptiX implementation of ray-scene intersection. */ + +#pragma once + +#include "kernel/bvh/types.h" +#include "kernel/bvh/util.h" + +#define OPTIX_DEFINE_ABI_VERSION_ONLY +#include <optix_function_table.h> + +CCL_NAMESPACE_BEGIN + +/* Utilities. */ + +template<typename T> ccl_device_forceinline T *get_payload_ptr_0() +{ + return pointer_unpack_from_uint<T>(optixGetPayload_0(), optixGetPayload_1()); +} +template<typename T> ccl_device_forceinline T *get_payload_ptr_2() +{ + return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3()); +} + +template<typename T> ccl_device_forceinline T *get_payload_ptr_6() +{ + return (T *)(((uint64_t)optixGetPayload_7() << 32) | optixGetPayload_6()); +} + +ccl_device_forceinline int get_object_id() +{ +#ifdef __OBJECT_MOTION__ + /* Always get the instance ID from the TLAS + * There might be a motion transform node between TLAS and BLAS which does not have one. */ + return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0)); +#else + return optixGetInstanceId(); +#endif +} + +/* Hit/miss functions. */ + +extern "C" __global__ void __miss__kernel_optix_miss() +{ + /* 'kernel_path_lamp_emission' checks intersection distance, so need to set it even on a miss. */ + optixSetPayload_0(__float_as_uint(optixGetRayTmax())); + optixSetPayload_5(PRIMITIVE_NONE); +} + +extern "C" __global__ void __anyhit__kernel_optix_local_hit() +{ +#if defined(__HAIR__) || defined(__POINTCLOUD__) + if (!optixIsTriangleHit()) { + /* Ignore curves and points. */ + return optixIgnoreIntersection(); + } +#endif + +#ifdef __BVH_LOCAL__ + const int object = get_object_id(); + if (object != optixGetPayload_4() /* local_object */) { + /* Only intersect with matching object. */ + return optixIgnoreIntersection(); + } + + const int prim = optixGetPrimitiveIndex(); + ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); + if (intersection_skip_self_local(ray->self, prim)) { + return optixIgnoreIntersection(); + } + + const uint max_hits = optixGetPayload_5(); + if (max_hits == 0) { + /* Special case for when no hit information is requested, just report that something was hit */ + optixSetPayload_5(true); + return optixTerminateRay(); + } + + int hit = 0; + uint *const lcg_state = get_payload_ptr_0<uint>(); + LocalIntersection *const local_isect = get_payload_ptr_2<LocalIntersection>(); + + if (lcg_state) { + for (int i = min(max_hits, local_isect->num_hits) - 1; i >= 0; --i) { + if (optixGetRayTmax() == local_isect->hits[i].t) { + return optixIgnoreIntersection(); + } + } + + hit = local_isect->num_hits++; + + if (local_isect->num_hits > max_hits) { + hit = lcg_step_uint(lcg_state) % local_isect->num_hits; + if (hit >= max_hits) { + return optixIgnoreIntersection(); + } + } + } + else { + if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) { + /* Record closest intersection only. + * Do not terminate ray here, since there is no guarantee about distance ordering in any-hit. + */ + return optixIgnoreIntersection(); + } + + local_isect->num_hits = 1; + } + + Intersection *isect = &local_isect->hits[hit]; + isect->t = optixGetRayTmax(); + isect->prim = prim; + isect->object = get_object_id(); + isect->type = kernel_data_fetch(objects, isect->object).primitive_type; + + const float2 barycentrics = optixGetTriangleBarycentrics(); + isect->u = 1.0f - barycentrics.y - barycentrics.x; + isect->v = barycentrics.x; + + /* Record geometric normal. */ + const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w; + const float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0); + const float3 tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1); + const float3 tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2); + local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); + + /* Continue tracing (without this the trace call would return after the first hit). */ + optixIgnoreIntersection(); +#endif +} + +extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() +{ +#ifdef __SHADOW_RECORD_ALL__ + int prim = optixGetPrimitiveIndex(); + const uint object = get_object_id(); +# ifdef __VISIBILITY_FLAG__ + const uint visibility = optixGetPayload_4(); + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + return optixIgnoreIntersection(); + } +# endif + + ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); + if (intersection_skip_self_shadow(ray->self, object, prim)) { + return optixIgnoreIntersection(); + } + + float u = 0.0f, v = 0.0f; + int type = 0; + if (optixIsTriangleHit()) { + const float2 barycentrics = optixGetTriangleBarycentrics(); + u = 1.0f - barycentrics.y - barycentrics.x; + v = barycentrics.x; + type = kernel_data_fetch(objects, object).primitive_type; + } +# ifdef __HAIR__ + else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) { + u = __uint_as_float(optixGetAttribute_0()); + v = __uint_as_float(optixGetAttribute_1()); + + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + type = segment.type; + prim = segment.prim; + +# if OPTIX_ABI_VERSION < 55 + /* Filter out curve endcaps. */ + if (u == 0.0f || u == 1.0f) { + return optixIgnoreIntersection(); + } +# endif + } +# endif + else { + type = kernel_data_fetch(objects, object).primitive_type; + u = 0.0f; + v = 0.0f; + } + +# ifndef __TRANSPARENT_SHADOWS__ + /* No transparent shadows support compiled in, make opaque. */ + optixSetPayload_5(true); + return optixTerminateRay(); +# else + const uint max_hits = optixGetPayload_3(); + const uint num_hits_packed = optixGetPayload_2(); + const uint num_recorded_hits = uint16_unpack_from_uint_0(num_hits_packed); + const uint num_hits = uint16_unpack_from_uint_1(num_hits_packed); + + /* If no transparent shadows, all light is blocked and we can stop immediately. */ + if (num_hits >= max_hits || + !(intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) { + optixSetPayload_5(true); + return optixTerminateRay(); + } + + /* Always use baked shadow transparency for curves. */ + if (type & PRIMITIVE_CURVE) { + float throughput = __uint_as_float(optixGetPayload_1()); + throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u); + optixSetPayload_1(__float_as_uint(throughput)); + optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits, num_hits + 1)); + + if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { + optixSetPayload_5(true); + return optixTerminateRay(); + } + else { + /* Continue tracing. */ + optixIgnoreIntersection(); + return; + } + } + + /* Record transparent intersection. */ + optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits + 1, num_hits + 1)); + + uint record_index = num_recorded_hits; + + const IntegratorShadowState state = optixGetPayload_0(); + + const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); + if (record_index >= max_record_hits) { + /* If maximum number of hits reached, find a hit to replace. */ + float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); + uint max_recorded_hit = 0; + + for (int i = 1; i < max_record_hits; i++) { + const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); + if (isect_t > max_recorded_t) { + max_recorded_t = isect_t; + max_recorded_hit = i; + } + } + + if (optixGetRayTmax() >= max_recorded_t) { + /* Accept hit, so that OptiX won't consider any more hits beyond the distance of the + * current hit anymore. */ + return; + } + + record_index = max_recorded_hit; + } + + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = optixGetRayTmax(); + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object; + INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type; + + /* Continue tracing. */ + optixIgnoreIntersection(); +# endif /* __TRANSPARENT_SHADOWS__ */ +#endif /* __SHADOW_RECORD_ALL__ */ +} + +extern "C" __global__ void __anyhit__kernel_optix_volume_test() +{ +#if defined(__HAIR__) || defined(__POINTCLOUD__) + if (!optixIsTriangleHit()) { + /* Ignore curves. */ + return optixIgnoreIntersection(); + } +#endif + + const uint object = get_object_id(); +#ifdef __VISIBILITY_FLAG__ + const uint visibility = optixGetPayload_4(); + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + return optixIgnoreIntersection(); + } +#endif + + if ((kernel_data_fetch(object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) { + return optixIgnoreIntersection(); + } + + const int prim = optixGetPrimitiveIndex(); + ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); + if (intersection_skip_self(ray->self, object, prim)) { + return optixIgnoreIntersection(); + } +} + +extern "C" __global__ void __anyhit__kernel_optix_visibility_test() +{ +#ifdef __HAIR__ +# if OPTIX_ABI_VERSION < 55 + if (optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE) { + /* Filter out curve endcaps. */ + const float u = __uint_as_float(optixGetAttribute_0()); + if (u == 0.0f || u == 1.0f) { + return optixIgnoreIntersection(); + } + } +# endif +#endif + + const uint object = get_object_id(); + const uint visibility = optixGetPayload_4(); +#ifdef __VISIBILITY_FLAG__ + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + return optixIgnoreIntersection(); + } +#endif + + const int prim = optixGetPrimitiveIndex(); + ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); + + if (visibility & PATH_RAY_SHADOW_OPAQUE) { + if (intersection_skip_self_shadow(ray->self, object, prim)) { + return optixIgnoreIntersection(); + } + else { + /* Shadow ray early termination. */ + return optixTerminateRay(); + } + } + else { + if (intersection_skip_self(ray->self, object, prim)) { + return optixIgnoreIntersection(); + } + } +} + +extern "C" __global__ void __closesthit__kernel_optix_hit() +{ + const int object = get_object_id(); + const int prim = optixGetPrimitiveIndex(); + + optixSetPayload_0(__float_as_uint(optixGetRayTmax())); /* Intersection distance */ + optixSetPayload_4(object); + + if (optixIsTriangleHit()) { + const float2 barycentrics = optixGetTriangleBarycentrics(); + optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x)); + optixSetPayload_2(__float_as_uint(barycentrics.x)); + optixSetPayload_3(prim); + optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type); + } + else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) { + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */ + optixSetPayload_2(optixGetAttribute_1()); + optixSetPayload_3(segment.prim); + optixSetPayload_5(segment.type); + } + else { + optixSetPayload_1(0); + optixSetPayload_2(0); + optixSetPayload_3(prim); + optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type); + } +} + +/* Custom primitive intersection functions. */ + +#ifdef __HAIR__ +ccl_device_inline void optix_intersection_curve(const int prim, const int type) +{ + const int object = get_object_id(); + +# ifdef __VISIBILITY_FLAG__ + const uint visibility = optixGetPayload_4(); + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + return; + } +# endif + + const float3 ray_P = optixGetObjectRayOrigin(); + const float3 ray_D = optixGetObjectRayDirection(); + const float ray_tmin = optixGetRayTmin(); + +# ifdef __OBJECT_MOTION__ + const float time = optixGetRayTime(); +# else + const float time = 0.0f; +# endif + + Intersection isect; + isect.t = optixGetRayTmax(); + + if (curve_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { + static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use"); + optixReportIntersection(isect.t, + type & PRIMITIVE_ALL, + __float_as_int(isect.u), /* Attribute_0 */ + __float_as_int(isect.v)); /* Attribute_1 */ + } +} + +extern "C" __global__ void __intersection__curve_ribbon() +{ + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, optixGetPrimitiveIndex()); + const int prim = segment.prim; + const int type = segment.type; + if (type & PRIMITIVE_CURVE_RIBBON) { + optix_intersection_curve(prim, type); + } +} + +#endif + +#ifdef __POINTCLOUD__ +extern "C" __global__ void __intersection__point() +{ + const int prim = optixGetPrimitiveIndex(); + const int object = get_object_id(); + const int type = kernel_data_fetch(objects, object).primitive_type; + +# ifdef __VISIBILITY_FLAG__ + const uint visibility = optixGetPayload_4(); + if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { + return; + } +# endif + + const float3 ray_P = optixGetObjectRayOrigin(); + const float3 ray_D = optixGetObjectRayDirection(); + const float ray_tmin = optixGetRayTmin(); + +# ifdef __OBJECT_MOTION__ + const float time = optixGetRayTime(); +# else + const float time = 0.0f; +# endif + + Intersection isect; + isect.t = optixGetRayTmax(); + + if (point_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { + static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use"); + optixReportIntersection(isect.t, type & PRIMITIVE_ALL); + } +} +#endif + +/* Scene intersection. */ + +ccl_device_intersect bool scene_intersect(KernelGlobals kg, + ccl_private const Ray *ray, + const uint visibility, + ccl_private Intersection *isect) +{ + uint p0 = 0; + uint p1 = 0; + uint p2 = 0; + uint p3 = 0; + uint p4 = visibility; + uint p5 = PRIMITIVE_NONE; + uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; + uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; + + uint ray_mask = visibility & 0xFF; + uint ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + } + else if (visibility & PATH_RAY_SHADOW_OPAQUE) { + ray_flags |= OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT; + } + + optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0, + ray->P, + ray->D, + ray->tmin, + ray->tmax, + ray->time, + ray_mask, + ray_flags, + 0, /* SBT offset for PG_HITD */ + 0, + 0, + p0, + p1, + p2, + p3, + p4, + p5, + p6, + p7); + + isect->t = __uint_as_float(p0); + isect->u = __uint_as_float(p1); + isect->v = __uint_as_float(p2); + isect->prim = p3; + isect->object = p4; + isect->type = p5; + + return p5 != PRIMITIVE_NONE; +} + +#ifdef __BVH_LOCAL__ +ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, + ccl_private const Ray *ray, + ccl_private LocalIntersection *local_isect, + int local_object, + ccl_private uint *lcg_state, + int max_hits) +{ + uint p0 = pointer_pack_to_uint_0(lcg_state); + uint p1 = pointer_pack_to_uint_1(lcg_state); + uint p2 = pointer_pack_to_uint_0(local_isect); + uint p3 = pointer_pack_to_uint_1(local_isect); + uint p4 = local_object; + uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; + uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; + + /* Is set to zero on miss or if ray is aborted, so can be used as return value. */ + uint p5 = max_hits; + + if (local_isect) { + local_isect->num_hits = 0; /* Initialize hit count to zero. */ + } + optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0, + ray->P, + ray->D, + ray->tmin, + ray->tmax, + ray->time, + 0xFF, + /* Need to always call into __anyhit__kernel_optix_local_hit. */ + OPTIX_RAY_FLAG_ENFORCE_ANYHIT, + 2, /* SBT offset for PG_HITL */ + 0, + 0, + p0, + p1, + p2, + p3, + p4, + p5, + p6, + p7); + + return p5; +} +#endif + +#ifdef __SHADOW_RECORD_ALL__ +ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, + IntegratorShadowState state, + ccl_private const Ray *ray, + uint visibility, + uint max_hits, + ccl_private uint *num_recorded_hits, + ccl_private float *throughput) +{ + uint p0 = state; + uint p1 = __float_as_uint(1.0f); /* Throughput. */ + uint p2 = 0; /* Number of hits. */ + uint p3 = max_hits; + uint p4 = visibility; + uint p5 = false; + uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; + uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + } + + optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0, + ray->P, + ray->D, + ray->tmin, + ray->tmax, + ray->time, + ray_mask, + /* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */ + OPTIX_RAY_FLAG_ENFORCE_ANYHIT, + 1, /* SBT offset for PG_HITS */ + 0, + 0, + p0, + p1, + p2, + p3, + p4, + p5, + p6, + p7); + + *num_recorded_hits = uint16_unpack_from_uint_0(p2); + *throughput = __uint_as_float(p1); + + return p5; +} +#endif + +#ifdef __VOLUME__ +ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, + ccl_private const Ray *ray, + ccl_private Intersection *isect, + const uint visibility) +{ + uint p0 = 0; + uint p1 = 0; + uint p2 = 0; + uint p3 = 0; + uint p4 = visibility; + uint p5 = PRIMITIVE_NONE; + uint p6 = ((uint64_t)ray) & 0xFFFFFFFF; + uint p7 = (((uint64_t)ray) >> 32) & 0xFFFFFFFF; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + } + + optixTrace(intersection_ray_valid(ray) ? kernel_data.device_bvh : 0, + ray->P, + ray->D, + ray->tmin, + ray->tmax, + ray->time, + ray_mask, + /* Need to always call into __anyhit__kernel_optix_volume_test. */ + OPTIX_RAY_FLAG_ENFORCE_ANYHIT, + 3, /* SBT offset for PG_HITV */ + 0, + 0, + p0, + p1, + p2, + p3, + p4, + p5, + p6, + p7); + + isect->t = __uint_as_float(p0); + isect->u = __uint_as_float(p1); + isect->v = __uint_as_float(p2); + isect->prim = p3; + isect->object = p4; + isect->type = p5; + + return p5 != PRIMITIVE_NONE; +} +#endif + +CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/device/optix/compat.h b/intern/cycles/kernel/device/optix/compat.h index aa4a6321a8b..1a11a533b7e 100644 --- a/intern/cycles/kernel/device/optix/compat.h +++ b/intern/cycles/kernel/device/optix/compat.h @@ -8,7 +8,6 @@ #include <optix.h> #define __KERNEL_GPU__ -#define __KERNEL_GPU_RAYTRACING__ #define __KERNEL_CUDA__ /* OptiX kernels are implicitly CUDA kernels too */ #define __KERNEL_OPTIX__ #define CCL_NAMESPACE_BEGIN diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index 204aa8182a1..6abb5aeacb9 100644 --- a/intern/cycles/kernel/device/optix/kernel.cu +++ b/intern/cycles/kernel/device/optix/kernel.cu @@ -20,34 +20,6 @@ #include "kernel/integrator/intersect_volume_stack.h" // clang-format on -#define OPTIX_DEFINE_ABI_VERSION_ONLY -#include <optix_function_table.h> - -template<typename T> ccl_device_forceinline T *get_payload_ptr_0() -{ - return pointer_unpack_from_uint<T>(optixGetPayload_0(), optixGetPayload_1()); -} -template<typename T> ccl_device_forceinline T *get_payload_ptr_2() -{ - return pointer_unpack_from_uint<T>(optixGetPayload_2(), optixGetPayload_3()); -} - -template<typename T> ccl_device_forceinline T *get_payload_ptr_6() -{ - return (T *)(((uint64_t)optixGetPayload_7() << 32) | optixGetPayload_6()); -} - -ccl_device_forceinline int get_object_id() -{ -#ifdef __OBJECT_MOTION__ - /* Always get the instance ID from the TLAS - * There might be a motion transform node between TLAS and BLAS which does not have one. */ - return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0)); -#else - return optixGetInstanceId(); -#endif -} - extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest() { const int global_index = optixGetLaunchIndex().x; @@ -84,396 +56,3 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_st integrator_intersect_volume_stack(nullptr, path_index); } -extern "C" __global__ void __miss__kernel_optix_miss() -{ - /* 'kernel_path_lamp_emission' checks intersection distance, so need to set it even on a miss. */ - optixSetPayload_0(__float_as_uint(optixGetRayTmax())); - optixSetPayload_5(PRIMITIVE_NONE); -} - -extern "C" __global__ void __anyhit__kernel_optix_local_hit() -{ -#if defined(__HAIR__) || defined(__POINTCLOUD__) - if (!optixIsTriangleHit()) { - /* Ignore curves and points. */ - return optixIgnoreIntersection(); - } -#endif - -#ifdef __BVH_LOCAL__ - const int object = get_object_id(); - if (object != optixGetPayload_4() /* local_object */) { - /* Only intersect with matching object. */ - return optixIgnoreIntersection(); - } - - const int prim = optixGetPrimitiveIndex(); - ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); - if (intersection_skip_self_local(ray->self, prim)) { - return optixIgnoreIntersection(); - } - - const uint max_hits = optixGetPayload_5(); - if (max_hits == 0) { - /* Special case for when no hit information is requested, just report that something was hit */ - optixSetPayload_5(true); - return optixTerminateRay(); - } - - int hit = 0; - uint *const lcg_state = get_payload_ptr_0<uint>(); - LocalIntersection *const local_isect = get_payload_ptr_2<LocalIntersection>(); - - if (lcg_state) { - for (int i = min(max_hits, local_isect->num_hits) - 1; i >= 0; --i) { - if (optixGetRayTmax() == local_isect->hits[i].t) { - return optixIgnoreIntersection(); - } - } - - hit = local_isect->num_hits++; - - if (local_isect->num_hits > max_hits) { - hit = lcg_step_uint(lcg_state) % local_isect->num_hits; - if (hit >= max_hits) { - return optixIgnoreIntersection(); - } - } - } - else { - if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) { - /* Record closest intersection only. - * Do not terminate ray here, since there is no guarantee about distance ordering in any-hit. - */ - return optixIgnoreIntersection(); - } - - local_isect->num_hits = 1; - } - - Intersection *isect = &local_isect->hits[hit]; - isect->t = optixGetRayTmax(); - isect->prim = prim; - isect->object = get_object_id(); - isect->type = kernel_data_fetch(objects, isect->object).primitive_type; - - const float2 barycentrics = optixGetTriangleBarycentrics(); - isect->u = 1.0f - barycentrics.y - barycentrics.x; - isect->v = barycentrics.x; - - /* Record geometric normal. */ - const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w; - const float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0); - const float3 tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1); - const float3 tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2); - local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); - - /* Continue tracing (without this the trace call would return after the first hit). */ - optixIgnoreIntersection(); -#endif -} - -extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() -{ -#ifdef __SHADOW_RECORD_ALL__ - int prim = optixGetPrimitiveIndex(); - const uint object = get_object_id(); -# ifdef __VISIBILITY_FLAG__ - const uint visibility = optixGetPayload_4(); - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - return optixIgnoreIntersection(); - } -# endif - - ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); - if (intersection_skip_self_shadow(ray->self, object, prim)) { - return optixIgnoreIntersection(); - } - - float u = 0.0f, v = 0.0f; - int type = 0; - if (optixIsTriangleHit()) { - const float2 barycentrics = optixGetTriangleBarycentrics(); - u = 1.0f - barycentrics.y - barycentrics.x; - v = barycentrics.x; - type = kernel_data_fetch(objects, object).primitive_type; - } -# ifdef __HAIR__ - else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) { - u = __uint_as_float(optixGetAttribute_0()); - v = __uint_as_float(optixGetAttribute_1()); - - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - type = segment.type; - prim = segment.prim; - -# if OPTIX_ABI_VERSION < 55 - /* Filter out curve endcaps. */ - if (u == 0.0f || u == 1.0f) { - return optixIgnoreIntersection(); - } -# endif - } -# endif - else { - type = kernel_data_fetch(objects, object).primitive_type; - u = 0.0f; - v = 0.0f; - } - -# ifndef __TRANSPARENT_SHADOWS__ - /* No transparent shadows support compiled in, make opaque. */ - optixSetPayload_5(true); - return optixTerminateRay(); -# else - const uint max_hits = optixGetPayload_3(); - const uint num_hits_packed = optixGetPayload_2(); - const uint num_recorded_hits = uint16_unpack_from_uint_0(num_hits_packed); - const uint num_hits = uint16_unpack_from_uint_1(num_hits_packed); - - /* If no transparent shadows, all light is blocked and we can stop immediately. */ - if (num_hits >= max_hits || - !(intersection_get_shader_flags(NULL, prim, type) & SD_HAS_TRANSPARENT_SHADOW)) { - optixSetPayload_5(true); - return optixTerminateRay(); - } - - /* Always use baked shadow transparency for curves. */ - if (type & PRIMITIVE_CURVE) { - float throughput = __uint_as_float(optixGetPayload_1()); - throughput *= intersection_curve_shadow_transparency(nullptr, object, prim, u); - optixSetPayload_1(__float_as_uint(throughput)); - optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits, num_hits + 1)); - - if (throughput < CURVE_SHADOW_TRANSPARENCY_CUTOFF) { - optixSetPayload_5(true); - return optixTerminateRay(); - } - else { - /* Continue tracing. */ - optixIgnoreIntersection(); - return; - } - } - - /* Record transparent intersection. */ - optixSetPayload_2(uint16_pack_to_uint(num_recorded_hits + 1, num_hits + 1)); - - uint record_index = num_recorded_hits; - - const IntegratorShadowState state = optixGetPayload_0(); - - const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE); - if (record_index >= max_record_hits) { - /* If maximum number of hits reached, find a hit to replace. */ - float max_recorded_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, t); - uint max_recorded_hit = 0; - - for (int i = 1; i < max_record_hits; i++) { - const float isect_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, i, t); - if (isect_t > max_recorded_t) { - max_recorded_t = isect_t; - max_recorded_hit = i; - } - } - - if (optixGetRayTmax() >= max_recorded_t) { - /* Accept hit, so that OptiX won't consider any more hits beyond the distance of the - * current hit anymore. */ - return; - } - - record_index = max_recorded_hit; - } - - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, u) = u; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, v) = v; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, t) = optixGetRayTmax(); - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, prim) = prim; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, object) = object; - INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, record_index, type) = type; - - /* Continue tracing. */ - optixIgnoreIntersection(); -# endif /* __TRANSPARENT_SHADOWS__ */ -#endif /* __SHADOW_RECORD_ALL__ */ -} - -extern "C" __global__ void __anyhit__kernel_optix_volume_test() -{ -#if defined(__HAIR__) || defined(__POINTCLOUD__) - if (!optixIsTriangleHit()) { - /* Ignore curves. */ - return optixIgnoreIntersection(); - } -#endif - - const uint object = get_object_id(); -#ifdef __VISIBILITY_FLAG__ - const uint visibility = optixGetPayload_4(); - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - return optixIgnoreIntersection(); - } -#endif - - if ((kernel_data_fetch(object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) { - return optixIgnoreIntersection(); - } - - const int prim = optixGetPrimitiveIndex(); - ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); - if (intersection_skip_self(ray->self, object, prim)) { - return optixIgnoreIntersection(); - } -} - -extern "C" __global__ void __anyhit__kernel_optix_visibility_test() -{ -#ifdef __HAIR__ -# if OPTIX_ABI_VERSION < 55 - if (optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE) { - /* Filter out curve endcaps. */ - const float u = __uint_as_float(optixGetAttribute_0()); - if (u == 0.0f || u == 1.0f) { - return optixIgnoreIntersection(); - } - } -# endif -#endif - - const uint object = get_object_id(); - const uint visibility = optixGetPayload_4(); -#ifdef __VISIBILITY_FLAG__ - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - return optixIgnoreIntersection(); - } -#endif - - const int prim = optixGetPrimitiveIndex(); - ccl_private Ray *const ray = get_payload_ptr_6<Ray>(); - - if (visibility & PATH_RAY_SHADOW_OPAQUE) { - if (intersection_skip_self_shadow(ray->self, object, prim)) { - return optixIgnoreIntersection(); - } - else { - /* Shadow ray early termination. */ - return optixTerminateRay(); - } - } - else { - if (intersection_skip_self(ray->self, object, prim)) { - return optixIgnoreIntersection(); - } - } -} - -extern "C" __global__ void __closesthit__kernel_optix_hit() -{ - const int object = get_object_id(); - const int prim = optixGetPrimitiveIndex(); - - optixSetPayload_0(__float_as_uint(optixGetRayTmax())); /* Intersection distance */ - optixSetPayload_4(object); - - if (optixIsTriangleHit()) { - const float2 barycentrics = optixGetTriangleBarycentrics(); - optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x)); - optixSetPayload_2(__float_as_uint(barycentrics.x)); - optixSetPayload_3(prim); - optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type); - } - else if ((optixGetHitKind() & (~PRIMITIVE_MOTION)) != PRIMITIVE_POINT) { - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */ - optixSetPayload_2(optixGetAttribute_1()); - optixSetPayload_3(segment.prim); - optixSetPayload_5(segment.type); - } - else { - optixSetPayload_1(0); - optixSetPayload_2(0); - optixSetPayload_3(prim); - optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type); - } -} - -#ifdef __HAIR__ -ccl_device_inline void optix_intersection_curve(const int prim, const int type) -{ - const int object = get_object_id(); - -# ifdef __VISIBILITY_FLAG__ - const uint visibility = optixGetPayload_4(); - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - return; - } -# endif - - const float3 ray_P = optixGetObjectRayOrigin(); - const float3 ray_D = optixGetObjectRayDirection(); - const float ray_tmin = optixGetRayTmin(); - -# ifdef __OBJECT_MOTION__ - const float time = optixGetRayTime(); -# else - const float time = 0.0f; -# endif - - Intersection isect; - isect.t = optixGetRayTmax(); - - if (curve_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { - static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use"); - optixReportIntersection(isect.t, - type & PRIMITIVE_ALL, - __float_as_int(isect.u), /* Attribute_0 */ - __float_as_int(isect.v)); /* Attribute_1 */ - } -} - -extern "C" __global__ void __intersection__curve_ribbon() -{ - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, optixGetPrimitiveIndex()); - const int prim = segment.prim; - const int type = segment.type; - if (type & PRIMITIVE_CURVE_RIBBON) { - optix_intersection_curve(prim, type); - } -} - -#endif - -#ifdef __POINTCLOUD__ -extern "C" __global__ void __intersection__point() -{ - const int prim = optixGetPrimitiveIndex(); - const int object = get_object_id(); - const int type = kernel_data_fetch(objects, object).primitive_type; - -# ifdef __VISIBILITY_FLAG__ - const uint visibility = optixGetPayload_4(); - if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) { - return; - } -# endif - - const float3 ray_P = optixGetObjectRayOrigin(); - const float3 ray_D = optixGetObjectRayDirection(); - const float ray_tmin = optixGetRayTmin(); - -# ifdef __OBJECT_MOTION__ - const float time = optixGetRayTime(); -# else - const float time = 0.0f; -# endif - - Intersection isect; - isect.t = optixGetRayTmax(); - - if (point_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) { - static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use"); - optixReportIntersection(isect.t, type & PRIMITIVE_ALL); - } -} -#endif diff --git a/intern/cycles/kernel/integrator/intersect_volume_stack.h b/intern/cycles/kernel/integrator/intersect_volume_stack.h index 9ba4a0a3964..b53bee11312 100644 --- a/intern/cycles/kernel/integrator/intersect_volume_stack.h +++ b/intern/cycles/kernel/integrator/intersect_volume_stack.h @@ -38,8 +38,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg, #ifdef __VOLUME_RECORD_ALL__ Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1]; - uint num_hits = scene_intersect_volume_all( - kg, &volume_ray, hits, 2 * volume_stack_size, visibility); + uint num_hits = scene_intersect_volume(kg, &volume_ray, hits, 2 * volume_stack_size, visibility); if (num_hits > 0) { Intersection *isect = hits; @@ -108,8 +107,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s #ifdef __VOLUME_RECORD_ALL__ Intersection hits[2 * MAX_VOLUME_STACK_SIZE + 1]; - uint num_hits = scene_intersect_volume_all( - kg, &volume_ray, hits, 2 * volume_stack_size, visibility); + uint num_hits = scene_intersect_volume(kg, &volume_ray, hits, 2 * volume_stack_size, visibility); if (num_hits > 0) { int enclosed_volumes[MAX_VOLUME_STACK_SIZE]; Intersection *isect = hits; diff --git a/intern/cycles/kernel/integrator/subsurface_random_walk.h b/intern/cycles/kernel/integrator/subsurface_random_walk.h index 7857673b271..e43bbb3c50a 100644 --- a/intern/cycles/kernel/integrator/subsurface_random_walk.h +++ b/intern/cycles/kernel/integrator/subsurface_random_walk.h @@ -377,7 +377,6 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg, hit = (ss_isect.num_hits > 0); if (hit) { - /* t is always in world space with OptiX and MetalRT. */ ray.tmax = ss_isect.hits[0].t; } diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h index 4f4b811a8e7..d809f6fc2bd 100644 --- a/intern/cycles/kernel/types.h +++ b/intern/cycles/kernel/types.h @@ -83,7 +83,6 @@ CCL_NAMESPACE_BEGIN #define __LAMP_MIS__ #define __CAMERA_MOTION__ #define __OBJECT_MOTION__ -#define __BAKING__ #define __PRINCIPLED__ #define __SUBSURFACE__ #define __VOLUME__ @@ -99,10 +98,6 @@ CCL_NAMESPACE_BEGIN # define __VOLUME_RECORD_ALL__ #endif /* __KERNEL_CPU__ */ -#ifdef __KERNEL_GPU_RAYTRACING__ -# undef __BAKING__ -#endif /* __KERNEL_GPU_RAYTRACING__ */ - /* MNEE currently causes "Compute function exceeds available temporary registers" * on Metal, disabled for now. */ #ifndef __KERNEL_METAL__ @@ -129,9 +124,6 @@ CCL_NAMESPACE_BEGIN # if !(__KERNEL_FEATURES & KERNEL_FEATURE_SUBSURFACE) # undef __SUBSURFACE__ # endif -# if !(__KERNEL_FEATURES & KERNEL_FEATURE_BAKING) -# undef __BAKING__ -# endif # if !(__KERNEL_FEATURES & KERNEL_FEATURE_PATCH_EVALUATION) # undef __PATCH_EVAL__ # endif |