Merge branch 'master' into sw/avx10

Author: stefanatwork

common/sys/vector.h

@@ -5,6 +5,7 @@
 
 #include "alloc.h"
 #include <algorithm>
+#include <type_traits>
 
 namespace embree
 {

common/tasking/taskschedulerinternal.h

@@ -14,6 +14,7 @@
 #include "../sys/atomic.h"
 #include "../math/range.h"
 
+#include <exception>
 #include <list>
 
 namespace embree

kernels/common/device.cpp

@@ -668,31 +668,35 @@ namespace embree
     bool ze_rtas_builder = false;
     for (uint32_t i=0; i<extensions.size(); i++)
     {
-      if (strncmp("ZE_experimental_rtas_builder",extensions[i].name,sizeof(extensions[i].name)) == 0)
+      if (strncmp("ZE_extension_rtas",extensions[i].name,sizeof(extensions[i].name)) == 0) {
         ze_rtas_builder = true;
+      }
+      if (strncmp("ZE_experimental_relaxed_allocation_limits", extensions[i].name, sizeof(extensions[i].name)) == 0) {
+        hasRelaxedAllocationLimits = true;
+      }
     }
     if (!ze_rtas_builder)
-      throw_RTCError(RTC_ERROR_LEVEL_ZERO_RAYTRACING_SUPPORT_MISSING, "ZE_experimental_rtas_builder extension not found. Please install a recent driver. On Linux, make sure that the package intel-level-zero-gpu-raytracing is installed");
+      throw_RTCError(RTC_ERROR_LEVEL_ZERO_RAYTRACING_SUPPORT_MISSING, "ZE_extension_rtas extension not found. Please install a recent driver. On Linux, make sure that the package intel-level-zero-gpu-raytracing is installed");
 
     result = ZeWrapper::initRTASBuilder(hDriver);
     if (result == ZE_RESULT_ERROR_DEPENDENCY_UNAVAILABLE) {
-      throw_RTCError(RTC_ERROR_LEVEL_ZERO_RAYTRACING_SUPPORT_MISSING, "cannot load ZE_experimental_rtas_builder extension. Please install a recent driver. On Linux, make sure that the package intel-level-zero-gpu-raytracing is installed");
+      throw_RTCError(RTC_ERROR_LEVEL_ZERO_RAYTRACING_SUPPORT_MISSING, "cannot load ZE_extension_rtas extension. Please install a recent driver. On Linux, make sure that the package intel-level-zero-gpu-raytracing is installed");
     }
     if (result != ZE_RESULT_SUCCESS)
-      throw_RTCError(RTC_ERROR_UNKNOWN, "cannot initialize ZE_experimental_rtas_builder extension");
+      throw_RTCError(RTC_ERROR_UNKNOWN, "cannot initialize ZE_extension_rtas extension");
 
     if (State::verbosity(1))
     {
       std::cout << "  Level Zero RTAS Builder" << std::endl;
     }
 
     /* check if extension library can get loaded */
-    ze_rtas_parallel_operation_exp_handle_t hParallelOperation;
-    result = ZeWrapper::zeRTASParallelOperationCreateExp(hDriver, &hParallelOperation);
+    ze_rtas_parallel_operation_ext_handle_t hParallelOperation;
+    result = ZeWrapper::zeRTASParallelOperationCreateExt(hDriver, &hParallelOperation);
     if (result == ZE_RESULT_ERROR_DEPENDENCY_UNAVAILABLE)
       throw_RTCError(RTC_ERROR_UNKNOWN, "Level Zero RTAS Build Extension cannot get loaded");
     if (result == ZE_RESULT_SUCCESS)
-      ZeWrapper::zeRTASParallelOperationDestroyExp(hParallelOperation);
+      ZeWrapper::zeRTASParallelOperationDestroyExt(hParallelOperation);
 
     gpu_maxWorkGroupSize = getGPUDevice().get_info<sycl::info::device::max_work_group_size>();
     gpu_maxComputeUnits  = getGPUDevice().get_info<sycl::info::device::max_compute_units>();    
@@ -707,7 +711,7 @@ namespace embree
       std::cout << std::endl;
     }
 
-    dispatchGlobalsPtr = zeRTASInitExp(gpu_device, gpu_context);
+    dispatchGlobalsPtr = zeRTASInit(gpu_device, gpu_context);
   }
 
   DeviceGPU::~DeviceGPU()

kernels/common/device.h

@@ -209,7 +209,8 @@ namespace embree
 
     unsigned int gpu_maxWorkGroupSize;
     unsigned int gpu_maxComputeUnits;
-
+    bool hasRelaxedAllocationLimits = false;
+ 
   public:
     void* dispatchGlobalsPtr = nullptr;
 
@@ -218,6 +219,8 @@ namespace embree
     inline sycl::context &getGPUContext() { return gpu_context; }    
 
     inline unsigned int getGPUMaxWorkGroupSize() { return gpu_maxWorkGroupSize; }
+    inline bool relaxedAllocationLimitsSupported() const { return hasRelaxedAllocationLimits; }
+
 
     void init_rthw_level_zero();
     void init_rthw_opencl();

kernels/common/rtcore.cpp

@@ -1405,7 +1405,7 @@ RTC_NAMESPACE_BEGIN;
     RTC_CATCH_END2(geometry);
   }
 
-  AffineSpace3fa loadTransform(RTCFormat format, const float* xfm)
+  static AffineSpace3fa loadTransform(RTCFormat format, const float* xfm)
   {
     AffineSpace3fa space = one;
     switch (format)

kernels/geometry/curve_intersector_distance.h

@@ -103,38 +103,49 @@ namespace embree
         const NativeCurve3fa curve3Di(v0,v1,v2,v3);
         const NativeCurve3fa curve3D = enlargeRadiusToMinWidth(context,geom,ray.org,curve3Di);
         const NativeCurve3fa curve2D = curve3D.xfm_pr(pre.ray_space,ray.org);
-      
-        /* evaluate the bezier curve */
-        vboolx valid = vfloatx(step) < vfloatx(float(N));
-        const Vec4vfx p0 = curve2D.template eval0<W>(0,N);
-        const Vec4vfx p1 = curve2D.template eval1<W>(0,N);
-
-        /* approximative intersection with cone */
-        const Vec4vfx v = p1-p0;
-        const Vec4vfx w = -p0;
-        const vfloatx d0 = madd(w.x,v.x,w.y*v.y);
-        const vfloatx d1 = madd(v.x,v.x,v.y*v.y);
-        const vfloatx u = clamp(d0*rcp(d1),vfloatx(zero),vfloatx(one));
-        const Vec4vfx p = madd(u,v,p0);
-        const vfloatx t = p.z*pre.depth_scale;
-        const vfloatx d2 = madd(p.x,p.x,p.y*p.y); 
-        const vfloatx r = p.w;
-        const vfloatx r2 = r*r;
-        valid &= (d2 <= r2) & (vfloatx(ray.tnear()) <= t) & (t <= vfloatx(ray.tfar));
-        if (EMBREE_CURVE_SELF_INTERSECTION_AVOIDANCE_FACTOR != 0.0f) 
-          valid &= t > float(EMBREE_CURVE_SELF_INTERSECTION_AVOIDANCE_FACTOR)*r*pre.depth_scale; // ignore self intersections
 
-        /* update hit information */
         bool ishit = false;
-        if (unlikely(any(valid))) {
-          DistanceCurveHit<NativeCurve3fa,W> hit(valid,u,0.0f,t,0,N,curve3D);
-          ishit = ishit | epilog(valid,hit);
+        int i = 0;
+
+#if !defined(__SYCL_DEVICE_ONLY__)
+        /* unrolled first iteration: on CPU SIMD this lets the compiler keep
+           the first chunk of segments in registers and skip the loop overhead.
+           On SYCL (W=1) it duplicates code without benefit and just bloats the
+           kernel, so the SYCL path falls straight into the unified loop. */
+        {
+          /* evaluate the bezier curve */
+          vboolx valid = vfloatx(step) < vfloatx(float(N));
+          const Vec4vfx p0 = curve2D.template eval0<W>(0,N);
+          const Vec4vfx p1 = curve2D.template eval1<W>(0,N);
+
+          /* approximative intersection with cone */
+          const Vec4vfx v = p1-p0;
+          const Vec4vfx w = -p0;
+          const vfloatx d0 = madd(w.x,v.x,w.y*v.y);
+          const vfloatx d1 = madd(v.x,v.x,v.y*v.y);
+          const vfloatx u = clamp(d0*rcp(d1),vfloatx(zero),vfloatx(one));
+          const Vec4vfx p = madd(u,v,p0);
+          const vfloatx t = p.z*pre.depth_scale;
+          const vfloatx d2 = madd(p.x,p.x,p.y*p.y);
+          const vfloatx r = p.w;
+          const vfloatx r2 = r*r;
+          valid &= (d2 <= r2) & (vfloatx(ray.tnear()) <= t) & (t <= vfloatx(ray.tfar));
+          if (EMBREE_CURVE_SELF_INTERSECTION_AVOIDANCE_FACTOR != 0.0f)
+            valid &= t > float(EMBREE_CURVE_SELF_INTERSECTION_AVOIDANCE_FACTOR)*r*pre.depth_scale; // ignore self intersections
+
+          /* update hit information */
+          if (unlikely(any(valid))) {
+            DistanceCurveHit<NativeCurve3fa,W> hit(valid,u,0.0f,t,0,N,curve3D);
+            ishit = ishit | epilog(valid,hit);
+          }
+          i = W;
         }
+#endif
 
-        if (unlikely(W < N)) 
+        if (unlikely(i < N))
         {
           /* process SIMD-size many segments per iteration */
-          for (int i=W; i<N; i+=W)
+          for (; i<N; i+=W)
           {
             /* evaluate the bezier curve */
             vboolx valid = vintx(i)+vintx(step) < vintx(N);

kernels/geometry/curve_intersector_sweep.h

@@ -82,10 +82,16 @@ namespace embree
         const Vec3fa dRdu = /*dQdu*/-dPdu;
         const Vec3fa dRdt = dQdt;//-dPdt;
 
-        const Vec3fa T = normalize(dPdu);
+        /* share a single rsqrt(dot(dPdu,dPdu)) between the normalize() and the
+           1/length(dPdu) needed for cos_err. Saves one transcendental per
+           Jacobian iteration (important on Xe iGPUs where sqrt/rsqrt are
+           ~14-20 cycle latency and there is no ILP to hide them). */
+        const float dPdu2 = dot(dPdu,dPdu);
+        const float rcp_len_dPdu = rsqrt(dPdu2);
+        const Vec3fa T = dPdu * rcp_len_dPdu;
         const Vec3fa dTdu = dnormalize(dPdu,ddPdu);
         //const Vec3fa dTdt = zero;
-        const float cos_err = P_err/length(dPdu);
+        const float cos_err = P_err * rcp_len_dPdu;
 
         /* Error estimate for dot(R,T):
 
@@ -348,8 +354,10 @@ namespace embree
         const Vec3ff dQ2 = abs(3.0f*(P3-P2) - W);
         const Vec3ff max_dQ = max(dQ0,dQ1,dQ2);
         const float m = max(max_dQ.x,max_dQ.y,max_dQ.z); //,max_dQ.w);
-        const float l = length(Vec3f(W));
-        const bool well_behaved = m < 0.2f*l;
+        /* compare squared values to avoid the sqrt in length(W); m is
+           non-negative since computed from abs() above, so this is exact. */
+        const float l2 = dot(Vec3f(W),Vec3f(W));
+        const bool well_behaved = m*m < (0.2f*0.2f)*l2;
 
         if (!well_behaved && stack.depth < max_depth) {
           stack.push();