Merge branch 'devel' into carsten/embree4_gpubuild_merge_devel

version.

.github/workflows/continuous.yml

@@ -201,20 +201,54 @@ jobs:
         scripts/test.py build --debug
         scripts/test.py build
 
-  linux-DG2-JIT-EMBARGO-test:
+  linux-DG2-JIT-GPU-BVH-BUILDER:
     secrets: inherit
     uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@devel
-    needs: ["linux-DG2-JIT-EMBARGO"]
+    needs: [setup_gfx, setup_dpcpp]
     with:
       image: embree/ubuntu:22.04
-      options: --device=/dev/dri:/dev/dri
-      runs-on: '[ "Linux", "docker", "dg2" ]'
+      runs-on: '[ "Linux", "docker", "build" ]'
       project: embree
       env-from-files: ./.github/workflows/dpcpp-sycl-nightly.env ./.github/workflows/gfx-ubuntu22-embargo.env
-      artifact-in: linux-DG2-JIT-EMBARGO
+      artifact-out: linux-DG2-JIT-GPU-BVH-BUILDER
+      artifact-path: ./build .ctest_conf
       cmd: |
-        scripts/test.py test --debug
-        scripts/test.py test
+        scripts/test.py configure platform:x64 compiler:dpcpp ispc:ispc1.18.1 isa:SSE2 build:RelWithAssert EMBREE_SYCL_SUPPORT:ON EMBREE_SYCL_GPU_BVH_BUILDER:ON USERGEOM:CONTEXT FILTER_FUNCTION:CONTEXT gfx:EMBARGO sycl:none tasking:TBB2021.7.0 intensity:4
+        scripts/test.py build --debug
+        scripts/test.py build
+
+  linux-DG2-JIT-PUBLIC-GPU-BVH-BUILDER:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@devel
+    needs: [setup_gfx, setup_dpcpp]
+    with:
+      image: embree/ubuntu:22.04
+      runs-on: '[ "Linux", "docker", "build" ]'
+      project: embree
+      env-from-files: ./.github/workflows/dpcpp-sycl-nightly.env ./.github/workflows/gfx-ubuntu22-public.env
+      install-gfx-driver: false      
+      artifact-out: linux-DG2-JIT-PUBLIC-GPU-BVH-BUILDER
+      artifact-path: ./build .ctest_conf
+      cmd: |
+        scripts/test.py configure platform:x64 compiler:dpcpp ispc:ispc1.18.1 isa:AVX512 build:RelWithAssert EMBREE_SYCL_SUPPORT:ON EMBREE_SYCL_GPU_BVH_BUILDER:ON USERGEOM:CONTEXT FILTER_FUNCTION:CONTEXT gfx:PUBLIC sycl:none tasking:TBB2021.7.0 intensity:4
+        scripts/test.py build --debug
+        scripts/test.py build
+
+
+ # linux-DG2-AOT-INTERNAL-test:
+ #   secrets: inherit
+ #   uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@devel
+ #   needs: ["linux-DG2-AOT-INTERNAL"]
+ #   with:
+ #     image: embree/ubuntu:22.04
+ #     options: --device=/dev/dri:/dev/dri
+ #     runs-on: '[ "Linux", "docker", "dg2" ]'
+ #     project: embree
+ #     env-from-files: ./.github/workflows/dpcpp-sycl-nightly.env ./.github/workflows/gfx-ubuntu22-internal.env
+ #     artifact-in: linux-DG2-AOT-INTERNAL
+ #     cmd: |
+ #       scripts/test.py test --debug
+ #       scripts/test.py test
 
   linux-DG2-JIT-PUBLIC-ICX:
     secrets: inherit
@@ -608,7 +642,7 @@ jobs:
       - linux-DG2-JIT-INTERNAL-test
       - linux-DG2-JIT-PUBLIC-test
       - linux-ATSM-JIT-PUBLIC-test
-      - linux-DG2-JIT-EMBARGO-test
+      - linux-DG2-JIT-EMBARGO
       - linux-DG2-JIT-PUBLIC-ICX-test
       - linux-PVC-1T-JIT-INTERNAL-test
       - linux-PVC-2T-JIT-INTERNAL-test

CMakeLists.txt

@@ -170,6 +170,12 @@ IF (EMBREE_SYCL_SUPPORT)
   ADD_DEFINITIONS("-DEMBREE_SYCL_SUPPORT")  # FIXME: only use define from rtcore_config.h
 ENDIF()
 
+OPTION(EMBREE_SYCL_GPU_BVH_BUILDER "Using fast GPU-based BVH builder" OFF)
+IF (EMBREE_SYCL_GPU_BVH_BUILDER)
+  SET(EMBREE_STATIC_LIB ON)  
+  ADD_DEFINITIONS("-DEMBREE_SYCL_GPU_BVH_BUILDER")
+ENDIF()
+
 CMAKE_DEPENDENT_OPTION(EMBREE_SYCL_RT_SIMULATION "Using hardware simulation" OFF "EMBREE_SYCL_SUPPORT" OFF)
 MARK_AS_ADVANCED(EMBREE_SYCL_RT_SIMULATION)
 IF (EMBREE_SYCL_RT_SIMULATION AND NOT EMBREE_SYCL_RT_VALIDATION_API AND EMBREE_SYCL_IMPLICIT_DISPATCH_GLOBALS)
@@ -180,6 +186,7 @@ IF (EMBREE_SYCL_RT_SIMULATION)
 ENDIF()
 
 CMAKE_DEPENDENT_OPTION(EMBREE_SYCL_RT_VALIDATION_API "Use rt_validation API instead of IGC provided rt_production API" OFF "EMBREE_SYCL_SUPPORT" OFF)
+
 IF (EMBREE_SYCL_RT_VALIDATION_API)
   ADD_DEFINITIONS("-DEMBREE_SYCL_RT_VALIDATION_API")
 ENDIF()

common/cmake/clang.cmake

@@ -109,7 +109,7 @@ ELSE()
     SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fexceptions")              # enable exceptions
     SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -pthread")                  # enable threads
     SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -msimd128")                 # enable SIMD intrinsics
-  ENDIF()
+  ENDIF()  
 
   SET(CMAKE_CXX_FLAGS_DEBUG "")
   SET(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} -g")              # generate debug information

common/cmake/embree-config.cmake

@@ -52,6 +52,8 @@ SET(EMBREE_CURVE_SELF_INTERSECTION_AVOIDANCE_FACTOR @EMBREE_CURVE_SELF_INTERSECT
 SET(EMBREE_DISC_POINT_SELF_INTERSECTION_AVOIDANCE @EMBREE_DISC_POINT_SELF_INTERSECTION_AVOIDANCE@)
 SET(EMBREE_MIN_WIDTH @EMBREE_MIN_WIDTH@)
 
+SET(EMBREE_SYCL_GPU_BVH_BUILDER @EMBREE_SYCL_GPU_BVH_BUILDER@)
+
 IF (EMBREE_STATIC_LIB AND (EMBREE_TASKING_SYSTEM STREQUAL "TBB"))
   INCLUDE(CMakeFindDependencyMacro)
   FIND_DEPENDENCY(TBB)

common/sys/platform.h

@@ -212,11 +212,15 @@
 #define PRINT2(x,y) embree_cout << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << embree_endl
 #define PRINT3(x,y,z) embree_cout << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << embree_endl
 #define PRINT4(x,y,z,w) embree_cout << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << ", " << STRING(w) << " = " << (w) << embree_endl
+#define PRINT5(x,y,z,w,a) embree_cout << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << ", " << STRING(w) << " = " << (w) << ", " << STRING(a) << " = " << (a) << embree_endl
+#define PRINT6(x,y,z,w,a,b) embree_cout << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << ", " << STRING(w) << " = " << (w) << ", " << STRING(a) << " = " << (a) << ", " << STRING(b) << " = " << (b) << embree_endl
 
 #define UPRINT(x) embree_cout_uniform << STRING(x) << " = " << (x) << embree_endl
 #define UPRINT2(x,y) embree_cout_uniform << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << embree_endl
 #define UPRINT3(x,y,z) embree_cout_uniform << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << embree_endl
 #define UPRINT4(x,y,z,w) embree_cout_uniform << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << ", " << STRING(w) << " = " << (w) << embree_endl
+#define UPRINT5(x,y,z,w,a) embree_cout_uniform << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << ", " << STRING(w) << " = " << (w) << ", " << STRING(a) << " = " << (a) << embree_endl
+#define UPRINT6(x,y,z,w,a,b) embree_cout_uniform << STRING(x) << " = " << (x) << ", " << STRING(y) << " = " << (y) << ", " << STRING(z) << " = " << (z) << ", " << STRING(w) << " = " << (w) << ", " << STRING(a) << " = " << (a) << ", " << STRING(b) << " = " << (b) << embree_endl
 
 #if defined(DEBUG) // only report file and line in debug mode
   #define THROW_RUNTIME_ERROR(str) \
@@ -361,9 +365,15 @@ __forceinline std::string toString(long long value) {
 /// SYCL specific
 ////////////////////////////////////////////////////////////////////////////////
 
+#if defined(__SYCL_DEVICE_ONLY__)
+#  define EMBREE_SYCL_SIMD(N) [[intel::reqd_sub_group_size(N)]]  
+#else
+#  define EMBREE_SYCL_SIMD(N)
+#endif
 
 #if defined(EMBREE_SYCL_SUPPORT) && defined(__SYCL_DEVICE_ONLY__)
 
+  
 #define sycl_printf0(format, ...) {               \
     static const CONSTANT char fmt[] = format;               \
     if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))       \

kernels/builders/primrefgen_presplit.h

@@ -38,6 +38,7 @@ namespace embree
 
       __forceinline bool split_pos(const PrimRef& prim, unsigned int& dim_o, float& fsplit_o) const
       {
+#if !defined(__SYCL_DEVICE_ONLY__)        
         /* compute morton code */
         const Vec3fa lower = prim.lower;
         const Vec3fa upper = prim.upper;
@@ -72,6 +73,7 @@ namespace embree
 
         dim_o = dim;
         fsplit_o = fsplit;
+#endif        
         return true;
       }
 
@@ -83,10 +85,11 @@ namespace embree
         const Vec3fa gupper = (upper-base)*Vec3fa(scale)-Vec3fa(0.2f);
         Vec3ia ilower(floor(glower));
         Vec3ia iupper(floor(gupper));
-        
+
+#if !defined(__SYCL_DEVICE_ONLY__)                
         /* this ignores dimensions that are empty */
         iupper = (Vec3ia)select(vint4(glower) >= vint4(gupper),vint4(ilower),vint4(iupper));
-        
+#endif        
         /* compute a morton code for the lower and upper grid coordinates. */
         const unsigned int lower_code = bitInterleave(ilower.x,ilower.y,ilower.z);
         const unsigned int upper_code = bitInterleave(iupper.x,iupper.y,iupper.z);

kernels/common/scene.cpp

@@ -685,7 +685,10 @@ namespace embree
   {
     /* select acceleration structures to build */
     unsigned int new_enabled_geometry_types = world.enabledGeometryTypesMask();
-
+// #if defined(EMBREE_SYCL_SUPPORT)
+//     if (!dynamic_cast<DeviceGPU*>(device)) // do not build software accel for GPU if not required
+// #endif
+      
     if (flags_modified || new_enabled_geometry_types != enabled_geometry_types)
     {
       accels_init();
@@ -723,6 +726,18 @@ namespace embree
     /* build all hierarchies of this scene */
     accels_build();
 
+    /* build acceleration structure for rthw */
+// #if defined(EMBREE_SYCL_SUPPORT)
+//     if (DeviceGPU* gpu_device = dynamic_cast<DeviceGPU*>(device))
+//       //if (gpu_device->rthw_support())
+//       //  bounds = LBBox<embree::Vec3fa>(rthwifBuild(this,quality_flags,hwaccel,device->gpu_build));
+//       if (gpu_device->rthw_support()) {
+//         const BBox3f aabb = rthwifBuild(this,quality_flags,hwaccel,device->gpu_build);
+//         bounds = LBBox<embree::Vec3fa>(aabb);
+//         hwaccel_bounds = aabb;
+//       }
+// #endif
+    
     /* make static geometry immutable */
     if (!isDynamicAccel()) {
       accels_immutable();

kernels/common/stat.cpp

@@ -3,6 +3,7 @@
 
 #include "stat.h"
 
+
 namespace embree
 {
   Stat Stat::instance; 
@@ -24,16 +25,16 @@ namespace embree
     //Counters::Data& data = instance.cntrs.active;
 
     /* print absolute numbers */
-    cout << "--------- ABSOLUTE ---------" << std::endl;
-    cout << "  #normal_travs   = " << float(data.normal.travs            )*1E-6 << "M" << std::endl;
-    cout << "    #nodes        = " << float(data.normal.trav_nodes       )*1E-6 << "M" << std::endl;
-    cout << "    #nodes_xfm    = " << float(data.normal.trav_xfm_nodes   )*1E-6 << "M" << std::endl;
-    cout << "    #leaves       = " << float(data.normal.trav_leaves      )*1E-6 << "M" << std::endl;
-    cout << "    #prims        = " << float(data.normal.trav_prims       )*1E-6 << "M" << std::endl;
-    cout << "    #prim_hits    = " << float(data.normal.trav_prim_hits   )*1E-6 << "M" << std::endl;
+    cout << "--------- ABSOLUTE ---------" << embree_endl;
+    cout << "  #normal_travs   = " << float(data.normal.travs            )*1E-6f << "M" << embree_endl;
+    cout << "    #nodes        = " << float(data.normal.trav_nodes       )*1E-6f << "M" << embree_endl;
+    cout << "    #nodes_xfm    = " << float(data.normal.trav_xfm_nodes   )*1E-6f << "M" << embree_endl;
+    cout << "    #leaves       = " << float(data.normal.trav_leaves      )*1E-6f << "M" << embree_endl;
+    cout << "    #prims        = " << float(data.normal.trav_prims       )*1E-6f << "M" << embree_endl;
+    cout << "    #prim_hits    = " << float(data.normal.trav_prim_hits   )*1E-6f << "M" << embree_endl;
 
-    cout << "    #stack nodes  = " << float(data.normal.trav_stack_nodes )*1E-6 << "M" << std::endl;
-    cout << "    #stack pop    = " << float(data.normal.trav_stack_pop )*1E-6 << "M" << std::endl;
+    cout << "    #stack nodes  = " << float(data.normal.trav_stack_nodes )*1E-6f << "M" << embree_endl;
+    cout << "    #stack pop    = " << float(data.normal.trav_stack_pop )*1E-6f << "M" << embree_endl;
 
     size_t normal_box_hits = 0;
     size_t weighted_box_hits = 0;
@@ -49,20 +50,20 @@ namespace embree
       cout << "[" << i << "] " << value << " ";
       average += (float)i*data.normal.trav_hit_boxes[i] / normal_box_hits;
     }
-    cout << "    average = " << average << std::endl;
+    cout << "    average = " << average << embree_endl;
     for (size_t i=0;i<SIZE_HISTOGRAM;i++) cout << "[" << i << "] " << 100.0f * data.normal.trav_hit_boxes[i]*i / weighted_box_hits << " ";
-    cout << std::endl;
+    cout << embree_endl;
 
     if (data.shadow.travs) {
-      cout << "  #shadow_travs = " << float(data.shadow.travs         )*1E-6 << "M" << std::endl;
-      cout << "    #nodes      = " << float(data.shadow.trav_nodes    )*1E-6 << "M" << std::endl;
-      cout << "    #nodes_xfm  = " << float(data.shadow.trav_xfm_nodes)*1E-6 << "M" << std::endl;
-      cout << "    #leaves     = " << float(data.shadow.trav_leaves   )*1E-6 << "M" << std::endl;
-      cout << "    #prims      = " << float(data.shadow.trav_prims    )*1E-6 << "M" << std::endl;
-      cout << "    #prim_hits  = " << float(data.shadow.trav_prim_hits)*1E-6 << "M" << std::endl;
+      cout << "  #shadow_travs = " << float(data.shadow.travs         )*1E-6f << "M" << embree_endl;
+      cout << "    #nodes      = " << float(data.shadow.trav_nodes    )*1E-6f << "M" << embree_endl;
+      cout << "    #nodes_xfm  = " << float(data.shadow.trav_xfm_nodes)*1E-6f << "M" << embree_endl;
+      cout << "    #leaves     = " << float(data.shadow.trav_leaves   )*1E-6f << "M" << embree_endl;
+      cout << "    #prims      = " << float(data.shadow.trav_prims    )*1E-6f << "M" << embree_endl;
+      cout << "    #prim_hits  = " << float(data.shadow.trav_prim_hits)*1E-6f << "M" << embree_endl;
 
-      cout << "    #stack nodes = " << float(data.shadow.trav_stack_nodes )*1E-6 << "M" << std::endl;
-      cout << "    #stack pop   = " << float(data.shadow.trav_stack_pop )*1E-6 << "M" << std::endl;
+      cout << "    #stack nodes = " << float(data.shadow.trav_stack_nodes )*1E-6f << "M" << embree_endl;
+      cout << "    #stack pop   = " << float(data.shadow.trav_stack_pop )*1E-6f << "M" << embree_endl;
 
       size_t shadow_box_hits = 0;
       size_t weighted_shadow_box_hits = 0;
@@ -73,14 +74,14 @@ namespace embree
       }
       cout << "    #hit_boxes    = ";
       for (size_t i=0;i<SIZE_HISTOGRAM;i++) cout << "[" << i << "] " << 100.0f * data.shadow.trav_hit_boxes[i] / shadow_box_hits << " ";
-      cout << std::endl;
+      cout << embree_endl;
       for (size_t i=0;i<SIZE_HISTOGRAM;i++) cout << "[" << i << "] " << 100.0f * data.shadow.trav_hit_boxes[i]*i / weighted_shadow_box_hits << " ";
-      cout << std::endl;
+      cout << embree_endl;
     }
-    cout << std::endl;
+    cout << embree_endl;
 
     /* print per traversal numbers */
-    cout << "--------- PER TRAVERSAL ---------" << std::endl;
+    cout << "--------- PER TRAVERSAL ---------" << embree_endl;
     float active_normal_travs       = float(cntrs.active.normal.travs      )/float(cntrs.all.normal.travs      );
     float active_normal_trav_nodes  = float(cntrs.active.normal.trav_nodes )/float(cntrs.all.normal.trav_nodes );
     float active_normal_trav_xfm_nodes  = float(cntrs.active.normal.trav_xfm_nodes )/float(cntrs.all.normal.trav_xfm_nodes );
@@ -89,13 +90,13 @@ namespace embree
     float active_normal_trav_prim_hits = float(cntrs.active.normal.trav_prim_hits  )/float(cntrs.all.normal.trav_prim_hits  );
     float active_normal_trav_stack_pop = float(cntrs.active.normal.trav_stack_pop  )/float(cntrs.all.normal.trav_stack_pop  );
 
-    cout << "  #normal_travs   = " << float(cntrs.code.normal.travs      )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_travs       << "% active" << std::endl;
-    cout << "    #nodes        = " << float(cntrs.code.normal.trav_nodes )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_nodes  << "% active" << std::endl;
-    cout << "    #node_xfm     = " << float(cntrs.code.normal.trav_xfm_nodes )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_xfm_nodes  << "% active" << std::endl;
-    cout << "    #leaves       = " << float(cntrs.code.normal.trav_leaves)/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_leaves << "% active" << std::endl;
-    cout << "    #prims        = " << float(cntrs.code.normal.trav_prims  )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_prims   << "% active" << std::endl;
-    cout << "    #prim_hits    = " << float(cntrs.code.normal.trav_prim_hits  )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_prim_hits   << "% active" << std::endl;
-    cout << "    #stack_pop    = " << float(cntrs.code.normal.trav_stack_pop  )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_stack_pop   << "% active" << std::endl;
+    cout << "  #normal_travs   = " << float(cntrs.code.normal.travs      )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_travs       << "% active" << embree_endl;
+    cout << "    #nodes        = " << float(cntrs.code.normal.trav_nodes )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_nodes  << "% active" << embree_endl;
+    cout << "    #node_xfm     = " << float(cntrs.code.normal.trav_xfm_nodes )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_xfm_nodes  << "% active" << embree_endl;
+    cout << "    #leaves       = " << float(cntrs.code.normal.trav_leaves)/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_leaves << "% active" << embree_endl;
+    cout << "    #prims        = " << float(cntrs.code.normal.trav_prims  )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_prims   << "% active" << embree_endl;
+    cout << "    #prim_hits    = " << float(cntrs.code.normal.trav_prim_hits  )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_prim_hits   << "% active" << embree_endl;
+    cout << "    #stack_pop    = " << float(cntrs.code.normal.trav_stack_pop  )/float(cntrs.code.normal.travs) << ", " << 100.0f*active_normal_trav_stack_pop   << "% active" << embree_endl;
 
     if (cntrs.all.shadow.travs) {
       float active_shadow_travs       = float(cntrs.active.shadow.travs      )/float(cntrs.all.shadow.travs      );
@@ -105,24 +106,24 @@ namespace embree
       float active_shadow_trav_prims   = float(cntrs.active.shadow.trav_prims  )/float(cntrs.all.shadow.trav_prims  );
       float active_shadow_trav_prim_hits = float(cntrs.active.shadow.trav_prim_hits  )/float(cntrs.all.shadow.trav_prim_hits  );
 
-      cout << "  #shadow_travs = " << float(cntrs.code.shadow.travs      )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_travs       << "% active" << std::endl;
-      cout << "    #nodes      = " << float(cntrs.code.shadow.trav_nodes )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_nodes  << "% active" << std::endl;
-      cout << "    #nodes_xfm  = " << float(cntrs.code.shadow.trav_xfm_nodes )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_xfm_nodes  << "% active" << std::endl;
-      cout << "    #leaves     = " << float(cntrs.code.shadow.trav_leaves)/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_leaves << "% active" << std::endl;
-      cout << "    #prims      = " << float(cntrs.code.shadow.trav_prims  )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_prims   << "% active" << std::endl;
-      cout << "    #prim_hits  = " << float(cntrs.code.shadow.trav_prim_hits  )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_prim_hits   << "% active" << std::endl;
+      cout << "  #shadow_travs = " << float(cntrs.code.shadow.travs      )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_travs       << "% active" << embree_endl;
+      cout << "    #nodes      = " << float(cntrs.code.shadow.trav_nodes )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_nodes  << "% active" << embree_endl;
+      cout << "    #nodes_xfm  = " << float(cntrs.code.shadow.trav_xfm_nodes )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_xfm_nodes  << "% active" << embree_endl;
+      cout << "    #leaves     = " << float(cntrs.code.shadow.trav_leaves)/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_leaves << "% active" << embree_endl;
+      cout << "    #prims      = " << float(cntrs.code.shadow.trav_prims  )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_prims   << "% active" << embree_endl;
+      cout << "    #prim_hits  = " << float(cntrs.code.shadow.trav_prim_hits  )/float(cntrs.code.shadow.travs) << ", " << 100.0f*active_shadow_trav_prim_hits   << "% active" << embree_endl;
 
     }
-    cout << std::endl;
+    cout << embree_endl;
 
      /* print user counters for performance tuning */
-    cout << "--------- USER ---------" << std::endl;
+    cout << "--------- USER ---------" << embree_endl;
     for (size_t i=0; i<10; i++)
-      cout << "#user" << i << " = " << float(cntrs.user[i])/float(cntrs.all.normal.travs+cntrs.all.shadow.travs) << " per traversal" << std::endl;
+      cout << "#user" << i << " = " << float(cntrs.user[i])/float(cntrs.all.normal.travs+cntrs.all.shadow.travs) << " per traversal" << embree_endl;
 
-    cout << "#user5/user3 " << 100.0f*float(cntrs.user[5])/float(cntrs.user[3]) << "%" << std::endl;
-    cout << "#user6/user3 " << 100.0f*float(cntrs.user[6])/float(cntrs.user[3]) << "%" << std::endl;
-    cout << "#user7/user3 " << 100.0f*float(cntrs.user[7])/float(cntrs.user[3]) << "%" << std::endl;
-    cout << std::endl;
+    cout << "#user5/user3 " << 100.0f*float(cntrs.user[5])/float(cntrs.user[3]) << "%" << embree_endl;
+    cout << "#user6/user3 " << 100.0f*float(cntrs.user[6])/float(cntrs.user[3]) << "%" << embree_endl;
+    cout << "#user7/user3 " << 100.0f*float(cntrs.user[7])/float(cntrs.user[3]) << "%" << embree_endl;
+    cout << embree_endl;
   }
 }

kernels/common/state.cpp

@@ -81,6 +81,7 @@ namespace embree
     object_accel_mb_min_leaf_size = 1;
     object_accel_mb_max_leaf_size = 1;
 
+    gpu_build = 0;    
     max_spatial_split_replications = 1.2f;
     useSpatialPreSplits = false;
 
@@ -420,6 +421,15 @@ namespace embree
         }
       }
 
+      else if (tok == Token::Id("gpu_build"))
+      {
+        if (cin->trySymbol("=")) {
+          Token flag = cin->get();
+          if      (flag == Token::Id("ploc"))          gpu_build = 1;
+          else if (flag == Token::Id("ploc_twolevel")) gpu_build = 2;                              
+        }        
+      }      
+
       else if (tok == Token::Id("max_spatial_split_replications") && cin->trySymbol("="))
         max_spatial_split_replications = cin->get().Float();
 

kernels/common/state.h

@@ -105,7 +105,12 @@ namespace embree
     std::string grid_accel_mb;           //!< acceleration structure to use for motion blur grids
     std::string grid_builder_mb;         //!< builder for motion blur grids
 
+    //public:
+    //std::string rthw_builder = "internal"; //!< ray tracing hardware builder to use
+
   public:
+    int gpu_build;                         //!< enabled if BVH build should happen on GPU    
+
     float max_spatial_split_replications;  //!< maximally replications*N many primitives in accel for spatial splits
     bool useSpatialPreSplits;              //!< use spatial pre-splits instead of the full spatial split builder
     size_t tessellation_cache_size;        //!< size of the shared tessellation cache 

kernels/rthwif/CMakeLists.txt

@@ -157,7 +157,7 @@ IF (EMBREE_SYCL_RT_VALIDATION_API)
   ADD_LIBRARY(embree_rthwif_sycl STATIC rttrace/rthwif_production.cpp)
   SET_PROPERTY(TARGET embree_rthwif_sycl APPEND PROPERTY COMPILE_FLAGS "-DEMBREE_SYCL_SUPPORT")
   SET_TARGET_PROPERTIES(embree_rthwif_sycl PROPERTIES COMPILE_FLAGS ${CMAKE_CXX_FLAGS_SYCL})
-
+  
   INSTALL(TARGETS embree_rthwif_sycl EXPORT embree_rthwif_sycl-targets ARCHIVE DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
   INSTALL(EXPORT embree_rthwif_sycl-targets DESTINATION "${EMBREE_CMAKEEXPORT_DIR}" COMPONENT devel)
 
@@ -168,15 +168,31 @@ ADD_LIBRARY(embree_rthwif ${RTHWIF_LIB_TYPE} rtbuild/rthwif_builder.cpp rtbuild/
 TARGET_LINK_LIBRARIES(embree_rthwif PUBLIC ${EMBREE_RTHWIF_SYCL} PRIVATE ${TBB_TARGET} simd sys)
 SET_TARGET_PROPERTIES(embree_rthwif PROPERTIES OUTPUT_NAME ${RTHWIF_NAME})
 IF (EMBREE_RTHWIF_STATIC_LIB)
-  TARGET_COMPILE_DEFINITIONS(embree_rthwif PUBLIC EMBREE_RTHWIF_STATIC_LIB)
+  TARGET_COMPILE_DEFINITIONS(embree_rthwif PUBLIC EMBREE_RTHWIF_STATIC_LIB)  
 ENDIF()
 
+IF (EMBREE_SYCL_GPU_BVH_BUILDER)
+ADD_LIBRARY(embree_rthwif_gpu STATIC rtbuild/rthwif_embree_builder_ploc.cpp rtbuild/qbvh6.cpp rtbuild/statistics.cpp)
+TARGET_LINK_LIBRARIES(embree_rthwif_gpu PUBLIC tbb simd sys -lze_loader)
+SET_TARGET_PROPERTIES(embree_rthwif_gpu PROPERTIES COMPILE_FLAGS ${CMAKE_CXX_FLAGS_SYCL})
+ENDIF()
+
+TARGET_LINK_LIBRARIES(embree_rthwif PUBLIC ${EMBREE_RTHWIF_SYCL} PRIVATE ${TBB_TARGET} simd sys)
+
+
 if (EMBREE_LEVEL_ZERO)
   TARGET_LINK_LIBRARIES(embree_rthwif PUBLIC -lze_loader)
   TARGET_COMPILE_DEFINITIONS(embree_rthwif PUBLIC EMBREE_LEVEL_ZERO)
   TARGET_COMPILE_DEFINITIONS(embree_rthwif PUBLIC EMBREE_SYCL_SUPPORT)
 endif()
 
+IF (EMBREE_SYCL_GPU_BVH_BUILDER)
+  TARGET_LINK_LIBRARIES(embree_rthwif PUBLIC embree_rthwif_gpu)
+ENDIF()
+
+SET_PROPERTY(TARGET embree_rthwif APPEND PROPERTY COMPILE_FLAGS "-DEMBREE_SYCL_SUPPORT")
+
+
 IF (EMBREE_STATIC_LIB OR NOT EMBREE_RTHWIF_STATIC_LIB)
   INSTALL(TARGETS embree_rthwif EXPORT ${RTHWIF_NAME}-targets
     LIBRARY NAMELINK_SKIP DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib
@@ -185,6 +201,11 @@ IF (EMBREE_STATIC_LIB OR NOT EMBREE_RTHWIF_STATIC_LIB)
   INSTALL(EXPORT ${RTHWIF_NAME}-targets DESTINATION "${EMBREE_CMAKEEXPORT_DIR}" COMPONENT devel)
 ENDIF()
 
+IF (EMBREE_SYCL_GPU_BVH_BUILDER)
+INSTALL(TARGETS embree_rthwif_gpu EXPORT embree_rthwif_gpu-targets ARCHIVE DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
+INSTALL(EXPORT embree_rthwif_gpu-targets DESTINATION "${EMBREE_CMAKEEXPORT_DIR}" COMPONENT devel)
+ENDIF()
+
 IF (EMBREE_LEVEL_ZERO OR RTHWIF_STANDALONE)
   ADD_SUBDIRECTORY(testing)
 ENDIF()

kernels/rthwif/rtbuild/gpu/AABB3f.h

@@ -0,0 +1,507 @@
+// Copyright 2009-2021 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+
+#pragma once
+
+#if defined(EMBREE_SYCL_SUPPORT)
+
+#include "common.h"
+
+namespace embree
+{
+  namespace gpu
+  {    
+    class AABB3f {
+    public:
+      float lower_x,lower_y,lower_z; 
+      float upper_x,upper_y,upper_z;
+      
+      __forceinline AABB3f() = default;
+
+      __forceinline AABB3f(const AABB3f &aabb) : lower_x(aabb.lower_x),
+                                                 lower_y(aabb.lower_y),
+                                                 lower_z(aabb.lower_z),
+                                                 upper_x(aabb.upper_x),
+                                                 upper_y(aabb.upper_y),
+                                                 upper_z(aabb.upper_z)
+      {}
+
+      __forceinline AABB3f(const float3 &v) : lower_x(v.x()),
+                                              lower_y(v.y()),
+                                              lower_z(v.z()),
+                                              upper_x(v.x()),
+                                              upper_y(v.y()),
+                                              upper_z(v.z())                                              
+      {}
+
+      __forceinline AABB3f(const float3 &lower, const float3 &upper) : lower_x(lower.x()),
+                                                                       lower_y(lower.y()),
+                                                                       lower_z(lower.z()),
+                                                                       upper_x(upper.x()),
+                                                                       upper_y(upper.y()),
+                                                                       upper_z(upper.z())
+                                                                       
+      {}
+
+      __forceinline AABB3f(const BBox3fa &aabb) : lower_x(aabb.lower.x),
+                                                  lower_y(aabb.lower.y),
+                                                  lower_z(aabb.lower.z),
+                                                  upper_x(aabb.upper.x),
+                                                  upper_y(aabb.upper.y),
+                                                  upper_z(aabb.upper.z)
+      {}
+
+      __forceinline AABB3f(const float lower_x, const float lower_y, const float lower_z, const float upper_x, const float upper_y, const float upper_z) : lower_x(lower_x),
+                                                                                                                                                           lower_y(lower_y),
+                                                                                                                                                           lower_z(lower_z),
+                                                                                                                                                           upper_x(upper_x),
+                                                                                                                                                           upper_y(upper_y),
+                                                                                                                                                           upper_z(upper_z)
+      {}
+      
+      __forceinline bool empty()
+      {
+        if (lower_x > upper_x ||
+            lower_y > upper_y ||
+            lower_z > upper_z) return true;
+        return false;
+      }
+      
+      __forceinline void init()
+      {
+        const float _pos_inf =  INFINITY;
+        const float _neg_inf = -INFINITY;
+        lower_x = _pos_inf;
+        lower_y = _pos_inf;
+        lower_z = _pos_inf;        
+        upper_x = _neg_inf;
+        upper_y = _neg_inf;
+        upper_z = _neg_inf;        
+      }
+
+      __forceinline void setZero()
+      {
+        lower_x = 0.0f;
+        lower_y = 0.0f;
+        lower_z = 0.0f;
+        upper_x = 0.0f;
+        upper_y = 0.0f;
+        upper_z = 0.0f;
+      }
+      
+      __forceinline bool encloses(const AABB3f &other)
+      {
+        if ( other.lower_x < lower_x ||
+             other.lower_y < lower_y ||
+             other.lower_z < lower_z )
+          return false;
+        if ( other.upper_x > upper_x ||
+             other.upper_y > upper_y ||
+             other.upper_z > upper_z )
+          return false;
+        return true;
+      }      
+      
+      __forceinline void extend(const class AABB3f &aabb)
+      {
+        lower_x = min(lower_x,aabb.lower_x);
+        lower_y = min(lower_y,aabb.lower_y);
+        lower_z = min(lower_z,aabb.lower_z);        
+        upper_x = max(upper_x,aabb.upper_x);
+        upper_y = max(upper_y,aabb.upper_y);
+        upper_z = max(upper_z,aabb.upper_z);        
+      }
+
+      
+      __forceinline void extend(const class AABB3f &aabb, const uint slot)
+      {
+        const uint subgroupLocalID = get_sub_group_local_id();
+        const bool cmp = subgroupLocalID == slot;        
+        lower_x = cselect(cmp,min(lower_x,aabb.lower_x),lower_x);
+        lower_y = cselect(cmp,min(lower_y,aabb.lower_y),lower_y);
+        lower_z = cselect(cmp,min(lower_z,aabb.lower_z),lower_z);        
+        upper_x = cselect(cmp,max(upper_x,aabb.upper_x),upper_x);
+        upper_y = cselect(cmp,max(upper_y,aabb.upper_y),upper_y);
+        upper_z = cselect(cmp,max(upper_z,aabb.upper_z),upper_z);        
+      }
+      
+
+      __forceinline void extend(const float3 &v)
+      {
+        lower_x = min(lower_x,v.x());
+        lower_y = min(lower_y,v.y());
+        lower_z = min(lower_z,v.z());        
+        upper_x = max(upper_x,v.x());
+        upper_y = max(upper_y,v.y());
+        upper_z = max(upper_z,v.z());        
+      }
+
+      __forceinline void extend(const float4 &v)
+      {
+        lower_x = min(lower_x,v.x());
+        lower_y = min(lower_y,v.y());
+        lower_z = min(lower_z,v.z());        
+        upper_x = max(upper_x,v.x());
+        upper_y = max(upper_y,v.y());
+        upper_z = max(upper_z,v.z());        
+      }
+      
+       __forceinline void extend(const BBox3fa &aabb)
+       {
+        lower_x = min(lower_x,aabb.lower.x);
+        lower_y = min(lower_y,aabb.lower.y);
+        lower_z = min(lower_z,aabb.lower.z);        
+        upper_x = max(upper_x,aabb.upper.x);
+        upper_y = max(upper_y,aabb.upper.y);
+        upper_z = max(upper_z,aabb.upper.z);        
+       }
+
+      __forceinline void enlarge(const float &v)
+      {
+        lower_x -= v;
+        lower_y -= v;
+        lower_z -= v;
+        upper_x += v;
+        upper_y += v;
+        upper_z += v;        
+      }      
+
+      __forceinline void add(const class AABB3f &aabb)
+      {
+        lower_x += aabb.lower_x;
+        lower_y += aabb.lower_y;
+        lower_z += aabb.lower_z;        
+        upper_x += aabb.upper_x;
+        upper_y += aabb.upper_y;
+        upper_z += aabb.upper_z;        
+      }
+      
+      __forceinline void mul(const float f)
+      {
+        lower_x *= f;
+        lower_y *= f;
+        lower_z *= f;        
+        upper_x *= f;
+        upper_y *= f;
+        upper_z *= f;                
+      }
+      
+      __forceinline float3 lower() const { return float3(lower_x,lower_y,lower_z); }
+      __forceinline float3 upper() const { return float3(upper_x,upper_y,upper_z); }
+      
+      __forceinline float3 size() const
+      {
+        return upper() - lower();
+      }
+
+      __forceinline float3 diag() const
+      {
+        return upper() - lower();
+      }
+      
+      __forceinline float3 centroid2() const
+      {
+        return upper() + lower();
+      }
+
+      __forceinline float3 centroid() const
+      {
+        return (upper() + lower())*0.5f;
+      }
+
+      __forceinline float length2() const
+      {
+        const float3 diag = size();
+        return  sycl::fma(diag.x(),diag.x(),
+                          sycl::fma(diag.y(),diag.y(),
+                                    diag.z() * diag.z()));
+        
+      }
+      
+      __forceinline float maxDiagDim() const
+      {
+        const float3 d = diag();
+        return max(max(d.x(),d.y()),d.z());        
+      }      
+      
+      __forceinline void atomic_merge_global(AABB3f &dest) const
+      {
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> _lower_x(dest.lower_x);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> _lower_y(dest.lower_y);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> _lower_z(dest.lower_z);
+
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> _upper_x(dest.upper_x);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> _upper_y(dest.upper_y);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> _upper_z(dest.upper_z);
+        
+        _lower_x.fetch_min(lower_x);
+        _lower_y.fetch_min(lower_y);
+        _lower_z.fetch_min(lower_z);
+
+        _upper_x.fetch_max(upper_x);
+        _upper_y.fetch_max(upper_y);
+        _upper_z.fetch_max(upper_z);        
+      }
+
+      __forceinline void atomic_merge_local(AABB3f &dest) const
+      {
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> _lower_x(dest.lower_x);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> _lower_y(dest.lower_y);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> _lower_z(dest.lower_z);
+
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> _upper_x(dest.upper_x);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> _upper_y(dest.upper_y);
+        sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> _upper_z(dest.upper_z);
+
+        _lower_x.fetch_min(lower_x);
+        _lower_y.fetch_min(lower_y);
+        _lower_z.fetch_min(lower_z);
+
+        _upper_x.fetch_max(upper_x);
+        _upper_y.fetch_max(upper_y);
+        _upper_z.fetch_max(upper_z);        
+      }
+      
+      __forceinline AABB3f sub_group_reduce() const
+      {
+        AABB3f result;
+        result.lower_x = embree::sub_group_reduce(lower_x, SYCL_EXT_ONEAPI::minimum<float>());
+        result.lower_y = embree::sub_group_reduce(lower_y, SYCL_EXT_ONEAPI::minimum<float>());
+        result.lower_z = embree::sub_group_reduce(lower_z, SYCL_EXT_ONEAPI::minimum<float>());
+        result.upper_x = embree::sub_group_reduce(upper_x, SYCL_EXT_ONEAPI::maximum<float>());
+        result.upper_y = embree::sub_group_reduce(upper_y, SYCL_EXT_ONEAPI::maximum<float>());
+        result.upper_z = embree::sub_group_reduce(upper_z, SYCL_EXT_ONEAPI::maximum<float>());        
+        return result;	
+      }
+
+      __forceinline AABB3f sub_group_broadcast(const sycl::id<1> &localID) const
+      {
+        AABB3f result;
+        result.lower_x = embree::sub_group_broadcast<float>(lower_x,localID);
+        result.lower_y = embree::sub_group_broadcast<float>(lower_y,localID);
+        result.lower_z = embree::sub_group_broadcast<float>(lower_z,localID);
+        result.upper_x = embree::sub_group_broadcast<float>(upper_x,localID);
+        result.upper_y = embree::sub_group_broadcast<float>(upper_y,localID);
+        result.upper_z = embree::sub_group_broadcast<float>(upper_z,localID);
+        return result;	
+      }
+
+      __forceinline AABB3f sub_group_shuffle(const sycl::id<1> &localID) const
+      {
+        AABB3f result;
+        result.lower_x = embree::sub_group_shuffle<float>(lower_x,localID);
+        result.lower_y = embree::sub_group_shuffle<float>(lower_y,localID);
+        result.lower_z = embree::sub_group_shuffle<float>(lower_z,localID);
+        result.upper_x = embree::sub_group_shuffle<float>(upper_x,localID);
+        result.upper_y = embree::sub_group_shuffle<float>(upper_y,localID);
+        result.upper_z = embree::sub_group_shuffle<float>(upper_z,localID);
+        return result;	
+      }
+      
+
+      __forceinline AABB3f sub_group_scan_exclusive_min_max() const
+      {
+        AABB3f result;
+        result.lower_x = embree::sub_group_exclusive_scan(lower_x, SYCL_EXT_ONEAPI::minimum<float>());
+        result.lower_y = embree::sub_group_exclusive_scan(lower_y, SYCL_EXT_ONEAPI::minimum<float>());
+        result.lower_z = embree::sub_group_exclusive_scan(lower_z, SYCL_EXT_ONEAPI::minimum<float>());
+        result.upper_x = embree::sub_group_exclusive_scan(upper_x, SYCL_EXT_ONEAPI::maximum<float>());
+        result.upper_y = embree::sub_group_exclusive_scan(upper_y, SYCL_EXT_ONEAPI::maximum<float>());
+        result.upper_z = embree::sub_group_exclusive_scan(upper_z, SYCL_EXT_ONEAPI::maximum<float>());
+        return result;	
+      }
+
+      __forceinline AABB3f sub_group_scan_inclusive_min_max() const
+	
+      {
+        AABB3f result;
+        result.lower_x = embree::sub_group_inclusive_scan(lower_x, SYCL_EXT_ONEAPI::minimum<float>());
+        result.lower_y = embree::sub_group_inclusive_scan(lower_y, SYCL_EXT_ONEAPI::minimum<float>());
+        result.lower_z = embree::sub_group_inclusive_scan(lower_z, SYCL_EXT_ONEAPI::minimum<float>());
+        result.upper_x = embree::sub_group_inclusive_scan(upper_x, SYCL_EXT_ONEAPI::maximum<float>());
+        result.upper_y = embree::sub_group_inclusive_scan(upper_y, SYCL_EXT_ONEAPI::maximum<float>());
+        result.upper_z = embree::sub_group_inclusive_scan(upper_z, SYCL_EXT_ONEAPI::maximum<float>());
+        return result;	
+      }
+
+      __forceinline AABB3f work_group_reduce() const
+      {
+        AABB3f result;
+        result.lower_x = work_group_reduce_min(lower_x);
+        result.lower_y = work_group_reduce_min(lower_y);
+        result.lower_z = work_group_reduce_min(lower_z);
+        result.upper_x = work_group_reduce_max(upper_x);
+        result.upper_y = work_group_reduce_max(upper_y);
+        result.upper_z = work_group_reduce_max(upper_z);
+        return result;	
+      }
+
+      __forceinline float halfArea() const
+      {
+        const float3 d = size();
+        return halfarea(d);
+      }
+
+      __forceinline float area() const
+      {
+        return halfArea() * 2.0f;
+      }
+      
+      __forceinline float maxAbsElement() const
+      {
+        const float3 abs_l = abs(lower());
+        const float3 abs_h = abs(upper());
+        const float max_abs_l = max(max(abs_l.x(),abs_l.y()),abs_l.z());
+        const float max_abs_h = max(max(abs_h.x(),abs_h.y()),abs_h.z());
+        return max(max_abs_l,max_abs_h);
+      }
+
+      __forceinline AABB3f sub_group_masked_reduce(const bool cmp) const
+      {
+        AABB3f result;
+        const float _pos_inf =  INFINITY;
+        const float _neg_inf = -INFINITY;
+        result.lower_x = cselect(cmp,lower_x,_pos_inf);
+        result.lower_y = cselect(cmp,lower_y,_pos_inf);
+        result.lower_z = cselect(cmp,lower_z,_pos_inf);
+        result.upper_x = cselect(cmp,upper_x,_neg_inf);
+        result.upper_y = cselect(cmp,upper_y,_neg_inf);
+        result.upper_z = cselect(cmp,upper_z,_neg_inf);
+        return result.sub_group_reduce();
+      }
+      
+      __forceinline AABB3f sub_group_invalidate_reduce_left(const uint pos) const
+      {
+        AABB3f result;
+        const uint subgroupLocalID = get_sub_group_local_id();
+        const bool cmp = subgroupLocalID <= pos;
+        const float _pos_inf =  INFINITY;
+        const float _neg_inf = -INFINITY;
+        result.lower_x = cselect(cmp,lower_x,_pos_inf);
+        result.lower_y = cselect(cmp,lower_y,_pos_inf);
+        result.lower_z = cselect(cmp,lower_z,_pos_inf);
+        result.upper_x = cselect(cmp,upper_x,_neg_inf);
+        result.upper_y = cselect(cmp,upper_y,_neg_inf);
+        result.upper_z = cselect(cmp,upper_z,_neg_inf);
+        return result.sub_group_reduce();
+      }
+
+      __forceinline AABB3f sub_group_invalidate_reduce_right(const uint pos) const
+      {
+        AABB3f result;
+        const uint subgroupLocalID = get_sub_group_local_id();
+        const bool cmp = subgroupLocalID > pos;
+        const float _pos_inf =  INFINITY;
+        const float _neg_inf = -INFINITY;        
+        result.lower_x = cselect(cmp,lower_x,_pos_inf);
+        result.lower_y = cselect(cmp,lower_y,_pos_inf);
+        result.lower_z = cselect(cmp,lower_z,_pos_inf);
+        result.upper_x = cselect(cmp,upper_x,_neg_inf);
+        result.upper_y = cselect(cmp,upper_y,_neg_inf);
+        result.upper_z = cselect(cmp,upper_z,_neg_inf);        
+        return result.sub_group_reduce();
+      }      
+
+      __forceinline AABB3f conservativeBounds() const {
+        const float ulps = 1.0f;
+        const float3 abs_l = abs(lower());
+        const float3 abs_h = abs(upper());
+        const float max_abs_l = max(max(abs_l.x(),abs_l.y()),abs_l.z());
+        const float max_abs_h = max(max(abs_h.x(),abs_h.y()),abs_h.z());             
+        const float err = ulps*std::numeric_limits<float>::epsilon() * max(max_abs_l,max_abs_h);
+        AABB3f bounds(lower(),upper());
+        bounds.enlarge(err);
+        return bounds;
+      }
+      
+      __forceinline BBox3f convert() const
+      {
+        return BBox3f(Vec3f(lower_x,lower_y,lower_z),Vec3f(upper_x,upper_y,upper_z));
+      }
+
+      __forceinline bool checkNumericalBounds() const
+      {
+        const static float FLT_LARGE = 1.844E18f;
+        if (lower_x < -FLT_LARGE || upper_x > FLT_LARGE) return false;
+        if (lower_y < -FLT_LARGE || upper_y > FLT_LARGE) return false;
+        if (lower_z < -FLT_LARGE || upper_z > FLT_LARGE) return false;
+        return true;
+      }
+
+      __forceinline uint numEqualDims() const
+      {
+        uint equal_dims = lower_x == upper_x ? 1 : 0;
+        equal_dims += lower_y == upper_y ? 1 : 0;
+        equal_dims += lower_z == upper_z ? 1 : 0;
+        return equal_dims;
+      }
+
+      
+      friend __forceinline embree_ostream operator<<(embree_ostream cout, const AABB3f &aabb)
+      {
+        cout << "AABB3f { ";
+        cout << "  lower = (" << aabb.lower_x << ", " << aabb.lower_y << ", " << aabb.lower_z << ") ";
+        cout << "  upper = (" << aabb.upper_x << ", " << aabb.upper_y << ", " << aabb.upper_z << ") ";
+        return cout << "}";
+      }
+    };
+
+      __forceinline AABB3f merge(const AABB3f &a, const AABB3f &b)
+      {
+        AABB3f bounds;
+        bounds.lower_x = min(a.lower_x,b.lower_x);
+        bounds.lower_y = min(a.lower_y,b.lower_y);
+        bounds.lower_z = min(a.lower_z,b.lower_z);        
+        bounds.upper_x = max(a.upper_x,b.upper_x);
+        bounds.upper_y = max(a.upper_y,b.upper_y);
+        bounds.upper_z = max(a.upper_z,b.upper_z);
+        return bounds;
+      }
+
+
+      __forceinline AABB3f intersect(const AABB3f &a, const AABB3f &b)
+      {
+        AABB3f bounds;
+        bounds.lower_x = max(a.lower_x,b.lower_x);
+        bounds.lower_y = max(a.lower_y,b.lower_y);
+        bounds.lower_z = max(a.lower_z,b.lower_z);        
+        bounds.upper_x = min(a.upper_x,b.upper_x);
+        bounds.upper_y = min(a.upper_y,b.upper_y);
+        bounds.upper_z = min(a.upper_z,b.upper_z);
+        return bounds;
+      }
+
+
+
+
+      __forceinline float distance(const AABB3f &a, const AABB3f &b)
+      {
+        const float3 dist_lower = a.lower() - b.lower();
+        const float3 dist_upper = a.upper() - b.upper();
+        return sycl::fma(dist_lower.x(),dist_lower.x(),
+                         sycl::fma(dist_lower.y(),dist_lower.y(),
+                                   sycl::fma(dist_lower.z(),dist_lower.z(),
+                                             sycl::fma(dist_upper.x(),dist_upper.x(),
+                                                       sycl::fma(dist_upper.y(), dist_upper.y(), dist_upper.z() * dist_upper.z())))));        
+      }
+    
+    __forceinline bool operator ==(const AABB3f& a, const AABB3f& b) {
+        if (a.lower_x != b.lower_x ||
+            a.lower_y != b.lower_y ||
+            a.lower_z != b.lower_z ||
+            a.upper_x != b.upper_x ||
+            a.upper_y != b.upper_y ||
+            a.upper_z != b.upper_z)
+          return false;
+        return true;
+    }
+    
+    __forceinline bool operator !=(const AABB3f& a, const AABB3f& b) {
+      return !(a==b);
+    }
+    
+    
+  }
+}
+
+#endif

kernels/rthwif/rtbuild/gpu/common.h

@@ -0,0 +1,348 @@
+// Copyright 2009-2021 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+
+#pragma once
+
+#include "../../../../common/sys/sycl.h"
+
+using sycl::float16;
+using sycl::float8;
+using sycl::float4;
+using sycl::float3;
+using sycl::float2;
+using sycl::int16;
+using sycl::int8;
+using sycl::int4;
+using sycl::int3;
+using sycl::int2;
+using sycl::uint16;
+using sycl::uint8;
+using sycl::uint4;
+using sycl::uint3;
+using sycl::uint2;
+using sycl::uchar16;
+using sycl::uchar8;
+using sycl::uchar4;
+using sycl::uchar3;
+using sycl::uchar2;
+using sycl::uchar;
+using sycl::ushort16;
+using sycl::ushort8;
+using sycl::ushort4;
+using sycl::ushort3;
+using sycl::ushort2;
+
+using sycl::fmax;
+using sycl::fmin;
+
+namespace embree
+{
+  namespace gpu
+  {
+    __forceinline size_t alignTo(const size_t size, const size_t alignment)
+    {
+      return ((size+alignment-1)/alignment)*alignment;
+    }
+
+    
+    template<typename T>
+      __forceinline T cfma(const T &a, const T &b, const T &c)
+      {
+	return sycl::fma(a,b,c);
+      }
+
+    template<typename T>
+      __forceinline T lerp(const T &a, const T &b, const float t) {      
+      return sycl::fma(T(1.0f-t),a,T(t)*b);
+    }
+      
+    __forceinline float halfarea(const sycl::float3 &d)
+    {
+      return sycl::fma(d.x(),d.y(),sycl::fma(d.x(),d.z(),d.y()*d.z()));
+    }
+
+    __forceinline float halfarea(const sycl::float4 &d)
+    {
+      return sycl::fma((float)d.x(),((float)d.y()+(float)d.z()),(float)d.y()*(float)d.z());
+    }
+    
+    __forceinline float area(const sycl::float3 &d)
+    {
+      return halfarea(d) * 2.0f;
+    }
+
+    __forceinline float area(const sycl::float4 &d)
+    {
+      return halfarea(d) * 2.0f;
+    }
+
+    __forceinline float dot3(const float3 &a,
+			     const float3 &b)
+    {
+      return sycl::dot(a,b);
+    }    
+
+
+    template<typename T>
+      static __forceinline uint atomic_add_global(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> counter(*dest);        
+        return counter.fetch_add(count);      
+      }
+
+    template<typename T>
+      static __forceinline uint atomic_min_global(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> counter(*dest);        
+        return counter.fetch_min(count);      
+      }
+
+    template<typename T>
+      static __forceinline uint atomic_max_global(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> counter(*dest);        
+        return counter.fetch_max(count);      
+      }
+
+    template<typename T>
+      static __forceinline uint atomic_or_global(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> counter(*dest);        
+        return counter.fetch_or(count);      
+      }
+    
+    template<typename T>
+      static __forceinline uint atomic_add_local(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> counter(*dest);        
+        return counter.fetch_add(count);      
+      }
+
+    template<typename T>
+      static __forceinline uint atomic_min_local(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> counter(*dest);        
+        return counter.fetch_min(count);      
+      }
+
+    template<typename T>
+      static __forceinline uint atomic_max_local(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> counter(*dest);        
+        return counter.fetch_max(count);      
+      }
+
+    template<typename T>
+      static __forceinline uint atomic_or_local(T *dest, const T count=1)
+      {
+        sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> counter(*dest);        
+        return counter.fetch_or(count);      
+      }
+    
+        
+    template<typename T>
+      static __forceinline uint as_uint(T t)
+      {
+	return __builtin_bit_cast(uint,t);
+      }
+
+    template<typename T>
+      static __forceinline uint as_int(T t)
+      {
+	return __builtin_bit_cast(int,t);
+      }
+    
+    template<typename T>
+      static __forceinline float as_float(T t)
+      {
+	return __builtin_bit_cast(float,t);
+      }
+
+    template<typename T>
+      static __forceinline void* as_void(T t)
+      {
+	return __builtin_bit_cast(void*,t);
+      }
+    
+    __forceinline uint bitInterleave3D(const uint3 &in)
+    {
+      uint x = in.x(), y = in.y(), z = in.z();
+      x = (x | (x << 16)) & 0x030000FF;
+      x = (x | (x << 8)) & 0x0300F00F;
+      x = (x | (x << 4)) & 0x030C30C3;
+      x = (x | (x << 2)) & 0x09249249;
+
+      y = (y | (y << 16)) & 0x030000FF;
+      y = (y | (y << 8)) & 0x0300F00F;
+      y = (y | (y << 4)) & 0x030C30C3;
+      y = (y | (y << 2)) & 0x09249249;
+
+      z = (z | (z << 16)) & 0x030000FF;
+      z = (z | (z << 8)) & 0x0300F00F;
+      z = (z | (z << 4)) & 0x030C30C3;
+      z = (z | (z << 2)) & 0x09249249;
+
+      return x | (y << 1) | (z << 2);
+    }
+
+
+    __forceinline uint bitInterleave4D(const uint4 &in)
+    {
+      uint x = in.x(), y = in.y(), z = in.z(), w = in.w();
+
+      x = x & 0x000000ff;
+      x = (x ^ (x << 16)) & 0x00c0003f;
+      x = (x ^ (x << 8)) & 0x00c03807;
+      x = (x ^ (x << 4)) & 0x08530853;
+      x = (x ^ (x << 2)) & 0x09090909;
+      x = (x ^ (x << 1)) & 0x11111111;
+
+      y = y & 0x000000ff;
+      y = (y ^ (y << 16)) & 0x00c0003f;
+      y = (y ^ (y << 8)) & 0x00c03807;
+      y = (y ^ (y << 4)) & 0x08530853;
+      y = (y ^ (y << 2)) & 0x09090909;
+      y = (y ^ (y << 1)) & 0x11111111;
+
+      z = z & 0x000000ff;
+      z = (z ^ (z << 16)) & 0x00c0003f;
+      z = (z ^ (z << 8)) & 0x00c03807;
+      z = (z ^ (z << 4)) & 0x08530853;
+      z = (z ^ (z << 2)) & 0x09090909;
+      z = (z ^ (z << 1)) & 0x11111111;
+
+      w = w & 0x000000ff;
+      w = (w ^ (w << 16)) & 0x00c0003f;
+      w = (w ^ (w << 8)) & 0x00c03807;
+      w = (w ^ (w << 4)) & 0x08530853;
+      w = (w ^ (w << 2)) & 0x09090909;
+      w = (w ^ (w << 1)) & 0x11111111;
+
+      return (x | (y << 1) | (z << 2) | (w << 3));
+    }
+
+    __forceinline uint64_t splitBy3(const uint a){
+      uint64_t x = a & 0x1fffff; 
+      x = (x | x << 32) & 0x1f00000000ffff; 
+      x = (x | x << 16) & 0x1f0000ff0000ff; 
+      x = (x | x << 8) & 0x100f00f00f00f00f;
+      x = (x | x << 4) & 0x10c30c30c30c30c3;
+      x = (x | x << 2) & 0x1249249249249249;
+      return x;
+    }
+
+    __forceinline uint64_t bitInterleave3D_64bits(const uint3 v){
+      return splitBy3(v.x()) | splitBy3(v.y()) << 1 | splitBy3(v.z()) << 2;
+    }
+
+    __forceinline uint64_t bitInterleave4D_64bits(const uint4 &v){
+      const uint low  = bitInterleave4D(v>>0);
+      const uint high = bitInterleave4D(v>>8);
+      return (((uint64_t)high)<<32) | low;      
+    }
+    
+    
+    __forceinline uint atomic_add_global_sub_group_shared(sycl::atomic_ref<uint, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> &counter, const uint add)
+    {
+      const uint subgroupLocalID = get_sub_group_local_id();    
+      const uint ballot = sub_group_ballot(true);
+      const uint first  = sycl::ctz(ballot);
+      const uint popc   = sycl::popcount(ballot);
+      const uint prefix = sycl::popcount(ballot & (((uint)1 << subgroupLocalID)-1));
+      uint index = 0;
+      if (subgroupLocalID == first)
+        index = counter.fetch_add(popc * add);
+      index = sub_group_broadcast(index,first);
+      return index + prefix;
+    }
+
+    __forceinline uint atomic_add_local_sub_group_shared(sycl::atomic_ref<uint, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> &counter, const uint add)
+    {
+      const uint subgroupLocalID = get_sub_group_local_id();    
+      const uint ballot = sub_group_ballot(true);
+      const uint first  = sycl::ctz(ballot);
+      const uint popc   = sycl::popcount(ballot);
+      const uint prefix = sycl::popcount(ballot & (((uint)1 << subgroupLocalID)-1));
+      uint index = 0;
+      if (subgroupLocalID == first)
+        index = counter.fetch_add(popc * add);
+      index = sub_group_broadcast(index,first);
+      return index + prefix;
+    }
+
+    template<typename T>        
+    __forceinline uint atomic_add_global_sub_group_varying(sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> &counter, const T add)
+    {
+      const uint subgroupLocalID = get_sub_group_local_id();    
+      const uint ballot = sub_group_ballot(true);
+      const uint first  = sycl::ctz(ballot);
+      const T total  = sub_group_reduce(add, SYCL_EXT_ONEAPI::plus<T>());
+      const T prefix = sub_group_exclusive_scan(add, SYCL_EXT_ONEAPI::plus<T>());
+      uint index = 0;
+      if (subgroupLocalID == first)
+        index = counter.fetch_add(total);
+      index = sub_group_broadcast(index,first);
+      return index + prefix;
+    }
+
+    template<typename T>    
+    static __forceinline T atomic_add_global_sub_group_varying(T *dest, const T count=1)
+    {
+      sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::global_space> counter(*dest);
+      return atomic_add_global_sub_group_varying(counter,count);
+    }
+
+    template<typename T>        
+    __forceinline uint atomic_add_local_sub_group_varying(sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::local_space> &counter, const T add)
+    {
+      const uint subgroupLocalID = get_sub_group_local_id();    
+      const uint ballot = sub_group_ballot(true);
+      const uint first  = sycl::ctz(ballot);
+      const T total  = sub_group_reduce(add, SYCL_EXT_ONEAPI::plus<T>());
+      const T prefix = sub_group_exclusive_scan(add, SYCL_EXT_ONEAPI::plus<T>());
+      uint index = 0;
+      if (subgroupLocalID == first)
+        index = counter.fetch_add(total);
+      index = sub_group_broadcast(index,first);
+      return index + prefix;
+    }
+
+    template<typename T>    
+    static __forceinline T atomic_add_local_sub_group_varying(T *dest, const T count=1)
+    {
+      sycl::atomic_ref<T, sycl::memory_order::relaxed, sycl::memory_scope::device,sycl::access::address_space::local_space> counter(*dest);
+      return sub_group_shared_varying_local_atomic(counter,count);
+    }
+    
+    __forceinline void waitOnQueueAndCatchException(sycl::queue &gpu_queue)
+    {
+      try {
+        gpu_queue.wait_and_throw();
+      } catch (sycl::exception const& e) {
+        std::cout << "Caught synchronous SYCL exception:\n"
+                  << e.what() << std::endl;
+        FATAL("SYCL Exception");     
+      }      
+    }
+
+    __forceinline void waitOnEventAndCatchException(sycl::event &event)
+    {
+      try {
+        event.wait_and_throw();
+      } catch (sycl::exception const& e) {
+        std::cout << "Caught synchronous SYCL exception:\n"
+                  << e.what() << std::endl;
+        FATAL("SYCL Exception");     
+      }      
+    }
+    
+    __forceinline double getDeviceExecutionTiming(sycl::event &queue_event)
+    {
+      const auto t0 = queue_event.template get_profiling_info<sycl::info::event_profiling::command_start>();
+      const auto t1 = queue_event.template get_profiling_info<sycl::info::event_profiling::command_end>();
+      return (t1-t0)*1E-6;      
+    }
+    
+  };
+};
+

kernels/rthwif/rtbuild/gpu/morton.h

@@ -0,0 +1,289 @@
+// Copyright 2009-2021 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+
+#pragma once
+
+#if defined(EMBREE_SYCL_SUPPORT)
+#include "common.h"
+
+namespace embree
+{
+  namespace gpu
+  {
+
+    struct Range {
+      uint start;
+      uint end;
+      __forceinline Range()
+      {}
+
+      __forceinline Range(const uint start, const uint end) : start(start), end(end)
+      {}
+
+      __forceinline uint size() const { return end-start; }
+
+      __forceinline uint center() const { return (end+start)>>1; }
+      
+    };
+
+
+    // ================================================================================================================================================================================
+    // ============================================================================= MORTON CODES =====================================================================================
+    // ================================================================================================================================================================================
+
+
+    struct __aligned(8) MortonCodePrimitive40x24Bits3D
+    {
+      static const uint GRID_SHIFT = 13; // * 3 = 30
+      
+      static const uint CODE_SHIFT = 64-(3*GRID_SHIFT+1);      
+      static const uint64_t INDEX_MASK = ((uint64_t)1 << CODE_SHIFT)-1;
+        
+      uint64_t index_code; // 64bit code + index combo
+
+      __forceinline MortonCodePrimitive40x24Bits3D() {}      
+      __forceinline MortonCodePrimitive40x24Bits3D(const uint64_t code, const uint index) { index_code = (code << CODE_SHIFT) | (uint64_t)index; }
+      __forceinline MortonCodePrimitive40x24Bits3D(const uint3 &v, const uint index)
+      {
+        const uint64_t code = gpu::bitInterleave3D(v);
+        MortonCodePrimitive40x24Bits3D(code,index);
+      }
+
+      __forceinline uint getIndex() const { return (uint)(index_code & INDEX_MASK); }
+      __forceinline uint64_t getCode()  const { return index_code; }
+
+      __forceinline uint64_t getMCode()  const { return index_code & (~INDEX_MASK); }      
+      
+      __forceinline operator const uint64_t&()  const { return index_code; }
+
+      __forceinline bool operator < (const MortonCodePrimitive40x24Bits3D& mc) const
+      {
+        return index_code < mc.index_code;
+      }
+    
+    };
+
+    struct __aligned(8) MortonCodePrimitive40x24Bits4D
+    {
+      static const uint GRID_SHIFT = 10; // * 4 = 40      
+      static const uint CODE_SHIFT = 64-4*GRID_SHIFT;      
+      static const uint64_t INDEX_MASK = ((uint64_t)1 << CODE_SHIFT)-1;
+        
+      uint64_t index_code; // 64bit code + index combo
+
+      __forceinline MortonCodePrimitive40x24Bits4D() {}
+      __forceinline MortonCodePrimitive40x24Bits4D(const uint64_t code, const uint index) { index_code = (code << CODE_SHIFT) | (uint64_t)index; }
+      __forceinline MortonCodePrimitive40x24Bits4D(const uint4 &v, const uint index)
+      {
+        const uint64_t code = gpu::bitInterleave4D(v);
+        MortonCodePrimitive40x24Bits4D(code,index);
+      }
+      
+      __forceinline uint getIndex() const { return (uint)(index_code & INDEX_MASK); }
+      __forceinline uint64_t getCode()  const { return index_code; }
+      __forceinline uint64_t getMCode()  const { return index_code & (~INDEX_MASK); }            
+      __forceinline operator const uint64_t&()  const { return index_code; }
+      __forceinline bool operator < (const MortonCodePrimitive40x24Bits4D& mc) const
+      {
+        return index_code < mc.index_code;
+      }
+    
+    };
+    
+
+    /* used by PLOC in a two phase process */
+    struct __aligned(8) MortonCodePrimitive64Bit_2x
+    {
+      static const uint GRID_SHIFT = 21; // * 3 = 63
+      
+      static const uint CODE_SHIFT = 32;      
+      static const uint64_t INDEX_MASK = ((uint64_t)1 << CODE_SHIFT)-1;
+        
+      uint64_t index_code; // 64bit code + index combo
+
+      __forceinline MortonCodePrimitive64Bit_2x() {}
+      __forceinline MortonCodePrimitive64Bit_2x(const uint64_t code, const uint index) { index_code = (code << CODE_SHIFT) | (uint64_t)index; }
+
+      __forceinline uint getIndex() const { return (uint)(index_code & INDEX_MASK); }
+      __forceinline uint64_t getCode()  const { return index_code; }
+
+      __forceinline uint64_t getMCode()  const { return index_code & (~INDEX_MASK); }      
+      
+      __forceinline operator const uint64_t&()  const { return index_code; }
+
+      __forceinline bool operator < (const MortonCodePrimitive64Bit_2x& mc) const
+      {
+        return index_code < mc.index_code;
+      }
+    
+    };
+
+
+    struct MortonCodePrimitive64x32Bits3D
+    {
+      static const uint GRID_SHIFT = 21; // * 3 = 63
+      static const uint KEY_BITS   = 64;
+      
+      uint64_t code;
+      uint index;
+
+      __forceinline MortonCodePrimitive64x32Bits3D() {}      
+      __forceinline MortonCodePrimitive64x32Bits3D(const uint64_t code, const uint index) : code(code), index(index) {}
+      __forceinline MortonCodePrimitive64x32Bits3D(const uint3 &v, const uint index) : code(gpu::bitInterleave3D_64bits(v)), index(index)
+      {
+      }
+      
+      __forceinline bool operator < (const MortonCodePrimitive64x32Bits3D& emc) const
+      {
+      
+        if (code < emc.code)
+          return true;
+        else if (code == emc.code)
+          return index < emc.index;
+        else
+          return false;
+      }
+
+      __forceinline uint64_t getCode() const { return code; }
+      __forceinline uint64_t getMCode()  const { return code; }            
+      __forceinline uint    getIndex() const { return index; }
+      __forceinline operator const uint64_t&()  const { return code; }    
+    };
+
+    struct MortonCodePrimitive64x32Bits4D
+    {
+      static const uint GRID_SHIFT = 16; // * 4 = 64
+      static const uint KEY_BITS   = 64;
+      
+      uint64_t code;
+      uint index;
+
+      __forceinline MortonCodePrimitive64x32Bits4D() {}
+      __forceinline MortonCodePrimitive64x32Bits4D(const uint64_t code, const uint index) : code(code), index(index) {}
+      __forceinline MortonCodePrimitive64x32Bits4D(const uint4 &v, const uint index) : code( gpu::bitInterleave4D_64bits(v)), index(index)
+      {
+      }
+      
+      __forceinline bool operator < (const MortonCodePrimitive64x32Bits4D& emc) const
+      {
+      
+        if (code < emc.code)
+          return true;
+        else if (code == emc.code)
+          return index < emc.index;
+        else
+          return false;
+      }
+
+      __forceinline uint64_t getCode() const { return code; }
+      __forceinline uint64_t getMCode()  const { return code; }            
+      __forceinline uint    getIndex() const { return index; }
+      __forceinline operator const uint64_t&()  const { return code; }    
+    };
+    
+
+
+    template<typename type>  
+      __forceinline unsigned int findSplit64Bit(const gpu::Range &current, const type *const morton)
+    {
+      const uint64_t code_start = morton[current.start].getCode();
+      const uint64_t code_end   = morton[current.end-1].getCode();
+
+      if (code_start != code_end)
+      {
+        uint64_t bitpos = sycl::clz(code_start^code_end);
+        
+        /* split the items at the topmost different morton code bit */
+        const uint64_t bitpos_diff = 63-bitpos;
+        const uint64_t bitmask = (uint64_t)1 << bitpos_diff;
+  
+        /* find location where bit differs using binary search */
+        unsigned int begin = current.start;
+        unsigned int end   = current.end;
+        while (begin + 1 != end) {
+          const unsigned int mid = (begin+end)/2;
+          const uint64_t bit = morton[mid].getCode() & bitmask;
+          if (bit == 0) begin = mid; else end = mid;
+        }
+        unsigned int center = end;
+        return center;        
+      }
+      else  /* if all items mapped to same morton code */
+      {
+        uint bitpos = sycl::clz(morton[current.start].getIndex()^morton[current.end-1].getIndex());
+        
+        /* split the items at the topmost different morton code bit */
+        const uint bitpos_diff = 31-bitpos;
+        const uint bitmask = (uint)1 << bitpos_diff;
+  
+        /* find location where bit differs using binary search */
+        unsigned int begin = current.start;
+        unsigned int end   = current.end;
+        while (begin + 1 != end) {
+          const unsigned int mid = (begin+end)/2;
+          const uint bit = morton[mid].getIndex() & bitmask;
+          if (bit == 0) begin = mid; else end = mid;
+        }
+        unsigned int center = end;
+        return center;
+      }
+    }
+
+    template<typename type>      
+      __forceinline uint delta(const type& m0, const type& m1)
+    {
+      return m0.getCode() != m1 .getCode() ? (m0.getCode()^m1.getCode()) : m0.getIndex()^m1.getIndex(); //sycl::ext::intel::ctz(m0.index_code^m1.index_code);
+    }
+    
+
+    template<typename type>  
+      __forceinline void splitRange(const gpu::Range &current, const type *const morton, gpu::Range &left, gpu::Range &right)
+    {
+      const uint split = findSplit64Bit(current,morton);
+      left.start  = current.start;
+      left.end    = split;
+      right.start = split;
+      right.end   = current.end;
+    }
+
+
+    template<typename type>          
+      __forceinline int sigma(const int i1, const int i2, const type *const morton, const uint numPrimitives) 
+    {
+      const int left  = min(i1, i2);
+      const int right = max(i1, i2);
+      if (left < 0 || right >= numPrimitives) return -1;
+      const ulong left_code  = morton[left].getCode();
+      const ulong right_code = morton[right].getCode();
+      if (left_code != right_code)
+        return sycl::clz(left_code ^ right_code);
+      else
+        return 64 + sycl::clz(morton[left].getIndex() ^ morton[right].getIndex());
+    }
+
+    template<typename type>          
+      __forceinline struct gpu::Range findSegment(const unsigned int idx, const type *const mc, uint const numPrimitives)
+    {
+      const int delta = sigma(idx, idx+1, mc, numPrimitives) >= sigma(idx, idx-1, mc, numPrimitives) ? 1 : -1;
+      const int sigma_min = sigma(idx, idx-delta, mc, numPrimitives);
+      int max_range = 2;
+      while (sigma(idx,idx + max_range * delta, mc, numPrimitives) > sigma_min)
+        max_range <<= 1;
+      int step  = 0;
+      int r = max_range;
+      do
+      {
+        r >>= 1;
+        if(sigma(idx, idx + (step + r)*delta, mc, numPrimitives) > sigma_min)
+          step += r;
+      }
+      while (r > 1);
+      struct gpu::Range range;
+      range.start = min(idx, idx + step*delta);
+      range.end   = max(idx, idx + step*delta) + 1;
+      return range;          
+    }    
+  }
+}
+
+#endif

kernels/rthwif/rtbuild/gpu/sort.h

@@ -0,0 +1,379 @@
+// Copyright 2009-2021 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+
+#pragma once
+
+#if defined(EMBREE_SYCL_SUPPORT)
+#include "common.h"
+
+#define RADIX_SORT_BINS        256    
+#define RADIX_SORT_MAX_NUM_DSS 128
+#define RADIX_SORT_WG_SIZE     512
+#define RADIX_ITERATIONS_64BIT 8
+
+namespace embree
+{
+  namespace gpu
+  {
+
+    __forceinline uint getNumWGsScratchSize(const ssize_t scratch_size)
+    {
+      return min(max( (int)floorf(scratch_size / (sizeof(uint)*RADIX_SORT_BINS))-1,(int)1),RADIX_SORT_MAX_NUM_DSS);
+    }
+    
+    __forceinline void localAtomicBallot(uint *const histogram, const uint ID, const uint add)
+    {
+      uint mask = sub_group_ballot(true);
+      while(mask)
+      {
+        const uint first = sycl::ctz(mask);
+        const uint index = sub_group_broadcast(ID,first);
+        const bool cmp = ID == index;
+        const uint cmp_mask = sub_group_ballot(cmp);
+        const uint reduced_count = sycl::popcount(cmp_mask) * add;
+        mask &= ~cmp_mask;
+        
+        sycl::atomic_ref<uint, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> hist(histogram[ID]);                                                                                    
+        if (get_sub_group_local_id() == first)
+          hist += reduced_count;
+      }        
+    }
+    
+
+    __forceinline void radix_sort_single_workgroup(sycl::queue &gpu_queue, uint64_t *_input, uint64_t *_output, const uint numPrimitives,  const uint start_iteration, const uint stop_iteration)
+    {
+      static const uint RADIX_SORT_SINGLE_WG_SIZE = 256; 
+
+      struct __aligned(RADIX_SORT_SINGLE_WG_SIZE/32 * sizeof(uint)) BinFlags
+      {
+        uint flags[RADIX_SORT_SINGLE_WG_SIZE/32];
+      };
+    
+      const sycl::nd_range<1> nd_range1(RADIX_SORT_SINGLE_WG_SIZE,sycl::range<1>(RADIX_SORT_SINGLE_WG_SIZE));          
+      sycl::event event = gpu_queue.submit([&](sycl::handler &cgh) {
+                                             sycl::local_accessor< uint, 1> histogram(sycl::range<1>(RADIX_SORT_SINGLE_WG_SIZE),cgh);
+                                             sycl::local_accessor< uint, 1> sums(sycl::range<1>(RADIX_SORT_SINGLE_WG_SIZE),cgh);
+                                             sycl::local_accessor< uint, 1> prefix_sums(sycl::range<1>(RADIX_SORT_SINGLE_WG_SIZE),cgh);
+                                             sycl::local_accessor< uint, 1> local_offset(sycl::range<1>(RADIX_SORT_SINGLE_WG_SIZE),cgh);
+                                             sycl::local_accessor< BinFlags, 1> bin_flags(sycl::range<1>(RADIX_SORT_BINS),cgh);
+                                                   
+                                             cgh.parallel_for(nd_range1,[=](sycl::nd_item<1> item) EMBREE_SYCL_SIMD(16)                                                                    
+                                                              {                                                                                                                                            
+                                                                const uint localID     = item.get_local_id(0);
+                                                                const uint step_local  = item.get_local_range().size();
+                                                                const uint subgroupID      = get_sub_group_id();
+                                                                const uint subgroupSize    = get_sub_group_size();
+                                                                const uint subgroupLocalID = get_sub_group_local_id();
+
+                                                                for (uint iter = start_iteration;iter<stop_iteration;iter++)
+                                                                {
+                                                                  item.barrier(sycl::access::fence_space::local_space);
+                                                                          
+                                                                  const uint shift = iter*8;
+
+                                                                  const uint64_t *const input  = (iter % 2) == 0 ? _input : _output;
+                                                                  uint64_t *const output = (iter % 2) == 1 ? _input : _output;
+                                                                        
+                                                                  // ==== bin key into global histogram =====
+                                                                      
+                                                                  if (localID < RADIX_SORT_BINS)
+                                                                    histogram[localID] = 0;
+
+                                                                  item.barrier(sycl::access::fence_space::local_space);
+                                                                    
+                                                                  for (uint ID = localID; ID < numPrimitives; ID += step_local)
+                                                                  {
+                                                                    const uint bin = ((uint)(input[ID] >> shift)) & (RADIX_SORT_BINS - 1);
+                                                                    //localAtomicBallot(histogram.get_pointer(),bin,1);
+                                                                    gpu::atomic_add_local((uint*)histogram.get_pointer() + bin,(uint)1);
+                                                                  }
+
+                                                                  item.barrier(sycl::access::fence_space::local_space);
+
+                                                                  // ==== reduce global histogram =====    
+                                                                                                                                              
+                                                                  SYCL_EXT_ONEAPI::sub_group sub_group = this_sub_group();
+                                                                  sub_group.barrier();
+
+                                                                  if (localID < RADIX_SORT_BINS)
+                                                                  {
+                                                                    const uint count = histogram[localID];
+                                                                    const uint sum = sub_group_reduce(count, std::plus<uint>());
+                                                                    const uint prefix_sum = sub_group_exclusive_scan(count, std::plus<uint>());
+
+                                                                    sums[localID] = sum;
+                                                                    prefix_sums[localID] = prefix_sum;
+                                                                  }
+                                                                        
+                                                                  item.barrier(sycl::access::fence_space::local_space);
+
+
+                                                                  if (subgroupID == 0)
+                                                                  {
+                                                                    uint off = 0;
+                                                                    for (int i = subgroupLocalID; i < RADIX_SORT_BINS; i += subgroupSize)
+                                                                    {
+                                                                      local_offset[i] = off + prefix_sums[i];
+                                                                      off += sums[i];
+                                                                    }
+                                                                  }
+
+                                                                  item.barrier(sycl::access::fence_space::local_space);                                                                          
+                                                                      
+                                                                  // ==== scatter key/value pairs according to global historgram =====      
+                                                                                                                         
+                                                                  const uint flags_bin = localID / 32;
+                                                                  const uint flags_bit = 1 << (localID % 32);                                                                      
+                                                                      
+                                                                  for (uint blockID = 0; blockID < numPrimitives; blockID += step_local)
+                                                                  {
+                                                                    item.barrier(sycl::access::fence_space::local_space);
+                                                                        
+                                                                    const uint ID = blockID + localID;                                                                            
+                                                                        
+                                                                    uint binID = 0;
+                                                                    uint binOffset = 0;
+
+                                                                    if (localID < RADIX_SORT_BINS)                                                                            
+                                                                      for (int i=0;i<RADIX_SORT_SINGLE_WG_SIZE/32;i++)
+                                                                        bin_flags[localID].flags[i] = 0;
+
+                                                                    item.barrier(sycl::access::fence_space::local_space);
+                                                                                 
+                                                                            
+                                                                    uint64_t in;
+                                                                    if (ID < numPrimitives)
+                                                                    {
+                                                                      in = input[ID];                                                                              
+                                                                      binID = (in >> shift) & (RADIX_SORT_BINS - 1);                                                                            
+                                                                      binOffset = local_offset[binID];
+                                                                              
+                                                                      sycl::atomic_ref<uint, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> bflags(bin_flags[binID].flags[flags_bin]);                                                                                    
+                                                                      bflags.fetch_add(flags_bit);
+
+                                                                    }
+
+                                                                          
+                                                                    item.barrier(sycl::access::fence_space::local_space);
+                                                                            
+                                                                    if (ID < numPrimitives)
+                                                                    {
+                                                                      uint prefix = 0;
+                                                                      uint count = 0;
+                                                                      for (uint i = 0; i < RADIX_SORT_SINGLE_WG_SIZE / 32; i++)
+                                                                      {
+                                                                        const uint bits = bin_flags[binID].flags[i];
+                                                                        const uint full_count    = sycl::popcount(bits);
+                                                                        const uint partial_count = sycl::popcount(bits & (flags_bit - 1));
+                                                                        prefix += (i  < flags_bin) ? full_count : 0;
+                                                                        prefix += (i == flags_bin) ? partial_count : 0;
+                                                                        count += full_count;
+                                                                      }
+
+                                                                      output[binOffset + prefix] = in;
+                                                                      if (prefix == count - 1)
+                                                                        local_offset[binID] += count;                                                                          
+                                                                    }
+                                                                          
+                                                                  }                                                                       
+                                                                }
+                                                              });
+                                                   
+                                           });
+
+                
+      try {
+        event.wait_and_throw();
+      } catch (sycl::exception const& e) {
+        std::cout << "Caught synchronous SYCL exception:\n"
+                  << e.what() << std::endl;
+        FATAL("SYCL Exception");     		
+      }
+    }
+
+    
+    template<typename sort_type>
+    sycl::event radix_sort_iteration_type(sycl::queue &gpu_queue, sycl::event &input_event, sort_type *input, sort_type *output, const uint primitives,  uint *global_histogram, const uint iter, const uint RADIX_SORT_NUM_DSS=256)
+    {
+      const uint shift = iter*8;
+
+      
+      // ==== bin key into global histogram =====
+      
+      const sycl::nd_range<1> nd_range_binning(sycl::range<1>(RADIX_SORT_WG_SIZE*RADIX_SORT_NUM_DSS),sycl::range<1>(RADIX_SORT_WG_SIZE));          
+      sycl::event bin_event = gpu_queue.submit([&](sycl::handler &cgh) {
+                                                 cgh.depends_on(input_event);
+                                                 sycl::local_accessor< uint, 1> histogram(sycl::range<1>(RADIX_SORT_BINS),cgh);                                                 
+                                                 cgh.parallel_for(nd_range_binning,[=](sycl::nd_item<1> item) EMBREE_SYCL_SIMD(16)
+                                                                  {
+                                                                    const uint localID     = item.get_local_id(0);
+                                                                    const uint step_local  = item.get_local_range().size();
+                                                                    const uint groupID     = item.get_group(0);
+
+                                                                    const uint startID = ((size_t)groupID + 0)*(size_t)primitives / RADIX_SORT_NUM_DSS;
+                                                                    const uint endID   = ((size_t)groupID + 1)*(size_t)primitives / RADIX_SORT_NUM_DSS;
+                                                                                                                                                                                                            
+                                                                    if (localID < RADIX_SORT_BINS)
+                                                                      histogram[localID] = 0;
+
+                                                                    item.barrier(sycl::access::fence_space::local_space);
+                                                                    
+                                                                    for (uint ID = startID + localID; ID < endID; ID += step_local)
+                                                                    {
+                                                                      const uint64_t key = input[ID];
+                                                                      const uint bin = ((uint)(key >> shift)) & (RADIX_SORT_BINS - 1);
+                                                                      //gpu::localAtomicBallot(histogram.get_pointer(),bin,1);
+                                                                      gpu::atomic_add_local((uint*)histogram.get_pointer() + bin,(uint)1);
+                                                                          
+                                                                    }
+
+                                                                    item.barrier(sycl::access::fence_space::local_space);
+    
+                                                                    if (localID < RADIX_SORT_BINS)
+                                                                      global_histogram[RADIX_SORT_BINS*groupID + localID] = histogram[localID];
+                                                                    
+                                                                  });
+                                                 
+                                               });
+            
+      // ==== scatter key/value pairs according to global historgram =====
+        
+      struct __aligned(64) BinFlags
+      {
+        uint flags[RADIX_SORT_WG_SIZE/32];
+      };
+      
+      const sycl::nd_range<1> nd_range_scatter(sycl::range<1>(RADIX_SORT_WG_SIZE*RADIX_SORT_NUM_DSS),sycl::range<1>(RADIX_SORT_WG_SIZE));          
+      sycl::event scatter_event = gpu_queue.submit([&](sycl::handler &cgh) {
+                                                     cgh.depends_on(bin_event);
+                                                     sycl::local_accessor< uint, 1> local_offset(sycl::range<1>(RADIX_SORT_BINS),cgh);
+                                                     sycl::local_accessor< uint, 1> sums(sycl::range<1>(RADIX_SORT_BINS),cgh);                                                     
+                                                     sycl::local_accessor< BinFlags, 1> bin_flags(sycl::range<1>(RADIX_SORT_BINS),cgh);
+                                                     cgh.parallel_for(nd_range_scatter,[=](sycl::nd_item<1> item) EMBREE_SYCL_SIMD(16)
+                                                                      {
+                                                                        const uint groupID     = item.get_group(0);
+                                                                        const uint localID     = item.get_local_id(0);
+                                                                        const uint step_local  = item.get_local_range().size();
+                                                                        const uint subgroupID      = get_sub_group_id();
+                                                                        const uint subgroupLocalID = get_sub_group_local_id();
+                                                                        
+
+                                                                        const uint startID = ((size_t)groupID + 0)*(size_t)primitives / RADIX_SORT_NUM_DSS;
+                                                                        const uint endID   = ((size_t)groupID + 1)*(size_t)primitives / RADIX_SORT_NUM_DSS;
+
+                                                                        
+                                                                        /* --- reduce global histogram --- */
+                                                                        uint local_hist = 0;
+                                                                        uint prefix_sum = 0;
+                                                                        
+                                                                        if (localID < RADIX_SORT_BINS)
+                                                                        {
+                                                                          uint t = 0;
+                                                                          for (uint j = 0; j < RADIX_SORT_NUM_DSS; j++)
+                                                                          {
+                                                                            const uint count = global_histogram[RADIX_SORT_BINS*j + localID];
+                                                                            local_hist = (j == groupID) ? t : local_hist;
+                                                                            t += count;
+                                                                          }
+                                                                        
+                                                                          const uint count = t;
+                                                                          const uint sum = sub_group_reduce(count, std::plus<uint>());
+                                                                          prefix_sum = sub_group_exclusive_scan(count, std::plus<uint>());
+
+                                                                          sums[subgroupID] = sum;
+                                                                        }
+                                                                      
+                                                                        item.barrier(sycl::access::fence_space::local_space);
+
+                                                                        if (localID < RADIX_SORT_BINS)
+                                                                        {                                                                        
+                                                                          const uint sums_prefix_sum = sub_group_broadcast(sub_group_exclusive_scan(sums[subgroupLocalID], std::plus<uint>()),subgroupID);
+                                                                          const uint global_hist = sums_prefix_sum + prefix_sum;
+                                                                          local_offset[localID] = global_hist + local_hist;
+                                                                        }
+
+                                                                        // === barrier comes later ===
+                                                                        
+                                                                        const uint flags_bin = localID / 32;
+                                                                        const uint flags_bit = 1 << (localID % 32);                                                                      
+                                                                      
+                                                                        for (uint blockID = startID; blockID < endID; blockID += step_local)
+                                                                        {
+                                                                        
+                                                                          const uint ID = blockID + localID;
+                                                                        
+                                                                          uint binID = 0;
+                                                                          uint binOffset = 0;
+
+#if 0                                                                          
+                                                                          if (localID < RADIX_SORT_BINS)
+                                                                            for (int i=0;i<RADIX_SORT_WG_SIZE/32;i++)
+                                                                              bin_flags[localID].flags[i] = 0;
+#else
+                                                                          uint *const bflags = (uint*)&bin_flags.get_pointer()[0];                                                                          
+                                                                          for (uint i=localID;i<RADIX_SORT_BINS*RADIX_SORT_WG_SIZE/32;i+=step_local)
+                                                                            bflags[i] = 0;
+#endif                                                                          
+
+                                                                          item.barrier(sycl::access::fence_space::local_space);
+
+                                                                          sort_type key_value;
+                                                                          uint64_t key;
+                                                                          if (ID < endID)
+                                                                          {
+                                                                            key_value = input[ID];
+                                                                            key = key_value;
+                                                                            binID = (key >> shift) & (RADIX_SORT_BINS - 1);
+                                                                            binOffset = local_offset[binID];
+                                                                            sycl::atomic_ref<uint, sycl::memory_order::relaxed, sycl::memory_scope::work_group,sycl::access::address_space::local_space> bflags(bin_flags[binID].flags[flags_bin]);                                                                            
+                                                                            bflags += flags_bit;
+                                                                          }
+
+                                                                          item.barrier(sycl::access::fence_space::local_space);
+                                                                        
+                                                                          if (ID < endID)
+                                                                          {
+                                                                            uint prefix = 0;
+                                                                            uint count = 0;
+                                                                            for (uint i = 0; i < RADIX_SORT_WG_SIZE / 32; i++)
+                                                                            {
+                                                                              const uint bits = bin_flags[binID].flags[i];
+                                                                              const uint full_count    = sycl::popcount(bits);
+                                                                              const uint partial_count = sycl::popcount(bits & (flags_bit - 1));
+                                                                              prefix += (i  < flags_bin) ? full_count : 0;
+                                                                              prefix += (i == flags_bin) ? partial_count : 0;
+                                                                              count += full_count;
+                                                                            }
+                                                                            output[binOffset + prefix] = key_value;
+                                                                            if (prefix == count - 1)
+                                                                              local_offset[binID] += count;                                                                          
+                                                                          }
+
+                                                                          item.barrier(sycl::access::fence_space::local_space);
+
+                                                                        }
+                                                                    
+                                                                      });
+                                                 
+                                                   });
+      return scatter_event;
+    }
+
+
+    template<typename sort_type>
+    sycl::event radix_sort_Nx8Bit(sycl::queue &gpu_queue, sort_type *input, sort_type *output, const uint items,  uint *global_histogram, const uint start_iteration, const uint end_iteration, sycl::event &initial, const uint RADIX_SORT_NUM_DSS=256)
+    {
+      sycl::event events[8]; // 8x8=64bit maximum
+      for (uint i=start_iteration;i<end_iteration;i++)
+      {
+        events[i] = radix_sort_iteration_type<sort_type>(gpu_queue,i == start_iteration ? initial : events[i-1],input,output,items,global_histogram,i,RADIX_SORT_NUM_DSS);
+        std::swap(input,output);
+      }
+      
+      return events[end_iteration-1];
+    }    
+    
+  };
+};
+
+#endif

kernels/rthwif/rtbuild/leaf.h

@@ -81,11 +81,13 @@ namespace embree
     PrimLeafDesc(uint32_t shaderIndex, uint32_t geomIndex, GeometryFlags gflags, uint32_t geomMask, Type type = TYPE_NONE)
     : shaderIndex(shaderIndex), geomMask(geomMask), geomIndex(geomIndex), type(type), geomFlags((uint32_t)gflags)
     {
+#if !defined(__SYCL_DEVICE_ONLY__)      
       if (shaderIndex > MAX_SHADER_INDEX)
         throw std::runtime_error("too large shader ID");
       
       if (geomIndex > MAX_GEOM_INDEX)
         throw std::runtime_error("too large geometry ID");
+#endif      
     }
 
     /* compares two PrimLeafDesc's for equality */
@@ -137,6 +139,7 @@ namespace embree
   public:
     uint32_t shaderIndex : 24;    // shader index used for shader record calculations
     uint32_t geomMask    : 8;     // geometry mask used for ray masking
+      
  
     uint32_t geomIndex      : 29; // the geometry index specifies the n'th geometry of the scene
     /*Type*/ uint32_t type  : 1;  // enable/disable culling for procedurals and instances
@@ -176,6 +179,9 @@ namespace embree
        * distance between them can be at most 0xFFFF as we use 16 bits
        * to encode that difference. */
       assert(primIndex0 <= primIndex1 && primIndex1 - primIndex0 < 0xFFFF);
+// #if !defined(__SYCL_DEVICE_ONLY__)
+//       print(std::cout,0);
+// #endif      
     }
 
     /* returns the i'th vertex */
@@ -277,22 +283,26 @@ namespace embree
     }
 
   public:
-    PrimLeafDesc leafDesc;  // the leaf header
-
-    uint32_t primIndex0;    // primitive index of first triangle
-    struct {
-      uint32_t primIndex1Delta : 16;  // delta encoded primitive index of second triangle
-      uint32_t j0              : 2;   // specifies first vertex of second triangle
-      uint32_t j1              : 2;   // specified second vertex of second triangle
-      uint32_t j2              : 2;   // specified third vertex of second triangle    
-      uint32_t last            : 1;   // true if the second triangle is the last triangle in a leaf list
-      uint32_t pad             : 9;   // unused bits
+    union {
+      struct {
+        PrimLeafDesc leafDesc;  // the leaf header
+        uint32_t primIndex0;    // primitive index of first triangle
+        struct {
+          uint32_t primIndex1Delta : 16;  // delta encoded primitive index of second triangle
+          uint32_t j0              : 2;   // specifies first vertex of second triangle
+          uint32_t j1              : 2;   // specified second vertex of second triangle
+          uint32_t j2              : 2;   // specified third vertex of second triangle    
+          uint32_t last            : 1;   // true if the second triangle is the last triangle in a leaf list
+          uint32_t pad             : 9;   // unused bits
+        };
+      };
+      uint header[4];
     };
     
     Vec3f v0;  // first vertex of first triangle
     Vec3f v1;  // second vertex of first triangle
     Vec3f v2;  // third vertex of first triangle
-    Vec3f v3;  // forth vertex used for second triangle
+    Vec3f v3;  // forth vertex used for second triangle    
   };
 
   static_assert(sizeof(QuadLeaf) == 64, "QuadLeaf must be 64 bytes large");

kernels/rthwif/rtbuild/qbvh6.cpp

@@ -6,7 +6,7 @@
 namespace embree
 {
   template<typename InternalNode>
-  void computeInternalNodeStatistics(BVHStatistics& stats, QBVH6::Node node, const BBox1f time_range, const float node_bounds_area, const float root_bounds_area)
+  void computeInternalNodeStatistics(BVHStatistics& stats, QBVH6::Node node, const BBox1f time_range, const float node_bounds_area, const float root_bounds_area, size_t depth)
   {
     InternalNode* inner = node.innerNode<InternalNode>();
 
@@ -16,7 +16,7 @@ namespace embree
       if (inner->valid(i))
       {
         size++;
-        computeStatistics(stats, inner->child(i), time_range, area(inner->bounds(i)), root_bounds_area, InternalNode::NUM_CHILDREN);
+        computeStatistics(stats, inner->child(i), time_range, area(inner->bounds(i)), root_bounds_area, InternalNode::NUM_CHILDREN, depth+1);
       }
     }
 
@@ -26,9 +26,10 @@ namespace embree
     stats.internalNode.numChildrenTotal += InternalNode::NUM_CHILDREN;
     stats.internalNode.nodeSAH += time_range.size() * node_bounds_area / root_bounds_area;
     stats.internalNode.numBytes += sizeof(InternalNode);
+    stats.depth = max(stats.depth, depth);
   }
 
-  void computeStatistics(BVHStatistics& stats, QBVH6::Node node, const BBox1f time_range, const float node_bounds_area, const float root_bounds_area, uint32_t numChildren)
+  void computeStatistics(BVHStatistics& stats, QBVH6::Node node, const BBox1f time_range, const float node_bounds_area, const float root_bounds_area, uint32_t numChildren, size_t depth)
   {
     switch (node.type)
     {
@@ -108,7 +109,7 @@ namespace embree
     }
     case NODE_TYPE_INTERNAL:
     {
-      computeInternalNodeStatistics<QBVH6::InternalNode6>(stats, node, time_range, node_bounds_area, root_bounds_area);
+      computeInternalNodeStatistics<QBVH6::InternalNode6>(stats, node, time_range, node_bounds_area, root_bounds_area, depth);
       break;
     }
     default:
@@ -119,14 +120,15 @@ namespace embree
   BVHStatistics QBVH6::computeStatistics() const
   {
     BVHStatistics stats;
-    if (empty()) return stats;
-    embree::computeStatistics(stats,root(),BBox1f(0,1),area(bounds),area(bounds),6);
+    if (empty()) return stats;    
+    embree::computeStatistics(stats,root(),BBox1f(0,1),area(bounds),area(bounds),6,1);
     return stats;
   }
 
   template<typename QInternalNode>
   void QBVH6::printInternalNodeStatistics(std::ostream& cout, QBVH6::Node node, uint32_t depth, uint32_t numChildren)
   {
+#if  !defined(__SYCL_DEVICE_ONLY__)    
     QInternalNode* inner = node.innerNode<QInternalNode>();
     inner->print(cout, depth, false);
     std::cout << std::endl;
@@ -138,10 +140,14 @@ namespace embree
     }
 
     cout << tab(depth) << "}" << std::endl;
+#endif    
   }
 
   void QBVH6::print( std::ostream& cout, QBVH6::Node node, uint32_t depth, uint32_t numChildren)
   {
+#if !defined(__SYCL_DEVICE_ONLY__)    
+    if (depth == 1) return;
+    
     switch (node.type)
     {
     case NODE_TYPE_INSTANCE: {
@@ -214,6 +220,7 @@ namespace embree
       std::cout << "{ INVALID_NODE }" << std::endl;
       //assert(false);
     }
+#endif    
   }
 
   unsigned* getBackPointersData(const QBVH6* base) { // FIXME: should be member function
@@ -234,7 +241,7 @@ namespace embree
 
   void QBVH6::print ( std::ostream& cout ) const
   {
-    
+#if  !defined(__SYCL_DEVICE_ONLY__)    
     cout << "QBVH @ "<< this <<" header: {\n";
     cout << "  rootNodeOffset = " << rootNodeOffset << std::endl;
     cout << "  bounds = " << bounds << std::endl;
@@ -261,6 +268,7 @@ namespace embree
       }
       cout << "}\n";
     }
+#endif    
   }
   
   template<typename QInternalNode>

kernels/rthwif/rtbuild/qbvh6_builder_sah.h

@@ -191,10 +191,12 @@ namespace embree
         __forceinline size_t end  () const { return prims.end(); }
         __forceinline size_t size () const { return prims.size(); }
         __forceinline bool   equalType() const { return type != UNKNOWN; }
-        
+
+#if !defined(__SYCL_DEVICE_ONLY__)        
         friend inline std::ostream& operator<<(std::ostream& cout, const BuildRecord& r) {
           return cout << "BuildRecord { depth = " << r.depth << ", pinfo = " << r.prims << ", type = " << r.type << " }";
         }
+#endif        
         
       public:
         size_t depth;        //!< Depth of the root of this subtree.
@@ -1011,6 +1013,9 @@ namespace embree
 
         ReductionTy build(uint32_t numGeometries, PrimInfo& pinfo_o, char* root)
         {
+          //PING;
+          //PRINT(numGeometries);
+          
           double t1 = verbose ? getSeconds() : 0.0;
 
           /* quadify all triangles */
@@ -1038,7 +1043,7 @@ namespace embree
           }, [](const PrimInfo& a, const PrimInfo& b) -> PrimInfo { return PrimInfo::merge(a,b); });
 
           double t3 = verbose ? getSeconds() : 0.0;
-          if (verbose) std::cout << "primrefgen   : " << std::setw(10) << (t3-t2)*1000.0 << "ms, " << std::setw(10) << 1E-6*double(numPrimitives)/(t3-t2) << " Mprims/s" << std::endl;
+          if (verbose) std::cout << "primrefgen   : " << numPrimitives << " " << std::setw(10) << (t3-t2)*1000.0 << "ms, " << std::setw(10) << 1E-6*double(numPrimitives)/(t3-t2) << " Mprims/s" << std::endl;
           
           /* if we need to filter out geometry, run again */
           if (pinfo.size() != numPrimitives)
@@ -1061,6 +1066,8 @@ namespace embree
           /* perform pre-splitting */
           if ((build_quality == RTHWIF_BUILD_QUALITY_HIGH) &&  numPrimitives)
           {
+            PRINT("PRE-SPLITTING DISABLED");
+            
             auto splitter = [this] (const PrimRef& prim, const size_t dim, const float pos, PrimRef& left_o, PrimRef& right_o) {
               splitTriangleOrQuad(prim,dim,pos,left_o,right_o);
             };
@@ -1168,6 +1175,14 @@ namespace embree
           qbvh->numTimeSegments = 1; 
           qbvh->dispatchGlobalsPtr = (uint64_t) dispatchGlobalsPtr;
 
+
+#if 0
+          BVHStatistics stats2 = qbvh->computeStatistics();
+          stats2.print(std::cout);
+          stats2.print_raw(std::cout);
+          qbvh->print();
+#endif
+          
 #if 0
           BVHStatistics stats = qbvh->computeStatistics();
           stats.print(std::cout);
@@ -1179,6 +1194,41 @@ namespace embree
           for (size_t i=0; i<bytes; i++) {
             if (i % 32 == 0) std::cout << std::endl << "  ";
             std::cout << "0x" << std::hex << std::setw(2) << std::setfill('0') << (unsigned)((unsigned char*)accel)[i] << ", ";
+=======
+              /* fill QBVH6 header */
+              allocator.clear();
+              QBVH6* qbvh = new (accel.data()) QBVH6(QBVH6::SizeEstimate());
+              qbvh->numPrims = numPrimitives;
+              uint64_t rootNodeOffset = QBVH6::Node((char*)(r0.node - (char*)qbvh), r0.type, r0.primRange.cur_prim);
+              assert(rootNodeOffset == QBVH6::rootNodeOffset);
+              qbvh->bounds = bounds;
+              qbvh->numTimeSegments = max(1u,maxTimeSegments);
+              qbvh->dispatchGlobalsPtr = (uint64_t) dispatchGlobalsPtr;
+              break;
+            }
+            catch (std::bad_alloc&)
+            {
+              if (verbose) {
+                allocator.cleanup();
+                FastAllocator::AllStatistics stats1(&allocator);
+                stats1.print(numPrimitives);
+              }
+              
+              if (bytes >= worst_case_bytes) {
+                throw std::runtime_error("BVH build failed");
+                return empty;
+              }
+            }
+          }
+          
+          /* print BVH statistics in verbose mode */
+          if (verbose)
+          {
+            QBVH6* qbvh = (QBVH6*) accel.data();
+            BVHStatistics stats = qbvh->computeStatistics();
+            stats.print(std::cout);
+            stats.print_raw(std::cout);
+            qbvh->print();
           }
           std::cout << std::endl << "};" << std::endl;*/
 #endif

kernels/rthwif/rtbuild/qnode.h

@@ -183,12 +183,12 @@ namespace embree
 
     NodeType nodeType;    // the type of the node    
     uint8_t pad;          // unused byte
-
+        
     int8_t exp_x;          // 2^exp_x is the size of the grid in x dimension
     int8_t exp_y;          // 2^exp_y is the size of the grid in y dimension
     int8_t exp_z;          // 2^exp_z is the size of the grid in z dimension
     uint8_t nodeMask;      // mask used for ray filtering
-
+    
     struct ChildData
     {
       uint8_t blockIncr : 2; // size of child in 64 byte blocks
@@ -196,12 +196,14 @@ namespace embree
       uint8_t pad : 2; // unused bits
     } childData[NUM_CHILDREN];
 
-    uint8_t lower_x[NUM_CHILDREN];  // the quantized lower bounds in x-dimension
-    uint8_t upper_x[NUM_CHILDREN];  // the quantized upper bounds in x-dimension
-    uint8_t lower_y[NUM_CHILDREN];  // the quantized lower bounds in y-dimension
-    uint8_t upper_y[NUM_CHILDREN];  // the quantized upper bounds in y-dimension
-    uint8_t lower_z[NUM_CHILDREN];  // the quantized lower bounds in z-dimension
-    uint8_t upper_z[NUM_CHILDREN];  // the quantized upper bounds in z-dimension
+    struct {
+      uint8_t lower_x[NUM_CHILDREN];  // the quantized lower bounds in x-dimension
+      uint8_t upper_x[NUM_CHILDREN];  // the quantized upper bounds in x-dimension
+      uint8_t lower_y[NUM_CHILDREN];  // the quantized lower bounds in y-dimension
+      uint8_t upper_y[NUM_CHILDREN];  // the quantized upper bounds in y-dimension
+      uint8_t lower_z[NUM_CHILDREN];  // the quantized lower bounds in z-dimension
+      uint8_t upper_z[NUM_CHILDREN];  // the quantized upper bounds in z-dimension
+    };
   };
 
   static_assert(sizeof(InternalNode6Data) == 64, "InternalNode6Data must be 64 bytes large");
@@ -370,6 +372,7 @@ namespace embree
     int64_t getChildOffset(uint32_t childID) const
     {
       int64_t ofs = this->childOffset;
+      if (ofs == 0) FATAL("OFFSET == 0");
       for (uint32_t j = 0; j < childID; j++)
         ofs += this->childData[j].blockIncr;
       return 64 * ofs;
@@ -485,7 +488,7 @@ namespace embree
       dst->setChildOffset((char*)this + getChildOffset(0));
     }
     
-#if !defined(__RTRT_GSIM)
+#if !defined(__RTRT_GSIM) &&  !defined(__SYCL_DEVICE_ONLY__)
     
     /* output of internal node */
     void print(std::ostream& cout, uint32_t depth, bool close) const
@@ -495,7 +498,6 @@ namespace embree
       cout << tab(depth) << "  childOffset = " << 64 * int64_t(this->childOffset) << std::endl;
       cout << tab(depth) << "  nodeType = " << NodeType(this->nodeType) << std::endl;
       cout << tab(depth) << "  nodeMask = " << std::bitset<8>(this->nodeMask) << std::endl;
-      
       for (uint32_t i = 0; i < NUM_CHILDREN; i++)
       {
         cout << tab(depth) << "  child" << i << " = { ";

kernels/rthwif/rtbuild/quadifier.h

@@ -5,6 +5,8 @@
 
 #include "../../common/default.h"
 
+#define WINDOW_SIZE 32
+
 namespace embree
 {
   enum QuadifierType : uint16_t
@@ -78,7 +80,7 @@ namespace embree
   }
 
   template<typename GetTriangleFunc>
-  __forceinline void merge_triangle_window( uint32_t geomID, static_deque<uint32_t,32>& triangleWindow, QuadifierType* quads_o, const GetTriangleFunc& getTriangle )
+  __forceinline void merge_triangle_window( uint32_t geomID, static_deque<uint32_t,WINDOW_SIZE>& triangleWindow, QuadifierType* quads_o, const GetTriangleFunc& getTriangle )
   {
     uint32_t primID0 = triangleWindow.pop_front();
     
@@ -118,7 +120,7 @@ namespace embree
   template<typename GetTriangleFunc>
   inline size_t pair_triangles( uint32_t geomID, QuadifierType* quads_o, uint32_t primID0, uint32_t primID1, const GetTriangleFunc& getTriangle ) 
   {
-    static_deque<uint32_t, 32> triangleWindow;
+    static_deque<uint32_t, WINDOW_SIZE> triangleWindow;
 
     size_t numTrianglePairs = 0;
     for (uint32_t primID=primID0; primID<primID1; primID++)
@@ -129,13 +131,11 @@ namespace embree
         merge_triangle_window(geomID, triangleWindow,quads_o,getTriangle);
         numTrianglePairs++;
       }
-    }
-    
+    }    
     while (triangleWindow.size()) {
       merge_triangle_window(geomID, triangleWindow,quads_o,getTriangle);
       numTrianglePairs++;
     }
-
     return numTrianglePairs;
   }
 }

kernels/rthwif/rtbuild/rthwif_builder.h

@@ -465,3 +465,10 @@ RTHWIF_API uint32_t rthwifGetParallelOperationMaxConcurrency( RTHWIF_PARALLEL_OP
  */
 
 RTHWIF_API RTHWIF_ERROR rthwifJoinParallelOperation( RTHWIF_PARALLEL_OPERATION parallelOperation );
+
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+RTHWIF_API RTHWIF_ERROR rthwifGetAccelSizeGPU(const RTHWIF_BUILD_ACCEL_ARGS& args_i, RTHWIF_ACCEL_SIZE& size_o, void *sycl_queue, uint verbose_level=0); 
+RTHWIF_API RTHWIF_ERROR rthwifBuildAccelGPU(const RTHWIF_BUILD_ACCEL_ARGS& args, void *sycl_queue, uint verbose_level=0);
+RTHWIF_API RTHWIF_ERROR rthwifPrefetchAccelGPU(const RTHWIF_BUILD_ACCEL_ARGS& args, void *sycl_queue, uint verbose_level=0);
+#endif 

kernels/rthwif/rtbuild/rthwif_embree_builder_ploc.cpp

@@ -0,0 +1,978 @@
+#include "../../sycl/rthwif_embree.h"
+#include "../../sycl/rthwif_embree_builder.h"
+#include "../rttrace/rthwif_internal.h"
+#include "rthwif_embree_builder_ploc.h"
+#include "qbvh6.h"
+#include "../common/algorithms/parallel_reduce.h"
+
+// === less than threshold, a single workgroup is used to perform all PLOC iterations in a single kernel launch ===
+#define SINGLE_WG_SWITCH_THRESHOLD            4*1024
+
+// === less than threshold, 40bits morton code + 24bits index are used, otherwise 64bit morton code + 32bit index ===
+#define FAST_MC_NUM_PRIMS_THRESHOLD           1024*1024
+
+// === max number of primitives fitting in 24bits ===
+#define FAST_MC_MAX_NUM_PRIMS                 ((uint)1<<24)
+
+// === less than threshold, a single workgroup is used for all radix sort iterations ===
+#define SMALL_SORT_THRESHOLD                  1024*4
+
+// === maximum number of workgroups with 1024 elements, DG2/PVC perform best with 64 ===
+#define MAX_LARGE_WGS                         256
+
+// === rebalance if BVH2 subtrees are degenerated ===
+#define BVH2_REBALANCE                        1
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)      
+
+namespace embree
+{
+  using namespace embree::isa;
+
+  __forceinline size_t estimateSizeInternalNodes(const size_t numQuads, const size_t numInstances, const size_t numProcedurals, const bool conservative)
+  {
+    const size_t N = numQuads + numInstances + numProcedurals; 
+    // === conservative estimate ===
+    size_t numFatLeaves = 0;
+    if (conservative)
+      numFatLeaves = ceilf( (float)N/2 ) + ceilf( (float)numInstances/2 ); // FIXME : better upper bound for instance case
+    else
+      numFatLeaves = ceilf( (float)N/3 ) + ceilf( (float)numInstances/2 ); // FIXME : better upper bound for instance case        
+    const size_t numInnerNodes = ceilf( (float)numFatLeaves/4 ); 
+    return gpu::alignTo(std::max( ((numFatLeaves + numInnerNodes) * 64) , N * 16),64);
+  }
+
+  __forceinline size_t estimateSizeLeafNodes(const size_t numQuads, const size_t numInstances, const size_t numProcedurals)
+  {
+    return (numQuads + numProcedurals + 2 * numInstances) * 64;  
+  }
+
+  __forceinline size_t estimateAccelBufferSize(const size_t numQuads, const size_t numInstances, const size_t numProcedurals, const bool conservative)
+  {
+    const size_t header              = 128;
+    const size_t node_size           = estimateSizeInternalNodes(numQuads,numInstances,numProcedurals,conservative);
+    const size_t leaf_size           = estimateSizeLeafNodes(numQuads,numInstances,numProcedurals); 
+    const size_t totalSize           = header + node_size + leaf_size;
+    return totalSize;
+  }
+
+  __forceinline size_t estimateScratchBufferSize(const size_t numPrimitives)
+  {
+    // === sizeof(size_t)*MAX_LARGE_WGS for prefix sums across large work groups ===
+    return sizeof(PLOCGlobals) + sizeof(size_t)*MAX_LARGE_WGS + numPrimitives * sizeof(LeafGenerationData);
+  }
+
+  uint getBVH2Depth(BVH2Ploc *bvh2, uint index, const uint numPrimitives)
+  {
+    if (BVH2Ploc::getIndex(index) < numPrimitives) //isLeaf 
+      return 1;
+    else
+      return 1 + std::max(getBVH2Depth(bvh2,bvh2[index].leftIndex(),numPrimitives),getBVH2Depth(bvh2,bvh2[index].rightIndex(),numPrimitives));
+  }
+
+  uint getNumLeaves(BVH2Ploc *bvh2, uint index, const uint numPrimitives)
+  {
+    if (BVH2Ploc::getIndex(index) < numPrimitives) //isLeaf 
+      return 1;
+    else
+      return getNumLeaves(bvh2,bvh2[index].leftIndex(),numPrimitives) + getNumLeaves(bvh2,bvh2[index].rightIndex(),numPrimitives);
+  }
+
+
+  uint getNumFatLeaves(BVH2Ploc *bvh2, uint index, const uint numPrimitives)
+  {
+    if (BVH2Ploc::isFatLeaf(index,numPrimitives)) //isLeaf 
+      return 1;
+    else
+      return getNumFatLeaves(bvh2,bvh2[index].leftIndex(),numPrimitives) + getNumFatLeaves(bvh2,bvh2[index].rightIndex(),numPrimitives);
+  }
+  
+
+
+  void printBVH2Path(BVH2Ploc *bvh2, uint index, const uint numPrimitives)
+  {
+    if (BVH2Ploc::getIndex(index) < numPrimitives) //isLeaf
+    {
+      PRINT2(index,"LEAF");
+    }
+    else
+    {
+      const uint depth = getBVH2Depth(bvh2,index,numPrimitives);
+      const uint leftIndex = bvh2[index].leftIndex();
+      const uint rightIndex = bvh2[index].rightIndex();
+      const bool isFatLeafLeft = BVH2Ploc::isFatLeaf(bvh2[index].left,numPrimitives);
+      const bool isFatLeafRight = BVH2Ploc::isFatLeaf(bvh2[index].right,numPrimitives);
+      const uint numLeavesLeft = getNumLeaves(bvh2,leftIndex,numPrimitives);
+      const uint numLeavesRight = getNumLeaves(bvh2,rightIndex,numPrimitives);      
+      PRINT6(index,depth,leftIndex,rightIndex,isFatLeafLeft,isFatLeafRight);
+      PRINT2(leftIndex,numLeavesLeft);
+      PRINT2(rightIndex,numLeavesRight);
+      
+      if (!isFatLeafLeft)
+        printBVH2Path(bvh2, leftIndex,numPrimitives);
+      if (!isFatLeafRight)      
+        printBVH2Path(bvh2,rightIndex,numPrimitives);      
+    }
+  }  
+  
+  
+  void checkBVH2PlocHW(BVH2Ploc *bvh2, uint index,uint &nodes,uint &leaves,float &nodeSAH, float &leafSAH, uint &maxDepth,const uint numPrimitives, const uint bvh2_max_allocations, const uint depth)
+  {
+    if (bvh2[index].bounds.empty()) {
+      PRINT2(index,bvh2[index]);
+      FATAL("invalid bounds in BVH2");
+    }
+    if (!bvh2[index].bounds.checkNumericalBounds())
+    {
+      PRINT2(index,bvh2[index]);      
+      FATAL("Numerical Bounds in BVH2");
+    }
+
+    if (BVH2Ploc::getIndex(index) < numPrimitives) //isLeaf 
+    {
+      leaves++;
+      leafSAH +=  bvh2[index].bounds.area();
+      assert(bvh2[index].getLeafIndex() < numPrimitives);
+    }
+    else
+    {
+      maxDepth = max(maxDepth,depth+1);
+      uint indices[BVH_BRANCHING_FACTOR];
+      const uint numChildren = openBVH2MaxAreaSortChildren(BVH2Ploc::getIndex(index),indices,bvh2,numPrimitives);
+      for (uint i=0;i<numChildren;i++)
+        if (BVH2Ploc::getIndex(indices[i]) > bvh2_max_allocations)
+          FATAL("OPENING ERROR");
+
+      nodes++;              
+      nodeSAH += bvh2[index].bounds.area();
+      
+      if (!bvh2[index].bounds.encloses( bvh2[ bvh2[index].leftIndex() ].bounds )) PRINT2("ENCLOSING ERROR LEFT",index);
+      checkBVH2PlocHW(bvh2,bvh2[index].leftIndex(),nodes,leaves,nodeSAH,leafSAH,maxDepth,numPrimitives,bvh2_max_allocations,depth+1);
+
+      if (!bvh2[index].bounds.encloses( bvh2[ bvh2[index].rightIndex() ].bounds )) PRINT2("ENCLOSING ERROR RIGHT",index);
+      checkBVH2PlocHW(bvh2,bvh2[index].rightIndex(),nodes,leaves,nodeSAH,leafSAH,maxDepth,numPrimitives,bvh2_max_allocations,depth+1);
+    }
+  }
+
+  struct BuildTimer {
+    enum Type {
+      PRE_PROCESS  = 0,
+      BUILD        = 1,
+      POST_PROCESS = 2,
+      ALLOCATION   = 3,            
+      TOTAL        = 4
+    };
+
+    double host_timers[TOTAL];
+    double device_timers[TOTAL];
+    
+    double t0,t1;
+
+    inline void reset()
+    {
+      for (uint i=0;i<TOTAL;i++)
+      {
+        host_timers[i] = 0.0;
+        device_timers[i] = 0.0;        
+      }        
+    }
+    
+    inline void start(const Type type)
+    {
+      t0 = getSeconds();
+    }
+
+    inline void stop(const Type type)
+    {
+      t1 = getSeconds();
+      host_timers[(int)type] += 1000.0*(t1-t0);
+    }
+
+    inline void add_to_device_timer(const Type type, double t)
+    {
+      device_timers[(int)type] += t;
+    }
+    
+    inline float get_accum_device_timer(const Type type) { return device_timers[(int)type]; }
+    inline float get_accum_host_timer  (const Type type) { return host_timers[(int)type]; }    
+    inline float get_host_timer() { return 1000.0*(t1-t0); }
+
+    inline float get_total_device_time()
+    {
+      double sum = 0.0;
+      for (uint i=0;i<ALLOCATION;i++) sum += device_timers[i];
+      return sum;
+    }
+
+    inline float get_total_host_time()
+    {
+      double sum = 0.0;
+      for (uint i=0;i<ALLOCATION;i++) sum += host_timers[i];
+      return sum;
+    }
+    
+    
+  };
+
+  __forceinline uint32_t getNumPrimitives(const RTHWIF_GEOMETRY_DESC* geom)
+  {
+    switch (geom->geometryType) {
+    case RTHWIF_GEOMETRY_TYPE_TRIANGLES  : return ((RTHWIF_GEOMETRY_TRIANGLES_DESC*)  geom)->triangleCount;
+    case RTHWIF_GEOMETRY_TYPE_AABBS_FPTR : return ((RTHWIF_GEOMETRY_AABBS_FPTR_DESC*) geom)->primCount;
+    case RTHWIF_GEOMETRY_TYPE_QUADS      : return ((RTHWIF_GEOMETRY_QUADS_DESC*)      geom)->quadCount;
+    case RTHWIF_GEOMETRY_TYPE_INSTANCE   : return 1;
+    default                              : return 0;
+    };
+  }
+
+  /* fill all arg members that app did not know of yet */
+  __forceinline RTHWIF_BUILD_ACCEL_ARGS rthwifPrepareBuildAccelArgs(const RTHWIF_BUILD_ACCEL_ARGS& args_i)
+  {
+    RTHWIF_BUILD_ACCEL_ARGS args;
+    memset(&args,0,sizeof(RTHWIF_BUILD_ACCEL_ARGS));
+    memcpy(&args,&args_i,std::min(sizeof(RTHWIF_BUILD_ACCEL_ARGS),args_i.structBytes));
+    args.structBytes = sizeof(RTHWIF_BUILD_ACCEL_ARGS);
+    return args;
+  } 
+  
+  __forceinline PrimitiveCounts countPrimitives(const RTHWIF_GEOMETRY_DESC** geometries, const uint numGeometries)
+  {
+    auto reduce = [&](const range<size_t>& r) -> PrimitiveCounts
+                  {
+                    PrimitiveCounts counts;
+                    for (size_t geomID = r.begin(); geomID < r.end(); geomID++)
+                    {
+                      const RTHWIF_GEOMETRY_DESC* geom = geometries[geomID];
+                      if (geom == nullptr) continue;    
+                      switch (geom->geometryType) {
+                      case RTHWIF_GEOMETRY_TYPE_TRIANGLES  :
+                      {
+                        counts.numTriangles   += ((RTHWIF_GEOMETRY_TRIANGLES_DESC*)  geom)->triangleCount;
+                        counts.numQuadBlocks  += (((RTHWIF_GEOMETRY_TRIANGLES_DESC *)geom)->triangleCount+TRIANGLE_QUAD_BLOCK_SIZE-1)/TRIANGLE_QUAD_BLOCK_SIZE;
+                        break;
+                      }
+                      case RTHWIF_GEOMETRY_TYPE_QUADS      :
+                      {
+                        counts.numQuads       += ((RTHWIF_GEOMETRY_QUADS_DESC*)  geom)->quadCount;
+                        counts.numQuadBlocks  += (((RTHWIF_GEOMETRY_QUADS_DESC *)geom)->quadCount+TRIANGLE_QUAD_BLOCK_SIZE-1)/TRIANGLE_QUAD_BLOCK_SIZE;                        
+                        break;
+                      }
+                      case RTHWIF_GEOMETRY_TYPE_AABBS_FPTR : counts.numProcedurals += ((RTHWIF_GEOMETRY_AABBS_FPTR_DESC*) geom)->primCount; break;
+                      case RTHWIF_GEOMETRY_TYPE_INSTANCE   : counts.numInstances   += 1; break;
+                      default: assert(false); break;        
+                      };                    
+                    };
+                    return counts;
+                  };
+
+    const uint COUNT_BLOCK_SIZE = 256;
+    const uint COUNT_PARALLEL_THRESHOLD = 256;
+    
+    const PrimitiveCounts primCounts = parallel_reduce((uint)0, numGeometries, COUNT_BLOCK_SIZE, COUNT_PARALLEL_THRESHOLD, PrimitiveCounts(), reduce,
+                                                       [&](const PrimitiveCounts& b0, const PrimitiveCounts& b1) -> PrimitiveCounts { return b0 + b1; });
+    return primCounts;
+  }
+
+  RTHWIF_ERROR createEmptyBVH(const RTHWIF_BUILD_ACCEL_ARGS& args, sycl::queue  &gpu_queue)
+  {
+    sycl::event queue_event =  gpu_queue.submit([&](sycl::handler &cgh) {
+                                                  cgh.single_task([=]() {
+                                                                    QBVH6* qbvh  = (QBVH6*)args.accelBuffer;       
+                                                                    qbvh->bounds = BBox3f(empty);
+                                                                    qbvh->numPrims       = 0;                                                                        
+                                                                    qbvh->nodeDataStart  = 2;
+                                                                    qbvh->nodeDataCur    = 3;
+                                                                    qbvh->leafDataStart  = 3;
+                                                                    qbvh->leafDataCur    = 3;        
+                                                                    new (qbvh->nodePtr(2)) QBVH6::InternalNode6(NODE_TYPE_INTERNAL);
+                                                                  });
+                                                });
+    gpu::waitOnEventAndCatchException(queue_event);    
+    if (args.accelBufferBytesOut) *args.accelBufferBytesOut = 128+64;
+    if (args.boundsOut)           { BBox3f geometryBounds (empty); *args.boundsOut = *(RTHWIF_AABB*)&geometryBounds; };
+    return RTHWIF_ERROR_NONE;    
+  }
+// =================================================================================================================================================================================
+// =================================================================================================================================================================================
+// =================================================================================================================================================================================
+
+  RTHWIF_API RTHWIF_ERROR rthwifGetAccelSizeGPU(const RTHWIF_BUILD_ACCEL_ARGS& args_i, RTHWIF_ACCEL_SIZE& size_o, void *sycl_queue, uint verbose_level=0)
+  {
+    double time0 = getSeconds();
+    
+    RTHWIF_BUILD_ACCEL_ARGS args = rthwifPrepareBuildAccelArgs(args_i);
+    const RTHWIF_GEOMETRY_DESC** geometries = args.geometries;
+    const uint numGeometries = args.numGeometries;
+    sycl::queue  &gpu_queue  = *(sycl::queue*)sycl_queue;
+
+    // =============================================================================    
+    // === GPU-based primitive count estimation including triangle quadification ===
+    // =============================================================================
+    
+    const PrimitiveCounts primCounts = getEstimatedPrimitiveCounts(gpu_queue,geometries,numGeometries,verbose_level >= 2);            
+
+    const uint numTriangles       = primCounts.numTriangles; // === original number of triangles ===
+    const uint numMergedTrisQuads = primCounts.numMergedTrisQuads;
+    const uint numQuads           = primCounts.numQuads;
+    const uint numProcedurals     = primCounts.numProcedurals;
+    const uint numInstances       = primCounts.numInstances;
+  
+    const uint numPrimitives = numMergedTrisQuads + numProcedurals + numInstances;
+
+    // =============================================    
+    // === allocation for empty scene is default ===
+    // =============================================
+    
+    size_t expectedBytes = 3*64; 
+    size_t worstCaseBytes = 4*64;
+
+    if (numPrimitives)
+    {    
+      expectedBytes  = estimateAccelBufferSize(     numMergedTrisQuads, numInstances, numProcedurals, false);
+      worstCaseBytes = estimateAccelBufferSize(numQuads + numTriangles, numInstances, numProcedurals, true);    
+    }
+
+    // ===============================================    
+    // === estimate accel and scratch buffer sizes ===
+    // ===============================================
+    
+    const size_t scratchBytes = estimateScratchBufferSize(std::max(numPrimitives,numGeometries));
+
+    if (verbose_level >= 2)
+    {
+      PRINT5(numMergedTrisQuads,numTriangles,numQuads,numProcedurals,numInstances);      
+      PRINT3(expectedBytes,worstCaseBytes,scratchBytes);
+    }
+        
+    /* return size to user */
+    RTHWIF_ACCEL_SIZE size;
+    memset(&size,0,sizeof(RTHWIF_ACCEL_SIZE));
+    size.accelBufferExpectedBytes = expectedBytes;
+    size.accelBufferWorstCaseBytes = worstCaseBytes;
+    size.scratchBufferBytes = scratchBytes;
+    size_t bytes_o = size_o.structBytes;
+    memset(&size_o,0,bytes_o);
+    memcpy(&size_o,&size,bytes_o);
+    size_o.structBytes = bytes_o;
+
+    double time1 = getSeconds();
+    if (verbose_level >= 1)
+      std::cout << "rthwifGetAccelSizeGPU time = " << (float)(time1-time0)*1000.0f << " ms" << std::endl;
+    
+    return RTHWIF_ERROR_NONE;
+  }
+
+  RTHWIF_API RTHWIF_ERROR rthwifPrefetchAccelGPU(const RTHWIF_BUILD_ACCEL_ARGS& args, void *sycl_queue, uint verbose_level=0)
+  {
+    double time0 = getSeconds();
+    
+    sycl::queue  &gpu_queue  = *(sycl::queue*)sycl_queue;
+
+#if 0
+    const RTHWIF_GEOMETRY_DESC** geometries = args.geometries;
+    const uint numGeometries                = args.numGeometries;  
+    
+    // ===================================    
+    // === prefetch builder scene data ===
+    // ===================================
+    
+    for (size_t geomID = 0; geomID < numGeometries; geomID++)
+    {
+      const RTHWIF_GEOMETRY_DESC* geom = geometries[geomID];
+      if (geom == nullptr) continue;    
+      switch (geom->geometryType) {
+      case RTHWIF_GEOMETRY_TYPE_TRIANGLES  :
+      {
+        RTHWIF_GEOMETRY_TRIANGLES_DESC *t = (RTHWIF_GEOMETRY_TRIANGLES_DESC*)geom;
+        if (t->vertexBuffer)   gpu_queue.prefetch(t->vertexBuffer,t->vertexCount*t->vertexStride);
+        if (t->triangleBuffer) gpu_queue.prefetch(t->triangleBuffer,t->triangleCount*t->triangleStride);      
+        gpu_queue.prefetch(t,sizeof(RTHWIF_GEOMETRY_TRIANGLES_DESC));
+        break;
+      }
+      case RTHWIF_GEOMETRY_TYPE_QUADS      :
+      {
+        RTHWIF_GEOMETRY_QUADS_DESC *q = (RTHWIF_GEOMETRY_QUADS_DESC*)geom;
+        if (q->vertexBuffer) gpu_queue.prefetch(q->vertexBuffer,q->vertexCount*q->vertexStride);
+        if (q->quadBuffer) gpu_queue.prefetch(q->quadBuffer,q->quadCount*q->quadStride);      
+        gpu_queue.prefetch(q,sizeof(RTHWIF_GEOMETRY_QUADS_DESC));      
+        break;
+      }
+      case RTHWIF_GEOMETRY_TYPE_AABBS_FPTR :
+      {
+        RTHWIF_GEOMETRY_AABBS_FPTR_DESC *a = (RTHWIF_GEOMETRY_AABBS_FPTR_DESC*)geom;
+        gpu_queue.prefetch(a,sizeof(RTHWIF_GEOMETRY_AABBS_FPTR_DESC));
+        break;
+      }
+      case RTHWIF_GEOMETRY_TYPE_INSTANCE   :
+      {
+        RTHWIF_GEOMETRY_INSTANCE_DESC *i = (RTHWIF_GEOMETRY_INSTANCE_DESC*)geom;
+        gpu_queue.prefetch(i->bounds,sizeof(RTHWIF_AABB));      
+        gpu_queue.prefetch(i,sizeof(RTHWIF_GEOMETRY_INSTANCE_DESC));
+        break;
+      }
+      default: assert(false); break;        
+      };                    
+    };
+
+    if (geometries) gpu_queue.prefetch(geometries,sizeof(RTHWIF_GEOMETRY_DESC*)*numGeometries);
+    if (args.accelBuffer)   gpu_queue.prefetch(args.accelBuffer  ,args.accelBufferBytes);
+    if (args.scratchBuffer) gpu_queue.prefetch(args.scratchBuffer,args.scratchBufferBytes);  
+#endif  
+    // ======================================================    
+    // === DUMMY KERNEL TO TRIGGER REMAINING USM TRANSFER ===
+    // ======================================================
+    
+    sycl::event queue_event =  gpu_queue.submit([&](sycl::handler &cgh) { cgh.single_task([=]() {}); });
+    gpu::waitOnEventAndCatchException(queue_event);
+
+    double time1 = getSeconds();
+    if (verbose_level >= 1)
+      std::cout << "rthwifPrefetchAccelGPU time = " << (float)(time1-time0)*1000.0f << " ms" << std::endl;
+    
+    return RTHWIF_ERROR_NONE;      
+  }
+
+  RTHWIF_API RTHWIF_ERROR rthwifBuildAccelGPU(const RTHWIF_BUILD_ACCEL_ARGS& args, void *sycl_queue, uint verbose_level=0)
+  {
+    BuildTimer timer;
+    timer.reset();
+
+    timer.start(BuildTimer::PRE_PROCESS);      
+    
+    // ================================    
+    // === GPU device/queue/context ===
+    // ================================
+  
+    sycl::queue  &gpu_queue  = *(sycl::queue*)sycl_queue;
+    const bool verbose1 = verbose_level >= 1;    
+    const bool verbose2 = verbose_level >= 2;
+    const uint gpu_maxComputeUnits  = gpu_queue.get_device().get_info<sycl::info::device::max_compute_units>();
+    const uint MAX_WGS = gpu_maxComputeUnits / 8;
+    
+    uint *host_device_tasks = (uint*)sycl::aligned_alloc(64,HOST_DEVICE_COMM_BUFFER_SIZE,gpu_queue.get_device(),gpu_queue.get_context(),sycl::usm::alloc::host);
+
+    
+    if (unlikely(verbose2))
+    {
+      const uint gpu_maxWorkGroupSize = gpu_queue.get_device().get_info<sycl::info::device::max_work_group_size>();
+      const uint gpu_maxLocalMemory   = gpu_queue.get_device().get_info<sycl::info::device::local_mem_size>();    
+      PRINT("PLOC++ GPU BVH BUILDER");            
+      PRINT( gpu_queue.get_device().get_info<sycl::info::device::global_mem_size>() );
+      PRINT(gpu_maxWorkGroupSize);
+      PRINT(gpu_maxComputeUnits);
+      PRINT(gpu_maxLocalMemory);
+    }
+
+    // =============================    
+    // === setup scratch pointer ===
+    // =============================
+    
+    PLOCGlobals *globals = (PLOCGlobals *)args.scratchBuffer;
+    uint *const sync_mem = (uint*)((char*)args.scratchBuffer + sizeof(PLOCGlobals));
+    uint *const scratch  = (uint*)((char*)args.scratchBuffer + sizeof(PLOCGlobals) + sizeof(uint)*MAX_LARGE_WGS);    
+  
+    // ======================          
+    // ==== init globals ====
+    // ======================
+    {
+      sycl::event queue_event =  gpu_queue.submit([&](sycl::handler &cgh) {
+                                                    cgh.single_task([=]() {
+                                                                      globals->reset();
+                                                                    });
+                                                  });
+      gpu::waitOnEventAndCatchException(queue_event);
+      if (unlikely(verbose1))
+      {
+        double dt = gpu::getDeviceExecutionTiming(queue_event);
+        timer.add_to_device_timer(BuildTimer::PRE_PROCESS,dt);
+        if (unlikely(verbose2)) std::cout << "=> Init Globals I: " << dt << " ms" << std::endl;        
+      }      
+    }  
+  
+    // ==============================================================================    
+    // === get primitive type count from geometries, compute quad blocks per geom ===
+    // ==============================================================================
+  
+    const RTHWIF_GEOMETRY_DESC** geometries = args.geometries;
+    uint numGeometries                = args.numGeometries;
+ 
+    double device_prim_counts_time = 0.0f;
+  
+    const PrimitiveCounts primCounts = countPrimitives(gpu_queue,geometries,numGeometries,globals,scratch,host_device_tasks,device_prim_counts_time,verbose1); 
+
+    // ================================================
+    
+    timer.stop(BuildTimer::PRE_PROCESS);
+    timer.add_to_device_timer(BuildTimer::PRE_PROCESS,device_prim_counts_time);                              
+    if (unlikely(verbose2)) std::cout << "=> Count Primitives from Geometries: " << timer.get_host_timer() << " ms (host) " << device_prim_counts_time << " ms (device) " << std::endl;      
+  
+    uint numQuads            = primCounts.numQuads + primCounts.numTriangles; // no quadification taken into account at this point
+    uint numProcedurals      = primCounts.numProcedurals;
+    uint numInstances        = primCounts.numInstances;
+    const uint numQuadBlocks = primCounts.numQuadBlocks;
+
+    const uint expected_numPrimitives = numQuads + numProcedurals + numInstances;    
+
+    // =================================================    
+    // === empty scene before removing invalid prims ===
+    // =================================================
+    
+    if (unlikely(expected_numPrimitives == 0)) createEmptyBVH(args,gpu_queue);
+        
+    if (numQuads)
+    {
+      // ==================================================
+      // === compute correct quadification using blocks === 
+      // ==================================================
+
+      timer.start(BuildTimer::PRE_PROCESS);      
+      double device_quadification_time = 0.0;
+      numQuads = countQuadsPerGeometryUsingBlocks(gpu_queue,globals,args.geometries,numGeometries,numQuadBlocks,scratch,scratch+numGeometries,host_device_tasks,device_quadification_time,verbose1);
+      timer.stop(BuildTimer::PRE_PROCESS);
+      timer.add_to_device_timer(BuildTimer::PRE_PROCESS,device_quadification_time);
+      if (unlikely(verbose2)) std::cout << "=> Count " << numQuads << " Quads " << timer.get_host_timer() << " ms (host) " << (float)device_quadification_time << " ms (device) " << std::endl;
+    }
+    
+    // ================================
+    // === estimate size of the BVH ===
+    // ================================
+
+    size_t numPrimitives             = numQuads + numInstances + numProcedurals;  // actual #prims can be lower due to invalid instances or procedurals but quads count is accurate at this point
+    const size_t allocated_size      = args.accelBufferBytes;
+    const size_t header              = 128;
+    const size_t leaf_size           = estimateSizeLeafNodes(numQuads,numInstances,numProcedurals);
+    const size_t node_size           = (header + leaf_size) <= allocated_size ? allocated_size - leaf_size - header : 0; 
+    const size_t node_data_start     = header;
+    const size_t leaf_data_start     = header + node_size;
+      
+    // =================================================================
+    // === if allocated accel buffer is too small, return with error ===
+    // =================================================================
+
+    const uint required_size = header + estimateSizeInternalNodes(numQuads,numInstances,numProcedurals,false) + leaf_size;
+    if (unlikely(allocated_size < required_size))
+    {
+      if (unlikely(verbose2))
+      {
+        PRINT2(required_size,allocated_size);
+        PRINT2(node_size,estimateSizeInternalNodes(numQuads,numInstances,numProcedurals,false));        
+        PRINT3("RETRY BVH BUILD DUE BECAUSE OF SMALL ACCEL BUFFER ALLOCATION!!!", args.accelBufferBytes,required_size );
+      }
+      if (args.accelBufferBytesOut) *args.accelBufferBytesOut = required_size;
+      if (host_device_tasks) sycl::free(host_device_tasks,gpu_queue.get_context());
+      return RTHWIF_ERROR_RETRY;
+    }
+
+    const size_t conv_mem_size = sizeof(numPrimitives)*numPrimitives;
+    const size_t NUM_ACTIVE_LARGE_WGS = min((numPrimitives+LARGE_WG_SIZE-1)/LARGE_WG_SIZE,(size_t)MAX_WGS);
+
+    // ===========================
+    // === set up all pointers ===
+    // ===========================
+    QBVH6* qbvh   = (QBVH6*)args.accelBuffer;
+    char *bvh_mem = (char*)qbvh + header;
+    char *const leaf_mem = (char*)qbvh + leaf_data_start;
+    BVH2Ploc *const bvh2 = (BVH2Ploc*)(leaf_mem);
+    typedef gpu::MortonCodePrimitive64Bit_2x MCPrim;
+    MCPrim *const mc0 = (MCPrim*)(bvh2 + numPrimitives);
+    MCPrim *const mc1 = mc0 + numPrimitives;     
+    MCPrim *const morton_codes[2] = { mc0, mc1 }; 
+    uint *const cluster_index     = (uint*) (bvh_mem + 0 * numPrimitives * sizeof(uint)); // * 2
+    BVH2SubTreeState *const bvh2_subtree_size = (BVH2SubTreeState*) (bvh_mem + 2 * numPrimitives * sizeof(uint)); // * 2        
+    uint *cluster_i[2] = { cluster_index + 0, cluster_index + numPrimitives };        
+    uint *const cluster_index_source = cluster_i[0];
+    uint *const   cluster_index_dest = cluster_i[1];
+    LeafGenerationData *leafGenData = (LeafGenerationData*)scratch;
+
+    // ==============================          
+    // ==== init globals phase 2 ====
+    // ==============================
+    {
+      sycl::event queue_event =  gpu_queue.submit([&](sycl::handler &cgh) {
+                                                    cgh.single_task([=]() {
+                                                                      globals->numPrimitives              = numPrimitives;
+                                                                      globals->node_mem_allocator_cur     = node_data_start/64;
+                                                                      globals->node_mem_allocator_start   = node_data_start/64;
+                                                                      globals->leaf_mem_allocator_cur     = leaf_data_start/64;
+                                                                      globals->leaf_mem_allocator_start   = leaf_data_start/64;
+                                                                      globals->bvh2_index_allocator       = numPrimitives; 
+                                                                    });
+                                                  });
+      gpu::waitOnEventAndCatchException(queue_event);
+      if (unlikely(verbose1))
+      {
+        double dt = gpu::getDeviceExecutionTiming(queue_event);
+        timer.add_to_device_timer(BuildTimer::PRE_PROCESS,dt);
+        if (unlikely(verbose2)) std::cout << "=> Init Globals II: " << dt << " ms" << std::endl;        
+      }      
+    }	    
+
+    timer.start(BuildTimer::PRE_PROCESS);        
+    
+    double create_primref_time = 0.0f;
+    // ===================================================          
+    // ==== merge triangles to quads, create primrefs ====
+    // ===================================================
+
+    if (numQuads)
+      createQuads_initPLOCPrimRefs(gpu_queue,globals,args.geometries,numGeometries,numQuadBlocks,scratch,bvh2,0,create_primref_time,verbose1);
+    
+    // ====================================          
+    // ==== create procedural primrefs ====
+    // ====================================
+
+    if (numProcedurals)
+      numProcedurals = createProcedurals_initPLOCPrimRefs(gpu_queue,args.geometries,numGeometries,sync_mem,NUM_ACTIVE_LARGE_WGS,bvh2,numQuads,host_device_tasks,create_primref_time,verbose1);
+
+    // ==================================          
+    // ==== create instance primrefs ====
+    // ==================================
+    
+    if (numInstances)
+      numInstances = createInstances_initPLOCPrimRefs(gpu_queue,args.geometries,numGeometries,sync_mem,NUM_ACTIVE_LARGE_WGS,bvh2,numQuads + numProcedurals,host_device_tasks,create_primref_time,verbose1);
+
+    // =================================================================================================    
+    // === recompute actual number of primitives after quadification and removing of invalid entries ===
+    // =================================================================================================
+    
+    numPrimitives = numQuads + numInstances + numProcedurals;
+
+    const GeometryTypeRanges geometryTypeRanges(numQuads,numProcedurals,numInstances);        
+    
+    if (unlikely(verbose2))
+    {
+      PRINT4(numPrimitives,numQuads,numInstances,numProcedurals);
+      PRINT3(node_size,leaf_size,args.accelBufferBytes);
+      PRINT2(node_size/64,leaf_size/64);      
+    }      
+  
+    // =================================================================================    
+    // === test for empty scene again after all final primitive counts are available ===
+    // =================================================================================
+
+    if (unlikely(numPrimitives == 0))
+    {
+      if (host_device_tasks) sycl::free(host_device_tasks,gpu_queue.get_context());      
+      return createEmptyBVH(args,gpu_queue);
+    }
+
+    timer.stop(BuildTimer::PRE_PROCESS);
+    timer.add_to_device_timer(BuildTimer::PRE_PROCESS,create_primref_time);    
+    if (unlikely(verbose2)) std::cout << "=> Create Quads/Procedurals/Instances etc, Init PrimRefs: " << timer.get_host_timer() << " ms (host) " << create_primref_time << " ms (device) " << std::endl;
+      
+    // ==========================================          
+    // ==== get centroid and geometry bounds ====
+    // ==========================================
+        
+    timer.start(BuildTimer::PRE_PROCESS);        
+    double device_compute_centroid_bounds_time = 0.0f;
+     
+    computeCentroidGeometryBounds(gpu_queue, &globals->geometryBounds, &globals->centroidBounds, bvh2, numPrimitives, device_compute_centroid_bounds_time, verbose1);
+    
+    timer.stop(BuildTimer::PRE_PROCESS);
+    timer.add_to_device_timer(BuildTimer::PRE_PROCESS,device_compute_centroid_bounds_time);
+
+
+    if (unlikely(verbose2))
+      std::cout << "=> Get Geometry and Centroid Bounds Phase: " << timer.get_host_timer() << " ms (host) " << device_compute_centroid_bounds_time << " ms (device) " << std::endl;		
+  
+    // ==============================          
+    // ==== compute morton codes ====
+    // ==============================
+
+    const bool fastMCMode = numPrimitives < FAST_MC_NUM_PRIMS_THRESHOLD || (args.quality == RTHWIF_BUILD_QUALITY_LOW && numPrimitives < FAST_MC_MAX_NUM_PRIMS);    
+    
+    timer.start(BuildTimer::PRE_PROCESS);        
+    double device_compute_mc_time = 0.0f;
+
+    if (!fastMCMode)
+      computeMortonCodes64Bit_SaveMSBBits(gpu_queue,&globals->centroidBounds,mc0,bvh2,(uint*)bvh2_subtree_size,numPrimitives,device_compute_mc_time,verbose1);
+    else
+      computeMortonCodes64Bit(gpu_queue,&globals->centroidBounds,(gpu::MortonCodePrimitive40x24Bits3D*)mc1,bvh2,numPrimitives,0,(uint64_t)-1,device_compute_mc_time,verbose1);
+            
+    timer.stop(BuildTimer::PRE_PROCESS);
+     
+    if (unlikely(verbose2))
+      std::cout << "=> Compute Morton Codes: " << timer.get_host_timer() << " ms (host) " << device_compute_mc_time << " ms (device) " << std::endl;		
+    
+    // ===========================          
+    // ==== sort morton codes ====
+    // ===========================
+
+    timer.start(BuildTimer::PRE_PROCESS);        
+    
+    if (!fastMCMode) // fastMCMode == 32bit key + 32bit value pairs, !fastMode == 64bit key + 32bit value pairs
+    {
+      const uint scratchMemWGs = gpu::getNumWGsScratchSize(conv_mem_size);
+      const uint nextPowerOf2 =  1 << (32 - sycl::clz(numPrimitives) - 1);
+      const uint sortWGs = min(max(min((int)nextPowerOf2/8192,(int)gpu_maxComputeUnits/4),1),(int)scratchMemWGs);
+
+      sycl::event initial = sycl::event();
+      sycl::event block0  = gpu::radix_sort_Nx8Bit(gpu_queue, morton_codes[0], morton_codes[1], numPrimitives, (uint*)scratch, 4, 8, initial, sortWGs);      
+      sycl::event restore = restoreMSBBits(gpu_queue,mc0,(uint*)bvh2_subtree_size,numPrimitives,block0,verbose1);      
+      sycl::event block1  = gpu::radix_sort_Nx8Bit(gpu_queue, morton_codes[0], morton_codes[1], numPrimitives, (uint*)scratch, 4, 8, restore, sortWGs);
+      gpu::waitOnEventAndCatchException(block1);      
+    }
+    else
+    {
+      if (numPrimitives < SMALL_SORT_THRESHOLD)
+        gpu::radix_sort_single_workgroup(gpu_queue, (uint64_t *)mc0, (uint64_t *)mc1, numPrimitives, 3,8);
+      else
+      {
+        const uint scratchMemWGs = gpu::getNumWGsScratchSize(conv_mem_size);        
+        const uint nextPowerOf2 =  1 << (32 - sycl::clz(numPrimitives) - 1);          
+        const uint sortWGs = min(max(min((int)nextPowerOf2/LARGE_WG_SIZE,(int)gpu_maxComputeUnits/4),1),(int)scratchMemWGs);
+        sycl::event initial = sycl::event();
+        sycl::event block0  = gpu::radix_sort_Nx8Bit(gpu_queue, (gpu::MortonCodePrimitive40x24Bits3D*)morton_codes[1], (gpu::MortonCodePrimitive40x24Bits3D*)morton_codes[0], numPrimitives, (uint*)scratch, 3, 8, initial, sortWGs);
+        gpu::waitOnEventAndCatchException(block0);              
+      }      
+    }
+    
+    timer.stop(BuildTimer::PRE_PROCESS);        
+    timer.add_to_device_timer(BuildTimer::PRE_PROCESS,timer.get_host_timer());
+            
+    if (unlikely(verbose2))
+      std::cout << "=> Sort Morton Codes: " << timer.get_host_timer() << " ms (host and device)" << std::endl;
+                      
+    // ===========================          
+    // ====== init clusters ======
+    // ===========================
+
+    
+    timer.start(BuildTimer::PRE_PROCESS);        
+    double device_init_clusters_time = 0.0f;
+
+    if (!fastMCMode)
+      initClusters(gpu_queue,mc0,bvh2,cluster_index,bvh2_subtree_size,numPrimitives,device_init_clusters_time,verbose1);
+    else
+      initClusters(gpu_queue,(gpu::MortonCodePrimitive40x24Bits3D*)mc0,bvh2,cluster_index,bvh2_subtree_size,numPrimitives,device_init_clusters_time,verbose1); 
+    
+    timer.stop(BuildTimer::PRE_PROCESS);        
+    timer.add_to_device_timer(BuildTimer::PRE_PROCESS,device_init_clusters_time);
+        
+    if (unlikely(verbose2))
+      std::cout << "=> Init Clusters: " << timer.get_host_timer() << " ms (host) " << device_init_clusters_time << " ms (device) " << std::endl;		
+
+    uint numPrims = numPrimitives;
+    // ===================================================================================================================================================
+    // ===================================================================================================================================================
+    // ===================================================================================================================================================
+
+    // === 8 or 16-wide search radius dependening on compiler flags ===
+    const uint SEARCH_RADIUS_SHIFT = args.quality == RTHWIF_BUILD_QUALITY_LOW ? 3 : 4;
+    
+    double device_ploc_iteration_time = 0.0f;
+        
+    uint iteration = 0;
+  
+    timer.start(BuildTimer::BUILD);        
+
+    // ========================            
+    // ==== clear sync mem ====
+    // ========================      
+
+    clearScratchMem(gpu_queue,sync_mem,0,NUM_ACTIVE_LARGE_WGS,device_ploc_iteration_time,verbose1);
+
+    float ratio = 100.0f;
+    for (;numPrims>1;iteration++)
+    {          
+      // ==================================================            
+      // ==== single kernel path if #prims < threshold ====
+      // ==================================================
+      
+      if (numPrims < SINGLE_WG_SWITCH_THRESHOLD)
+      {
+        double singleWG_time = 0.0f;
+        singleWGBuild(gpu_queue, globals, bvh2, cluster_index_source, cluster_index_dest, bvh2_subtree_size, numPrims, SEARCH_RADIUS_SHIFT, singleWG_time, verbose1);
+        timer.add_to_device_timer(BuildTimer::BUILD,singleWG_time);
+        numPrims = 1;
+      }
+      else  
+      {            
+        // ===================================================================================
+        // ==== nearest neighbor search, merge clusters and create bvh2 nodes (fast path) ====
+        // ===================================================================================
+        device_ploc_iteration_time = 0.0f;
+        iteratePLOC(gpu_queue,globals,bvh2,cluster_index_source,cluster_index_dest,bvh2_subtree_size,sync_mem,numPrims,NUM_ACTIVE_LARGE_WGS,host_device_tasks,SEARCH_RADIUS_SHIFT,device_ploc_iteration_time,ratio < BOTTOM_LEVEL_RATIO,verbose1);
+        timer.add_to_device_timer(BuildTimer::BUILD,device_ploc_iteration_time);
+      
+        const uint new_numPrims = *host_device_tasks;
+        assert(new_numPrims < numPrims);
+        ratio = (float)(numPrims-new_numPrims) / numPrims * 100.0f;
+        numPrims = new_numPrims;                      
+        // ==========================            
+      }        
+      if (unlikely(verbose2))
+        PRINT5(iteration,numPrims,ratio,(float)device_ploc_iteration_time,(float)timer.get_accum_device_timer(BuildTimer::BUILD));
+    }
+  
+    timer.stop(BuildTimer::BUILD);        
+
+    if (unlikely(verbose2))
+      std::cout << "=> PLOC phase: " <<  timer.get_host_timer() << " ms (host) " << (float)timer.get_accum_device_timer(BuildTimer::BUILD) << " ms (device) " << std::endl;    
+
+
+
+#if BVH2_REBALANCE == 1
+        // ===============================================================================================================
+        // ========================================== rebalance BVH2 if degenerated ======================================
+        // ===============================================================================================================
+#if 0
+        uint maxDepth = 0;
+        for (uint i=0;i<numPrims;i++)
+        {
+          const uint depth = getBVH2Depth(bvh2,cluster_index_source[i],numPrimitives);
+          maxDepth = max(maxDepth,depth);
+        }
+        PRINT(maxDepth);
+#endif
+        
+        double rebalanceBVH2_time = 0.0f;
+        rebalanceBVH2(gpu_queue,bvh2,bvh2_subtree_size,numPrimitives,rebalanceBVH2_time,verbose1);
+        if (unlikely(verbose2))
+          PRINT(rebalanceBVH2_time);
+        timer.add_to_device_timer(BuildTimer::BUILD,rebalanceBVH2_time);
+
+#if 0
+        maxDepth = 0;
+        for (uint i=0;i<numPrims;i++)
+        {
+          const uint depth = getBVH2Depth(bvh2,cluster_index_source[i],numPrimitives);
+          maxDepth = max(maxDepth,depth);
+        }
+        PRINT(maxDepth);        
+#endif
+        // ===============================================================================================================
+        // ===============================================================================================================
+        // ===============================================================================================================
+#endif        
+                
+    // =====================================                
+    // === check and convert BVH2 (host) ===
+    // =====================================
+
+    if (unlikely(verbose2))
+        PRINT2(globals->bvh2_index_allocator,2*numPrimitives);        
+    
+    if (unlikely(verbose2))
+    {
+      PRINT2(globals->bvh2_index_allocator,2*numPrimitives);              
+      if (globals->bvh2_index_allocator >= 2*numPrimitives)
+        FATAL("BVH2 construction, allocator");
+      PRINT(globals->rootIndex);      
+      uint nodes = 0;
+      uint leaves = 0;
+      float nodeSAH = 0;
+      float leafSAH = 0;
+      uint maxDepth = 0;
+      checkBVH2PlocHW(bvh2,globals->rootIndex,nodes,leaves,nodeSAH,leafSAH,maxDepth,numPrimitives,globals->bvh2_index_allocator,0);
+      nodeSAH /= globals->geometryBounds.area();
+      leafSAH /= globals->geometryBounds.area();                
+      PRINT5(nodes,leaves,nodeSAH,leafSAH,maxDepth);
+ 
+      /* --- dummy kernel to trigger USM transfer again to not screw up device timings --- */
+      sycl::event queue_event =  gpu_queue.submit([&](sycl::handler &cgh) {
+                                                    cgh.single_task([=]() {
+                                                                    });
+                                                  });
+      gpu::waitOnEventAndCatchException(queue_event);
+    }
+    
+    // =============================    
+    // === convert BVH2 to QBVH6 ===
+    // =============================
+    timer.start(BuildTimer::PRE_PROCESS);    
+    float conversion_device_time = 0.0f;
+    const bool convert_success = convertBVH2toQBVH6(gpu_queue,globals,host_device_tasks,args.geometries,qbvh,bvh2,leafGenData,numPrimitives,numInstances != 0,geometryTypeRanges,conversion_device_time,verbose1);
+
+    /* --- init final QBVH6 header --- */        
+    {     
+      sycl::event queue_event =  gpu_queue.submit([&](sycl::handler &cgh) {
+                                                    cgh.single_task([=]() {
+                                                                      /* init qbvh */
+                                                                      qbvh->bounds.lower.x = globals->geometryBounds.lower_x;
+                                                                      qbvh->bounds.lower.y = globals->geometryBounds.lower_y;
+                                                                      qbvh->bounds.lower.z = globals->geometryBounds.lower_z;
+                                                                      qbvh->bounds.upper.x = globals->geometryBounds.upper_x;
+                                                                      qbvh->bounds.upper.y = globals->geometryBounds.upper_y;
+                                                                      qbvh->bounds.upper.z = globals->geometryBounds.upper_z;
+                                                                      qbvh->numPrims       = numPrimitives;                                                                        
+                                                                      qbvh->nodeDataStart  = globals->node_mem_allocator_start;
+                                                                      qbvh->nodeDataCur    = globals->node_mem_allocator_cur;
+                                                                      qbvh->leafDataStart  = globals->leaf_mem_allocator_start;
+                                                                      qbvh->leafDataCur    = globals->leaf_mem_allocator_cur;
+                                                                      *(gpu::AABB3f*)host_device_tasks = globals->geometryBounds;
+                                                                    });
+                                                  });
+      gpu::waitOnEventAndCatchException(queue_event);
+    }	    
+    
+    if (args.boundsOut) *args.boundsOut = *(RTHWIF_AABB*)host_device_tasks;
+  
+    timer.stop(BuildTimer::POST_PROCESS);
+    timer.add_to_device_timer(BuildTimer::POST_PROCESS,conversion_device_time);
+
+    if (unlikely(verbose2))
+      std::cout << "=> BVH2 -> QBVH6 Flattening: " <<  timer.get_host_timer() << " ms (host) " << conversion_device_time << " ms (device) " << std::endl;
+
+    // ==========================================================    
+    // ==========================================================
+    // ==========================================================
+
+    if (unlikely(verbose2))
+    {
+      // === memory allocation and usage stats ===
+      const uint nodes_used   = globals->node_mem_allocator_cur-globals->node_mem_allocator_start;
+      const uint leaves_used  = globals->leaf_mem_allocator_cur-globals->leaf_mem_allocator_start;
+      const float nodes_util  = 100.0f * (float)(globals->node_mem_allocator_cur-globals->node_mem_allocator_start) / (node_size/64);
+      const float leaves_util = 100.0f * (float)(globals->leaf_mem_allocator_cur-globals->leaf_mem_allocator_start) / (leaf_size/64);
+      PRINT4(globals->node_mem_allocator_start,globals->node_mem_allocator_cur,nodes_used,nodes_util);
+      PRINT4(globals->leaf_mem_allocator_start,globals->leaf_mem_allocator_cur,leaves_used,leaves_util);      
+      PRINT(globals->numLeaves);
+    }
+
+    if (unlikely(convert_success == false))
+    {
+      if (args.accelBufferBytesOut) *args.accelBufferBytesOut = estimateAccelBufferSize(numQuads, numInstances, numProcedurals, true); 
+      if (host_device_tasks) sycl::free(host_device_tasks,gpu_queue.get_context());      
+      return RTHWIF_ERROR_RETRY;
+    }
+    
+#if defined(EMBREE_SYCL_ALLOC_DISPATCH_GLOBALS)
+    {
+      sycl::event queue_event =  gpu_queue.submit([&](sycl::handler &cgh) {
+                                                    cgh.single_task([=]() {
+                                                                      HWAccel* hwaccel = (HWAccel*)args.accelBuffer;  
+                                                                      hwaccel->dispatchGlobalsPtr = (uint64_t)args.dispatchGlobalsPtr;                                                                      
+                                                                    });
+                                                  });
+      gpu::waitOnEventAndCatchException(queue_event);
+    }
+#endif      
+
+    if (args.accelBufferBytesOut)
+      *args.accelBufferBytesOut = args.accelBufferBytes;
+
+#if 1
+    if (verbose2)
+    {
+      gpu::waitOnQueueAndCatchException(gpu_queue);
+      
+      qbvh->print(std::cout,qbvh->root(),0,6);
+      BVHStatistics stats = qbvh->computeStatistics();      
+      stats.print(std::cout);
+      stats.print_raw(std::cout);
+      PRINT("VERBOSE STATS DONE");
+    }        
+#endif
+
+    if (host_device_tasks) sycl::free(host_device_tasks,gpu_queue.get_context());      
+    
+    if (unlikely(verbose1))
+      std::cout << "=> BVH build time: host = " << timer.get_total_host_time() << " ms , device = " << timer.get_total_device_time() << " ms , numPrimitives (original) = " << expected_numPrimitives << " , numPrimitives (build) = " << numPrimitives << std::endl;
+
+    return RTHWIF_ERROR_NONE;    
+  }
+}
+
+#endif    

kernels/rthwif/rtbuild/statistics.cpp

@@ -74,7 +74,8 @@ namespace embree
     cout << "  numBuildPrimitivesPostSplit = " << numBuildPrimitivesPostSplit << std::endl;
     cout << "  primRefSplits               = " << std::setprecision(2) << percent(numBuildPrimitivesPostSplit,numBuildPrimitives) << "%" << std::endl;
     cout << "  numBVHPrimitives            = " << totalPrimitives << std::endl;
-    cout << "  spatialSplits               = " << std::setprecision(2) << percent(totalPrimitives,numScenePrimitives) << "%" << std::endl;    
+    cout << "  spatialSplits               = " << std::setprecision(2) << percent(totalPrimitives,numScenePrimitives) << "%" << std::endl;
+    cout << "  maxDepth                    = " << depth << std::endl;
     cout << std::endl;
      
     cout << "                      #nodes     SAH   total       bytes     used    total   b/node  b/child   b/prim  #child     fill" << std::endl;

kernels/rthwif/rtbuild/statistics.h

@@ -98,7 +98,7 @@ namespace embree
     };
 
     BVHStatistics ()
-    : numScenePrimitives(0), numBuildPrimitives(0), numBuildPrimitivesPostSplit(0) {}
+      : numScenePrimitives(0), numBuildPrimitives(0), numBuildPrimitivesPostSplit(0), depth(1) {}
         
     void print    (std::ostream& cout) const;
     void print_raw(std::ostream& cout) const;
@@ -106,6 +106,7 @@ namespace embree
     size_t numScenePrimitives;
     size_t numBuildPrimitives;
     size_t numBuildPrimitivesPostSplit;
+    size_t depth;
     NodeStat internalNode;
     LeafStat quadLeaf;
     LeafStat proceduralLeaf;

kernels/rthwif/testing/CMakeLists.txt

@@ -18,6 +18,14 @@ IF (NOT EMBREE_SYCL_IMPLICIT_DISPATCH_GLOBALS)
   ADD_DEFINITIONS("-DEMBREE_SYCL_ALLOC_DISPATCH_GLOBALS")
 ENDIF()
 
+#ADD_EXECUTABLE(rthwif_cornell_box rthwif_cornell_box.cpp ../rthwif_production.cpp)
+#SET_PROPERTY(TARGET rthwif_cornell_box APPEND PROPERTY COMPILE_FLAGS "-fsycl -fsycl-targets=spir64 -DEMBREE_SYCL_SUPPORT")
+#SET_PROPERTY(TARGET rthwif_cornell_box APPEND PROPERTY LINK_FLAGS "-fsycl -lze_loader")
+
+#ADD_EXECUTABLE(rthwif_test rthwif_test.cpp ../rthwif_production.cpp)
+#TARGET_LINK_LIBRARIES(rthwif_test embree_rthwif sys tbb)
+#TARGET_LINK_LIBRARIES(rthwif_test embree_rthwif_gpu sys tbb)
+
 IF (NOT DEFINED EMBREE_SYCL_RT_VALIDATION_API)
   OPTION(EMBREE_SYCL_RT_VALIDATION_API "Use rt_validation API instead of IGC provided rt_production API" OFF)
 ENDIF()
@@ -35,8 +43,21 @@ TARGET_LINK_LIBRARIES(embree_rthwif_cornell_box embree_rthwif sys simd ${TBB_TAR
 SET_PROPERTY(TARGET embree_rthwif_cornell_box APPEND PROPERTY COMPILE_FLAGS "-fsycl -fsycl-targets=spir64 -DEMBREE_SYCL_SUPPORT")
 SET_PROPERTY(TARGET embree_rthwif_cornell_box APPEND PROPERTY LINK_FLAGS "-fsycl -lze_loader -fsycl-targets=spir64 -Xsycl-target-backend=spir64 \" -cl-intel-greater-than-4GB-buffer-required \"")
 
+#ADD_EXECUTABLE(rthwif_test rthwif_test.cpp)
+#TARGET_LINK_LIBRARIES(rthwif_test embree_rthwif sys simd tbb)
+#IF (EMBREE_SYCL_GPU_BVH_BUILDER)
+#TARGET_LINK_LIBRARIES(rthwif_test embree_rthwif_gpu sys tbb)
+#ENDIF()
+
+#SET_PROPERTY(TARGET rthwif_test APPEND PROPERTY COMPILE_FLAGS "-fsycl -fsycl-targets=spir64 -DEMBREE_SYCL_SUPPORT")
+#SET_PROPERTY(TARGET rthwif_test APPEND PROPERTY LINK_FLAGS "-fsycl -lze_loader -fsycl-targets=spir64 -Xsycl-target-backend=spir64 \" -cl-intel-greater-than-4GB-buffer-required \"")
+
 ADD_EXECUTABLE(embree_rthwif_test rthwif_test.cpp)
 TARGET_LINK_LIBRARIES(embree_rthwif_test embree_rthwif sys simd ${TBB_TARGET} ${RT_SIM_LIBRARY})
+IF (EMBREE_SYCL_GPU_BVH_BUILDER)
+TARGET_LINK_LIBRARIES(embree_rthwif_test embree_rthwif_gpu sys tbb)
+ENDIF()
+
 SET_PROPERTY(TARGET embree_rthwif_test APPEND PROPERTY COMPILE_FLAGS "-fsycl -fsycl-targets=spir64 -DEMBREE_SYCL_SUPPORT")
 SET_PROPERTY(TARGET embree_rthwif_test APPEND PROPERTY LINK_FLAGS "-fsycl -lze_loader -fsycl-targets=spir64 -Xsycl-target-backend=spir64 \" -cl-intel-greater-than-4GB-buffer-required \"")
 

kernels/rthwif/testing/rthwif_test.cpp

@@ -27,6 +27,8 @@
 #include <map>
 #include <iostream>
 
+#define VERBOSE 0
+
 namespace embree {
   double getSeconds();
 }
@@ -944,6 +946,19 @@ void free_accel_buffer(void* ptr, sycl::context context)
 #endif
 }
 
+  
+  void exception_handler(sycl::exception_list exceptions)
+  {
+    for (std::exception_ptr const& e : exceptions) {
+      try {
+        std::rethrow_exception(e);
+      } catch(sycl::exception const& e) {
+        std::cout << "Caught asynchronous SYCL exception: " << e.what() << std::endl;
+      }
+    }
+  };
+
+
 struct Scene
 {
   typedef InstanceGeometryT<Scene> InstanceGeometry;
@@ -1073,7 +1088,7 @@ struct Scene
         geometries[j]->transform(world2local);
         scene->geometries.push_back(geometries[j]);
       }
-
+      
       //std::shared_ptr<InstanceGeometry> instance = std::make_shared<InstanceGeometry>(local2world,scene,procedural);
       uint32_t instUserID = RandomSampler_getUInt(rng);
       std::shared_ptr<InstanceGeometry> instance(new InstanceGeometry(local2world,scene,procedural,instUserID));
@@ -1101,14 +1116,22 @@ struct Scene
       std::swap(geometries[g],geometries[k]);
     }
   }
-
+  
   void buildAccel(sycl::device& device, sycl::context& context, BuildMode buildMode, bool benchmark = false)
   {
      RTHWIF_BUILD_QUALITY quality = (RTHWIF_BUILD_QUALITY) (RandomSampler_getUInt(rng) % 3);
      
     /* fill geometry descriptor buffer */
+    const size_t numGeometries = size();        
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    sycl::queue queue(device, exception_handler);
+    
+    GEOMETRY_DESC *desc = (GEOMETRY_DESC*)sycl::aligned_alloc(64,sizeof(GEOMETRY_DESC)*numGeometries,device,context,sycl::usm::alloc::shared);
+    RTHWIF_GEOMETRY_DESC **geom = (RTHWIF_GEOMETRY_DESC**)sycl::aligned_alloc(64,numGeometries*sizeof(RTHWIF_GEOMETRY_DESC*),device,context,sycl::usm::alloc::shared);
+#else        
     std::vector<GEOMETRY_DESC> desc(size());
     std::vector<const RTHWIF_GEOMETRY_DESC*> geom(size());
+#endif    
     size_t numPrimitives = 0;
     for (size_t geomID=0; geomID<size(); geomID++)
     {
@@ -1123,7 +1146,7 @@ struct Scene
       numPrimitives += g->getNumPrimitives();
       g->buildAccel(device,context,buildMode,quality);
       g->getDesc(&desc[geomID]);
-      geom[geomID] = (const RTHWIF_GEOMETRY_DESC*) &desc[geomID];
+      geom[geomID] = (RTHWIF_GEOMETRY_DESC*) &desc[geomID];
     }
 
     /* estimate accel size */
@@ -1132,15 +1155,23 @@ struct Scene
     RTHWIF_BUILD_ACCEL_ARGS args;
     memset(&args,0,sizeof(args));
     args.structBytes = sizeof(args);
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    args.geometries = (const RTHWIF_GEOMETRY_DESC**) geom;    
+#else    
     args.geometries = (const RTHWIF_GEOMETRY_DESC**) geom.data();
-    args.numGeometries = geom.size();
+#endif    
+    args.numGeometries = numGeometries; //geom.size();
     args.accelBuffer = nullptr;
     args.accelBufferBytes = 0;
     args.scratchBuffer = nullptr;
     args.scratchBufferBytes = 0;
     args.quality = quality;
     args.flags = RTHWIF_BUILD_FLAG_NONE;
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    args.parallelOperation = nullptr;
+#else    
     args.parallelOperation = parallelOperation;
+#endif    
     args.boundsOut = &bounds;
     args.accelBufferBytesOut = &accelBufferBytesOut;
     args.buildUserPtr = nullptr;
@@ -1151,7 +1182,11 @@ struct Scene
     RTHWIF_ACCEL_SIZE size;
     memset(&size,0,sizeof(RTHWIF_ACCEL_SIZE));
     size.structBytes = sizeof(RTHWIF_ACCEL_SIZE);
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    RTHWIF_ERROR err = rthwifGetAccelSizeGPU(args,size,&queue,VERBOSE);    
+#else    
     RTHWIF_ERROR err = rthwifGetAccelSize(args,size);
+#endif    
     if (err != RTHWIF_ERROR_NONE)
       throw std::runtime_error("BVH size estimate failed");
 
@@ -1160,11 +1195,20 @@ struct Scene
 
     /* allocate scratch buffer */
     size_t sentinelBytes = 1024; // add that many zero bytes to catch buffer overruns
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+      // === scratch buffer === 
+    char *scratchBuffer  = (char*)sycl::aligned_alloc(64,size.scratchBufferBytes+sentinelBytes,device,context,sycl::usm::alloc::shared);
+    assert(scratchBuffer);    
+    memset(scratchBuffer,0,size.scratchBufferBytes+sentinelBytes);
+    args.scratchBuffer = scratchBuffer;
+    args.scratchBufferBytes = size.scratchBufferBytes;           
+#else    
     std::vector<char> scratchBuffer(size.scratchBufferBytes+sentinelBytes);
     memset(scratchBuffer.data(),0,scratchBuffer.size());
     args.scratchBuffer = scratchBuffer.data();
     args.scratchBufferBytes = size.scratchBufferBytes;
-
+#endif
+    
     accel = nullptr;
     size_t accelBytes = 0;
     
@@ -1177,23 +1221,35 @@ struct Scene
       accel = alloc_accel_buffer(accelBytes+sentinelBytes,device,context);
       memset(accel,0,accelBytes+sentinelBytes);
 
+      
+      size_t numIterations = benchmark ? 16 : 1;
+
+      args.numGeometries = numGeometries; //geom.size();
+      args.accelBuffer = accel;
+      args.accelBufferBytes = size.accelBufferWorstCaseBytes;
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)      
+      rthwifPrefetchAccelGPU(args,&queue,VERBOSE);      
+#endif
+      
       /* build accel */
       double t0 = embree::getSeconds();
-      size_t numIterations = benchmark ? 16 : 1;
 
       for (size_t i=0; i<numIterations; i++)
       {
-        args.numGeometries = geom.size();
-        args.accelBuffer = accel;
-        args.accelBufferBytes = size.accelBufferWorstCaseBytes;
+        
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+        err = rthwifBuildAccelGPU(args,&queue,VERBOSE);        
+#else                
         err = rthwifBuildAccel(args);
+#endif        
 
         if (args.parallelOperation) {
           assert(err == RTHWIF_ERROR_PARALLEL_OPERATION);
           uint32_t maxThreads = rthwifGetParallelOperationMaxConcurrency(parallelOperation);
           tbb::parallel_for(0u, maxThreads, 1u, [&](uint32_t) {
-            err = rthwifJoinParallelOperation(parallelOperation);
-          });
+                                                  err = rthwifJoinParallelOperation(parallelOperation);
+                                                });
         }
       
         if (err != RTHWIF_ERROR_NONE)
@@ -1222,22 +1278,33 @@ struct Scene
         memset(accel,0,accelBytes+sentinelBytes);
 
         /* build accel */
-        args.numGeometries = geom.size();
+        args.numGeometries = numGeometries; //geom.size();
         args.accelBuffer = accel;
         args.accelBufferBytes = accelBytes;
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+        err = rthwifBuildAccelGPU(args,&queue,VERBOSE);
+#else        
         err = rthwifBuildAccel(args);
+#endif        
 
         if (args.parallelOperation) {
           assert(err == RTHWIF_ERROR_PARALLEL_OPERATION);
           uint32_t maxThreads = rthwifGetParallelOperationMaxConcurrency(parallelOperation);
           tbb::parallel_for(0u, maxThreads, 1u, [&](uint32_t) {
-            err = rthwifJoinParallelOperation(parallelOperation);
-          });
+                                                  err = rthwifJoinParallelOperation(parallelOperation);
+                                                });
         }
         
         if (err != RTHWIF_ERROR_RETRY)
           break;
 
+        // PRINT(accelBufferBytesOut);
+        // PRINT(bytes);
+        // PRINT(size.accelBufferExpectedBytes);        
+        // PRINT(size.accelBufferWorstCaseBytes);
+        // PRINT(accelBufferBytesOut);
+        
         if (accelBufferBytesOut < bytes || size.accelBufferWorstCaseBytes < accelBufferBytesOut )
           throw std::runtime_error("failed build returned wrong new estimate");
 
@@ -1271,6 +1338,12 @@ struct Scene
         throw std::runtime_error("wrong acceleration structure size returned");
       }
     }
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    sycl::free(scratchBuffer,context);    
+    sycl::free(geom    ,context);
+    sycl::free(desc   ,context);    
+#endif                      
   }
   
   void buildTriMap(Transform local_to_world, std::vector<uint32_t> id_stack, uint32_t instUserID, bool procedural_instance, std::vector<Hit>& tri_map)
@@ -1311,17 +1384,6 @@ struct Scene
   void* accel;
 };
 
-void exception_handler(sycl::exception_list exceptions)
-{
-  for (std::exception_ptr const& e : exceptions) {
-    try {
-      std::rethrow_exception(e);
-    } catch(sycl::exception const& e) {
-      std::cout << "Caught asynchronous SYCL exception: " << e.what() << std::endl;
-    }
-  }
-};
-
 void render(uint32_t i, const TestInput& in, TestOutput& out, intel_raytracing_acceleration_structure_t accel)
 {
   intel_raytracing_ext_flag_t flags = intel_get_raytracing_ext_flag();
@@ -1786,7 +1848,6 @@ uint32_t executeTest(sycl::device& device, sycl::queue& queue, sycl::context& co
     scene->createInstances(scene->size(),3, inst == InstancingType::SW_INSTANCING);
 
   scene->addNullGeometries(16);
-
   scene->buildAccel(device,context,BuildMode::BUILD_EXPECTED_SIZE,false);
 
   /* calculate test input and expected output */
@@ -1912,6 +1973,9 @@ uint32_t executeBuildTest(sycl::device& device, sycl::queue& queue, sycl::contex
     const uint32_t numPrimitives = i>10 ? i*i : i;
     std::cout << "testing " << numPrimitives << " primitives" << std::endl;
     numErrors += executeBuildTest(device,queue,context,test,buildMode,numPrimitives,i);
+    if (VERBOSE)
+      if (numErrors)
+        exit(1);
   }
   return numErrors;
 }
@@ -2075,7 +2139,7 @@ int main(int argc, char* argv[])
     
   /* initialize SYCL device */
   device = sycl::device(sycl::gpu_selector_v);
-  sycl::queue queue = sycl::queue(device,exception_handler);
+  sycl::queue queue = sycl::queue(device, exception_handler);
   context = queue.get_context();
 
   dispatchGlobalsPtr = allocDispatchGlobals(device,context);
@@ -2087,16 +2151,21 @@ int main(int argc, char* argv[])
   
   uint32_t numErrors = 0;
   if (test >= TestType::BENCHMARK_TRIANGLES)
+  {
     numErrors = executeBenchmark(device,queue,context,test);
+  }
   else if (test >= TestType::BUILD_TEST_TRIANGLES)
+  {
     numErrors = executeBuildTest(device,queue,context,test,buildMode);
+  }
   else
+  {
     numErrors = executeTest(device,queue,context,inst,test);
+  }
 
   sycl::free(dispatchGlobalsPtr, context);
 
   rthwifExit();
-  
   return numErrors ? 1 : 0;
 }
 

kernels/sycl/rthwif_embree_builder.cpp

@@ -17,6 +17,9 @@
 
 namespace embree
 {
+
+  BBox3fa rthwifBuildPloc(Scene* scene, RTCBuildQuality quality_flags, Device::avector<char,64>& buffer_o, const bool two_level=false);
+
   using namespace embree::isa;
 
   enum Flags : uint32_t {
@@ -31,7 +34,7 @@ namespace embree
   struct DispatchGlobals
   {
     uint64_t rtMemBasePtr;               // base address of the allocated stack memory
-    uint64_t callStackHandlerKSP;             // this is the KSP of the continuation handler that is invoked by BTD when the read KSP is 0
+    uint64_t callStackHandlerKSP;        // this is the KSP of the continuation handler that is invoked by BTD when the read KSP is 0
     uint32_t asyncStackSize;             // async-RT stack size in 64 byte blocks
     uint32_t numDSSRTStacks : 16;        // number of stacks per DSS
     uint32_t syncRayQueryCount : 4;      // number of ray queries in the sync-RT stack: 0-15 mapped to: 1-16
@@ -388,6 +391,7 @@ namespace embree
     out->transform = (float*) &out->xfmdata;
     out->bounds = (RTHWIF_AABB*) &dynamic_cast<Scene*>(geom->object)->hwaccel_bounds;
     out->xfmdata = *(RTHWIF_TRANSFORM_FLOAT4X4_COLUMN_MAJOR*) &local2world;
+    
     EmbreeHWAccel* hwaccel = (EmbreeHWAccel*) dynamic_cast<Scene*>(geom->object)->hwaccel.data();
     out->accel = hwaccel->AccelTable[0];
   }
@@ -439,8 +443,33 @@ namespace embree
     return (RTHWIF_BUILD_FLAGS) result;
   }  
 
-  BBox3f rthwifBuild(Scene* scene, AccelBuffer& accel)
+  
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+  
+  void exception_handler(sycl::exception_list exceptions)
   {
+    for (std::exception_ptr const& e : exceptions) {
+      try {
+        std::rethrow_exception(e);
+      } catch(sycl::exception const& e) {
+        std::cout << "Caught asynchronous SYCL exception: " << e.what() << std::endl;
+      }
+    }
+  };
+
+#endif
+  
+  
+  BBox3f rthwifBuild(Scene* scene, AccelBuffer& accel)
+  {    
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    DeviceGPU *gpu_device = dynamic_cast<DeviceGPU*>(scene->device);    
+    sycl::device &sycl_device = gpu_device->getGPUDevice();
+    sycl::property_list PropList;
+    if (gpu_device->verbose) PropList = { sycl::property::queue::enable_profiling() };
+    sycl::queue sycl_queue(sycl_device, exception_handler, PropList );
+#endif
+    
     auto getType = [&](unsigned int geomID) -> GEOMETRY_TYPE
     {
       /* no HW support for MB yet */
@@ -521,12 +550,22 @@ namespace embree
       totalBytes += sizeof_RTHWIF_GEOMETRY(type);
     }
 
-    /* fill geomdesc buffers */
-    std::vector<RTHWIF_GEOMETRY_DESC*> geomDescr(scene->size());
+    const size_t numGeometries = scene->size();
+    
+    /* fill geomdesc buffers */    
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    const size_t geomDescrBytes = sizeof(RTHWIF_GEOMETRY_DESC*)*numGeometries;
+    RTHWIF_GEOMETRY_DESC** geomDescr = (RTHWIF_GEOMETRY_DESC**)sycl::aligned_alloc_shared(64,geomDescrBytes,gpu_device->getGPUDevice(),gpu_device->getGPUContext(),sycl::ext::oneapi::property::usm::device_read_only());
+    assert(geomDescr);        
+    char *geomDescrData = (char*)sycl::aligned_alloc_shared(64,totalBytes,gpu_device->getGPUDevice(),gpu_device->getGPUContext(),sycl::ext::oneapi::property::usm::device_read_only());
+    assert(geomDescrData);
+#else    
+    std::vector<RTHWIF_GEOMETRY_DESC*> geomDescr(numGeometries);
     std::vector<char> geomDescrData(totalBytes);
-
+#endif
+    
     size_t offset = 0;
-    for (size_t geomID=0; geomID<scene->size(); geomID++)
+    for (size_t geomID=0; geomID<numGeometries; geomID++) // FIXME: slow if numGeometries is large
     {
       geomDescr[geomID] = nullptr;     
       Geometry* geom = scene->get(geomID);
@@ -537,24 +576,37 @@ namespace embree
       createGeometryDesc(&geomDescrData[offset],scene,scene->get(geomID),type);
       geomDescr[geomID] = (RTHWIF_GEOMETRY_DESC*) &geomDescrData[offset];
       offset += sizeof_RTHWIF_GEOMETRY(type);
+#if !defined(EMBREE_SYCL_GPU_BVH_BUILDER)      
       assert(offset <= geomDescrData.size());
+#endif      
     }
 
-    RTHWIF_PARALLEL_OPERATION parallelOperation = rthwifNewParallelOperation();
-
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    RTHWIF_PARALLEL_OPERATION parallelOperation = nullptr;    
+#else
+    RTHWIF_PARALLEL_OPERATION parallelOperation = rthwifNewParallelOperation();    
+#endif    
     /* estimate static accel size */
     BBox1f time_range(0,1);
     RTHWIF_AABB bounds;
     RTHWIF_BUILD_ACCEL_ARGS args;
     memset(&args,0,sizeof(args));
     args.structBytes = sizeof(args);
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    args.geometries = (const RTHWIF_GEOMETRY_DESC**) geomDescr;    
+#else    
     args.geometries = (const RTHWIF_GEOMETRY_DESC**) geomDescr.data();
-    args.numGeometries = geomDescr.size();
+#endif    
+    args.numGeometries = numGeometries;
     args.accelBuffer = nullptr;
     args.accelBufferBytes = 0;
     args.scratchBuffer = nullptr;
     args.scratchBufferBytes = 0;
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)    
+    args.quality = gpu_device->quality_flags == RTC_BUILD_QUALITY_LOW ? RTHWIF_BUILD_QUALITY_LOW : RTHWIF_BUILD_QUALITY_MEDIUM;
+#else      
     args.quality = convertBuildQuality(scene->quality_flags);
+#endif    
     args.flags = convertBuildFlags(scene->scene_flags);
     args.parallelOperation = parallelOperation;
     args.boundsOut = &bounds;
@@ -566,15 +618,28 @@ namespace embree
     RTHWIF_ACCEL_SIZE sizeTotal;
     memset(&sizeTotal,0,sizeof(RTHWIF_ACCEL_SIZE));
     sizeTotal.structBytes = sizeof(RTHWIF_ACCEL_SIZE);
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    RTHWIF_ERROR err = rthwifGetAccelSizeGPU(args,sizeTotal,&sycl_queue,gpu_device->verbose);    
+#else    
     RTHWIF_ERROR err = rthwifGetAccelSize(args,sizeTotal);
+#endif    
     if (err != RTHWIF_ERROR_NONE)
       throw_RTCError(RTC_ERROR_UNKNOWN,"BVH size estimate failed");
 
     /* allocate scratch buffer */
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    // === scratch buffer === 
+    char *scratchBuffer  = (char*)sycl::aligned_alloc(64,sizeTotal.scratchBufferBytes,gpu_device->getGPUDevice(),gpu_device->getGPUContext(),gpu_device->verbose > 1 ? sycl::usm::alloc::shared : sycl::usm::alloc::device);
+    assert(scratchBuffer);
+    args.scratchBuffer = scratchBuffer;
+    args.scratchBufferBytes = sizeTotal.scratchBufferBytes;
+#else        
     std::vector<char> scratchBuffer(sizeTotal.scratchBufferBytes);
     args.scratchBuffer = scratchBuffer.data();
     args.scratchBufferBytes = scratchBuffer.size();
-
+#endif
+    
     size_t headerBytes = sizeof(EmbreeHWAccel) + std::max(1u,maxTimeSegments)*8;
     align(headerBytes,128);
 
@@ -587,36 +652,55 @@ namespace embree
       size.accelBufferExpectedBytes  = maxTimeSegments*sizeTotal.accelBufferExpectedBytes;
       size.accelBufferWorstCaseBytes = maxTimeSegments*sizeTotal.accelBufferWorstCaseBytes;
       size_t bytes = headerBytes+size.accelBufferExpectedBytes;
-        
+      
       /* allocate BVH data */
       if (accel.size() < bytes) accel.resize(bytes);
+      
+#if !defined(EMBREE_SYCL_GPU_BVH_BUILDER)      
       memset(accel.data(),0,accel.size()); // FIXME: not required
-
+#endif
+      
       /* build BVH for each time segment */
       for (uint32_t i=0; i<maxTimeSegments; i++)
       {
         const float t0 = float(i+0)/float(maxTimeSegments);
         const float t1 = float(i+1)/float(maxTimeSegments);
         time_range = BBox1f(t0,t1);
-        
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+        args.geometries = (const RTHWIF_GEOMETRY_DESC**) geomDescr;
+#else        
         args.geometries = (const RTHWIF_GEOMETRY_DESC**) geomDescr.data();
-        args.numGeometries = geomDescr.size();
-        args.accelBuffer = accel.data() + headerBytes + i*sizeTotal.accelBufferExpectedBytes;
-        args.accelBufferBytes = sizeTotal.accelBufferExpectedBytes;
-        bounds = { { INFINITY, INFINITY, INFINITY }, { -INFINITY, -INFINITY, -INFINITY } };
+#endif        
+        args.numGeometries = numGeometries; 
+        args.accelBuffer = accel.data() + headerBytes + i*sizeTotal.accelBufferExpectedBytes; // why + headerBytes?
         
+        args.accelBufferBytes = sizeTotal.accelBufferExpectedBytes;
+        bounds = { { INFINITY, INFINITY, INFINITY }, { -INFINITY, -INFINITY, -INFINITY } };  // why does the host initializes the bounds
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+        err = rthwifPrefetchAccelGPU(args,&sycl_queue,gpu_device->verbose); //triggers additional USM transfers
+#endif
+        
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)        
+        err = rthwifBuildAccelGPU(args,&sycl_queue,gpu_device->verbose);
+#else
         err = rthwifBuildAccel(args);
+        
+#endif        
+        
         if (args.parallelOperation) {
           assert(err == RTHWIF_ERROR_PARALLEL_OPERATION);
           uint32_t maxThreads = rthwifGetParallelOperationMaxConcurrency(parallelOperation);
           parallel_for(maxThreads, [&](uint32_t) { err = rthwifJoinParallelOperation(parallelOperation); });
         }
-        
-        fullBounds.extend(*(BBox3f*) args.boundsOut);
+
+        if (err == RTHWIF_ERROR_NONE) // added if
+          fullBounds.extend(*(BBox3f*) args.boundsOut);
 
         if (err == RTHWIF_ERROR_RETRY)
         {
-          if (sizeTotal.accelBufferExpectedBytes == sizeTotal.accelBufferWorstCaseBytes)
+          if (sizeTotal.accelBufferExpectedBytes == sizeTotal.accelBufferWorstCaseBytes) 
             throw_RTCError(RTC_ERROR_UNKNOWN,"build error");
           
           sizeTotal.accelBufferExpectedBytes = std::min(sizeTotal.accelBufferWorstCaseBytes,(size_t(1.2*sizeTotal.accelBufferExpectedBytes)+127)&-128);
@@ -631,12 +715,32 @@ namespace embree
     if (err != RTHWIF_ERROR_NONE)
       throw_RTCError(RTC_ERROR_UNKNOWN,"build error");
 
+    // === moving this to device code prevents USM down and up transfer of accel ===
+
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    EmbreeHWAccel hwaccel;
+    hwaccel.numTimeSegments = maxTimeSegments;
+    for (size_t i=0; i<maxTimeSegments; i++)
+      hwaccel.AccelTable[i] = (char*)accel.data() + headerBytes + i*sizeTotal.accelBufferExpectedBytes;    
+    sycl::event queue_event =  sycl_queue.memcpy(accel.data(),&hwaccel,sizeof(EmbreeHWAccel));
+    queue_event.wait();
+#else    
     EmbreeHWAccel* hwaccel = (EmbreeHWAccel*) accel.data();
     hwaccel->numTimeSegments = maxTimeSegments;
 
     for (size_t i=0; i<maxTimeSegments; i++)
       hwaccel->AccelTable[i] = (char*)hwaccel + headerBytes + i*sizeTotal.accelBufferExpectedBytes;
+#endif
+    
+    // =============================================================================
+    
+#if defined(EMBREE_SYCL_GPU_BVH_BUILDER)
+    sycl::free(geomDescr        ,gpu_device->getGPUContext());
+    sycl::free(geomDescrData    ,gpu_device->getGPUContext());    
+    sycl::free(scratchBuffer    ,gpu_device->getGPUContext());
+#endif
 
     return fullBounds;
   }
+
 }

tutorials/common/scenegraph/CMakeLists.txt

@@ -7,6 +7,9 @@ ADD_LIBRARY(scenegraph STATIC
     xml_writer.cpp
     obj_loader.cpp
     ply_loader.cpp
+#    rtas_loader.cpp
+#    RTASFile.cpp
+#    RTASToEmbree.cpp    
     corona_loader.cpp
     texture.cpp
     scenegraph.cpp

tutorials/common/scenegraph/scenegraph.cpp

@@ -7,6 +7,7 @@
 #include "obj_loader.h"
 #include "ply_loader.h"
 #include "corona_loader.h"
+//#include "rtas_loader.h"
 
 namespace embree
 {
@@ -20,6 +21,7 @@ namespace embree
     else if (toLowerCase(filename.ext()) == std::string("ply" )) return loadPLY(filename);
     else if (toLowerCase(filename.ext()) == std::string("xml" )) return loadXML(filename);
     else if (toLowerCase(filename.ext()) == std::string("scn" )) return loadCorona(filename);
+    //else if (toLowerCase(filename.ext()) == std::string("rtas" )) return loadRTAS(filename);    
     else throw std::runtime_error("unknown scene format: " + filename.ext());
   }
 

tutorials/common/tutorial/tutorial.cpp

@@ -25,6 +25,7 @@
 #include "../scenegraph/geometry_creation.h"
 #include "../scenegraph/obj_loader.h"
 #include "../scenegraph/xml_loader.h"
+//#include "../scenegraph/rtas_loader.h"
 #include "../image/image.h"
 
 #if defined(EMBREE_SYCL_SUPPORT) && defined(EMBREE_SYCL_TUTORIAL)
@@ -644,7 +645,12 @@ namespace embree
     initRayStats();
     
     for (unsigned int i=0; i<numFrames; i++)
+    {
+      double t0 = getSeconds();
       render(pixels,width,height,render_time,ispccamera);
+      double t1 = getSeconds();
+      std::cout << "Render time = " << (t1-t0)*1000.0 << " ms => " << 1.0 / (t1-t0) << " fps" << std::endl;
+    }
     
     Ref<Image> image = new Image4uc(width, height, (Col4uc*)pixels);
     storeImage(image, fileName);
@@ -1226,6 +1232,10 @@ namespace embree
         scene->add(loadOBJ(keyFramesFilenames[i],subdiv_mode != "",true));
       else if (keyFramesFilenames[i].ext() != "")
         scene->add(SceneGraph::load(keyFramesFilenames[i]));
+
+      //if (toLowerCase(keyFramesFilenames[i].ext()) == std::string("rtas"))
+      //  scene->add(loadRTAS(keyFramesFilenames[i],true));
+
       
       if (verbosity >= 1) 
         std::cout << " [DONE]" << std::endl << std::flush;

tutorials/common/tutorial/tutorial_device.h

@@ -40,6 +40,7 @@ inline sycl::nd_range<2> make_nd_range(unsigned int size0, unsigned int size1)
 
   return sycl::nd_range(sycl::range(size0,size1),wg_size);
 }
+
 #endif
 
 enum Shader { 

tutorials/dynamic_scene/dynamic_scene_device.cpp

@@ -264,12 +264,74 @@ void animateSphere (int id, float time)
     v->z = pos.z+r*sin(f*phif)*cos(thetaf);
   }
 #endif
-
   /* commit mesh */
   rtcUpdateGeometryBuffer(geom,RTC_BUFFER_TYPE_VERTEX,0);
   rtcCommitGeometry(geom);
 }
 
+#if defined(EMBREE_SYCL_TUTORIAL)
+
+struct State {
+  Vertex* vertices;
+  float pos_x, pos_y, pos_z;
+  float r,f;
+};
+
+__forceinline size_t alignTo(const size_t size, const size_t alignment) { return ((size+alignment-1)/alignment)*alignment; }
+
+void animateSphereGPU (const uint numSpheres, float time)
+{
+  State state[NUM_SPHERES];
+
+  for (uint id=0;id<numSpheres;id++)
+  {
+    /* animate vertices */
+    RTCGeometry geom = rtcGetGeometry(data.g_scene,id);
+    state[id].vertices = (Vertex*) rtcGetGeometryBufferData(geom,RTC_BUFFER_TYPE_VERTEX,0);
+    state[id].pos_x = Vec3fa(data.position[id]).x;
+    state[id].pos_y = Vec3fa(data.position[id]).y;
+    state[id].pos_z = Vec3fa(data.position[id]).z;    
+    state[id].r = data.radius[id];
+    state[id].f = 2.0f*(1.0f+0.5f*sinf(time+id));
+  }
+  
+  const float rcpNumTheta = rcp((float)data.numTheta);
+  const float rcpNumPhi   = rcp((float)data.numPhi);
+  const int data_numPhi   = data.numPhi;  
+  const int data_numTheta = data.numTheta;
+  const float PI = float(pi);
+  
+  sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
+                                                 const sycl::nd_range<3> nd_range(sycl::range<3>(alignTo(data_numPhi+1,4),alignTo(data_numTheta,4),alignTo(numSpheres,4)),sycl::range<3>(4,4,1));
+                                                       cgh.parallel_for(nd_range,[=](sycl::nd_item<3> item) {                                
+                                                                        const unsigned int phi   = item.get_global_id(0); 
+                                                                        const unsigned int theta = item.get_global_id(1);
+                                                                        const unsigned int id    = item.get_global_id(2);
+                                                                        const State &s = state[id];
+                                                                        if (phi < data_numPhi+1 && theta < data_numTheta && id < numSpheres)
+                                                                        {
+                                                                          Vertex* v = &s.vertices[phi*data_numTheta+theta];
+                                                                          const float phif   = phi*PI*rcpNumPhi;
+                                                                          const float thetaf = theta*2.0f*PI*rcpNumTheta;
+                                                                          v->x = s.pos_x+s.r*sinf(s.f*phif)*sinf(thetaf);
+                                                                          v->y = s.pos_y+s.r*cosf(phif);
+                                                                          v->z = s.pos_z+s.r*sinf(s.f*phif)*cosf(thetaf);
+                                                                        }
+                                                                      });
+                                                     });
+  
+  event.wait_and_throw();
+
+  for (uint id=0;id<numSpheres;id++)  
+  {
+    /* commit mesh */
+    RTCGeometry geom = rtcGetGeometry(data.g_scene,id);    
+    rtcUpdateGeometryBuffer(geom,RTC_BUFFER_TYPE_VERTEX,0);
+    rtcCommitGeometry(geom);
+  }  
+}
+#endif
+
 extern "C" void renderFrameStandard (int* pixels,
                           const unsigned int width,
                           const unsigned int height,
@@ -313,9 +375,13 @@ extern "C" void device_render (int* pixels,
                            const float time,
                            const ISPCCamera& camera)
 {
-  /* animate sphere */
+  /* animate sphere */  
+#if defined(EMBREE_SYCL_TUTORIAL)
+  animateSphereGPU(data.numSpheres,time);
+#else  
   for (int i=0; i<data.numSpheres; i++)
     animateSphere(i,time+i);
+#endif  
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);