From 5ec249555038fece3e0f580efa67235d7e7fa986 Mon Sep 17 00:00:00 2001
From: Xavier Hallade <xavier.hallade@intel.com>
Date: Thu, 11 May 2023 18:47:25 +0200
Subject: [PATCH] Cycles: oneAPI: enable Hardware Raytracing for Raytrace/MNEE
 kernels

We do so if Embree 4.1+ is present.
---
 intern/cycles/device/oneapi/device_impl.cpp   |  6 +++-
 intern/cycles/kernel/device/oneapi/kernel.cpp | 30 +++++++++++--------
 2 files changed, 22 insertions(+), 14 deletions(-)

diff --git a/intern/cycles/device/oneapi/device_impl.cpp b/intern/cycles/device/oneapi/device_impl.cpp
index 264063c5676..c6218697589 100644
--- a/intern/cycles/device/oneapi/device_impl.cpp
+++ b/intern/cycles/device/oneapi/device_impl.cpp
@@ -122,8 +122,12 @@ bool OneapiDevice::check_peer_access(Device * /*peer_device*/)
 
 bool OneapiDevice::can_use_hardware_raytracing_for_features(uint requested_features) const
 {
-  /* MNEE and Ray-trace kernels currently don't work correctly with HWRT. */
+  /* MNEE and Raytrace kernels work correctly with Hardware Raytracing starting with Embree 4.1. */
+#  if defined(RTC_VERSION) && RTC_VERSION < 40100
   return !(requested_features & (KERNEL_FEATURE_MNEE | KERNEL_FEATURE_NODE_RAYTRACE));
+#  else
+  return true;
+#  endif
 }
 
 BVHLayoutMask OneapiDevice::get_bvh_layout_mask(uint requested_features) const
diff --git a/intern/cycles/kernel/device/oneapi/kernel.cpp b/intern/cycles/kernel/device/oneapi/kernel.cpp
index f57a2ee9255..92a37755ee5 100644
--- a/intern/cycles/kernel/device/oneapi/kernel.cpp
+++ b/intern/cycles/kernel/device/oneapi/kernel.cpp
@@ -174,27 +174,30 @@ bool oneapi_kernel_is_required_for_features(const std::string &kernel_name,
   return true;
 }
 
-bool oneapi_kernel_is_raytrace_or_mnee(const std::string &kernel_name)
+bool oneapi_kernel_is_compatible_with_hardware_raytracing(const std::string &kernel_name)
 {
-  return (kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE)) !=
-          std::string::npos) ||
+  /* MNEE and Raytrace kernels work correctly with Hardware Raytracing starting with Embree 4.1.
+   */
+#  if defined(RTC_VERSION) && RTC_VERSION < 40100
+  return (kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE)) ==
+          std::string::npos) &&
          (kernel_name.find(device_kernel_as_string(
-              DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE)) != std::string::npos);
+              DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE)) == std::string::npos);
+#  else
+  return true;
+#  endif
 }
 
-bool oneapi_kernel_is_using_embree(const std::string &kernel_name)
+bool oneapi_kernel_has_intersections(const std::string &kernel_name)
 {
-#  ifdef WITH_EMBREE_GPU
-  /* MNEE and Ray-trace kernels aren't yet enabled to use Embree. */
   for (int i = 0; i < (int)DEVICE_KERNEL_NUM; i++) {
     DeviceKernel kernel = (DeviceKernel)i;
     if (device_kernel_has_intersection(kernel)) {
       if (kernel_name.find(device_kernel_as_string(kernel)) != std::string::npos) {
-        return !oneapi_kernel_is_raytrace_or_mnee(kernel_name);
+        return true;
       }
     }
   }
-#  endif
   return false;
 }
 
@@ -217,7 +220,8 @@ bool oneapi_load_kernels(SyclQueue *queue_,
         const std::string &kernel_name = kernel_id.get_name();
 
         if (!oneapi_kernel_is_required_for_features(kernel_name, kernel_features) ||
-            !oneapi_kernel_is_using_embree(kernel_name))
+            !(oneapi_kernel_has_intersections(kernel_name) &&
+              oneapi_kernel_is_compatible_with_hardware_raytracing(kernel_name)))
         {
           continue;
         }
@@ -260,14 +264,14 @@ bool oneapi_load_kernels(SyclQueue *queue_,
       /* In case HWRT is on, compilation of kernels using Embree is already handled in previous
        * block. */
       if (!oneapi_kernel_is_required_for_features(kernel_name, kernel_features) ||
-          (use_hardware_raytracing && oneapi_kernel_is_using_embree(kernel_name)))
+          (use_hardware_raytracing && oneapi_kernel_has_intersections(kernel_name) &&
+           oneapi_kernel_is_compatible_with_hardware_raytracing(kernel_name)))
       {
         continue;
       }
 
 #  ifdef WITH_EMBREE_GPU
-      if (oneapi_kernel_is_using_embree(kernel_name) ||
-          oneapi_kernel_is_raytrace_or_mnee(kernel_name)) {
+      if (oneapi_kernel_has_intersections(kernel_name)) {
         sycl::kernel_bundle<sycl::bundle_state::input> one_kernel_bundle_input =
             sycl::get_kernel_bundle<sycl::bundle_state::input>(queue->get_context(), {kernel_id});
         one_kernel_bundle_input