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.

16 files changed, 2528 insertions, 2326 deletions

diff --git a/intern/cycles/bvh/embree.cpp b/intern/cycles/bvh/embree.cpp
index eed7ae19965..be5785de473 100644
--- a/intern/cycles/bvh/embree.cpp
+++ b/intern/cycles/bvh/embree.cpp
@@ -21,13 +21,9 @@
 
 #  include "bvh/embree.h"
 
-/* Kernel includes are necessary so that the filter function for Embree can access the packed BVH.
- */
-#  include "kernel/bvh/embree.h"
-#  include "kernel/bvh/util.h"
+#  include "kernel/device/cpu/bvh.h"
 #  include "kernel/device/cpu/compat.h"
 #  include "kernel/device/cpu/globals.h"
-#  include "kernel/sample/lcg.h"
 
 #  include "scene/hair.h"
 #  include "scene/mesh.h"
@@ -46,265 +42,6 @@ static_assert(Object::MAX_MOTION_STEPS <= RTC_MAX_TIME_STEP_COUNT,
 static_assert(Object::MAX_MOTION_STEPS == Geometry::MAX_MOTION_STEPS,
               "Object and Geometry max motion steps inconsistent");
 
-#  define IS_HAIR(x) (x & 1)
-
-/* 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.
- */
-static void rtc_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.
- */
-static void rtc_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, &current_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, &current_isect, ctx->local_object_id);
-      }
-      else {
-        kernel_embree_convert_hit(kg, ray, hit, &current_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 (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, &current_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;
-  }
-}
-
-static void rtc_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;
-  }
-}
-
-static void rtc_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;
-  }
-
-  rtc_filter_occluded_func(args);
-}
-
 static size_t unaccounted_mem = 0;
 
 static bool rtc_memory_monitor_func(void *userPtr, const ssize_t bytes, const bool)
@@ -535,8 +272,8 @@ void BVHEmbree::add_triangles(const Object *ob, const Mesh *mesh, int i)
   set_tri_vertex_buffer(geom_id, mesh, false);
 
   rtcSetGeometryUserData(geom_id, (void *)prim_offset);
-  rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
-  rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
+  rtcSetGeometryOccludedFilterFunction(geom_id, kernel_embree_filter_occluded_func);
+  rtcSetGeometryIntersectFilterFunction(geom_id, kernel_embree_filter_intersection_func);
   rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
 
   rtcCommitGeometry(geom_id);
@@ -739,8 +476,8 @@ void BVHEmbree::add_points(const Object *ob, const PointCloud *pointcloud, int i
   set_point_vertex_buffer(geom_id, pointcloud, false);
 
   rtcSetGeometryUserData(geom_id, (void *)prim_offset);
-  rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_backface_cull);
-  rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_backface_cull);
+  rtcSetGeometryIntersectFilterFunction(geom_id, kernel_embree_filter_func_backface_cull);
+  rtcSetGeometryOccludedFilterFunction(geom_id, kernel_embree_filter_occluded_func_backface_cull);
   rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
 
   rtcCommitGeometry(geom_id);
@@ -799,12 +536,13 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, int i)
 
   rtcSetGeometryUserData(geom_id, (void *)prim_offset);
   if (hair->curve_shape == CURVE_RIBBON) {
-    rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_intersection_func);
-    rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func);
+    rtcSetGeometryIntersectFilterFunction(geom_id, kernel_embree_filter_intersection_func);
+    rtcSetGeometryOccludedFilterFunction(geom_id, kernel_embree_filter_occluded_func);
   }
   else {
-    rtcSetGeometryIntersectFilterFunction(geom_id, rtc_filter_func_backface_cull);
-    rtcSetGeometryOccludedFilterFunction(geom_id, rtc_filter_occluded_func_backface_cull);
+    rtcSetGeometryIntersectFilterFunction(geom_id, kernel_embree_filter_func_backface_cull);
+    rtcSetGeometryOccludedFilterFunction(geom_id,
+                                         kernel_embree_filter_occluded_func_backface_cull);
   }
   rtcSetGeometryMask(geom_id, ob->visibility_for_tracing());
 
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, &current_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, &current_isect, ctx->local_object_id);
+      }
+      else {
+        kernel_embree_convert_hit(kg, ray, hit, &current_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, &current_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