Update rtcore_scene.h (#555)

Add missing forward declaration for RTCGeometry in rtcore_scene.h

.github/workflows/dpcpp-none.env

@@ -1 +0,0 @@
-DPCPP_VERSION=

.github/workflows/dpcpp-sycl-nightly.env

@@ -1 +1 @@
-DPCPP_VERSION=sycl-nightly/20230417
+DPCPP_VERSION=intel-llvm/v6.2.0

.github/workflows/dpcpp-sycl-release.env

@@ -1 +0,0 @@
-DPCPP_VERSION=sycl-nightly/20230417

.github/workflows/gfx-ubuntu22-embargo.env

@@ -1 +0,0 @@
-GFX_DRIVER_VERSION=gfx-driver-builds/ci/master/gfx-driver-ci-master-13324/artifacts/Linux/Ubuntu/22.04/Release

.github/workflows/gfx-ubuntu22-internal.env

@@ -1 +1 @@
-GFX_DRIVER_VERSION=neo-builds/ci/master/ci-neo-master-025905/artifacts/linux/ubuntu/22.04
+GFX_DRIVER_VERSION=neo-builds/ci/master/ci-neo-master-033051/artifacts/linux/ubuntu/22.04

.github/workflows/gfx-ubuntu22-public.env

@@ -1 +1 @@
-GFX_DRIVER_VERSION=neo-builds/releases/releases_23.05/neo-releases-23.05.25593.11/artifacts/linux/ubuntu/22.04
+GFX_DRIVER_VERSION=linux-latest

.github/workflows/gfx-ubuntu22-release.env

@@ -1 +0,0 @@
-GFX_DRIVER_VERSION=neo-builds/releases/releases_23.05/neo-releases-23.05.25593.11/artifacts/linux/ubuntu/20.04

.github/workflows/gfx-ubuntu24-internal.env

@@ -0,0 +1 @@
+GFX_DRIVER_VERSION=neo-builds/ci/master/ci-neo-master-033048/artifacts/linux/ubuntu/24.04

.github/workflows/gfx-ubuntu24-public.env

@@ -0,0 +1 @@
+GFX_DRIVER_VERSION=linux-latest

.github/workflows/gfx-windows-internal.env

@@ -1 +1 @@
-GFX_DRIVER_VERSION=gfx-driver-builds/ci/master/gfx-driver-ci-master-13324/artifacts/Windows
+GFX_DRIVER_VERSION=gfx-driver-ci-master-20127

.github/workflows/gfx-windows-public.env

@@ -1 +1 @@
-GFX_DRIVER_VERSION=gfx-driver-builds/prod/releases_23ww08-ci-master-13370/prod-hini-releases_23ww08-ci-master-13370-revenue-pr-1014250/artifacts/Windows/
+GFX_DRIVER_VERSION=windows-latest

.github/workflows/gfx-windows-release.env

@@ -1 +0,0 @@
-GFX_DRIVER_VERSION=gfx-driver-builds/prod/releases_23ww08-ci-master-13370/prod-hini-releases_23ww08-ci-master-13370-revenue-pr-1014250/artifacts/Windows/

.github/workflows/performance.yml

@@ -1,155 +0,0 @@
-## Copyright 2022 Intel Corporation
-## SPDX-License-Identifier: Apache-2.0
-
-name: performance
-
-on: 
-  schedule:
-    - cron:  '00 01 * * *'
-  workflow_dispatch:
-
-jobs:
-
-  ########################################
-  # Linux build
-  ########################################
-
-  linux-build-SYCL-DPCPP:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@devel
-    with:
-      image: embree/ubuntu:20.04
-      options: --device=/dev/dri:/dev/dri
-      runs-on: '[ "Linux", "docker", "dg2" ]'
-      project: embree
-      env-from-files: ./.github/workflows/dpcpp-none.env ./.github/workflows/gfx-ubuntu22-release.env
-      artifact-out: linux-SYCL-DPCPP
-      artifact-path: ./build/*.tar.gz
-      cmd: |
-        source scripts/cmake-presets/linux-env.sh
-        source ${NAS_LINUX}/intel/oneAPI/compiler/2023.0.0/env/vars.sh
-        $CMAKE_EXE --preset package-linux-icx-sycl-ispc1_19_0-tbb2021_9_0-perf
-        $CMAKE_EXE --build build --config Release --target package
-        $CMAKE_EXE --build build --target post_package
-
-  linux-build-x86-ICX-ISPC:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/embree_docker.yml@main
-    with:
-      image: embree/rockylinux:8.7
-      runs-on: '[ "Linux", "docker", "avx512" ]'
-      artifact-out: linux-build-x86-ICX-ISPC
-      artifact-path: ./build/*.tar.gz
-      cmd: |
-        source scripts/cmake-presets/linux-env.sh
-        source $NAS_LINUX/intel/oneAPI/compiler/2023.1.0/env/vars.sh
-        $CMAKE_EXE --preset package-linux-icx-ispc1_19_0-tbb2021_9_0-perf
-        $CMAKE_EXE --build build --config Release --target package
-        $CMAKE_EXE --build build --target post_package
-
-  linux-test-DG2-A770-perf:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@main
-    needs: ["linux-build-SYCL-DPCPP"]
-    with:
-      image: embree/ubuntu:22.04
-      options: --device=/dev/dri:/dev/dri
-      runs-on: '[ "Linux", "docker", "dg2", "perf", "A770" ]'
-      project: embree
-      env-from-files: ./.github/workflows/dpcpp-none.env ./.github/workflows/gfx-ubuntu22-release.env
-      artifact-in: linux-SYCL-DPCPP
-      setup-benny: true
-      cmd: |
-        source <(scripts/test.py env --debug)
-        cd build
-        tar -xf embree-?.?.?.sycl.x86_64.linux.tar.gz
-        cd bin
-        chmod +x ../embree-vars.sh
-        ../embree-vars.sh
-        $GITHUB_WORKSPACE/scripts/run-perf.sh gpu
-
-
-  linux-test-PVC-perf:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@main
-    needs: ["linux-build-SYCL-DPCPP"]
-    with:
-      image: embree/ubuntu:22.04
-      options: --device=/dev/dri:/dev/dri
-      runs-on: '[ "Linux", "docker", "pvc", "perf" ]'
-      project: embree
-      env-from-files: ./.github/workflows/dpcpp-none.env ./.github/workflows/gfx-ubuntu22-release.env
-      artifact-in: linux-SYCL-DPCPP
-      setup-benny: true
-      cmd: |
-        export NEOReadDebugKeys=1
-        # Run it on single tile (disable implicit scaling)
-        export EnableImplicitScaling=0
-        source <(scripts/test.py env --debug)
-        cd build
-        tar -xf embree-?.?.?.sycl.x86_64.linux.tar.gz
-        cd bin
-        chmod +x ../embree-vars.sh
-        ../embree-vars.sh
-        $GITHUB_WORKSPACE/scripts/run-perf.sh gpu
-
-
-  linux-test-i9-12900k:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/perf.yml@devel
-    needs: ["linux-build-x86-ICX-ISPC"]
-    with:
-      runs-on: '[ "Linux", "perf", "vis-perf-i9-12900k-1" ]'
-      artifact-in: linux-build-x86-ICX-ISPC
-      setup-benny: true
-      cmd: |
-        cd build
-        tar -xf embree-?.?.?.x86_64.linux.tar.gz
-        cd bin
-        $GITHUB_WORKSPACE/scripts/run-perf.sh cpu
-
-
-  linux-test-x8380-1:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/perf.yml@main
-    needs: ["linux-build-x86-ICX-ISPC"]
-    with:
-      runs-on: '[ "Linux", "perf", "vis-perf-x8380-1" ]'
-      artifact-in: linux-build-x86-ICX-ISPC
-      setup-benny: true
-      cmd: |
-        cd build
-        tar -xf embree-?.?.?.x86_64.linux.tar.gz
-        cd bin
-        $GITHUB_WORKSPACE/scripts/run-perf.sh cpu
-
-
-  linux-test-x8280-1:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/perf.yml@main
-    needs: ["linux-build-x86-ICX-ISPC"]
-    with:
-      runs-on: '[ "Linux", "perf", "vis-perf-x8280-1" ]'
-      artifact-in: linux-build-x86-ICX-ISPC
-      setup-benny: true
-      cmd: |
-        cd build
-        tar -xf embree-?.?.?.x86_64.linux.tar.gz
-        cd bin
-        $GITHUB_WORKSPACE/scripts/run-perf.sh cpu
-
-
-  linux-test-a3970x-1:
-    secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/perf.yml@main
-    needs: ["linux-build-x86-ICX-ISPC"]
-    with:
-      runs-on: '[ "Linux", "perf", "vis-perf-a3970x-1" ]'
-      artifact-in: linux-build-x86-ICX-ISPC
-      setup-benny: true
-      cmd: |
-        cd build
-        tar -xf embree-?.?.?.x86_64.linux.tar.gz
-        cd bin
-        $GITHUB_WORKSPACE/scripts/run-perf.sh cpu
-

.github/workflows/public.continuous.yml

@@ -0,0 +1,116 @@
+## Copyright 2024 Intel Corporation
+## SPDX-License-Identifier: Apache-2.0
+
+name: continuous (public)
+
+on:
+  push:
+  pull_request:
+  workflow_dispatch:
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.ref }}
+  cancel-in-progress: true
+
+permissions: read-all
+
+jobs:
+  rocky-8:
+    runs-on: ubuntu-latest
+    container:
+      image: rockylinux:8
+    steps:
+      - name: Install packages (rocky 8)
+        run: |
+          yum -y install epel-release
+          yum -y install gcc-c++ clang tar git-lfs make tbb-devel freeglut-devel glfw-devel libXmu-devel libXi-devel python3
+          yum -y install cmake
+          yum -y install libXcursor-devel libXi-devel libXinerama-devel libXrandr-devel
+          yum -y install ncurses-compat-libs 
+          yum -y install environment-modules
+          dnf makecache --refresh
+          dnf config-manager --set-enabled powertools
+          dnf -y update
+          dnf -y install ninja-build
+
+      - name: Checkout Repository
+        uses: actions/checkout@v4
+
+      - name: Build and Run
+        run: |
+          mkdir build
+          cd build
+          cmake -DBUILD_TESTING=ON -DEMBREE_TUTORIALS=ON -DEMBREE_ISPC_SUPPORT=OFF -DEMBREE_TESTING_INTENSITY=3 ..
+          make -j$(nproc)
+          ctest
+
+  ubuntu-24-tbb-and-ispc-from-os:
+    runs-on: ubuntu-latest
+    container:
+      image: ubuntu:24.04
+    steps:
+      - name: Install packages (ubuntu 24.04)
+        run: |
+          apt-get update
+          apt-get -y upgrade
+          apt-get -y install g++ cmake git-lfs freeglut3-dev libglfw3-dev libxmu-dev libxi-dev unzip xorg-dev pip libtbb-dev ispc
+          apt-get -y install libopenimageio-dev
+
+      - name: Checkout Repository
+        uses: actions/checkout@v4
+
+      - name: Build and Run
+        run: |
+          mkdir build
+          cd build
+          cmake -DBUILD_TESTING=ON -DEMBREE_TUTORIALS=ON -DEMBREE_ISPC_SUPPORT=ON -DEMBREE_TESTING_INTENSITY=3 ..
+          make -j$(nproc)
+          ctest
+
+  ubuntu-24-tbb-and-ispc-prebuilt:
+    runs-on: ubuntu-latest
+    container:
+      image: ubuntu:24.04
+    steps:
+      - name: Install packages (ubuntu 24.04)
+        run: |
+          apt-get update
+          apt-get -y upgrade
+          apt-get -y install g++ cmake git-lfs freeglut3-dev libglfw3-dev libxmu-dev libxi-dev unzip xorg-dev pip
+          apt-get -y install libopenimageio-dev
+
+      - name: Checkout Repository
+        uses: actions/checkout@v4
+
+      - name: Build and Run
+        run: |
+          mkdir build
+          cd build
+          cmake -D EMBREE_EXTRA_OPTIONS="-DBUILD_TESTING=ON -DEMBREE_TUTORIALS=ON -DEMBREE_ISPC_SUPPORT=ON -DEMBREE_TESTING_INTENSITY=3" ../superbuild
+          cmake --build .
+          cd embree/build
+          ctest
+
+  ubuntu-24-tbb-src-ispc-prebuilt:
+    runs-on: ubuntu-latest
+    container:
+      image: ubuntu:24.04
+    steps:
+      - name: Install packages (ubuntu 24.04)
+        run: |
+          apt-get update
+          apt-get -y upgrade
+          apt-get -y install g++ cmake git-lfs freeglut3-dev libglfw3-dev libxmu-dev libxi-dev unzip xorg-dev pip
+          apt-get -y install libopenimageio-dev
+
+      - name: Checkout Repository
+        uses: actions/checkout@v4
+
+      - name: Build and Run
+        run: |
+          mkdir build
+          cd build
+          cmake -D BUILD_TBB_FROM_SOURCE=ON -D EMBREE_EXTRA_OPTIONS="-DBUILD_TESTING=ON -DEMBREE_TUTORIALS=ON -DEMBREE_ISPC_SUPPORT=ON -DEMBREE_TESTING_INTENSITY=3" ../superbuild
+          cmake --build .
+          cd embree/build
+          ctest

.github/workflows/release.yml

@@ -3,190 +3,379 @@
 
 name: release
 
+permissions: read-all
+
 on:
   schedule:
-    - cron:  '00 00 * * *'
+    - cron:  '01 00 * * *'
   workflow_dispatch:
 
 jobs:
 
-#####################################################################
-# Release Builds
 ####################################################################
-  release-windows-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP:
+# CPU Builds and Tests
+####################################################################
+  build-windows:
     secrets: inherit
     uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/windows.yml@main
     with:
-      runs-on: '[ "Windows", "NAS", "build", "avx512" ]'
-      artifact-out: release-windows-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      runs-on: '[ "Windows", "build"]'
+      artifact-out: build-windows
       artifact-path: ./build/embree-*.zip
-      shell: cmd
+      dpcpp-version: intel-llvm/v6.2.0
       cmd: |
-        call "C:\\Program Files (x86)\\Intel\\oneAPI\\compiler\\2023.1.0\\env\\vars.bat"
-        cmake --preset package-windows-icx-ispc1_19_0-tbb2021_9_0
-        cmake --build build --config Release --target package
-        cmake --build build --target post_package
-        cmake --build build --target test_package
+        $opts  = "-DCMAKE_BUILD_TYPE=Release"
+        $opts += " -DBUILD_TESTING=ON"
+        $opts += " -DEMBREE_TUTORIALS=ON"
+        $opts += " -DEMBREE_ISPC_SUPPORT=ON"
+        $opts += " -DEMBREE_TESTING_INTENSITY=3"
+        $opts += " -DEMBREE_MAX_ISA=AVX512"
+        $opts += " -DEMBREE_INSTALL_DEPENDENCIES=ON"
+        $opts += " -DEMBREE_TESTING_INSTALL_TESTS=ON"
+        $opts += " -DCMAKE_INSTALL_TESTDIR=testing"
+        $opts += " -DEMBREE_ZIP_MODE=ON"
 
+        cmake -S superbuild -B build -G Ninja -D EMBREE_EXTRA_OPTIONS="$opts"
+        cmake --build build --config Release
+        cmake --build build/embree/build --config Release --target build
+        Copy-Item build/embree/build/embree-*.zip build/ -Force
 
-  release-macosx-x64-Release-CLANG-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP:
+  windows-test:
     secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/macos.yml@devel
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/windows.yml@main
+    needs: ["build-windows"]
+    strategy:
+      fail-fast: false
+      matrix:
+        platform: [avx2, avx512]
+    with:
+      runs-on: '[ "Windows", "${{ matrix.platform }}" ]'
+      artifact-in: build-windows
+      cmd: |
+        $EmbreeDir = .\\scripts\\prepare-test-files.ps1 -TestType ctest
+        .\\scripts\\run-tests.ps1 -TestType ctest -EmbreeDir $EmbreeDir -CtestConfig Release
+
+  build-macosx-x86_64:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/macos.yml@main
     with:
       runs-on: '[ "macOS", "build", "x86_64" ]'
-      artifact-out: release-macosx-x64-Release-CLANG-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-path: ./build/*.zip
       cmd: |
-        source scripts/cmake-presets/macos-env.sh
-        $CMAKE_EXE --preset package-macos-gcc-ispc1_19_0-tbb2021_9_0
-        $CMAKE_EXE --build build --config Release --target package
-        $CMAKE_EXE --build build --target post_package
+        opts="-DCMAKE_BUILD_TYPE=Release"
+        opts+=" -DBUILD_TESTING=ON"
+        opts+=" -DEMBREE_TUTORIALS=ON"
+        opts+=" -DEMBREE_ISPC_SUPPORT=ON"
+        opts+=" -DEMBREE_TESTING_INTENSITY=3"
+        opts+=" -DCMAKE_CXX_COMPILER=clang++"
+        opts+=" -DCMAKE_C_COMPILER=clang"
+        opts+=" -DCMAKE_INSTALL_BINDIR=bin"
+        opts+=" -DCMAKE_INSTALL_DOCDIR=doc"
+        opts+=" -DCMAKE_INSTALL_INCLUDEDIR=include"
+        opts+=" -DCMAKE_INSTALL_TESTDIR=testing"
+        opts+=" -DCMAKE_MACOSX_RPATH=ON"
+        opts+=" -DCMAKE_SKIP_INSTALL_RPATH=OFF"
+        opts+=" -DEMBREE_BUILD_GLFW_FROM_SOURCE=ON"
+        opts+=" -DEMBREE_INSTALL_DEPENDENCIES=ON"
+        opts+=" -DEMBREE_MAX_ISA=AVX512"
+        opts+=" -DEMBREE_PATCH_GLFW_SOURCE=ON"
+        opts+=" -DEMBREE_TASKING_SYSTEM=TBB"
+        opts+=" -DEMBREE_TESTING_INSTALL_TESTS=ON"
+        opts+=" -DEMBREE_ZIP_MODE=ON"
+        opts+=" -DCPACK_PRE_BUILD_SCRIPTS=$CPACK_PRE_BUILD_SIGNING_SCRIPT"
 
+        cmake -S superbuild -B build -D EMBREE_EXTRA_OPTIONS="$opts"
+        cmake --build build --config Release
+        cmake --build build/embree/build --config Release --target build
+        cp build/embree/build/embree-*.zip build/
+      notarize-package: ./build/embree-*.x86_64.macosx.zip
+      artifact-out: build-macosx-x86_64
+      artifact-path: ./build/embree-*.zip
 
-  release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP:
+  test-macosx-x86_64:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/macos.yml@main
+    needs: ["build-macosx-x86_64"]
+    strategy:
+      fail-fast: false
+      matrix:
+        platform: [arm, x86_64]
+    with:
+      runs-on: '[ "macOS", "${{ matrix.platform }}" ]'
+      artifact-in: build-macosx-x86_64
+      cmd: |
+        EMBREE_PATH=$(./scripts/prepare-test-files-macos.sh --type ctest)
+        ./scripts/run-tests.sh --type ctest --embree-dir "$EMBREE_PATH" --ctest-config Release
+
+  build-macosx-arm:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/macos.yml@main
+    with:
+      runs-on: '[ "macOS", "build", "arm" ]'
+      cmd: |
+        opts="-DCMAKE_BUILD_TYPE=Release"
+        opts+=" -DBUILD_TESTING=ON"
+        opts+=" -DEMBREE_TUTORIALS=ON"
+        opts+=" -DEMBREE_ISPC_SUPPORT=ON"
+        opts+=" -DEMBREE_TESTING_INTENSITY=3"
+        opts+=" -DCMAKE_CXX_COMPILER=clang++"
+        opts+=" -DCMAKE_C_COMPILER=clang"
+        opts+=" -DCMAKE_INSTALL_BINDIR=bin"
+        opts+=" -DCMAKE_INSTALL_DOCDIR=doc"
+        opts+=" -DCMAKE_INSTALL_INCLUDEDIR=include"
+        opts+=" -DCMAKE_INSTALL_TESTDIR=testing"
+        opts+=" -DCMAKE_MACOSX_RPATH=ON"
+        opts+=" -DCMAKE_SKIP_INSTALL_RPATH=OFF"
+        opts+=" -DEMBREE_BUILD_GLFW_FROM_SOURCE=ON"
+        opts+=" -DEMBREE_INSTALL_DEPENDENCIES=ON"
+        opts+=" -DEMBREE_MAX_ISA=NEON2X"
+        opts+=" -DEMBREE_PATCH_GLFW_SOURCE=ON"
+        opts+=" -DEMBREE_TASKING_SYSTEM=TBB"
+        opts+=" -DEMBREE_TESTING_INSTALL_TESTS=ON"
+        opts+=" -DEMBREE_ZIP_MODE=ON"
+        opts+=" -DCPACK_PRE_BUILD_SCRIPTS=$CPACK_PRE_BUILD_SIGNING_SCRIPT"
+
+        cmake -S superbuild -B build -D BUILD_TBB_FROM_SOURCE=ON -D EMBREE_EXTRA_OPTIONS="$opts"
+        cmake --build build --config Release
+        cmake --build build/embree/build --config Release --target build
+        cp build/embree/build/embree-*.zip build/
+      notarize-package: build/embree-*.arm64.macosx.zip
+      artifact-out: build-macosx-arm
+      artifact-path: ./build/embree-*.zip
+
+  test-macosx-arm:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/macos.yml@main
+    needs: ["build-macosx-arm"]
+    strategy:
+      fail-fast: false
+      matrix:
+        platform: [arm]
+    with:
+      runs-on: '[ "macOS", "${{ matrix.platform }}" ]'
+      artifact-in: build-macosx-arm
+      cmd: |
+        EMBREE_PATH=$(./scripts/prepare-test-files-macos.sh --type ctest)
+        ./scripts/run-tests.sh --type ctest --embree-dir "$EMBREE_PATH" --ctest-config Release
+
+  build-linux:
     secrets: inherit
     uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker.yml@main
     with:
-      image: embree/rockylinux:8.7
-      runs-on: '[ "Linux", "docker", "avx512" ]'
-      artifact-out: release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      image: rockylinux:8.7
+      runs-on: '[ "Linux", "build" ]'
+      artifact-out: build-linux
       artifact-path: ./build/*.tar.gz
+      dpcpp-version: intel-llvm/v6.2.0
+      pre-cmd: |
+        dnf -y install epel-release
+        dnf -y config-manager --set-enabled powertools
+        dnf -y install ninja-build
       cmd: |
-        source scripts/cmake-presets/linux-env.sh
-        source $NAS_LINUX/intel/oneAPI/compiler/2023.1.0/env/vars.sh
-        $CMAKE_EXE --preset package-linux-icx-ispc1_19_0-tbb2021_9_0
-        $CMAKE_EXE --build build --config Release --target package
-        $CMAKE_EXE --build build --target post_package
-        $CMAKE_EXE --build build --target test_package
+        opts="-DCMAKE_BUILD_TYPE=Release"
+        opts+=" -DBUILD_TESTING=ON"
+        opts+=" -DEMBREE_TUTORIALS=ON"
+        opts+=" -DEMBREE_ISPC_SUPPORT=ON"
+        opts+=" -DEMBREE_TESTING_INTENSITY=3"
+        opts+=" -DEMBREE_MAX_ISA=AVX512"
+        opts+=" -DEMBREE_INSTALL_DEPENDENCIES=ON"
+        opts+=" -DEMBREE_TESTING_INSTALL_TESTS=ON"
+        opts+=" -DCMAKE_INSTALL_TESTDIR=testing"
+        opts+=" -DEMBREE_ZIP_MODE=ON"
+
+        cmake -S superbuild -B build -D EMBREE_EXTRA_OPTIONS="$opts"
+        cmake --build build --config Release
+        cmake --build build/embree/build --config Release --target build
+        cp build/embree/build/embree-*.tar.gz build/
+
+  linux-test:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker.yml@main
+    needs: ["build-linux"]
+    strategy:
+      fail-fast: false
+      matrix:
+        platform: [avx512, avx2]
+    with:
+      image: ubuntu:24.04
+      runs-on: '[ "Linux", "${{ matrix.platform }}" ]'
+      artifact-in: build-linux
+      cmd: |
+        EMBREE_PATH=$(./scripts/prepare-test-files.sh --type ctest | tail -1)
+        ./scripts/run-tests.sh --type ctest --embree-dir "$EMBREE_PATH" --ctest-config Release
 
 
-#####################################################################
-# DPC++ Release Builds
 ####################################################################
-  release-windows-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP:
+# SYCL Builds and Tests
+####################################################################
+  build-windows-sycl:
     secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/windows_gpu.yml@devel
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/windows.yml@main
     with:
-      runs-on: '[ "Windows", "NAS", "dg2" ]'
-      project: embree
-      env-from-files: ./.github/workflows/dpcpp-none.env ./.github/workflows/gfx-windows-internal.env
-      artifact-out: release-windows-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      runs-on: '[ "Windows", "build" ]'
+      dpcpp-version: intel-llvm/v6.2.0
+      artifact-out: build-windows-sycl
       artifact-path: ./build/*.zip
-      shell: cmd
       cmd: |
-        call "C:\\Program Files (x86)\\Intel\\oneAPI\\compiler\\2023.1.0\\env\\vars.bat"
-        cmake --preset package-windows-icx-sycl-ispc1_19_0-tbb2021_9_0
-        cmake --build build --config Release --target package
-        cmake --build build --target post_package
-        cmake --build build --target test_package
+        $opts  = "-DCMAKE_BUILD_TYPE=Release"
+        $opts += " -DBUILD_TESTING=ON"
+        $opts += " -DEMBREE_TUTORIALS=ON"
+        $opts += " -DEMBREE_ISPC_SUPPORT=ON"
+        $opts += " -DEMBREE_TESTING_INTENSITY=3"
+        $opts += " -DEMBREE_SYCL_SUPPORT=ON"
+        $opts += " -DEMBREE_SYCL_AOT_DEVICES=none"
+        $opts += " -DEMBREE_MAX_ISA=AVX512"
+        $opts += " -DEMBREE_INSTALL_DEPENDENCIES=ON"
+        $opts += " -DEMBREE_TESTING_INSTALL_TESTS=ON"
+        $opts += " -DCMAKE_INSTALL_TESTDIR=testing"
+        $opts += " -DEMBREE_ZIP_MODE=ON"
 
+        cmake -S superbuild -B build -G Ninja -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -D EMBREE_EXTRA_OPTIONS="$opts"
+        cmake --build build --config Release
+        cmake --build build/embree/build --config Release --target build
+        Copy-Item build/embree/build/embree-*.zip build/ -Force
 
-  # FIXME: remove this AOT release once JIT properly tested in CI
-  #release-windows-x64-Release-SYCL-ICX2023_0_0-AOT-ISPC1_18_1-TBB2021_7_0-PACKAGE-ZIP:
-  #  secrets: inherit
-  #  uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/embree_windows.yml@main
-  #  with:
-  #    runs-on: '[ "Windows", "NAS", "dg2" ]'
-  #    artifact-out: release-windows-x64-Release-SYCL-ICX2023_0_0-AOT-ISPC1_18_1-TBB2021_7_0-PACKAGE-ZIP
-  #    artifact-path: ./build/*.zip
-  #    cmd: |
-  #      python scripts/test.py configure platform:x64 build:Release compiler:dpcpp EMBREE_SYCL_SUPPORT:ON sycl:dg2 gfx:RELEASE ispc:ispc1.18.1 isas:SSE2-SSE42-AVX-AVX2-AVX512 intensity:4 package:ZIP
-  #      python scripts/test.py build
-  #      python scripts/test.py test
-
-
-  release-linux-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP:
+  build-linux-sycl:
     secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@devel
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker.yml@main
     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-none.env ./.github/workflows/gfx-ubuntu22-release.env
-      artifact-out: release-linux-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      image: rockylinux:8.7
+      runs-on: '[ "Linux", "build" ]'
+      dpcpp-version: intel-llvm/v6.2.0
+      artifact-out: build-linux-sycl
       artifact-path: ./build/*.tar.gz
+      pre-cmd: |
+        dnf -y install epel-release
+        dnf -y config-manager --set-enabled powertools
+        dnf -y install ninja-build
       cmd: |
-        source scripts/cmake-presets/linux-env.sh
-        source ${NAS_LINUX}/intel/oneAPI/compiler/2023.0.0/env/vars.sh
-        $CMAKE_EXE --preset package-linux-icx-sycl-ispc1_19_0-tbb2021_9_0
-        $CMAKE_EXE --build build --config Release --target package
-        $CMAKE_EXE --build build --target post_package
-        $CMAKE_EXE --build build --target test_package
+        opts="-DCMAKE_BUILD_TYPE=Release"
+        opts+=" -DBUILD_TESTING=ON"
+        opts+=" -DEMBREE_TUTORIALS=ON"
+        opts+=" -DEMBREE_ISPC_SUPPORT=ON"
+        opts+=" -DEMBREE_TESTING_INTENSITY=3"
+        opts+=" -DEMBREE_SYCL_SUPPORT=ON"
+        opts+=" -DEMBREE_SYCL_AOT_DEVICES=none"
+        opts+=" -DEMBREE_MAX_ISA=AVX512"
+        opts+=" -DEMBREE_INSTALL_DEPENDENCIES=ON"
+        opts+=" -DEMBREE_TESTING_INSTALL_TESTS=ON"
+        opts+=" -DCMAKE_INSTALL_TESTDIR=testing"
+        opts+=" -DEMBREE_ZIP_MODE=ON"
+
+        cmake -S superbuild -B build -G Ninja -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -D EMBREE_EXTRA_OPTIONS="$opts"
+        cmake --build build --config Release
+        cmake --build build/embree/build --config Release --target build
+        cp build/embree/build/embree-*.tar.gz build/
+
+  windows-sycl-test:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/windows.yml@main
+    needs: ["build-windows-sycl"]
+    strategy:
+      fail-fast: false
+      matrix:
+        platform: [dg2, bmg, lnl]
+    with:
+      runs-on: '[ "Windows", "${{ matrix.platform }}" ]'
+      env-from-files: ./.github/workflows/gfx-windows-public.env
+      artifact-in: build-windows-sycl
+      cmd: |
+        $EmbreeDir = .\\scripts\\prepare-test-files.ps1 -TestType ctest
+        .\\scripts\\run-tests.ps1 -TestType ctest -EmbreeDir $EmbreeDir -CtestConfig Release
+
+  linux-sycl-test:
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker.yml@main
+    needs: ["build-linux-sycl"]
+    strategy:
+      fail-fast: false
+      matrix:
+        platform: [dg2, pvc, atsm]
+    with:
+      image: ubuntu:24.04
+      options: --device=/dev/dri:/dev/dri
+      runs-on: '[ "Linux", "${{ matrix.platform }}" ]'
+      env-from-files: ./.github/workflows/gfx-ubuntu24-public.env
+      artifact-in: build-linux-sycl
+      cmd: |
+        EMBREE_PATH=$(./scripts/prepare-test-files.sh --type ctest | tail -1)
+        ./scripts/run-tests.sh --type ctest --embree-dir "$EMBREE_PATH" --ctest-config Release
 
 
-#####################################################################
+####################################################################
 # Release binaries scans
-#####################################################################
+####################################################################
   binary-analysis:
     needs:
-      - release-windows-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      - release-macosx-x64-Release-CLANG-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      - release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      - build-windows
+      - build-macosx-x86_64
+      - build-linux
     secrets: inherit
     uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/binary_analysis.yml@main
     with:
       project: embree
-      artifact-in-windows: release-windows-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-in-macos: release-macosx-x64-Release-CLANG-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-in-linux: release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      artifact-in-windows: build-windows
+      artifact-in-macos: build-macosx-x86_64
+      artifact-in-linux: build-linux
       path: "build/*.zip build/*.gz"
 
   antivirus-scan:
     needs:
-      - release-windows-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      - release-macosx-x64-Release-CLANG-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      - release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      - build-windows
+      - build-macosx-x86_64
+      - build-linux
     secrets: inherit
     uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/antivirus_scan.yml@main
     with:
       project: embree
-      artifact-in-windows: release-windows-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-in-macos: release-macosx-x64-Release-CLANG-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-in-linux: release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      artifact-in-windows: build-windows
+      artifact-in-macos: build-macosx-x86_64
+      artifact-in-linux: build-linux
       path: "build/*.zip build/*.gz"
 
-  binary-analysis-dpcpp:
+  binary-analysis-sycl:
     needs:
-      - release-windows-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      - release-linux-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      - build-windows-sycl
+      - build-linux-sycl
     secrets: inherit
     uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/binary_analysis.yml@main
     with:
       project: embree
-      artifact-in-windows: release-windows-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-in-linux: release-linux-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      artifact-in-windows: build-windows-sycl
+      artifact-in-linux: build-linux-sycl
       path: "build/*.zip build/*.gz"
-      artifact-out-name: "binary-analysis-dpcpp"
+      artifact-out-name: "binary-analysis-sycl"
 
-  antivirus-scan-dpcpp:
+  antivirus-scan-sycl:
     needs:
-      - release-windows-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      - release-linux-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      - build-windows-sycl
+      - build-linux-sycl
     secrets: inherit
     uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/antivirus_scan.yml@main
     with:
       project: embree
-      artifact-in-windows: release-windows-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-in-linux: release-linux-x64-Release-SYCL-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
+      artifact-in-windows: build-windows-sycl
+      artifact-in-linux: build-linux-sycl
       path: "build/*.zip build/*.gz"
-      artifact-out-name: "antivirus-scan-dpcpp"
+      artifact-out-name: "antivirus-scan-sycl"
 
-
-#####################################################################
-# Testing package
-#####################################################################
-  test-package:
+  binary-analysis-arm:
+    needs:
+      - build-macosx-arm
     secrets: inherit
-    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/docker_gpu.yml@devel
-    needs: [release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP]
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/binary_analysis.yml@main
     with:
-      runs-on: '[ "Linux", "docker", "dg2" ]'
       project: embree
-      env-from-files: ./.github/workflows/dpcpp-none.env ./.github/workflows/gfx-ubuntu22-release.env
-      artifact-in: release-linux-x64-Release-ICX2023_1_0-ISPC1_19_0-TBB2021_9_0-PACKAGE-ZIP
-      artifact-out: release-testing-PACKAGE-ZIP
-      artifact-path: ./build/*.tar.gz
-      cmd: |
-        rm build/embree-4.2.0.x86_64.linux.tar.gz
-        mv build/embree-4.2.0.x86_64.linux-testing.tar.gz build/embree-4.2.0-testing.tar.gz
+      artifact-in-macos: build-macosx-arm
+      path: "build/*.zip"
+      artifact-out-name: "binary-analysis-arm"
+
+  antivirus-scan-arm:
+    needs:
+      - build-macosx-arm
+    secrets: inherit
+    uses: intel-innersource/libraries.devops.renderkit.workflows/.github/workflows/antivirus_scan.yml@main
+    with:
+      project: embree
+      artifact-in-macos: build-macosx-arm
+      path: "build/*.zip"
+      artifact-out-name: "antivirus-scan-arm"

.gitignore

@@ -9,3 +9,5 @@ kernels/config.h
 kernels/hash.h
 include/embree4/rtcore_config.h
 build
+build*
+out

.gitlab-ci.yml

@@ -1,63 +0,0 @@
-## Copyright 2009-2021 Intel Corporation
-## SPDX-License-Identifier: Apache-2.0
-
-variables:
-  COMMON: "platform:x64 compiler:dpcpp ispc:ispc1.16.1 isa:SSE2 tasking:TBB2019.9"
-  # Retry mechanism knobs in case of sporadic errors
-  GET_SOURCES_ATTEMPTS: 10
-  ARTIFACT_DOWNLOAD_ATTEMPTS: 10
-  EXECUTOR_JOB_SECTION_ATTEMPTS: 3
-
-
-####################################################################
-# Stages
-####################################################################
-
-stages:
-  - benchmark
-
-
-####################################################################
-# Preliminary benchmark setup
-####################################################################
-
-.benchmark:
-  stage:  benchmark
-  tags:
-    - vis-perf-x8280-1
-  script:
-    - "scripts/test.py configure platform:x64 build:Release compiler:GCC ispc:ispc1.17.0 isas:AVX512 tasking:TBB2021.2.0 intensity:0 EMBREE_USE_GOOGLE_BENCHMARK:ON EMBREE_GOOGLE_BENCHMARK_DIR:/NAS/packages/apps/google-benchmark/vis-perf-x8280-1/lib64/cmake/benchmark"
-    - "scripts/test.py build"
-    - "scripts/test.py test"
-    - "scripts/test.py run-benchmark.sh"
-  needs: []
-
-benchmark_nightly:
-  extends: .benchmark
-  only:   [web, schedules]
-
-benchmark_manual:
-  extends: .benchmark
-  except:   [schedules]
-  when: manual
-
-.benchmark_gpu:
-  stage:  benchmark
-  image: $DOCKER_REGISTRY/embree/ubuntu:20.04
-  tags:   [docker-gpu, dg2, perf]
-  before_script:
-    - git config --global --add safe.directory $CI_PROJECT_DIR
-  script:
-    - "scripts/run-benchmark-gpu.sh"
-    - "scripts/record-benchmark-gpu.sh"
-  needs: []
-
-benchmark_gpu_nightly:
-  extends: .benchmark_gpu
-  only:   [schedules]
-
-benchmark_gpu_manual:
-  extends: .benchmark_gpu
-  except:   [schedules]
-  when: manual
-

CHANGELOG.md

@@ -1,6 +1,62 @@
 Version History
 ---------------
 
+### Embree 4.4.1
+
+-  Enabled support for Intel® Core™ Ultra series 3 integrated GPUs.
+-  Added specialization constants for curves, motion blur, and motion blurred instances to reduce code complexity.
+-  Enabled AVX512 support with MSVC compiler.
+-  Fixed numerous potential uninitialized variables, integer over/underflows, and out-of-bounds accesses.
+-  Fixed bounds checking for buffer stride and grid dimensions.
+-  Fixed rare hang in rtcNewDevice when using single thread mode (start_threads=1).
+-  Fixed morton builder for instance arrays.
+-  Improved compiler warning handling for C++26, GCC, and ARM builds.
+-  Improved AVX512 flag consistency between GCC and Clang.
+-  Updated TinyEXR dependency to fix resource leaks.
+-  Reduced memory copies and improved move semantics usage.
+-  Bumped CMake minimum required version to 3.10.
+
+### Embree 4.4.0
+
+-  Added support for passing geometry data to Embree using explicit host and SYCL device memory (see `rtcSetSharedGeometryBufferHostDevice`, `rtcNewBufferHostDevice`, and other API calls with `HostDevice` suffix).
+-  Embree does not use SYCL shared memory anymore internally on systems without host unified memory (i.e., discrete GPUs). Therefore, memory transfers are triggered by specific Embree API calls (e.g. `rtcCommitScene`, `rtcCommitBuffer`).
+-  Objects of type `RTCScene` are not accessible on a SYCL device anymore. Ray queries on SYCL devices must be performed using `RTCTraversable` objects and the `rtcTraversableIntersect` and `rtcTraversableOccluded` API calls.
+-  Embree does not query the availability of RDRAND for its ISA detection anymore, which caused issues on some older AMD CPUs.
+-  Performance improvements on GPU for the two level instancing case (RTC_MAX_INSTANCE_LEVEL_COUNT 2).
+
+### Embree 4.3.3
+-   Added RTCError RTC_ERROR_LEVEL_ZERO_RAYTRACING_SUPPORT_MISSING which can indicate a GPU driver that is too old or not installed properly.
+-   Added the API function rtcGetDeviceLastErrorMessage to query additional information about the last RTCError returned by rtcGetDeviceError. This can be used in case device creation failed and a rtcErrorFunction could not be set up for this purpose.
+-   Added the API function rtcGetErrorString which returns a string representation of a given RTCError error code. This is purely meant for convenient error information reporting on the user application side.
+-   Performance improvements on GPU for the one level instancing case (RTC_MAX_INSTANCE_LEVEL_COUNT 1).
+-   Reduced the number of unnecessary GPU-CPU USM back-migrations which can increase build performance for scene with many instances on GPU.
+-   Started adding public CI tests for streamlining integration of external pull requests.
+-   Work-around for problem with unsigned Windows binaries.
+
+### Embree 4.3.2
+-   Embree now uses level zero raytracing extension to build BVH which enables forward-compatibility. On Linux, the package intel-level-zero-gpu-raytracing has to be installed in addition to the other packages listed here [https://dgpu-docs.intel.com/](https://dgpu-docs.intel.com/).
+-   MacOS universal binary compilation now works.
+-   Some bugfixes for AVX512 on MacOS x86 machines.
+-   Known issue: It is recommended to run Embree on Intel® Data Center GPU Max Series (e.g. Intel® Data Center GPU Max 1550) with the following environment settings: NEOReadDebugKeys=1 UseKmdMigration=0
+-   Known issue: ISPC version of tutorials will not successfully build with MacOS universal binary compilation.
+
+### Embree 4.3.1
+-   Add missing EMBREE_GEOMETRY types to embree-config.cmake
+-   User defined thread count now takes precedence for internal task scheduler
+-   Fixed static linking issue with ze_wrapper library
+-   Better error reporting for SYCL platform and driver problems in embree_info and tutorial apps.
+-   Patch to glfw source is not applied by default anymore.
+-   Known issue: Running Embree on Intel® Data Center GPU Max Series with 2 tiles (e.g. Intel® Data Center GPU Max 1550) requires setting the environment variable ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE.
+-   Known issue: Embree build using Apple Clang 15 and ARM support (via the SEE2NEON library) may cause "EXEC_BAD_INSTRUCTION" runtime exceptions. Please use Apple Clang <= 14 on macOS.
+
+### Embree 4.3.0
+-   Added instance array primitive for reducing memory requirements in scenes
+    with large amounts of similar instances.
+-   Properly checks driver if L0 RTAS extension can get loaded.
+-   Added varying version of rtcGetGeometryTransform for ISPC.
+-   Fixed signature of RTCMemoryMonitorFunction for ISPC.
+-   Add support for ARM64 Windows platform in CMake.
+
 ### Embree 4.2.0
 -   SYCL version of Embree with GPU support is no longer in beta phase.
 -   Improved BVH build performance on many core machines for applications that oversubscribe threads.

CMakeLists.txt

@@ -1,11 +1,11 @@
 ## Copyright 2009-2021 Intel Corporation
 ## SPDX-License-Identifier: Apache-2.0
 
-CMAKE_MINIMUM_REQUIRED(VERSION 3.1.0)
+CMAKE_MINIMUM_REQUIRED(VERSION 3.10)
 
 SET(EMBREE_VERSION_MAJOR 4)
-SET(EMBREE_VERSION_MINOR 2)
-SET(EMBREE_VERSION_PATCH 0)
+SET(EMBREE_VERSION_MINOR 4)
+SET(EMBREE_VERSION_PATCH 1)
 SET(EMBREE_VERSION_NOTE "")
 
 SET(EMBREE_VERSION ${EMBREE_VERSION_MAJOR}.${EMBREE_VERSION_MINOR}.${EMBREE_VERSION_PATCH})
@@ -57,18 +57,9 @@ IF(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/.git")
 ENDIF()
 
 IF(COMMAND cmake_policy)
-  if (POLICY CMP0003)
-    cmake_policy(SET CMP0003 NEW)
-  endif()
-  if (POLICY CMP0042)
-    cmake_policy(SET CMP0042 NEW)
-  endif()
   if(POLICY CMP0072)
     cmake_policy(SET CMP0072 NEW)
   endif()
-  if(POLICY CMP0022)
-    cmake_policy(SET CMP0022 NEW)
-  endif()
   if(POLICY CMP0074)
     cmake_policy(SET CMP0074 NEW)
   endif()
@@ -205,10 +196,11 @@ OPTION(EMBREE_GEOMETRY_CURVE "Enables support for curve geometries." ON)
 OPTION(EMBREE_GEOMETRY_SUBDIVISION "Enables support for subdiv geometries." ON)
 OPTION(EMBREE_GEOMETRY_USER "Enables support for user geometries." ON)
 OPTION(EMBREE_GEOMETRY_INSTANCE "Enables support for instances." ON)
+OPTION(EMBREE_GEOMETRY_INSTANCE_ARRAY "Enables support for instance arrays." ON)
 SET(EMBREE_MAX_INSTANCE_LEVEL_COUNT 1 CACHE STRING "Maximum number of instance levels.")
 
-IF (NOT EMBREE_GEOMETRY_INSTANCE AND NOT EMBREE_MAX_INSTANCE_LEVEL_COUNT EQUAL 1)
-  message(FATAL_ERROR "EMBREE_MAX_INSTANCE_LEVEL_COUNT must be 1 when EMBREE_GEOMETRY_INSTANCE is disabled")
+IF (NOT EMBREE_GEOMETRY_INSTANCE AND NOT EMBREE_GEOMETRY_INSTANCE_ARRAY AND NOT EMBREE_MAX_INSTANCE_LEVEL_COUNT EQUAL 1)
+  message(FATAL_ERROR "EMBREE_MAX_INSTANCE_LEVEL_COUNT must be 1 when EMBREE_GEOMETRY_INSTANCE and EMBREE_GEOMETRY_INSTANCE_ARRAY are disabled")
 ENDIF()
 
 OPTION(EMBREE_GEOMETRY_GRID "Enables support for grid geometries." ON)
@@ -226,10 +218,10 @@ OPTION(EMBREE_MIN_WIDTH "Enables min-width feature to enlarge curve and point th
 ##############################################################
 
 # detect ARM compilation
-IF (APPLE AND CMAKE_SYSTEM_NAME STREQUAL "Darwin" AND (CMAKE_SYSTEM_PROCESSOR STREQUAL "arm64" OR CMAKE_OSX_ARCHITECTURES MATCHES "arm64"))
+IF (APPLE AND CMAKE_SYSTEM_NAME STREQUAL "Darwin" AND (CMAKE_SYSTEM_PROCESSOR STREQUAL "arm64" AND CMAKE_OSX_ARCHITECTURES STREQUAL "") OR ("arm64" IN_LIST CMAKE_OSX_ARCHITECTURES))
   MESSAGE(STATUS "Building for Apple silicon")
   SET(EMBREE_ARM ON)
-ELSEIF(CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64")
+ELSEIF(CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64" OR CMAKE_SYSTEM_PROCESSOR STREQUAL "ARM64")
   MESSAGE(STATUS "Building for AArch64")
   SET(EMBREE_ARM ON)
 ENDIF()
@@ -285,6 +277,9 @@ message(DEBUG "CMAKE_GENERATOR_TOOLSET: ${CMAKE_GENERATOR_TOOLSET}")
 message(DEBUG "CMAKE_CXX_COMPILER_ID: ${CMAKE_CXX_COMPILER_ID}")
 message(DEBUG "CMAKE_CXX_COMPILER: ${CMAKE_CXX_COMPILER}")
 
+include(CheckCXXCompilerFlag)
+check_cxx_compiler_flag(-fsycl COMPILER_HAS_SYCL_SUPPORT)
+
 IF (WIN32)
   IF (NOT DEFAULT_CMAKE_CONFIGURATION_TYPES_SET)
     SET(CMAKE_CONFIGURATION_TYPES "${CONFIGURATION_TYPES}" CACHE STRING "List of generated configurations." FORCE)
@@ -293,8 +288,7 @@ IF (WIN32)
   SET_PROPERTY(GLOBAL PROPERTY USE_FOLDERS ON)
 
   IF ((${CMAKE_CXX_COMPILER_ID} MATCHES "IntelLLVM") OR
-      (${CMAKE_CXX_COMPILER_ID} MATCHES "Clang" AND ${CMAKE_CXX_COMPILER} MATCHES "icx") OR
-      (${CMAKE_CXX_COMPILER} MATCHES "dpcpp"))
+      (${CMAKE_CXX_COMPILER_ID} MATCHES "Clang" AND COMPILER_HAS_SYCL_SUPPORT))
     MESSAGE("-- DPCPP compiler detected")
     INCLUDE(dpcpp)
   ELSEIF(${CMAKE_GENERATOR_TOOLSET} MATCHES "^LLVM")
@@ -311,9 +305,6 @@ IF (WIN32)
     INCLUDE(gnu)
   ELSE()
     set(EMBREE_SYCL_SUPPORT OFF)
-    IF (EMBREE_ISA_AVX512)
-      MESSAGE(FATAL_ERROR "Microsoft Visual C++ Compiler does not support AVX512. Please use Intel Compiler or Clang.")
-    ENDIF()
     MESSAGE("-- MSVC detected")
     INCLUDE (msvc)
   ENDIF()
@@ -397,7 +388,7 @@ IF (EMBREE_MAX_ISA STREQUAL "NONE")
     OPTION(EMBREE_ISA_SSE42 "Enables SSE4.2 ISA." ON)
     OPTION(EMBREE_ISA_AVX "Enables AVX ISA." ${COMPILER_SUPPORTS_AVX})
     OPTION(EMBREE_ISA_AVX2 "Enables AVX2 ISA." ${COMPILER_SUPPORTS_AVX2})
-    IF (WIN32 OR APPLE)
+    IF (APPLE)
       OPTION(EMBREE_ISA_AVX512 "Enables AVX512 ISA." OFF)
     ELSE()
       OPTION(EMBREE_ISA_AVX512 "Enables AVX512 ISA." ${COMPILER_SUPPORTS_AVX512})

CMakePresets.json

@@ -1,220 +0,0 @@
-{
-  "version": 4,
-  "cmakeMinimumRequired": {
-    "major": 3,
-    "minor": 11,
-    "patch": 0
-  },
-  "include": [
-    "scripts/cmake-presets/os.json", 
-    "scripts/cmake-presets/package.json",
-    "scripts/cmake-presets/compiler.json",
-    "scripts/cmake-presets/tbb.json"
-  ],
-  "configurePresets": [
-    {
-      "name": "package-linux-gcc-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "gcc", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-linux-icx-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "icx", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-    {
-      "name": "package-linux-icx-ispc1_19_0-tbb2021_9_0-perf",
-      "inherits": ["package-linux", "env", "icx", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512",
-        "EMBREE_USE_GOOGLE_BENCHMARK": "ON",
-        "EMBREE_BUILD_GOOGLE_BENCHMARK_FROM_SOURCE": "ON"
-      }
-    },
-
-
-    {
-      "name": "package-linux-clang-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "clang", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-static-linux-gcc-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "gcc", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-static-linux-icx-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "icx", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-static-linux-clang-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "clang", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-linux-icx-sycl-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "icx", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_SYCL_SUPPORT": "ON",
-        "EMBREE_SYCL_AOT_DEVICES": "none",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-    {
-      "name": "package-linux-icx-sycl-ispc1_19_0-tbb2021_9_0-perf",
-      "inherits": ["package-linux", "env", "icx", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_SYCL_SUPPORT": "ON",
-        "EMBREE_SYCL_AOT_DEVICES": "none",
-        "EMBREE_MAX_ISA": "AVX512",
-        "EMBREE_USE_GOOGLE_BENCHMARK": "ON",
-        "EMBREE_BUILD_GOOGLE_BENCHMARK_FROM_SOURCE": "ON"
-      }
-    },
-
-    {
-      "name": "package-linux-dpcpp-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-linux", "env", "dpcpp", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_SYCL_SUPPORT": "ON",
-        "EMBREE_SYCL_AOT_DEVICES": "none",
-        "EMBREE_MAX_ISA": "SSE2"
-      }
-    },
-
-
-
-    {
-      "name": "package-macos-gcc-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-macos", "env", "gcc", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-macos-clang-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-macos", "env", "clang", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-
-
-    {
-      "name": "package-windows-icx-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-windows", "env", "icx-windows", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-windows-v142-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-windows", "env", "v142", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "SSE2"
-      }
-    },
-
-    {
-      "name": "package-windows-v141-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-windows", "env", "v141", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "SSE2"
-      }
-    },
-
-    {
-      "name": "package-windows-v140-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-windows", "env", "v140", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "SSE2"
-      }
-    },
-    
-    {
-      "name": "package-static-windows-icx-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-static-windows", "env", "icx-windows", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    },
-
-    {
-      "name": "package-static-windows-v142-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-static-windows", "env", "v142", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_MAX_ISA": "SSE2"
-      }
-    },
-
-    {
-      "name": "package-windows-icx-sycl-ispc1_19_0-tbb2021_9_0",
-      "inherits": ["package-windows", "env", "icx-windows", "ispc1_19_0", "tbb2021_9_0"],
-      "binaryDir": "${sourceDir}/build",
-      "cacheVariables": {
-        "CMAKE_BUILD_TYPE": "Release",
-        "EMBREE_SYCL_SUPPORT": "ON",
-        "EMBREE_SYCL_AOT_DEVICES": "none",
-        "EMBREE_MAX_ISA": "AVX512"
-      }
-    }
-  ]
-}

common/algorithms/parallel_any_of.h

@@ -12,7 +12,8 @@ namespace embree
   template<typename Index, class UnaryPredicate>
     __forceinline bool parallel_any_of (Index first, Index last, UnaryPredicate pred)
   {
-    bool ret = false;
+    std::atomic_bool ret;
+    ret = false;
     
 #if defined(TASKING_TBB)
 #if TBB_INTERFACE_VERSION >= 12002

common/cmake/FindTBB.cmake

@@ -37,7 +37,7 @@
 #===============================================================================
 
 # We use INTERFACE libraries, which are only supported in 3.x
-cmake_minimum_required(VERSION 3.1)
+cmake_minimum_required(VERSION 3.10)
 
 # These two are used to automatically find the root and include directories.
 set(_TBB_INCLUDE_SUBDIR "include")

common/cmake/clang.cmake

@@ -8,10 +8,19 @@ MACRO(_SET_IF_EMPTY VAR VALUE)
 ENDMACRO()
 
 IF (EMBREE_ARM)
-  SET(FLAGS_SSE2 "-D__SSE__ -D__SSE2__")
-  SET(FLAGS_SSE42 "-D__SSE4_2__  -D__SSE4_1__")
-  SET(FLAGS_AVX "-D__AVX__ -D__SSE4_2__  -D__SSE4_1__  -D__BMI__ -D__BMI2__ -D__LZCNT__")
-  SET(FLAGS_AVX2 "-D__AVX2__ -D__AVX__ -D__SSE4_2__  -D__SSE4_1__  -D__BMI__ -D__BMI2__ -D__LZCNT__")
+  IF ("x86_64" IN_LIST CMAKE_OSX_ARCHITECTURES)
+    # set ARM an x86 flags for macOS universal binary build
+    SET(FLAGS_SSE2 "-D__SSE__ -D__SSE2__ -msse -msse2 -mno-sse4.2")
+    SET(FLAGS_SSE42 "-D__SSE4_2__  -D__SSE4_1__ -msse4.2")
+    SET(FLAGS_AVX "-D__AVX__ -D__SSE4_2__  -D__SSE4_1__  -D__BMI__ -D__BMI2__ -D__LZCNT__ -mavx")
+    SET(FLAGS_AVX2 "-D__AVX2__ -D__AVX__ -D__SSE4_2__  -D__SSE4_1__  -D__BMI__ -D__BMI2__ -D__LZCNT__ -mf16c -mavx2 -mfma -mlzcnt -mbmi -mbmi2")
+    _SET_IF_EMPTY(FLAGS_AVX512 "-march=skx")
+  ELSE ()
+    SET(FLAGS_SSE2 "-D__SSE__ -D__SSE2__")
+    SET(FLAGS_SSE42 "-D__SSE4_2__  -D__SSE4_1__")
+    SET(FLAGS_AVX "-D__AVX__ -D__SSE4_2__  -D__SSE4_1__  -D__BMI__ -D__BMI2__ -D__LZCNT__")
+    SET(FLAGS_AVX2 "-D__AVX2__ -D__AVX__ -D__SSE4_2__  -D__SSE4_1__  -D__BMI__ -D__BMI2__ -D__LZCNT__")
+  ENDIF ()
 ELSE ()
   # for `thread` keyword
   _SET_IF_EMPTY(FLAGS_SSE2  "-msse -msse2 -mno-sse4.2")

common/cmake/dpcpp.cmake

@@ -72,6 +72,7 @@ IF (EMBREE_SYCL_SUPPORT)
            ${SYCL_COMPILER_LIB_DIR}/sycl[0-9].lib
            ${SYCL_COMPILER_LIB_DIR}/sycl[0-9][0-9].lib)
     ENDIF()
+    LIST(GET SYCL_LIB 0 SYCL_LIB)
     GET_FILENAME_COMPONENT(SYCL_LIB_NAME ${SYCL_LIB} NAME_WE)
   ELSE()
     SET(SYCL_LIB_NAME "sycl")
@@ -143,10 +144,16 @@ IF (EMBREE_SYCL_SUPPORT)
   SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-pessimizing-move") # disabled: warning: moving a temporary object prevents copy elision [-Wpessimizing-move]
 
   IF (SYCL_ONEAPI_ICX AND WIN32)
+    IF (${CMAKE_CXX_COMPILER_VERSION} VERSION_GREATER_EQUAL 2024.0)
+      SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -I\"${SYCL_COMPILER_DIR}/../opt/compiler/include/sycl\" -I\"${SYCL_COMPILER_DIR}/../opt/compiler/include/sycl/sycl\"")       # disable warning from SYCL header
+    ENDIF()
     SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -I\"${SYCL_COMPILER_DIR}/../include/sycl\" -I\"${SYCL_COMPILER_DIR}/../include/\"")       # disable warning from SYCL header
     SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Qstd=c++17")
   ELSE()
     SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++17")
+    IF (SYCL_ONEAPI_ICX AND ${CMAKE_CXX_COMPILER_VERSION} VERSION_GREATER_EQUAL 2024.0)
+      SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -isystem \"${SYCL_COMPILER_DIR}/../opt/compiler/include/sycl\" -isystem \"${SYCL_COMPILER_DIR}/../opt/compiler/include/sycl/sycl\"")       # disable warning from SYCL header
+    ENDIF()
     SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -isystem \"${SYCL_COMPILER_DIR}/../include/sycl\" -isystem \"${SYCL_COMPILER_DIR}/../include/\"")       # disable warning from SYCL header
   ENDIF()
 
@@ -166,6 +173,8 @@ IF(SYCL_ONEAPI_ICX)
     SET(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} /Qimf-use-svml:false")
     SET(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /Qno-intel-lib")
     SET(CMAKE_C_LINK_FLAGS "${CMAKE_C_LINK_FLAGS} /Qno-intel-lib")
+    SET(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} /Qoption,link,/DEPENDENTLOADFLAG:0x2000")
+    SET(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} /Qoption,link,/DEPENDENTLOADFLAG:0x2000")
   ELSE()
     SET(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} -static-intel")
     SET(CMAKE_C_LINK_FLAGS "${CMAKE_C_LINK_FLAGS} -static-intel")

common/cmake/embree-config.cmake

@@ -45,6 +45,9 @@ SET(EMBREE_GEOMETRY_QUAD @EMBREE_GEOMETRY_QUAD@)
 SET(EMBREE_GEOMETRY_CURVE @EMBREE_GEOMETRY_CURVE@)
 SET(EMBREE_GEOMETRY_SUBDIVISION @EMBREE_GEOMETRY_SUBDIVISION@)
 SET(EMBREE_GEOMETRY_USER @EMBREE_GEOMETRY_USER@)
+SET(EMBREE_GEOMETRY_INSTANCE @EMBREE_GEOMETRY_INSTANCE@)
+SET(EMBREE_GEOMETRY_INSTANCE_ARRAY @EMBREE_GEOMETRY_INSTANCE_ARRAY@)
+SET(EMBREE_GEOMETRY_GRID @EMBREE_GEOMETRY_GRID@)
 SET(EMBREE_GEOMETRY_POINT @EMBREE_GEOMETRY_POINT@)
 
 SET(EMBREE_RAY_PACKETS @EMBREE_RAY_PACKETS@)
@@ -94,6 +97,7 @@ IF (EMBREE_SYCL_SUPPORT)
   ENDIF()
   IF (EMBREE_STATIC_LIB)
     INCLUDE("${EMBREE_ROOT_DIR}/@EMBREE_CMAKEEXPORT_DIR@/embree_rthwif-targets.cmake")
+    INCLUDE("${EMBREE_ROOT_DIR}/@EMBREE_CMAKEEXPORT_DIR@/ze_wrapper-targets.cmake")
   ENDIF()
   INCLUDE("${EMBREE_ROOT_DIR}/@EMBREE_CMAKEEXPORT_DIR@/embree_sycl-targets.cmake")
 
@@ -102,6 +106,7 @@ ENDIF()
 INCLUDE("${EMBREE_ROOT_DIR}/@EMBREE_CMAKEEXPORT_DIR@/embree-targets.cmake")
 
 
+SET(EMBREE_TESTING_ONLY_SYCL_TESTS @EMBREE_TESTING_ONLY_SYCL_TESTS@)
 SET(EMBREE_TESTING_INTENSITY @EMBREE_TESTING_INTENSITY@)
 SET(EMBREE_TESTING_MEMCHECK @EMBREE_TESTING_MEMCHECK@)
-SET(EMBREE_TESTING_BENCHMARK @EMBREE_TESTING_BENCHMARK@)
+SET(EMBREE_TESTING_BENCHMARK @EMBREE_TESTING_BENCHMARK@)

common/cmake/gnu.cmake

@@ -17,7 +17,7 @@ ELSE ()
   _SET_IF_EMPTY(FLAGS_SSE42 "-msse4.2")
   _SET_IF_EMPTY(FLAGS_AVX   "-mavx")
   _SET_IF_EMPTY(FLAGS_AVX2  "-mf16c -mavx2 -mfma -mlzcnt -mbmi -mbmi2")
-  _SET_IF_EMPTY(FLAGS_AVX512 "-mavx512f -mavx512dq -mavx512cd -mavx512bw -mavx512vl -mf16c -mavx2 -mfma -mlzcnt -mbmi -mbmi2 -mprefer-vector-width=256")
+  _SET_IF_EMPTY(FLAGS_AVX512 "-march=skylake-avx512")
 ENDIF ()
 
 OPTION(EMBREE_IGNORE_CMAKE_CXX_FLAGS "When enabled Embree ignores default CMAKE_CXX_FLAGS." ON)

common/cmake/ispc.cmake

@@ -39,8 +39,6 @@ IF (NOT EMBREE_ISPC_EXECUTABLE)
   FIND_PROGRAM(EMBREE_ISPC_EXECUTABLE ispc PATHS ${ISPC_DIR_HINT} DOC "Path to the ISPC executable.")
   IF (NOT EMBREE_ISPC_EXECUTABLE)
     MESSAGE(FATAL_ERROR "Intel SPMD Compiler (ISPC) not found. Disable EMBREE_ISPC_SUPPORT or install ISPC.")
-  ELSE()
-    MESSAGE(STATUS "Found Intel SPMD Compiler (ISPC): ${EMBREE_ISPC_EXECUTABLE}")
   ENDIF()
 ENDIF()
 
@@ -51,6 +49,8 @@ EXECUTE_PROCESS(COMMAND ${EMBREE_ISPC_EXECUTABLE} --version
 
 IF (NOT ${ISPC_RESULT} STREQUAL "0")
   MESSAGE(FATAL_ERROR "Error executing ISPC executable '${EMBREE_ISPC_EXECUTABLE}': ${ISPC_RESULT}")
+ELSE()
+  MESSAGE(STATUS "Found working Intel SPMD Compiler (ISPC): ${EMBREE_ISPC_EXECUTABLE}")
 ENDIF()
 
 STRING(REGEX MATCH "([0-9]+[.][0-9]+[.][0-9]+)" DUMMY "${ISPC_OUTPUT}")

common/cmake/msvc.cmake

@@ -5,6 +5,7 @@ SET(FLAGS_SSE2  "/D__SSE__ /D__SSE2__")
 SET(FLAGS_SSE42 "${FLAGS_SSE2} /D__SSE3__ /D__SSSE3__ /D__SSE4_1__ /D__SSE4_2__")
 SET(FLAGS_AVX   "${FLAGS_SSE42} /arch:AVX")
 SET(FLAGS_AVX2  "${FLAGS_SSE42} /arch:AVX2")
+SET(FLAGS_AVX512  "${FLAGS_AVX2} /arch:AVX512")
 
 SET(COMMON_CXX_FLAGS "")
 SET(COMMON_CXX_FLAGS "${COMMON_CXX_FLAGS} /EHsc")        # catch C++ exceptions only and extern "C" functions never throw a C++ exception
@@ -22,6 +23,9 @@ MACRO(DISABLE_STACK_PROTECTOR_FOR_FILE file)
   ENDIF()
 ENDMACRO()
 
+SET(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} /DEPENDENTLOADFLAG:0x2000")
+SET(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} /DEPENDENTLOADFLAG:0x2000")
+
 SET(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} ${COMMON_CXX_FLAGS}")
 STRING(REPLACE "/RTC1" "" CMAKE_CXX_FLAGS_DEBUG ${CMAKE_CXX_FLAGS_DEBUG})         # disable native runtime checks
 SET(CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} /DDEBUG")                     # enables assertions

common/cmake/package.cmake

@@ -23,12 +23,11 @@ ENDIF()
 # Install SYCL specific files
 ##############################################################
 
-# SYCL library
 IF (EMBREE_SYCL_SUPPORT AND EMBREE_INSTALL_DEPENDENCIES)
   GET_FILENAME_COMPONENT(DPCPP_COMPILER_DIR ${CMAKE_CXX_COMPILER} PATH)
-  
+
   IF (WIN32)
-    
+
     FILE(GLOB LIB_SYCL_LIB_FILES LIST_DIRECTORIES FALSE
       "${DPCPP_COMPILER_DIR}/../lib/sycl.lib"
       "${DPCPP_COMPILER_DIR}/../lib/sycl?.lib")
@@ -38,48 +37,81 @@ IF (EMBREE_SYCL_SUPPORT AND EMBREE_INSTALL_DEPENDENCIES)
       "${DPCPP_COMPILER_DIR}/../bin/sycl.dll"
       "${DPCPP_COMPILER_DIR}/../bin/sycl?.dll")
     INSTALL(FILES ${LIB_SYCL_DLL_FILES} DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
-    
-    INSTALL(FILES "${DPCPP_COMPILER_DIR}/../bin/pi_level_zero.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
-    INSTALL(FILES "${DPCPP_COMPILER_DIR}/../bin/win_proxy_loader.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib OPTIONAL)
 
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/svml_dispmd.dll")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/svml_dispmd.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
+    # copy level zero dll
+    IF (EXISTS "${DPCPP_COMPILER_DIR}/pi_level_zero.dll")
+      INSTALL(FILES "${DPCPP_COMPILER_DIR}/pi_level_zero.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
     ENDIF()
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/libmmd.dll")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/libmmd.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
+    IF (EXISTS "${DPCPP_COMPILER_DIR}/ur_adapter_level_zero.dll")
+      INSTALL(FILES "${DPCPP_COMPILER_DIR}/ur_adapter_level_zero.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
     ENDIF()
-    
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../bin/win_proxy_loader.dll")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../bin/win_proxy_loader.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
+
+    # copy unified runtime loader dll
+    IF(EXISTS "${DPCPP_COMPILER_DIR}/ur_loader.dll")
+      INSTALL(FILES "${DPCPP_COMPILER_DIR}/ur_loader.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib OPTIONAL)
     ENDIF()
+
+    # copy proxy loader dll
+    IF(EXISTS "${DPCPP_COMPILER_DIR}/win_proxy_loader.dll")
+      INSTALL(FILES "${DPCPP_COMPILER_DIR}/win_proxy_loader.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib OPTIONAL)
+    ENDIF()
+    IF (EXISTS "${DPCPP_COMPILER_DIR}/pi_win_proxy_loader.dll")
+      INSTALL(FILES "${DPCPP_COMPILER_DIR}/pi_win_proxy_loader.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
+    ENDIF()
+    IF(EXISTS "${DPCPP_COMPILER_DIR}/ur_win_proxy_loader.dll")
+      INSTALL(FILES "${DPCPP_COMPILER_DIR}/ur_win_proxy_loader.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib OPTIONAL)
+    ENDIF()
+
+    # copy additional oneAPI dlls
+    IF(SYCL_ONEAPI_ICX)
+      IF (EXISTS "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/libmmd.dll")
+        INSTALL(FILES "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/libmmd.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
+      ENDIF()
+      IF (EXISTS "${DPCPP_COMPILER_DIR}/libmmd.dll")
+        INSTALL(FILES "${DPCPP_COMPILER_DIR}/libmmd.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
+      ENDIF()
+      IF (EXISTS "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/svml_dispmd.dll")
+        INSTALL(FILES "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/svml_dispmd.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
+      ENDIF()
+    ENDIF()
+
   ELSE()
-    
+
     FILE(GLOB LIB_SYCL_FILES LIST_DIRECTORIES FALSE
       "${DPCPP_COMPILER_DIR}/../lib/libsycl.so"
       "${DPCPP_COMPILER_DIR}/../lib/libsycl.so.?"
       "${DPCPP_COMPILER_DIR}/../lib/libsycl.so.?.?"
-      "${DPCPP_COMPILER_DIR}/../lib/libsycl.so.?.?.?")
+      "${DPCPP_COMPILER_DIR}/../lib/libsycl.so.?.?.?"
+      "${DPCPP_COMPILER_DIR}/../lib/libsycl.so.?.?.?-?"
+      "${DPCPP_COMPILER_DIR}/../lib/libur_adapter_level_zero.so*"
+      "${DPCPP_COMPILER_DIR}/../lib/libur_loader.so*"
+      "${DPCPP_COMPILER_DIR}/../lib/libumf.so*")
     INSTALL(FILES ${LIB_SYCL_FILES} DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    
-    INSTALL(FILES "${DPCPP_COMPILER_DIR}/../lib/libpi_level_zero.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
 
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libsvml.so")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libsvml.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
+    # copy additional oneAPI runtime libraries
+    IF(SYCL_ONEAPI_ICX)
+      IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libsvml.so")
+        INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libsvml.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
+      ENDIF()
+
+      IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libirng.so")
+        INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libirng.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
+      ENDIF()
+
+      IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libimf.so")
+        INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libimf.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
+      ENDIF()
+
+      IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so")
+        FILE(GLOB LIB_SYCL_FILES LIST_DIRECTORIES FALSE
+          "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so"
+          "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so.?")
+        INSTALL(FILES ${LIB_SYCL_FILES} DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
+      ENDIF()
     ENDIF()
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libirng.so")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libirng.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    ENDIF()
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libimf.so")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libimf.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    ENDIF()
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so")
-      FILE(GLOB LIB_SYCL_FILES LIST_DIRECTORIES FALSE
-        "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so"
-        "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so.?")
-      INSTALL(FILES ${LIB_SYCL_FILES} DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    ENDIF()
-   
+
   ENDIF()
+
 ENDIF()
 
 ##############################################################
@@ -87,15 +119,6 @@ ENDIF()
 ##############################################################
 
 IF (WIN32)
-  IF(SYCL_ONEAPI_ICX AND EMBREE_INSTALL_DEPENDENCIES)
-    GET_FILENAME_COMPONENT(DPCPP_COMPILER_DIR ${CMAKE_CXX_COMPILER} PATH)
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/libmmd.dll")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/libmmd.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
-    ENDIF()
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/svml_dispmd.dll")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../redist/intel64_win/compiler/svml_dispmd.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
-    ENDIF()
-  ENDIF()
 
   SET(CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS_SKIP TRUE)
   SET(CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS_NO_WARNINGS TRUE)
@@ -103,32 +126,6 @@ IF (WIN32)
   LIST(FILTER CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS INCLUDE REGEX ".*msvcp[0-9]+\.dll|.*vcruntime[0-9]+\.dll|.*vcruntime[0-9]+_[0-9]+\.dll")
   INSTALL(FILES ${CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS} DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib)
 
-ELSE()
-
-  IF(SYCL_ONEAPI_ICX AND EMBREE_INSTALL_DEPENDENCIES)
-
-    GET_FILENAME_COMPONENT(DPCPP_COMPILER_DIR ${CMAKE_CXX_COMPILER} PATH)
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libsvml.so")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libsvml.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    ENDIF()
-
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libirng.so")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libirng.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    ENDIF()
-
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libimf.so")
-      INSTALL(FILES "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libimf.so" DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    ENDIF()
-
-    IF (EXISTS "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so")
-      FILE(GLOB LIB_SYCL_FILES LIST_DIRECTORIES FALSE
-        "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so"
-        "${DPCPP_COMPILER_DIR}/../compiler/lib/intel64_lin/libintlc.so.?")
-      INSTALL(FILES ${LIB_SYCL_FILES} DESTINATION "${CMAKE_INSTALL_LIBDIR}" COMPONENT lib)
-    ENDIF()
-
-  ENDIF()
-
 ENDIF()
 
 ##############################################################
@@ -241,7 +238,9 @@ IF (EMBREE_SYCL_SUPPORT)
 ENDIF()
 #SET(CPACK_PACKAGE_ICON "${PROJECT_SOURCE_DIR}/embree-doc/images/icon.png")
 #SET(CPACK_PACKAGE_RELOCATABLE TRUE)
-SET(CPACK_STRIP_FILES TRUE)
+IF(NOT WIN32)
+  SET(CPACK_STRIP_FILES TRUE)
+ENDIF()
 
 SET(CPACK_PACKAGE_VERSION_MAJOR ${EMBREE_VERSION_MAJOR})
 SET(CPACK_PACKAGE_VERSION_MINOR ${EMBREE_VERSION_MINOR})
@@ -281,7 +280,6 @@ SET(CPACK_RESOURCE_FILE_LICENSE "${PROJECT_SOURCE_DIR}/LICENSE.txt")
 
 # Windows specific settings
 IF(WIN32)
-
   IF (CMAKE_SIZEOF_VOID_P EQUAL 8)
     SET(ARCH x64)
     SET(CPACK_PACKAGE_NAME "${CPACK_PACKAGE_NAME} x64")
@@ -293,19 +291,7 @@ IF(WIN32)
   SET(CPACK_GENERATOR ZIP)
   SET(CPACK_PACKAGE_FILE_NAME "${CPACK_PACKAGE_FILE_NAME}.${ARCH}.windows")
   SET(PACKAGE_BASE_NAME "${CPACK_PACKAGE_FILE_NAME}")
-  #SET(CPACK_MONOLITHIC_INSTALL 1)
-  IF (EMBREE_TESTING_PACKAGE)
-    SET(PACKAGE_SCRIPT "${PROJECT_SOURCE_DIR}/scripts/package_win.bat")
-    ADD_TEST(NAME "BuildPackage" WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}" COMMAND ${PACKAGE_SCRIPT} ${CMAKE_BUILD_TYPE} "${PACKAGE_BASE_NAME}" "${ARCH}")
-  ENDIF()
-
-  add_custom_target(
-    post_package "${PROJECT_SOURCE_DIR}/scripts/package_post_build_win.bat" "${PACKAGE_BASE_NAME}" "${EMBREE_SIGN_FILE}"
-  )
-
-  add_custom_target(
-    test_package "${PROJECT_SOURCE_DIR}/scripts/package_test_win.bat" "${PACKAGE_BASE_NAME}"
-  )
+  SET(PACKAGE_EXT "zip")
 
 # MacOSX specific settings
 ELSEIF(APPLE)
@@ -314,41 +300,52 @@ ELSEIF(APPLE)
   SET(CPACK_RESOURCE_FILE_README "${PROJECT_BINARY_DIR}/README.txt")
 
   SET(CPACK_GENERATOR ZIP)
-  SET(PACKAGE_BASE_NAME "${CPACK_PACKAGE_FILE_NAME}")
-  SET(CPACK_PACKAGE_FILE_NAME "${CPACK_PACKAGE_FILE_NAME}.x86_64.macosx")
-  #SET(CPACK_MONOLITHIC_INSTALL 1)
-  IF (EMBREE_TESTING_PACKAGE)
-    ADD_TEST(NAME "BuildPackage" WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}" COMMAND "${PROJECT_SOURCE_DIR}/scripts/package_macosx.sh" ${CMAKE_BUILD_TYPE} "${PACKAGE_BASE_NAME}" "${EMBREE_SIGN_FILE}")
+  IF (EMBREE_ARM)
+    SET(CPACK_PACKAGE_FILE_NAME "${CPACK_PACKAGE_FILE_NAME}.arm64.macosx")
+  ELSE()
+    SET(CPACK_PACKAGE_FILE_NAME "${CPACK_PACKAGE_FILE_NAME}.x86_64.macosx")
   ENDIF()
+  SET(PACKAGE_BASE_NAME "${CPACK_PACKAGE_FILE_NAME}")
+  SET(PACKAGE_EXT "zip")
 
   add_custom_target(
-    post_package "${PROJECT_SOURCE_DIR}/scripts/package_post_build_macosx.sh" ${PACKAGE_BASE_NAME} ${EMBREE_SIGN_FILE}
-  )
-
-  add_custom_target(
-    test_package "${PROJECT_SOURCE_DIR}/scripts/package_test_macosx.sh" ${PACKAGE_BASE_NAME}
+    post_package_notarize "${PROJECT_SOURCE_DIR}/scripts/package_post_build_notarize_macosx.sh" ${PACKAGE_BASE_NAME} ${EMBREE_SIGN_FILE}
   )
 
 # Linux specific settings
 ELSE()
-
   SET(CPACK_GENERATOR TGZ)
   SET(CPACK_PACKAGE_FILE_NAME "${CPACK_PACKAGE_FILE_NAME}.x86_64.linux")
   SET(PACKAGE_BASE_NAME "${CPACK_PACKAGE_FILE_NAME}")
+  SET(PACKAGE_EXT "tar.gz")
   IF (EMBREE_SYCL_SUPPORT)
     SET(EMBREE_VERSION_SYCL_SUFFIX ".sycl")
   ENDIF()
-
-  add_custom_target(
-    post_package "${PROJECT_SOURCE_DIR}/scripts/package_post_build_linux.sh" ${PACKAGE_BASE_NAME}
-  )
-
-  add_custom_target(
-    test_package "${PROJECT_SOURCE_DIR}/scripts/package_test_linux.sh" ${PACKAGE_BASE_NAME}
-  )
-
 ENDIF()
 
-IF (EMBREE_TESTING_PACKAGE)
-  SET_TESTS_PROPERTIES(BuildPackage PROPERTIES TIMEOUT 1200)
-ENDIF()
+
+add_custom_target(
+  build ${CMAKE_COMMAND} --build . --config ${CMAKE_BUILD_TYPE} --target package -j8
+  COMMAND ${CMAKE_COMMAND} -DPACKAGE_BASENAME=${PACKAGE_BASE_NAME} -DPACKAGE_EXT=${PACKAGE_EXT} -P ${PROJECT_SOURCE_DIR}/scripts/package_build.cmake
+)
+
+add_custom_target(
+  test_package ${CMAKE_COMMAND} -DWHAT="UNPACK" -DPACKAGE_BASENAME=${PACKAGE_BASE_NAME} -DPACKAGE_EXT=${PACKAGE_EXT} -P ${PROJECT_SOURCE_DIR}/scripts/package_test.cmake
+  COMMAND cd embree_install/testing && ${CMAKE_COMMAND} -B build -DEMBREE_TESTING_INTENSITY=${EMBREE_TESTING_INTENSITY}
+  COMMAND ctest --test-dir ${CMAKE_CURRENT_BINARY_DIR}/embree_install/testing/build -VV -C ${CMAKE_BUILD_TYPE} --output-log ctest.output 
+  COMMAND ${CMAKE_COMMAND} -DWHAT="CHECK" -P ${PROJECT_SOURCE_DIR}/scripts/package_test.cmake
+)
+
+
+if(WIN32)
+  set(INTEGRATE_BINARY "./build/Release/test.exe")
+else()
+  set(INTEGRATE_BINARY "./build/test")
+endif()
+
+add_custom_target(
+  test_integration ${CMAKE_COMMAND} -DWHAT="UNPACK" -DPACKAGE_BASENAME=${PACKAGE_BASE_NAME} -DPACKAGE_EXT=${PACKAGE_EXT} -P ${PROJECT_SOURCE_DIR}/scripts/package_test.cmake 
+  COMMAND cd ${PROJECT_SOURCE_DIR}/tests/integration/test_embree_release && ${CMAKE_COMMAND} -B build --preset ${EMBREE_TESTING_INTEGRATION_PRESET} -Dembree_DIR="${CMAKE_CURRENT_BINARY_DIR}/embree_install/lib/cmake/embree-${EMBREE_VERSION}"
+  COMMAND cd ${PROJECT_SOURCE_DIR}/tests/integration/test_embree_release && ${CMAKE_COMMAND} --build build --config Release
+  COMMAND cd ${PROJECT_SOURCE_DIR}/tests/integration/test_embree_release && ${INTEGRATE_BINARY}
+)

common/cmake/test.cmake

@@ -90,7 +90,7 @@ if (BUILD_TESTING OR EMBREE_TESTING_INSTALL_TESTS)
                COMMAND ${executable}_sycl ${ARGN})
       SET(testnamedef EMBREE_TEST_${testname}_sycl_DEFINED)
       SET(${testnamedef} 1 CACHE INTERNAL "")
-      SET_TESTS_PROPERTIES(${testname}_sycl PROPERTIES TIMEOUT 50)
+      SET_TESTS_PROPERTIES(${testname}_sycl PROPERTIES TIMEOUT 150)
     ENDIF()
   ENDMACRO()
 
@@ -296,7 +296,6 @@ if (BUILD_TESTING OR EMBREE_TESTING_INSTALL_TESTS)
     #ADD_EMBREE_TEST_ECS(verify_memcheck verify NO_REFERENCE CONDITIONS "EMBREE_TESTING_MEMCHECK == ON")       # 8) same as 7) but with extra condition
 
   FUNCTION (ADD_EMBREE_TEST_ECS testname executable)
-
     if (EMBREE_TESTING_FILTER_TESTNAMES)
       list(FIND EMBREE_TESTING_FILTER_TESTNAMES "${testname}" i)
       if (${i} EQUAL -1)
@@ -442,6 +441,7 @@ if (BUILD_TESTING OR EMBREE_TESTING_INSTALL_TESTS)
       endforeach()
 
 
+
       if (memcheck)
         list(PREPEND args "${MY_PROJECT_BINARY_DIR}/${executable}")
         list(PREPEND args "${EMBREE_MEMORYCHECK_COMMAND_OPTIONS}")
@@ -464,13 +464,11 @@ if (BUILD_TESTING OR EMBREE_TESTING_INSTALL_TESTS)
         if (no_postfix)
           ADD_EMBREE_GENERIC_TEST(${testname} ${executable} ${args})
         else()
-          if((NOT ${EMBREE_SYCL_SUPPORT}) OR (NOT ${EMBREE_TESTING_ONLY_SYCL_TESTS}))
-            if (NOT(no_cpp))
-              ADD_EMBREE_GENERIC_CPP_TEST(${testname} ${executable} ${args})
-            endif()
-            if (NOT(no_ispc))
-              ADD_EMBREE_GENERIC_ISPC_TEST(${testname} ${executable} ${args})
-            endif()
+          if (NOT(no_cpp))
+            ADD_EMBREE_GENERIC_CPP_TEST(${testname} ${executable} ${args})
+          endif()
+          if (NOT(no_ispc))
+            ADD_EMBREE_GENERIC_ISPC_TEST(${testname} ${executable} ${args})
           endif()
           if (NOT(no_sycl))
             ADD_EMBREE_GENERIC_SYCL_TEST(${testname} ${executable} ${args}) 
@@ -555,17 +553,17 @@ if (BUILD_TESTING OR EMBREE_TESTING_INSTALL_TESTS)
     endif()
   ENDFUNCTION()
 
-  # add klocwork test
-  IF (EMBREE_TESTING_KLOCWORK)
-    ADD_TEST(NAME Klocwork-Build WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}" COMMAND "${PROJECT_SOURCE_DIR}/scripts/klocwork_build.sh")
-    ADD_TEST(NAME Klocwork-Check WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}" COMMAND "${PROJECT_SOURCE_DIR}/scripts/klocwork_check.sh")
-    SET_TESTS_PROPERTIES(Klocwork-Build PROPERTIES TIMEOUT 2400)
-    SET_TESTS_PROPERTIES(Klocwork-Check PROPERTIES TIMEOUT 300)
-  ENDIF()
-
 else()
   FUNCTION(ADD_EMBREE_TEST_ECS testname executable)
   ENDFUNCTION()
+  FUNCTION(ADD_EMBREE_GENERIC_TEST testname executable)
+  ENDFUNCTION()
+  FUNCTION(ADD_EMBREE_GENERIC_CPP_TEST testname executable)
+  ENDFUNCTION()
+  FUNCTION(ADD_EMBREE_GENERIC_SYCL_TEST testname executable)
+  ENDFUNCTION()
+  FUNCTION(ADD_EMBREE_GENERIC_ISPC_TEST testname executable)
+  ENDFUNCTION()
   FUNCTION(SET_EMBREE_TEST_PROPERTIES testname)
   ENDFUNCTION()
 endif()

common/math/emath.h

@@ -48,6 +48,9 @@ namespace embree
   __forceinline int   toInt  (const float& a) { return int(a); }
   __forceinline float toFloat(const int&   a) { return float(a); }
 
+  __forceinline int   asInt  (const float& a) { return *((int*)&a); }
+  __forceinline float asFloat(const int&   a) { return *((float*)&a); }
+
 #if defined(__WIN32__)
   __forceinline bool finite ( const float x ) { return _finite(x) != 0; }
 #endif

common/math/math_sycl.h

@@ -75,8 +75,6 @@ namespace embree
   __forceinline float floor( const float x ) { return sycl::floor(x); }
   __forceinline float ceil ( const float x ) { return sycl::ceil(x); }
   __forceinline float frac ( const float x ) { return x-floor(x); }
-  __forceinline float sincos ( const float phi, float* sin, float* cos ) { return *sin = sycl::sincos(phi, cos); }
-  __forceinline float sincosf ( const float phi, float* sin, float* cos ) { return *sin = sycl::sincos(phi, cos); }
 
   //__forceinline double abs  ( const double x ) { return ::fabs(x); }
   //__forceinline double sign ( const double x ) { return x<0?-1.0:1.0; }

common/math/vec2fa_sycl.h

@@ -95,8 +95,8 @@ namespace embree
   __forceinline Vec2fa abs  ( const Vec2fa& a ) { return Vec2fa(sycl::fabs(a.x),sycl::fabs(a.y)); }
   __forceinline Vec2fa sign ( const Vec2fa& a ) { return Vec2fa(sycl::sign(a.x),sycl::sign(a.y)); }
 
-   //__forceinline Vec2fa rcp  ( const Vec2fa& a ) { return Vec2fa(sycl::recip(a.x),sycl::recip(a.y)); }
-  __forceinline Vec2fa rcp  ( const Vec2fa& a ) { return Vec2fa(__sycl_std::__invoke_native_recip<float>(a.x),__sycl_std::__invoke_native_recip<float>(a.y)); }
+   //__forceinline Vec2fa rcp  ( const Vec2fa& a ) { return Vec2fa(sycl::native::recip(a.x),sycl::native::recip(a.y)); }
+  __forceinline Vec2fa rcp  ( const Vec2fa& a ) { return Vec2fa(sycl::native::recip(a.x),sycl::native::recip(a.y)); }
   __forceinline Vec2fa sqrt ( const Vec2fa& a ) { return Vec2fa(sycl::sqrt(a.x),sycl::sqrt(a.y)); }
   __forceinline Vec2fa sqr  ( const Vec2fa& a ) { return Vec2fa(a.x*a.x,a.y*a.y); }
   

common/math/vec3fa_sycl.h

@@ -92,7 +92,7 @@ namespace embree
   __forceinline Vec3fa sign ( const Vec3fa& a ) { return Vec3fa(sycl::sign(a.x),sycl::sign(a.y),sycl::sign(a.z)); }
 
   //__forceinline Vec3fa rcp  ( const Vec3fa& a ) { return Vec3fa(sycl::recip(a.x),sycl::recip(a.y),sycl::recip(a.z)); }
-  __forceinline Vec3fa rcp  ( const Vec3fa& a ) { return Vec3fa(__sycl_std::__invoke_native_recip<float>(a.x),__sycl_std::__invoke_native_recip<float>(a.y),__sycl_std::__invoke_native_recip<float>(a.z)); }
+  __forceinline Vec3fa rcp  ( const Vec3fa& a ) { return Vec3fa(sycl::native::recip(a.x),sycl::native::recip(a.y),sycl::native::recip(a.z)); }
   __forceinline Vec3fa sqrt ( const Vec3fa& a ) { return Vec3fa(sycl::sqrt(a.x),sycl::sqrt(a.y),sycl::sqrt(a.z)); }
   __forceinline Vec3fa sqr  ( const Vec3fa& a ) { return Vec3fa(a.x*a.x,a.y*a.y,a.z*a.z); }
 
@@ -393,7 +393,7 @@ namespace embree
   __forceinline Vec3fx sign ( const Vec3fx& a ) { return Vec3fx(sycl::sign(a.x),sycl::sign(a.y),sycl::sign(a.z),sycl::sign(a.z)); }
 
   //__forceinline Vec3fx rcp  ( const Vec3fx& a ) { return Vec3fx(sycl::recip(a.x),sycl::recip(a.y),sycl::recip(a.z)); }
-  __forceinline Vec3fx rcp  ( const Vec3fx& a ) { return Vec3fx(__sycl_std::__invoke_native_recip<float>(a.x),__sycl_std::__invoke_native_recip<float>(a.y),__sycl_std::__invoke_native_recip<float>(a.z),__sycl_std::__invoke_native_recip<float>(a.w)); }
+  __forceinline Vec3fx rcp  ( const Vec3fx& a ) { return Vec3fx(sycl::native::recip(a.x),sycl::native::recip(a.y),sycl::native::recip(a.z),sycl::native::recip(a.w)); }
   __forceinline Vec3fx sqrt ( const Vec3fx& a ) { return Vec3fx(sycl::sqrt(a.x),sycl::sqrt(a.y),sycl::sqrt(a.z),sycl::sqrt(a.w)); }
   __forceinline Vec3fx sqr  ( const Vec3fx& a ) { return Vec3fx(a.x*a.x,a.y*a.y,a.z*a.z,a.w*a.w); }
 
@@ -614,4 +614,12 @@ namespace embree
 
   //__forceinline Vec3ia::Vec3ia( const Vec3fx& a )
   //  : x((int)a.x), y((int)a.y), z((int)a.z) {}
+
 }
+
+#if __SYCL_COMPILER_VERSION >= 20210801
+namespace sycl {
+  template<> struct is_device_copyable<embree::Vec3fa> : std::true_type {};
+  template<> struct is_device_copyable<const embree::Vec3fa> : std::true_type {};
+}
+#endif

common/simd/arm/emulation.h

@@ -34,6 +34,7 @@ __forceinline __m128 _mm_broadcast_ss (float const * mem_addr)
 #define _MM_SET_EXCEPTION_MASK(x)
 // #define _MM_SET_FLUSH_ZERO_MODE(x)
 
+/*
 __forceinline int _mm_getcsr()
 {
   return 0;
@@ -43,6 +44,7 @@ __forceinline void _mm_mfence()
 {
   __sync_synchronize();
 }
+*/
 
 __forceinline __m128i _mm_load4epu8_epi32(__m128i *ptr)
 {

common/simd/vboolf4_sse2.h

@@ -119,7 +119,7 @@ namespace embree
 #if defined(__aarch64__)
   template<int i0, int i1, int i2, int i3>
   __forceinline vboolf4 shuffle(const vboolf4& v) {
-    return vreinterpretq_f32_u8(vqtbl1q_u8( vreinterpretq_u8_s32(v), _MN_SHUFFLE(i0, i1, i2, i3)));
+    return vreinterpretq_f32_u8(vqtbl1q_u8( vreinterpretq_u8_s32((int32x4_t)v.v), _MN_SHUFFLE(i0, i1, i2, i3)));
   }
 
   template<int i0, int i1, int i2, int i3>

common/simd/vfloat16_avx512.h

@@ -316,6 +316,17 @@ namespace embree
     return madd(t,b-a,a);
   }
 
+  __forceinline bool isvalid (const vfloat16& v) {
+    return all((v > vfloat16(-FLT_LARGE)) & (v < vfloat16(+FLT_LARGE)));
+  }
+
+  __forceinline void xchg(vboolf16 m, vfloat16& a, vfloat16& b)
+  {
+    vfloat16 c = a;
+    a = select(m,b,a);
+    b = select(m,c,b); 
+  }
+
   ////////////////////////////////////////////////////////////////////////////////
   /// Rounding Functions
   ////////////////////////////////////////////////////////////////////////////////

common/sys/alloc.cpp

@@ -12,51 +12,6 @@
   
 namespace embree
 {
-  size_t total_allocations = 0;
-
-#if defined(EMBREE_SYCL_SUPPORT)
-  
-  __thread sycl::context* tls_context_tutorial = nullptr;
-  __thread sycl::device* tls_device_tutorial = nullptr;
-  
-  __thread sycl::context* tls_context_embree = nullptr;
-  __thread sycl::device* tls_device_embree = nullptr;
-  
-  void enableUSMAllocEmbree(sycl::context* context, sycl::device* device)
-  {
-    if (tls_context_embree != nullptr) throw std::runtime_error("USM allocation already enabled");
-    if (tls_device_embree != nullptr) throw std::runtime_error("USM allocation already enabled");
-    tls_context_embree = context;
-    tls_device_embree = device;
-  }
-
-  void disableUSMAllocEmbree()
-  {
-    if (tls_context_embree  == nullptr) throw std::runtime_error("USM allocation not enabled");
-    if (tls_device_embree  == nullptr) throw std::runtime_error("USM allocation not enabled");
-    tls_context_embree = nullptr;
-    tls_device_embree = nullptr;
-  }
-
-  void enableUSMAllocTutorial(sycl::context* context, sycl::device* device)
-  {
-    //if (tls_context_tutorial != nullptr) throw std::runtime_error("USM allocation already enabled");
-    //if (tls_device_tutorial != nullptr) throw std::runtime_error("USM allocation already enabled");
-    tls_context_tutorial = context;
-    tls_device_tutorial = device;
-  }
-
-  void disableUSMAllocTutorial()
-  {
-    if (tls_context_tutorial  == nullptr) throw std::runtime_error("USM allocation not enabled");
-    if (tls_device_tutorial  == nullptr) throw std::runtime_error("USM allocation not enabled");
-    
-    tls_context_tutorial = nullptr;
-    tls_device_tutorial = nullptr;
-  }
-
-#endif
-  
   void* alignedMalloc(size_t size, size_t align)
   {
     if (size == 0)
@@ -71,8 +26,9 @@ namespace embree
 
   void alignedFree(void* ptr)
   {
-    if (ptr)
+    if (ptr) {
       _mm_free(ptr);
+    }
   }
 
 #if defined(EMBREE_SYCL_SUPPORT)
@@ -86,65 +42,67 @@ namespace embree
       return nullptr;
 
     assert((align & (align-1)) == 0);
-    total_allocations++;    
 
     void* ptr = nullptr;
-    if (mode == EMBREE_USM_SHARED_DEVICE_READ_ONLY)
+    if (mode == EmbreeUSMMode::DEVICE_READ_ONLY)
       ptr = sycl::aligned_alloc_shared(align,size,*device,*context,sycl::ext::oneapi::property::usm::device_read_only());
     else
       ptr = sycl::aligned_alloc_shared(align,size,*device,*context);
-      
+
+    if (size != 0 && ptr == nullptr)
+      throw std::bad_alloc();
+
+    return ptr;
+  }
+
+  void* alignedSYCLMalloc(sycl::context* context, sycl::device* device, size_t size, size_t align, EmbreeUSMMode mode, EmbreeMemoryType type)
+  {
+    assert(context);
+    assert(device);
+    
+    if (size == 0)
+      return nullptr;
+
+    assert((align & (align-1)) == 0);
+
+    void* ptr = nullptr;
+    if (type == EmbreeMemoryType::USM_SHARED) {
+      if (mode == EmbreeUSMMode::DEVICE_READ_ONLY)
+        ptr = sycl::aligned_alloc_shared(align,size,*device,*context,sycl::ext::oneapi::property::usm::device_read_only());
+      else
+        ptr = sycl::aligned_alloc_shared(align,size,*device,*context);
+    }
+    else if (type == EmbreeMemoryType::USM_HOST) {
+      ptr = sycl::aligned_alloc_host(align,size,*context);
+    }
+    else if (type == EmbreeMemoryType::USM_DEVICE) {
+      ptr = sycl::aligned_alloc_device(align,size,*device,*context);
+    }
+    else {
+      ptr = alignedMalloc(size,align);
+    }
+
     if (size != 0 && ptr == nullptr)
       throw std::bad_alloc();
 
     return ptr;
   }
   
-  static MutexSys g_alloc_mutex;
-  
-  void* alignedSYCLMalloc(size_t size, size_t align, EmbreeUSMMode mode)
-  {
-    if (tls_context_tutorial) return alignedSYCLMalloc(tls_context_tutorial, tls_device_tutorial, size, align, mode);
-    if (tls_context_embree  ) return alignedSYCLMalloc(tls_context_embree,   tls_device_embree,   size, align, mode);
-    return nullptr;
-  }
-
   void alignedSYCLFree(sycl::context* context, void* ptr)
   {
     assert(context);
     if (ptr) {
-      sycl::free(ptr,*context);
+      sycl::usm::alloc type = sycl::get_pointer_type(ptr, *context);
+      if (type == sycl::usm::alloc::host || type == sycl::usm::alloc::device || type == sycl::usm::alloc::shared)
+        sycl::free(ptr,*context);
+      else {
+        alignedFree(ptr);
+      }
     }
   }
 
-  void alignedSYCLFree(void* ptr)
-  {
-    if (tls_context_tutorial) return alignedSYCLFree(tls_context_tutorial, ptr);
-    if (tls_context_embree  ) return alignedSYCLFree(tls_context_embree, ptr);
-  }
-
 #endif
 
-  void* alignedUSMMalloc(size_t size, size_t align, EmbreeUSMMode mode)
-  {
-#if defined(EMBREE_SYCL_SUPPORT)
-    if (tls_context_embree || tls_context_tutorial)
-      return alignedSYCLMalloc(size,align,mode);
-    else
-#endif
-      return alignedMalloc(size,align);
-  }
-
-  void alignedUSMFree(void* ptr)
-  {
-#if defined(EMBREE_SYCL_SUPPORT)
-    if (tls_context_embree || tls_context_tutorial)
-      return alignedSYCLFree(ptr);
-    else
-#endif
-      return alignedFree(ptr);
-  }
-
   static bool huge_pages_enabled = false;
   static MutexSys os_init_mutex;
 

common/sys/alloc.h

@@ -9,71 +9,42 @@
 
 namespace embree
 {
-#if defined(EMBREE_SYCL_SUPPORT)
-
-  /* enables SYCL USM allocation */
-  void enableUSMAllocEmbree(sycl::context* context, sycl::device* device);
-  void enableUSMAllocTutorial(sycl::context* context, sycl::device* device);
-
-  /* disables SYCL USM allocation */
-  void disableUSMAllocEmbree();
-  void disableUSMAllocTutorial();
-
-#endif
-  
 #define ALIGNED_STRUCT_(align)                                            \
   void* operator new(size_t size) { return alignedMalloc(size,align); }   \
   void operator delete(void* ptr) { alignedFree(ptr); }                   \
   void* operator new[](size_t size) { return alignedMalloc(size,align); } \
   void operator delete[](void* ptr) { alignedFree(ptr); }
   
-#define ALIGNED_STRUCT_USM_(align)                                          \
-  void* operator new(size_t size) { return alignedUSMMalloc(size,align); }   \
-  void operator delete(void* ptr) { alignedUSMFree(ptr); }                   \
-  void* operator new[](size_t size) { return alignedUSMMalloc(size,align); } \
-  void operator delete[](void* ptr) { alignedUSMFree(ptr); }
-  
 #define ALIGNED_CLASS_(align)                                          \
  public:                                                               \
     ALIGNED_STRUCT_(align)                                             \
  private:
 
-#define ALIGNED_CLASS_USM_(align)                                          \
- public:                                                               \
-    ALIGNED_STRUCT_USM_(align)                                             \
- private:
-
-  enum EmbreeUSMMode {
-    EMBREE_USM_SHARED = 0,
-    EMBREE_USM_SHARED_DEVICE_READ_WRITE = 0,
-    EMBREE_USM_SHARED_DEVICE_READ_ONLY = 1
-  };
-  
   /*! aligned allocation */
   void* alignedMalloc(size_t size, size_t align);
   void alignedFree(void* ptr);
 
-  /*! aligned allocation using SYCL USM */
-  void* alignedUSMMalloc(size_t size, size_t align = 16, EmbreeUSMMode mode = EMBREE_USM_SHARED_DEVICE_READ_ONLY);
-  void alignedUSMFree(void* ptr);
+
+  enum class EmbreeUSMMode {
+    DEFAULT = 0,
+    DEVICE_READ_WRITE = 0,
+    DEVICE_READ_ONLY = 1
+  };
+
+  enum class EmbreeMemoryType {
+    USM_HOST = 0,
+    USM_DEVICE = 1,
+    USM_SHARED = 2,
+    MALLOC = 3
+  };
 
 #if defined(EMBREE_SYCL_SUPPORT)
-  
+
   /*! aligned allocation using SYCL USM */
   void* alignedSYCLMalloc(sycl::context* context, sycl::device* device, size_t size, size_t align, EmbreeUSMMode mode);
+  void* alignedSYCLMalloc(sycl::context* context, sycl::device* device, size_t size, size_t align, EmbreeUSMMode mode, EmbreeMemoryType type);
   void alignedSYCLFree(sycl::context* context, void* ptr);
 
-  // deleter functor to use as deleter in std unique or shared pointers that
-  // capture raw pointers created by sycl::malloc and it's variants
-  template<typename T>
-  struct sycl_deleter
-  {
-    void operator()(T const* ptr)
-    {
-      alignedUSMFree((void*)ptr);
-    }
-  };
-
 #endif
   
   /*! allocator that performs aligned allocations */

common/sys/barrier.cpp

@@ -101,7 +101,7 @@ namespace embree
     __forceinline void wait()
     {
       mutex.lock();
-      count++;
+      count = count + 1;
       
       if (count == barrierSize) {
         count = 0;

common/sys/filename.cpp

@@ -35,18 +35,6 @@ namespace embree
       filename.resize(filename.size()-1);
   }
   
-  /*! returns path to home folder */
-  FileName FileName::homeFolder() 
-  {
-#ifdef __WIN32__
-    const char* home = getenv("UserProfile");
-#else
-    const char* home = getenv("HOME");
-#endif
-    if (home) return home;
-    return "";
-  }
-
   /*! returns path to executable */
   FileName FileName::executableFolder() {
     return FileName(getExecutableFileName()).path();

common/sys/filename.h

@@ -20,9 +20,6 @@ namespace embree
 
     /*! create a valid filename from a string */
     FileName (const std::string& filename);
-    
-    /*! returns path to home folder */
-    static FileName homeFolder();
 
     /*! returns path to executable */
     static FileName executableFolder();

common/sys/intrinsics.h

@@ -223,7 +223,7 @@ namespace embree
 
 #if defined(EMBREE_SYCL_SUPPORT) && defined(__SYCL_DEVICE_ONLY__)
   __forceinline unsigned int   bsf(unsigned v) {
-    return SYCL_CTZ::ctz(v);
+    return sycl::ctz(v);
   }
 
 #else
@@ -248,7 +248,7 @@ namespace embree
   
 #if defined(EMBREE_SYCL_SUPPORT) && defined(__SYCL_DEVICE_ONLY__)
   __forceinline size_t bsf(size_t v) {
-    return SYCL_CTZ::ctz(v);
+    return sycl::ctz(v);
   }
 #else
   

common/sys/platform.h

@@ -34,13 +34,6 @@
 
 #pragma clang diagnostic pop
 
-#define SYCL_ONEAPI sycl
-#define SYCL_EXT_ONEAPI sycl::ext::oneapi
-#define SYCL_SUBGROUP sycl
-#define SYCL_EXPERIMENTAL sycl::ext::oneapi::experimental
-#define SYCL_INTEL sycl::ext::intel
-#define SYCL_CTZ sycl
-
 #include "sycl.h"
 
 #if defined(EMBREE_SYCL_SUPPORT) && defined(__SYCL_DEVICE_ONLY__)
@@ -367,38 +360,38 @@ __forceinline std::string toString(long long value) {
 
 #define sycl_printf0(format, ...) {               \
     static const CONSTANT char fmt[] = format;               \
-    if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))       \
-      SYCL_EXT_ONEAPI::experimental::printf(fmt, __VA_ARGS__ );  \
+    if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))       \
+      sycl::ext::oneapi::experimental::printf(fmt, __VA_ARGS__ );  \
   }
 
 #define sycl_printf0_(format) {               \
     static const CONSTANT char fmt[] = format;               \
-    if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))       \
-      SYCL_EXT_ONEAPI::experimental::printf(fmt);                \
+    if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))       \
+      sycl::ext::oneapi::experimental::printf(fmt);                \
   }
 
 #else
 
 #define sycl_printf0(format, ...) {                          \
     static const CONSTANT char fmt[] = format;               \
-    SYCL_EXT_ONEAPI::experimental::printf(fmt, __VA_ARGS__ );    \
+    sycl::ext::oneapi::experimental::printf(fmt, __VA_ARGS__ );    \
   }
 
 #define sycl_printf0_(format) {                              \
     static const CONSTANT char fmt[] = format;               \
-    SYCL_EXT_ONEAPI::experimental::printf(fmt);                  \
+    sycl::ext::oneapi::experimental::printf(fmt);                  \
   }
 
 #endif
 
 #define sycl_printf(format, ...) {               \
     static const CONSTANT char fmt[] = format;               \
-    SYCL_EXT_ONEAPI::experimental::printf(fmt, __VA_ARGS__ );    \
+    sycl::ext::oneapi::experimental::printf(fmt, __VA_ARGS__ );    \
   }
 
 #define sycl_printf_(format) {               \
     static const CONSTANT char fmt[] = format;               \
-    SYCL_EXT_ONEAPI::experimental::printf(fmt);                  \
+    sycl::ext::oneapi::experimental::printf(fmt);                  \
   }
 
 #if defined(EMBREE_SYCL_SUPPORT) && defined(__SYCL_DEVICE_ONLY__)
@@ -419,7 +412,7 @@ namespace embree
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, int   i)
   {
     if (cout.uniform) {
-      if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+      if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
         sycl_printf("%i",i);
     }
     else
@@ -431,7 +424,7 @@ namespace embree
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, unsigned int i)
   {
     if (cout.uniform) {
-      if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+      if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
         sycl_printf("%u",i);
     } else
       sycl_printf("%u ",i);
@@ -442,7 +435,7 @@ namespace embree
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, float f)
   {
     if (cout.uniform) {
-      if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+      if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
         sycl_printf("%f",f);
     } else
       sycl_printf("%f ",f);
@@ -453,7 +446,7 @@ namespace embree
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, double d)
   {
     if (cout.uniform) {
-      if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+      if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
         sycl_printf("%f",d);
     } else
       sycl_printf("%f ",d);
@@ -464,7 +457,7 @@ namespace embree
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, uint64_t l)
   {
     if (cout.uniform) {
-      if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+      if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
         sycl_printf("%lu",l);
     } else
       sycl_printf("%lu ",l);
@@ -475,7 +468,7 @@ namespace embree
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, long l)
   {
     if (cout.uniform) {
-      if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+      if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
         sycl_printf("%l",l);
     } else
       sycl_printf("%l ",l);
@@ -487,7 +480,7 @@ namespace embree
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, void* p)
   {
     if (cout.uniform) {
-      if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+      if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
         sycl_printf("%p",p);
     } else
       sycl_printf("%p ",p);
@@ -497,14 +490,14 @@ namespace embree
 
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, const char* c)
   {
-     if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+     if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
        sycl_printf("%s",c);
      return cout;
   }
 
   inline sycl_ostream_ operator <<(sycl_ostream_ cout, sycl_endl_)
   {
-    if (get_sub_group_local_id() == SYCL_CTZ::ctz(intel_sub_group_ballot(true)))
+    if (get_sub_group_local_id() == sycl::ctz(intel_sub_group_ballot(true)))
       sycl_printf_("\n");
     return cout;
   }

common/sys/sycl.h

@@ -31,9 +31,6 @@ using sycl::ushort4;
 using sycl::ushort3;
 using sycl::ushort2;
 
-using sycl::fmax;
-using sycl::fmin;
-
 #ifdef __SYCL_DEVICE_ONLY__
 #define GLOBAL __attribute__((opencl_global))
 #define LOCAL  __attribute__((opencl_local))
@@ -184,8 +181,15 @@ namespace embree
     return sycl::select(b,a,mask);
   }
   
-  __forceinline const SYCL_EXT_ONEAPI::sub_group this_sub_group() {
-    return SYCL_EXPERIMENTAL::this_sub_group(); 
+#define XSTR(x) STR(x)
+#define STR(x) #x
+
+  __forceinline const sycl::sub_group this_sub_group() {
+#if __LIBSYCL_MAJOR_VERSION >= 8
+    return sycl::ext::oneapi::this_work_item::get_sub_group();
+#else
+    return sycl::ext::oneapi::experimental::this_sub_group();
+#endif
   }
   
   __forceinline const uint32_t get_sub_group_local_id() {
@@ -209,23 +213,23 @@ namespace embree
   }
 
   __forceinline bool sub_group_all_of(bool pred) {
-    return SYCL_ONEAPI::all_of_group(this_sub_group(),pred);
+    return sycl::all_of_group(this_sub_group(),pred);
   }
 
   __forceinline bool sub_group_any_of(bool pred) {
-    return SYCL_ONEAPI::any_of_group(this_sub_group(),pred);
+    return sycl::any_of_group(this_sub_group(),pred);
   }
   
   __forceinline bool sub_group_none_of(bool pred) {
-    return SYCL_ONEAPI::none_of_group(this_sub_group(),pred);
+    return sycl::none_of_group(this_sub_group(),pred);
   }
 
   template <typename T> __forceinline T sub_group_broadcast(T x, sycl::id<1> local_id) {
-    return SYCL_SUBGROUP::group_broadcast<SYCL_EXT_ONEAPI::sub_group>(this_sub_group(),x,local_id);
+    return sycl::group_broadcast<sycl::sub_group>(this_sub_group(),x,local_id);
   }
   
   template <typename T> __forceinline T sub_group_make_uniform(T x) {
-    return sub_group_broadcast(x,SYCL_CTZ::ctz(intel_sub_group_ballot(true)));
+    return sub_group_broadcast(x,sycl::ctz(intel_sub_group_ballot(true)));
   }
 
   __forceinline void assume_uniform_array(void* ptr) {
@@ -235,59 +239,47 @@ namespace embree
   }
 
   template <typename T, class BinaryOperation> __forceinline T sub_group_reduce(T x, BinaryOperation binary_op) {
-    return SYCL_SUBGROUP::reduce_over_group<SYCL_EXT_ONEAPI::sub_group>(this_sub_group(),x,binary_op);
+    return sycl::reduce_over_group<sycl::sub_group>(this_sub_group(),x,binary_op);
   }
 
   template <typename T, class BinaryOperation> __forceinline T sub_group_reduce(T x, T init, BinaryOperation binary_op) {
-    return SYCL_SUBGROUP::reduce_over_group<SYCL_EXT_ONEAPI::sub_group>(this_sub_group(),x,init,binary_op);
+    return sycl::reduce_over_group<sycl::sub_group>(this_sub_group(),x,init,binary_op);
   }
   
   template <typename T> __forceinline T sub_group_reduce_min(T x, T init) {
-    return sub_group_reduce(x, init, SYCL_EXT_ONEAPI::minimum<T>());
+    return sub_group_reduce(x, init, sycl::ext::oneapi::minimum<T>());
   }
 
   template <typename T> __forceinline T sub_group_reduce_min(T x) {
-    return sub_group_reduce(x, SYCL_EXT_ONEAPI::minimum<T>());
+    return sub_group_reduce(x, sycl::ext::oneapi::minimum<T>());
   }
 
   template <typename T> __forceinline T sub_group_reduce_max(T x) {
-    return sub_group_reduce(x, SYCL_EXT_ONEAPI::maximum<T>());
+    return sub_group_reduce(x, sycl::ext::oneapi::maximum<T>());
   }
   
   template <typename T> __forceinline T sub_group_reduce_add(T x) {
-    return sub_group_reduce(x, SYCL_EXT_ONEAPI::plus<T>());
+    return sub_group_reduce(x, sycl::ext::oneapi::plus<T>());
   }
 
   template <typename T, class BinaryOperation> __forceinline T sub_group_exclusive_scan(T x, BinaryOperation binary_op) {
-    return SYCL_SUBGROUP::exclusive_scan_over_group(this_sub_group(),x,binary_op);
+    return sycl::exclusive_scan_over_group(this_sub_group(),x,binary_op);
   }
 
   template <typename T, class BinaryOperation> __forceinline T sub_group_exclusive_scan_min(T x) {
-    return sub_group_exclusive_scan(x,SYCL_EXT_ONEAPI::minimum<T>());
+    return sub_group_exclusive_scan(x,sycl::ext::oneapi::minimum<T>());
   }
 
   template <typename T, class BinaryOperation> __forceinline T sub_group_exclusive_scan(T x, T init, BinaryOperation binary_op) {
-    return SYCL_SUBGROUP::exclusive_scan_over_group(this_sub_group(),x,init,binary_op);
+    return sycl::exclusive_scan_over_group(this_sub_group(),x,init,binary_op);
   }
 
   template <typename T, class BinaryOperation> __forceinline T sub_group_inclusive_scan(T x, BinaryOperation binary_op) {
-    return SYCL_SUBGROUP::inclusive_scan_over_group(this_sub_group(),x,binary_op);
+    return sycl::inclusive_scan_over_group(this_sub_group(),x,binary_op);
   }
 
   template <typename T, class BinaryOperation> __forceinline T sub_group_inclusive_scan(T x, BinaryOperation binary_op, T init) {
-    return SYCL_SUBGROUP::inclusive_scan_over_group(this_sub_group(),x,binary_op,init);
-  }
-
-  template <typename T> __forceinline T sub_group_shuffle(T x, sycl::id<1> local_id) {
-    return this_sub_group().shuffle(x, local_id);
-  }
-
-  template <typename T> __forceinline T sub_group_shuffle_down(T x, uint32_t delta) {
-    return this_sub_group().shuffle_down(x, delta);
-  }
-  
-  template <typename T> __forceinline T sub_group_shuffle_up(T x, uint32_t delta) {
-    return this_sub_group().shuffle_up(x, delta);
+    return sycl::inclusive_scan_over_group(this_sub_group(),x,binary_op,init);
   }
 
   template <typename T> __forceinline T sub_group_load(const void* src) {

common/sys/sysinfo.cpp

@@ -344,6 +344,19 @@ namespace embree
     if (cpuid_leaf_7[EBX] & CPU_FEATURE_BIT_AVX512VL  ) cpu_features |= CPU_FEATURE_AVX512VL;
     if (cpuid_leaf_7[ECX] & CPU_FEATURE_BIT_AVX512VBMI) cpu_features |= CPU_FEATURE_AVX512VBMI;
 
+#if defined(__MACOSX__)
+    if (   (cpu_features & CPU_FEATURE_AVX512F)
+        || (cpu_features & CPU_FEATURE_AVX512DQ)
+        || (cpu_features & CPU_FEATURE_AVX512CD)
+        || (cpu_features & CPU_FEATURE_AVX512BW)
+        || (cpu_features & CPU_FEATURE_AVX512VL) )
+      {
+        // on macOS AVX512 will be enabled automatically by the kernel when the first AVX512 instruction is called
+        // see https://github.com/apple/darwin-xnu/blob/0a798f6738bc1db01281fc08ae024145e84df927/osfmk/i386/fpu.c#L176
+        // therefore we ignore the state of XCR0
+        cpu_features |= CPU_FEATURE_ZMM_ENABLED;
+      }
+#endif
     return cpu_features;
 
 #elif defined(__ARM_NEON) || defined(__EMSCRIPTEN__)

common/sys/sysinfo.h

@@ -158,7 +158,7 @@ namespace embree
   static const int SSE41  = SSSE3 | CPU_FEATURE_SSE41;
   static const int SSE42  = SSE41 | CPU_FEATURE_SSE42 | CPU_FEATURE_POPCNT;
   static const int AVX    = SSE42 | CPU_FEATURE_AVX | CPU_FEATURE_YMM_ENABLED;
-  static const int AVXI   = AVX | CPU_FEATURE_F16C | CPU_FEATURE_RDRAND;
+  static const int AVXI   = AVX | CPU_FEATURE_F16C;
   static const int AVX2   = AVXI | CPU_FEATURE_AVX2 | CPU_FEATURE_FMA3 | CPU_FEATURE_BMI1 | CPU_FEATURE_BMI2 | CPU_FEATURE_LZCNT;
   static const int AVX512 = AVX2 | CPU_FEATURE_AVX512F | CPU_FEATURE_AVX512DQ | CPU_FEATURE_AVX512CD | CPU_FEATURE_AVX512BW | CPU_FEATURE_AVX512VL | CPU_FEATURE_ZMM_ENABLED;
   static const int NEON = CPU_FEATURE_NEON | CPU_FEATURE_SSE | CPU_FEATURE_SSE2;

common/sys/vector.h

@@ -5,6 +5,7 @@
 
 #include "alloc.h"
 #include <algorithm>
+#include <type_traits>
 
 namespace embree
 {
@@ -124,8 +125,10 @@ namespace embree
 
       __forceinline       T* data()       { return items; };
       __forceinline const T* data() const { return items; };
+      
+      /* dangerous only use if you know what you're doing */
+      __forceinline void setDataPtr(T* data) { items = data; }
 
-     
       /******************** Modifiers **************************/
 
       __forceinline void push_back(const T& nt) 

common/tasking/taskschedulerinternal.cpp

@@ -8,8 +8,6 @@
 
 namespace embree
 {
-  RTC_NAMESPACE_BEGIN
-  
   static MutexSys g_mutex;
   size_t TaskScheduler::g_numThreads = 0;
   __thread TaskScheduler* TaskScheduler::g_instance = nullptr;
@@ -150,7 +148,8 @@ namespace embree
   {
     Lock<MutexSys> lock(g_mutex);
     assert(newNumThreads);
-    newNumThreads = min(newNumThreads, (size_t) getNumberOfLogicalThreads());
+    if (newNumThreads == std::numeric_limits<size_t>::max())
+      newNumThreads = (size_t) getNumberOfLogicalThreads();
 
     numThreads = newNumThreads;
     if (!startThreads && !running) return;
@@ -230,7 +229,8 @@ namespace embree
   TaskScheduler::TaskScheduler()
     : threadCounter(0), anyTasksRunning(0), hasRootTask(false)
   {
-    threadLocal.resize(2*getNumberOfLogicalThreads()); // FIXME: this has to be 2x as in the compatibility join mode with rtcCommitScene the worker threads also join. When disallowing rtcCommitScene to join a build we can remove the 2x.
+    assert(threadPool);
+    threadLocal.resize(2 * TaskScheduler::threadCount()); // FIXME: this has to be 2x as in the compatibility join mode with rtcCommitScene the worker threads also join. When disallowing rtcCommitScene to join a build we can remove the 2x.
     for (size_t i=0; i<threadLocal.size(); i++)
       threadLocal[i].store(nullptr);
   }
@@ -397,6 +397,4 @@ namespace embree
   dll_export void TaskScheduler::removeScheduler(const Ref<TaskScheduler>& scheduler) {
     threadPool->remove(scheduler);
   }
-
-  RTC_NAMESPACE_END
 }

common/tasking/taskschedulerinternal.h

@@ -3,6 +3,7 @@
 
 #pragma once
 
+#include "../../include/embree4/rtcore.h"
 #include "../sys/platform.h"
 #include "../sys/alloc.h"
 #include "../sys/barrier.h"
@@ -12,17 +13,12 @@
 #include "../sys/ref.h"
 #include "../sys/atomic.h"
 #include "../math/range.h"
-#include "../../include/embree4/rtcore.h"
 
+#include <exception>
 #include <list>
 
 namespace embree
 {
-
-  /* The tasking system exports some symbols to be used by the tutorials. Thus we 
-     hide is also in the API namespace when requested. */
-  RTC_NAMESPACE_BEGIN
-
   struct TaskScheduler : public RefCount
   {
     ALIGNED_STRUCT_(64);
@@ -375,10 +371,4 @@ namespace embree
     static __thread Thread* thread_local_thread;
     static ThreadPool* threadPool;
   };
-
-  RTC_NAMESPACE_END
-
-#if defined(RTC_NAMESPACE)
-    using RTC_NAMESPACE::TaskScheduler;
-#endif
 }

common/tasking/taskschedulertbb.cpp

@@ -65,7 +65,7 @@ namespace embree
       BarrierSys barrier(numThreads);
       tbb::parallel_for(size_t(0), size_t(numThreads), size_t(1), [&] ( size_t i ) {
           barrier.wait();
-        });
+        }, tbb::static_partitioner());
     }
   }
 

doc/CMakeLists.txt

@@ -22,5 +22,5 @@ FetchContent_Declare(IntelStyle
 
 FetchContent_GetProperties(IntelStyle)
 if(NOT IntelStyle_POPULATED)
-  FetchContent_Populate(IntelStyle)
+  FetchContent_MakeAvailable(IntelStyle)
 endif()

doc/Makefile

@@ -10,7 +10,7 @@ webimages := $(addprefix www/, $(images_jpg))
 docfiles := $(addprefix doc/, README.md CHANGELOG.md readme.pdf)
 spec := embree-spec.rst
 man_src := $(wildcard src/api/*.md)
-manfiles := $(patsubst src/api/%.md,man/man3/%.4embree4,$(man_src))
+manfiles := $(patsubst src/api/%.md,man/man3/%.3embree4,$(man_src))
 
 all: www images man doc spec
 www: $(webpages) $(webimages)
@@ -175,7 +175,7 @@ doc/readme.pdf: tmp/readme.tex tmp/readme_all.tex $(images_jpg) $(images_pdf)
 ### manpages
 ########################################################################
 
-man/man3/%.4embree4: src/api/%.md
+man/man3/%.3embree4: src/api/%.md
 	$(PANDOC) -s -t man $< -o $@
 
 ########################################################################

doc/src/api-ref.md

@@ -43,6 +43,16 @@
 ```
 \pagebreak
 
+## rtcGetDeviceLastErrorMessage
+``` {include=src/api/rtcGetDeviceLastErrorMessage.md}
+```
+\pagebreak
+
+## rtcGetErrorString
+``` {include=src/api/rtcGetErrorString.md}
+```
+\pagebreak
+
 ## rtcSetDeviceErrorFunction
 ``` {include=src/api/rtcSetDeviceErrorFunction.md}
 ```
@@ -103,6 +113,11 @@
 ```
 \pagebreak
 
+## rtcCommitSceneWithQueue
+``` {include=src/api/rtcCommitSceneWithQueue.md}
+```
+\pagebreak
+
 ## rtcJoinCommitScene
 ``` {include=src/api/rtcJoinCommitScene.md}
 ```
@@ -139,6 +154,11 @@
 ```
 \pagebreak
 
+## rtcGetSceneTraversable
+``` {include=src/api/rtcGetSceneTraversable.md}
+```
+\pagebreak
+
 ## rtcNewGeometry
 ``` {include=src/api/rtcNewGeometry.md}
 ```
@@ -184,6 +204,11 @@
 ```
 \pagebreak
 
+## RTC_GEOMETRY_TYPE_INSTANCE_ARRAY
+``` {include=src/api/RTC_GEOMETRY_TYPE_INSTANCE_ARRAY.md}
+```
+\pagebreak
+
 ## RTCCurveFlags
 ``` {include=src/api/RTCCurveFlags.md}
 ```
@@ -254,11 +279,21 @@
 ```
 \pagebreak
 
+## rtcSetSharedGeometryBufferHostDevice
+``` {include=src/api/rtcSetSharedGeometryBufferHostDevice.md}
+```
+\pagebreak
+
 ## rtcSetNewGeometryBuffer
 ``` {include=src/api/rtcSetNewGeometryBuffer.md}
 ```
 \pagebreak
 
+## rtcSetNewGeometryBufferHostDevice
+``` {include=src/api/rtcSetNewGeometryBufferHostDevice.md}
+```
+\pagebreak
+
 ## RTCFormat
 ``` {include=src/api/RTCFormat.md}
 ```
@@ -274,6 +309,11 @@
 ```
 \pagebreak
 
+## rtcGetGeometryBufferDataDevice
+``` {include=src/api/rtcGetGeometryBufferDataDevice.md}
+```
+\pagebreak
+
 ## rtcUpdateGeometryBuffer
 ``` {include=src/api/rtcUpdateGeometryBuffer.md}
 ```
@@ -319,6 +359,11 @@
 ```
 \pagebreak
 
+## rtcGetGeometryUserDataFromTraversable
+``` {include=src/api/rtcGetGeometryUserDataFromScenTraversable
+```
+\pagebreak
+
 ## rtcSetGeometryUserPrimitiveCount
 ``` {include=src/api/rtcSetGeometryUserPrimitiveCount.md}
 ```
@@ -355,6 +400,11 @@
 ```
 \pagebreak
 
+## rtcSetGeometryInstancedScenes
+``` {include=src/api/rtcSetGeometryInstancedScenes.md}
+```
+\pagebreak
+
 ## rtcSetGeometryTransform
 ``` {include=src/api/rtcSetGeometryTransform.md}
 ```
@@ -370,11 +420,21 @@
 ```
 \pagebreak
 
+## rtcGetGeometryTransformEx
+``` {include=src/api/rtcGetGeometryTransformEx.md}
+```
+\pagebreak
+
 ## rtcGetGeometryTransformFromScene
 ``` {include=src/api/rtcGetGeometryTransformFromScene.md}
 ```
 \pagebreak
 
+## rtcGetGeometryTransformFromTraversable
+``` {include=src/api/rtcGetGeometryTransformFromTraversable.md}
+```
+\pagebreak
+
 ## rtcSetGeometryTessellationRate
 ``` {include=src/api/rtcSetGeometryTessellationRate.md}
 ```
@@ -435,17 +495,26 @@
 ```
 \pagebreak
 
-
 ## rtcNewBuffer
 ``` {include=src/api/rtcNewBuffer.md}
 ```
 \pagebreak
 
+## rtcNewBufferHostDevice
+``` {include=src/api/rtcNewBufferHostDevice.md}
+```
+\pagebreak
+
 ## rtcNewSharedBuffer
 ``` {include=src/api/rtcNewSharedBuffer.md}
 ```
 \pagebreak
 
+## rtcNewSharedBufferHostDevice
+``` {include=src/api/rtcNewSharedBufferHostDevice.md}
+```
+\pagebreak
+
 ## rtcRetainBuffer
 ``` {include=src/api/rtcRetainBuffer.md}
 ```
@@ -461,6 +530,21 @@
 ```
 \pagebreak
 
+## rtcGetBufferDataDevice
+``` {include=src/api/rtcGetBufferDataDevice.md}
+```
+\pagebreak
+
+## rtcCommitBuffer
+``` {include=src/api/rtcCommitBuffer.md}
+```
+\pagebreak
+
+## rtcCommitBufferWithQueue
+``` {include=src/api/rtcCommitBufferWithQueue.md}
+```
+\pagebreak
+
 ## RTCRay
 ``` {include=src/api/RTCRay.md}
 ```
@@ -516,41 +600,81 @@
 ```
 \pagebreak
 
+## rtcTraversableIntersect1
+``` {include=src/api/rtcTraversableIntersect1.md}
+```
+\pagebreak
+
 ## rtcOccluded1
 ``` {include=src/api/rtcOccluded1.md}
 ```
 \pagebreak
 
+## rtcTraversableOccluded1
+``` {include=src/api/rtcTraversableOccluded1.md}
+```
+\pagebreak
+
 ## rtcIntersect4/8/16
 ``` {include=src/api/rtcIntersect4.md}
 ```
 \pagebreak
 
+## rtcTraversableIntersect4/8/16
+``` {include=src/api/rtcTraversableIntersect4.md}
+```
+\pagebreak
+
 ## rtcOccluded4/8/16
 ``` {include=src/api/rtcOccluded4.md}
 ```
 \pagebreak
 
+## rtcTraversableOccluded4/8/16
+``` {include=src/api/rtcTraversableOccluded4.md}
+```
+\pagebreak
+
 ## rtcForwardIntersect1
 ``` {include=src/api/rtcForwardIntersect1.md}
 ```
 \pagebreak
 
+## rtcTraversableForwardIntersect1
+``` {include=src/api/rtcTraversableForwardIntersect1.md}
+```
+\pagebreak
+
 ## rtcForwardOccluded1
 ``` {include=src/api/rtcForwardOccluded1.md}
 ```
 \pagebreak
 
+## rtcTraversableForwardOccluded1
+``` {include=src/api/rtcTraversableForwardOccluded1.md}
+```
+\pagebreak
+
 ## rtcForwardIntersect4/8/16
 ``` {include=src/api/rtcForwardIntersect4.md}
 ```
 \pagebreak
 
+## rtcTraversableForwardIntersect4/8/16
+``` {include=src/api/rtcTraversableForwardIntersect4.md}
+```
+\pagebreak
+
 ## rtcForwardOccluded4/8/16
 ``` {include=src/api/rtcForwardOccluded4.md}
 ```
 \pagebreak
 
+## rtcTraversableForwardOccluded4/8/16
+``` {include=src/api/rtcTraversableForwardOccluded4.md}
+```
+\pagebreak
+
 ## rtcInitPointQueryContext
 ``` {include=src/api/rtcInitPointQueryContext.md}
 ```
@@ -561,6 +685,21 @@
 ```
 \pagebreak
 
+## rtcTraversablePointQuery
+``` {include=src/api/rtcTraversablePointQuery.md}
+```
+\pagebreak
+
+## rtcPointQuery4/8/16
+``` {include=src/api/rtcPointQuery4.md}
+```
+\pagebreak
+
+## rtcTraversablePointQuery4/8/16
+``` {include=src/api/rtcTraversablePointQuery4.md}
+```
+\pagebreak
+
 ## rtcCollide
 ``` {include=src/api/rtcCollide.md}
 ```

doc/src/api.md

@@ -46,6 +46,13 @@ CPU with such a device. To render on the CPU and GPU in parallel, the
 user has to create a second Embree device and create a second
 scene to be used on the CPU.
 
+Starting with Embree 4.4 scene objects (`RTCScene` types) are
+not valid handles on SYCL devices anymore and therefore can not
+be used for Embree API calls in a SYCL kernel. Instead, Embree API
+calls on a SYCL kernel have a variation which use traversable objects
+(`RTCTraversable` type). To get a traversable object for a scene object
+the application can call `rtcGetSceneTraversable`.
+
 Files containing SYCL code, have to get compiled with the Intel®
 oneAPI DPC++ compiler. Please see section [Linux SYCL Compilation] and
 [Windows SYCL Compilation] for supported compilers. The DPC++ compiler
@@ -53,18 +60,29 @@ performs a two-phase compilation, where host code is compiled in a
 first phase, and device code compiled in a second compilation phase.
 
 Standard Embree API functions for scene construction can get used on
-the host but not the device. Data buffers that are shared with Embree
-(e.g. for vertex of index buffers) have to get allocated as SYCL
-unified shared memory (USM memory), using the `sycl::malloc` or
+the host but not the device. 
+
+Before version 4.4, Embree made heavy use of unified shared memory (USM)
+shared memory which simplifies memory management with SYCL devices
+by letting the SYCL runtime transfer data from host to device implicitly.
+However, some applications require more control over when and how data
+is migrated from CPU to GPU. Embree 4.4 allows to use explicit host and device
+memory allocations. See for example `rtcSetNewGeometryBufferHostDevice`,
+`rtcSetSharedGeometryBufferHostDevice`, `rtcNewBufferHostDevice`, and
+`rtcNewSharedBufferHostDevice`. It is still possible to share
+data buffers with Embree using SYCL USM shared memory by using
+the API calls without the `HostDevice` suffix.
+
+The easiest way to share data buffers with Embree
+(e.g. for vertex of index buffers) is to allocate the
+data as USM shared memory, using the `sycl::malloc` or
 `sycl::aligned_alloc` calls with `sycl::usm::alloc::shared` property,
 or the sycl::aligned_alloc_shared call, e.g:
 
     void* ptr = sycl::aligned_alloc(16, bytes, queue, sycl::usm::alloc::shared);
 
 These shared allocations have to be valid during rendering, as Embree
-may access contained data when tracing rays. Embree does not support
-device-only memory allocations, as the BVH builder implemented on the
-CPU relies on reading the data buffers.
+may access contained data when tracing rays.
 
 Device side rendering can get invoked by submitting a SYCL
 `parallel_for` to the SYCL queue:
@@ -72,6 +90,8 @@ Device side rendering can get invoked by submitting a SYCL
     const sycl::specialization_id<RTCFeatureFlags> feature_mask;
 
     RTCFeatureFlags required_features = RTC_FEATURE_FLAG_TRIANGLE;
+    
+    RTCTraversable traversable = rtcGetSceneTraversable(scene);
 
     queue.submit([=](sycl::handler& cgh)
     {
@@ -99,7 +119,7 @@ Device side rendering can get invoked by submitting a SYCL
         rayhit.hit.geomID = RTC_INVALID_GEOMETRY_ID;
         rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID;
   
-        rtcIntersect1(scene, &rayhit, &args);
+        rtcTraversableIntersect1(traversable, &rayhit, &args);
   
         result->geomID = rayhit.hit.geomID;
         result->primID = rayhit.hit.primID;
@@ -109,21 +129,23 @@ Device side rendering can get invoked by submitting a SYCL
     queue.wait_and_throw();
 
 
-This example passes a feature mask using a specialization contant to
-the `rtcIntersect1` function, which is recommended for GPU
+This example passes a feature mask using a specialization constant to
+the `rtcTraversableIntersect1` function, which is recommended for GPU
 rendering. For best performance, this feature mask should get used to
 enable only features required by the application to render the scene,
 e.g. just triangles in this example.
 
 Inside the SYCL `parallel_for` loop you can use rendering related functions,
-such as the `rtcIntersect1` and `rtcOccluded1` functions to trace rays,
-`rtcForwardIntersect1` and `rtcForwardOccluded1` to continue object
-traversal from inside a user geometry callback, 
-and `rtcGetGeometryUserDataFromScene` to get the user data pointer of some
+such as the `rtcTraversableIntersect1` and `rtcTraversableOccluded1` functions 
+to trace rays, `rtcTraversableForwardIntersect1/Ex` and `rtcTraversableForwardOccluded1/Ex`
+to continue object traversal from inside a user geometry callback,
+and `rtcGetGeometryUserDataFromTraversable` to get the user data pointer of some
 geometry.
 
-Have a look at the [Minimal] tutorial for a minimal SYCL example.
-
+Have a look at the [Minimal] tutorial for a minimal SYCL example and the
+[Host Device Memory] tutorial shows four different ways in which
+data buffers can be created by or shared with Embree using explicit
+host/device data buffers.
 
 SYCL JIT caching
 -----------------
@@ -169,7 +191,7 @@ device. To render on the CPU just use the standard C99 API without
 relying on SYCL.
 
 The SYCL language spec puts some restrictions to device functions,
-such as disallowing: global variable access, malloc, invokation of
+such as disallowing: global variable access, malloc, invocation of
 virtual functions, function pointers, runtime type information,
 exceptions, recursion, etc. See Section `5.4. Language Restrictions
 for device functions` of the [SYCL
@@ -183,10 +205,16 @@ reasons.
 Some features are not supported by the Embree SYCL API thus cannot
 get used on the GPU:
 
-- The packet tracing functions `rtcIntersect4/8/16` and
-  `rtcOccluded4/8/16` are not supported in SYCL
+- Since Embree 4.4, all the ray query functions
+  that take an `RTCScene` object as argument cannot get used in
+  SYCL device side code. Instead, the API functions
+  taking a `RTCTraversable` object (e.g. `rtcTraversableIntersect1`)
+  have to be used.
+
+- The packet tracing functions `rtcTraversableIntersect4/8/16` and
+  `rtcTraversableOccluded4/8/16` are not supported in SYCL
   device side code. Using these functions makes no sense for SYCL, as
-  the programming model is implicitely executed in SIMT mode on the
+  the programming model is implicitly executed in SIMT mode on the
   GPU anyway.
 
 - Filter and user geometry callbacks stored inside the geometry
@@ -203,10 +231,10 @@ get used on the GPU:
   interpolating over, thus its implementation on the GPU would contain
   a large switch statement for all potential geometry types.
 
-- Tracing rays using `rtcIntersect1` and `rtcOccluded1` functions from
-  user geometry callbacks is not supported in SYCL. Please use the
-  tail recursive `rtcForwardIntersect1` and `rtcForwardOccluded1`
-  calls instead.
+- Tracing rays using `rtcTraversableIntersect1` and `rtcTraversableOccluded1`
+  functions from user geometry callbacks is not supported in SYCL. 
+  Please use the tail recursive `rtcTraversableForwardIntersect1` 
+  and `rtcTraversableForwardOccluded1` calls instead.
 
 - Subdivision surfaces are not supported for Embree SYCL devices.
 
@@ -220,12 +248,8 @@ get used on the GPU:
 Embree SYCL Known Issues
 ------------------------
 
-- The SYCL support of Embree is in beta phase. Current functionality,
-  quality, and GPU performance may not reflect that of the final
-  product.
-  
-- Compilation with build configuration "debug" is currently not working on
-  Windows.
+- Compilation with build configuration "debug" is currently not feasible because
+  compilation times are very long.
 
 
 Upgrading from Embree 3 to Embree 4
@@ -303,20 +327,21 @@ required:
   instantiated object. In Embree 4 using `rtcIntersect` recursively is
   disallowed on the GPU but still supported on the CPU. To properly
   continue a ray inside an instantiated object use the new
-  `rtcForwardIntersect1` and `rtcForwardOccluded1` functions.
+  `rtc(Traversable)ForwardIntersect1` and `rtc(Traversable)ForwardOccluded1` functions.
 
-- The geometry object of Embree 4 is a host side only object, thus
+- The geometry object and scene object of Embree 4 are a host side only objects, thus
   accessing it during rendering from the GPU is not allowed. Thus all
-  API functions that take an RTCGeometry object as argument cannot get
-  used during rendering. Thus in particular the
-  `rtcGetGeometryUserData(RTCGeometry)` call cannot get used, but
-  there is an alternative function
-  `rtcGetGeometryUserDataFromScene(RTCScene scene,uint geomID)` that
-  should get used instead.
+  API functions that take an RTCGeometry object or RTCScene object as argument cannot get
+  used during rendering. In particular the `rtcGetGeometryUserData(RTCGeometry)`
+  call cannot get used, but there is an alternative function
+  `rtcGetGeometryUserDataFromTraversable(RTCTraversable traversable,uint geomID)` that
+  should get used instead. To perform ray queries on the GPU (e.g. `rtcTraversableIntersect1`)
+  the application has to get a `RTCTraversable` object first (using
+  `rtcGetSceneTraversable`) and pass it to the SYCL kernel.
 
 - The user geometry callback and filter callback functions should get
   passed through the intersection and occlusion argument structures to
-  the `rtcIntersect1` and `rtcOccluded1` functions directly to allow
+  the `rtcTraversableIntersect1` and `rtcTraversableOccluded1` functions directly to allow
   inlining. The experimental geometry version of the callbacks is
   disabled in SYCL and should not get used.
 

doc/src/api/RTCBufferType.md

@@ -26,6 +26,8 @@
       RTC_BUFFER_TYPE_VERTEX_CREASE_INDEX  = 20,
       RTC_BUFFER_TYPE_VERTEX_CREASE_WEIGHT = 21,
       RTC_BUFFER_TYPE_HOLE                 = 22,
+      
+      RTC_BUFFER_TYPE_TRANSFORM            = 23,
     
       RTC_BUFFER_TYPE_FLAGS = 32
     };
@@ -60,6 +62,10 @@ The `RTC_BUFFER_TYPE_FACE`, `RTC_BUFFER_TYPE_LEVEL`,
 special buffers required to create subdivision meshes (see
 [RTC_GEOMETRY_TYPE_SUBDIVISION]).
 
+The `RTC_BUFFER_TYPE_TRANSFORM` buffer is used to provide instance
+transformation information for instance array geometries (see
+[RTC_GEOMETRY_TYPE_INSTANCE_ARRAY]).
+
 The `RTC_BUFFER_TYPE_FLAGS` can get used to add additional flag per
 primitive of a geometry, and is currently only used for linear curves.
 

doc/src/api/RTCFormat.md

@@ -39,6 +39,8 @@
       RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR,
 
       RTC_FORMAT_GRID,
+
+      RTC_FORMAT_QUATERNION_DECOMPOSITION
     };
 
 #### DESCRIPTION
@@ -58,14 +60,16 @@ of (size 2,3,4, etc.). This format is typcally used to specify to
 format of vertex buffers, e.g. the `RTC_FORMAT_FLOAT3` type for vertex
 buffers of triangle meshes.
 
-The `RTC_FORMAT_FLOAT3X4_ROW_MAJOR` and
-`RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR` formats, specify a 3x4 floating
-point matrix layed out either row major or column major. The
-`RTC_FORMAT_FLOAT4X4_ROW_MAJOR` and `RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR`
-formats, specify a 4x4 floating point matrix layed out either row
-major or column major. These matrix formats are used in the
-[rtcSetGeometryTransform] function in order to set a transformation
-matrix for geometries.
+The `RTC_FORMAT_FLOAT3X4_ROW_MAJOR` and `RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR`
+formats, specify a 3x4 floating point matrix layed out either row major or
+column major. The `RTC_FORMAT_FLOAT4X4_ROW_MAJOR` and
+`RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR` formats, specify a 4x4 floating point matrix
+layed out either row major or column major. The
+`RTC_FORMAT_QUATERNION_DECOMPOSITION` format specifies a structure that
+represents a quaternion decomposition (see [RTCQuaternionDecomposition]) of an affine
+transformation. These formats are used in the [rtcSetGeometryTransform]
+function or in geometry buffers with type `RTC_BUFFER_TYPE_TRANSFORM` in order
+to set a transformation matrix for instance and instance array geometries.
 
 The `RTC_FORMAT_GRID` is a special data format used to specify grid
 primitives of layout RTCGrid when creating grid geometries
@@ -77,3 +81,4 @@ primitives of layout RTCGrid when creating grid geometries
 
 [rtcSetGeometryBuffer], [rtcSetSharedGeometryBuffer],
 [rtcSetNewGeometryBuffer], [rtcSetGeometryTransform]
+[RTCQuaternionDecomposition]

doc/src/api/RTC_GEOMETRY_TYPE_INSTANCE_ARRAY.md

@@ -0,0 +1,79 @@
+% RTC_GEOMETRY_TYPE_INSTANCE_ARRAY(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    RTC_GEOMETRY_TYPE_INSTANCE_ARRAY - instance array geometry type
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    RTCGeometry geometry =
+       rtcNewGeometry(device, RTC_GEOMETRY_TYPE_INSTANCE_ARRAY);
+
+#### DESCRIPTION
+
+Embree supports instance arrays, which is a more memory efficient way to
+represent large amounts of instances of the same or a small set of (sub)scenes.
+The main difference to regular Embree instances is that Embree instance arrays
+have a buffer of transformations (either affine transformations or quaternion
+decompositions [RTCQuaternionDecomposition]) that can be allocated by
+Embree or a shared buffer, similar to vertex buffers for triangle meshes.
+Optionally, instead of instancing only one scene, an instance array can instance
+multiple scenes by passing an array of scenes and a corresponding
+index buffer that specifies which instance of the instance array instances
+which of the scenes in the scenes array.
+
+Instance arrays are created by passing `RTC_GEOMETRY_TYPE_INSTANCE_ARRAY` to
+the `rtcNewGeometry` function call. The instanced scene can be either be set
+using the `rtcSetGeometryInstancedScene` call, or if multiple scenes should be
+instanced by passing an array of scenes using `rtcSetGeometryInstancedScenes`.
+The latter also requires to specify an index buffer using
+`rtcSetNewGeometryBuffer` or `rtcSetSharedGeometryBuffer` with
+`RTC_BUFFER_TYPE_INDEX` as the buffer type.
+
+Because the transformation information can become large for a large amount of
+instances, the instance array allows to share the transformation buffer between
+the user application and Embree. It can be either stored in a buffer created by
+Embree with `rtcSetNewGeometryBuffer` or an already existing buffer can be
+shared using `rtcSetSharedGeometryBuffer`. In either case, the buffer type has
+to be `RTC_BUFFER_TYPE_TRANSFORM` and the allowed formats are
+`RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR`, `RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR`,
+`RTC_FORMAT_FLOAT3X4_ROW_MAJOR`, and `RTC_FORMAT_QUATERNION_DECOMPOSITION`.
+Embree will not modify the data in the transformation buffer.
+
+Embree instance arrays support both single-level instancing and multi-level instancing.
+The maximum instance nesting depth is `RTC_MAX_INSTANCE_LEVEL_COUNT`; it
+can be configured at compile-time using the constant `EMBREE_MAX_INSTANCE_LEVEL_COUNT`.
+Users should adapt this constant to their needs: instances nested any deeper are silently
+ignored in release mode, and cause assertions in debug mode.
+
+Please note that `rtcCommitScene` on the instanced scene(s) should be
+called first, followed by `rtcCommitGeometry` on the instance array,
+followed by `rtcCommitScene` for the top-level scene containing the
+instance array.
+
+If a ray hits the instance, the `geomID` and `primID` members of the hit are
+set to the geometry ID and primitive ID of the hit primitive in the instanced
+scene. The `instID` member of the hit is set to the geometry ID of the instance
+array in the top-level scene and the `instPrimID` member is set to the index of
+the hit instance of the instance array.
+
+For multi-segment motion blur, the number of time steps must be first specified
+using the `rtcSetGeometryTimeStepCount` function. Then a transformation for
+each time step can be specified using the `rtcSetNewGeometryBuffer` or
+`rtcSetSharedGeometryBuffer` function and passing the time step as the `slot`
+parameter of these calls.
+
+See the [Instance Array Geometry] tutorial for an example of how to use instance arrays.
+
+#### EXIT STATUS
+
+On failure `NULL` is returned and an error code is set that can be
+queried using `rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcNewGeometry], [rtcSetGeometryInstancedScene],
+[rtcSetGeometryInstancedScenes], [rtcSetNewGeometryBuffer],
+[rtcSetSharedGeometryBuffer], [rtcGetGeometryTransformEx]

doc/src/api/rtcCommitBuffer.md

@@ -0,0 +1,32 @@
+% rtcCommitBuffer(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcCommitBuffer - commits buffer content from host to device
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcCommitBuffer(RTCBuffer buffer);
+
+#### DESCRIPTION
+
+If the buffer was created using `rtcNewBufferHostDevice` the
+`rtcCommitBuffer` function commits changes of the host buffer data
+to the device. This call is not necessary if the buffer was created
+with a non SYCL Embree device.
+
+The call to `rtcCommitBuffer` will internally use a temporary SYCL
+queue and wait for the memory copy to finish. The function
+`rtcCommitBufferWithQueue` can be used to asyncronously copy the
+data to the device.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcCommitBufferWithQueue] [rtcNewBufferHostDevice]

doc/src/api/rtcCommitBufferWithQueue.md

@@ -0,0 +1,39 @@
+% rtcCommitBufferWithQueue(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcCommitBufferWithQueue - commits buffer content from host to device using a given SYCL queue
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcCommitBufferWithQueue(RTCBuffer buffer,
+        sycl::queue queue, sycl::event* event);
+
+#### DESCRIPTION
+
+If the buffer was created using `rtcNewBufferHostDevice` the 
+`rtcCommitBufferWithQueue` function commits changes of the host buffer data 
+to the device. If the buffer was created with a non SYCL Embree device 
+or the SYCL device has host unified memory, this call has no effect.
+
+The call to `rtcCommitBufferWithQueue` will use the given SYCL queue
+to copy the memory asynchronously. If the SYCL event argument `event`
+is a valid pointer, Embree will use this pointer to return a copy of 
+the SYCL event associated to the memory copy. The parameter `event` 
+is optional and will be ignored if it is a null pointer.
+
+The user is responsible for synchronization using the SYCL queue or 
+the optional SYCL event.
+
+This function is only avaiable on Embree versions with enabled SYCL support.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcCommitBuffer] [rtcNewBufferHostDevice]

doc/src/api/rtcCommitScene.md

@@ -46,6 +46,18 @@ e.g. primitives containing NaNs, INFs, or values greater
 than 1.844E18f (as no reasonable calculations can be performed with
 such values without causing overflows).
 
+In case the RTCDevice associated with the `scene` is a SYCL device,
+`rtcCommitScene` will internally create a temporary SYCL queue to
+issue memory transfers from host to device memory. In this case,
+the call to `rtcCommitScene` will be blocking and return only after
+the memory transfers are completed and SYCL kernels can traverse
+the scene savely on the device. The function `rtcCommitSceneWithQueue`
+can be used to commit all changes for the specified scene asynchronously.
+Handles to traversal objects (`RTCTraversable` type) that have been
+aquired before a call to `rtcCommitScene` will be invalidated. 
+Use `rtcGetSceneTraversable` to get a valid handle from the scene after
+every scene commit.
+
 #### EXIT STATUS
 
 On failure an error code is set that can be queried using
@@ -53,4 +65,4 @@ On failure an error code is set that can be queried using
 
 #### SEE ALSO
 
-[rtcJoinCommitScene]
+[rtcCommitSceneWithQueue], [rtcJoinCommitScene]

doc/src/api/rtcCommitSceneWithQueue.md

@@ -0,0 +1,51 @@
+% rtcCommitSceneWithQueue(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcCommitSceneWithQueue - commits scene changes using a given
+      SYCL queue for host-to-device memory transfers
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcCommitScene(RTCScene scene,
+      sycl::queue queue,
+      sycl::event* event);
+
+#### DESCRIPTION
+
+The `rtcCommitSceneWithQueue` function commits all changes for the specified
+scene (`scene` argument). This internally triggers building of a
+spatial acceleration structure for the scene using all available
+worker threads. Ray queries can be performed only after committing
+all scene changes.
+
+The `rtcCommitSceneWithQueue` function differs from `rtcCommitScene` only
+in that the former uses a SYCL queue passed by the application for asynchronous
+host-to-device memory transfers and the latter uses an internal and temporary
+SYCL queue and is blocking, i.e. the function returns after all memory
+transfers are completed.
+
+The application is responsible for synchronization when using `rtcCommitSceneWithQueue`
+by either calling `wait()` on the `queue` or use the optional `event` argument
+to receive a copy of the sycl::event associated with the last memcpy command
+that was submitted to the queue. The argument `event` can also be a null pointer
+in which case the argument will be ignored internally.
+
+A SYCL kernel calling rtcTraversableIntersect or rtcTraversableOccluded
+using a RTCTraversable associated with the RTCScene `scene` has a dependency on
+the host-to-device memory transfers and therefore has to be synchronized properly.
+
+This function is only available when SYCL support is enabled. In case the
+RTCDevice associated with the `scene` is not a SYCL device, `rtcCommitSceneWithQueue`
+will behave the same as `rtcCommitScene`.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcCommitScene]

doc/src/api/rtcForwardIntersect1.md

@@ -1,8 +1,8 @@
-% rtcForwardIntersect1(3) | Embree Ray Tracing Kernels 4
+% rtcForwardIntersect1/Ex(3) | Embree Ray Tracing Kernels 4
 
 #### NAME
 
-    rtcForwardIntersect1 - forwards a single ray to new scene
+    rtcForwardIntersect1/Ex - forwards a single ray to new scene
       from user geometry callback
 
 #### SYNOPSIS
@@ -16,16 +16,23 @@
       unsigned int instID
     );
 
+    void rtcForwardIntersect1Ex(
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay* ray,
+      unsigned int instID,
+      unsigned int instPrimID,
+    );
+
 #### DESCRIPTION
 
-The `rtcForwardIntersect1` function forwards the traversal of a
-transformed ray (`ray` argument) into a scene (`scene` argument) from
-a user geometry callback. The function can only get invoked from a
-user geometry callback for a ray traversal initiated with the
-`rtcIntersect1` function. The callback arguments structure of the
-callback invokation has to get passed to the ray forwarding (`args`
-argument). The user geometry callback should instantly terminate after
-invoking the `rtcForwardIntersect1` function.
+The `rtcForwardIntersect1` and `rtcForwardIntersect1Ex` functions forward the
+traversal of a transformed ray (`ray` argument) into a scene (`scene` argument)
+from a user geometry callback. The function can only get invoked from a user
+geometry callback for a ray traversal initiated with the `rtcIntersect1`
+function. The callback arguments structure of the callback invokation has to
+get passed to the ray forwarding (`args` argument). The user geometry callback
+should instantly terminate after invoking the `rtcForwardIntersect1/Ex` function.
 
 Only the ray origin and ray direction members of the ray
 argument are used for forwarding, all additional ray properties are
@@ -42,10 +49,15 @@ This function can get used to implement user defined instancing using
 user geometries, e.g. by transforming the ray in a special way, and/or
 selecting between different scenes to instantiate.
 
+For user defined instance arrays, the `rtcForwardIntersect1Ex` variant has an
+additional `instPrimID` argument which is pushed to the instance primitive ID
+stack. Instance primitive IDs identify which instance of an instance array was
+hit.
+
 When using Embree on the CPU it is possible to recursively invoke
 `rtcIntersect1` directly from a user geometry callback. However, when
 SYCL is used, recursively tracing rays is not directly supported, and
-the `rtcForwardIntersect1` function must be used.
+the `rtcForwardIntersect1/Ex` functions must be used.
 
 The ray structure must be aligned to 16 bytes.
 

doc/src/api/rtcForwardIntersect4.md

@@ -1,8 +1,8 @@
-% rtcForwardIntersect4/8/16(3) | Embree Ray Tracing Kernels 4
+% rtcForwardIntersect4/8/16/Ex(3) | Embree Ray Tracing Kernels 4
 
 #### NAME
 
-    rtcForwardIntersect4/8/16 - forwards a ray packet to new scene
+    rtcForwardIntersect4/8/16/Ex - forwards a ray packet to new scene
       from user geometry callback
 
 #### SYNOPSIS
@@ -17,11 +17,11 @@
       unsigned int instID
     );
 
-    void rtcForwardIntersect4(
+    void rtcForwardIntersect8(
       void int* valid,
       const struct RTCIntersectFunctionNArguments* args,
       RTCScene scene,
-      struct RTCRay4* ray,
+      struct RTCRay8* ray,
       unsigned int instID
     );
 
@@ -30,19 +30,47 @@
       const struct RTCIntersectFunctionNArguments* args,
       RTCScene scene,
       struct RTCRay16* ray,
-      unsigned int instID
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcForwardIntersect4Ex(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay4* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcForwardIntersect8Ex(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay8* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcForwardIntersect16Ex(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay16* ray,
+      unsigned int instID,
+      unsigned int instPrimID
     );
 
 #### DESCRIPTION
 
-The `rtcForwardIntersect4/8/16` functions forward the traversal of a
-transformed ray packet (`ray` argument) into a scene (`scene` argument) from
-a user geometry callback. The function can only get invoked from a
-user geometry callback for a ray traversal initiated with the
-`rtcIntersect4/8/16` function. The callback arguments structure of the
-callback invokation has to get passed to the ray forwarding (`args`
-argument). The user geometry callback should instantly terminate after
-invoking the `rtcForwardIntersect4/8/16` function.
+The `rtcForwardIntersect4/8/16` and `rtcForwardIntersect4/8/16Ex` functions
+forward the traversal of a transformed ray packet (`ray` argument) into a scene
+(`scene` argument) from a user geometry callback. The function can only get
+invoked from a user geometry callback for a ray traversal initiated with the
+`rtcIntersect4/8/16` function. The callback arguments structure of the callback
+invokation has to get passed to the ray forwarding (`args` argument). The user
+geometry callback should instantly terminate after invoking the
+`rtcForwardIntersect4/8/16/Ex` function.
 
 Only the ray origin and ray direction members of the ray
 argument are used for forwarding, all additional ray properties are
@@ -59,6 +87,11 @@ This function can get used to implement user defined instancing using
 user geometries, e.g. by transforming the ray in a special way, and/or
 selecting between different scenes to instantiate.
 
+For user defined instance arrays, the `rtcForwardIntersect4/8/16Ex` variant has an
+additional `instPrimID` argument which is pushed to the instance primitive ID
+stack. Instance primitive IDs identify which instance of an instance array was
+hit.
+
 When using Embree on the CPU it is possible to recursively invoke
 `rtcIntersect4/8/16` directly from a user geometry callback. However, when
 SYCL is used, recursively tracing rays is not directly supported, and

doc/src/api/rtcForwardOccluded1.md

@@ -1,8 +1,8 @@
-% rtcForwardOccluded1(3) | Embree Ray Tracing Kernels 4
+% rtcForwardOccluded1/Ex(3) | Embree Ray Tracing Kernels 4
 
 #### NAME
 
-    rtcForwardOccluded1 - forwards a single ray to new scene
+    rtcForwardOccluded1/Ex - forwards a single ray to new scene
       from user geometry callback
 
 #### SYNOPSIS
@@ -16,16 +16,23 @@
       unsigned int instID
     );
 
+    void rtcForwardOccluded1Ex(
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
 #### DESCRIPTION
 
-The `rtcForwardOccluded1` function forwards the traversal of a
-transformed ray (`ray` argument) into a scene (`scene` argument) from
-a user geometry callback. The function can only get invoked from a
-user geometry callback for a ray traversal initiated with the
-`rtcOccluded1` function. The callback arguments structure of the
-callback invokation has to get passed to the ray forwarding (`args`
-argument). The user geometry callback should instantly terminate after
-invoking the `rtcForwardOccluded1` function.
+The `rtcForwardOccluded1` and `rtcForwardOccluded1Ex` functions forward the
+traversal of a transformed ray (`ray` argument) into a scene (`scene` argument)
+from a user geometry callback. The function can only get invoked from a user
+geometry callback for a ray traversal initiated with the `rtcOccluded1`
+function. The callback arguments structure of the callback invokation has to
+get passed to the ray forwarding (`args` argument). The user geometry callback
+should instantly terminate after invoking the `rtcForwardOccluded1/Ex` function.
 
 Only the ray origin and ray direction members of the ray
 argument are used for forwarding, all additional ray properties are
@@ -42,10 +49,15 @@ This function can get used to implement user defined instancing using
 user geometries, e.g. by transforming the ray in a special way, and/or
 selecting between different scenes to instantiate.
 
+For user defined instance arrays, the `rtcForwardIntersect1Ex` variant has an
+additional `instPrimID` argument which is pushed to the instance primitive ID
+stack. Instance primitive IDs identify which instance of an instance array was
+hit.
+
 When using Embree on the CPU it is possible to recursively invoke
 `rtcOccluded1` directly from a user geometry callback. However, when
 SYCL is used, recursively tracing rays is not directly supported, and
-the `rtcForwardOccluded1` function must be used.
+the `rtcForwardOccluded1/Ex` function must be used.
 
 The ray structure must be aligned to 16 bytes.
 

doc/src/api/rtcForwardOccluded4.md

@@ -1,8 +1,8 @@
-% rtcForwardOccluded4/8/16(3) | Embree Ray Tracing Kernels 4
+% rtcForwardOccluded4/8/16/Ex(3) | Embree Ray Tracing Kernels 4
 
 #### NAME
 
-    rtcForwardOccluded4/8/16 - forwards a ray packet to new scene
+    rtcForwardOccluded4/8/16/Ex - forwards a ray packet to new scene
       from user geometry callback
 
 #### SYNOPSIS
@@ -17,11 +17,11 @@
       unsigned int instID
     );
 
-    void rtcForwardOccluded4(
+    void rtcForwardOccluded8(
       void int* valid,
       const struct RTCOccludedFunctionNArguments* args,
       RTCScene scene,
-      struct RTCRay4* ray,
+      struct RTCRay8* ray,
       unsigned int instID
     );
 
@@ -33,16 +33,43 @@
       unsigned int instID
     );
 
+    void rtcForwardOccluded4Ex(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay4* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcForwardOccluded8Ex(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay8* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcForwardOccluded16Ex(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCScene scene,
+      struct RTCRay16* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
 #### DESCRIPTION
 
-The `rtcForwardOccluded4/8/16` functions forward the traversal of a
-transformed ray packet (`ray` argument) into a scene (`scene` argument) from
-a user geometry callback. The function can only get invoked from a
-user geometry callback for a ray traversal initiated with the
-`rtcOccluded4/8/16` function. The callback arguments structure of the
-callback invokation has to get passed to the ray forwarding (`args`
-argument). The user geometry callback should instantly terminate after
-invoking the `rtcForwardOccluded4/8/16` function.
+The `rtcForwardOccluded4/8/16` and `rtcForwardOccluded4/8/16Ex` functions
+forward the traversal of a transformed ray packet (`ray` argument) into a scene
+(`scene` argument) from a user geometry callback. The function can only get
+invoked from a user geometry callback for a ray traversal initiated with the
+`rtcOccluded4/8/16` function. The callback arguments structure of the callback
+invokation has to get passed to the ray forwarding (`args` argument). The user
+geometry callback should instantly terminate after invoking the
+`rtcForwardOccluded4/8/16/Ex` function.
 
 Only the ray origin and ray direction members of the ray
 argument are used for forwarding, all additional ray properties are
@@ -59,6 +86,11 @@ This function can get used to implement user defined instancing using
 user geometries, e.g. by transforming the ray in a special way, and/or
 selecting between different scenes to instantiate.
 
+For user defined instance arrays, the `rtcForwardIntersect4/8/16Ex` variant has
+an additional `instPrimID` argument which is pushed to the instance primitive
+ID stack. Instance primitive IDs identify which instance of an instance array
+was hit.
+
 When using Embree on the CPU it is possible to recursively invoke
 `rtcOccluded4/8/16` directly from a user geometry callback. However, when
 SYCL is used, recursively tracing rays is not directly supported, and

doc/src/api/rtcGetBufferData.md

@@ -13,7 +13,10 @@
 #### DESCRIPTION
 
 The `rtcGetBufferData` function returns a pointer to the buffer data
-of the specified buffer object (`buffer` argument).
+of the specified buffer object (`buffer` argument). If the buffer
+was created using `rtcNewBufferHostDevice` and the SYCL device has no
+host unified memory, this pointer is only valid on the host. To get
+a device pointer in this case, use `rtcGetBufferDataDevice`.
 
 #### EXIT STATUS
 
@@ -22,4 +25,4 @@ On failure an error code is set that can be queried using
 
 #### SEE ALSO
 
-[rtcNewBuffer]
+[rtcGetBufferDataDevice], [rtcNewBuffer], [rtcNewBufferHostDevice]

doc/src/api/rtcGetBufferDataDevice.md

@@ -0,0 +1,28 @@
+% rtcGetBufferData(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetBufferDataDevice - gets a device pointer to the buffer data
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void* rtcGetBufferDataDevice(RTCBuffer buffer);
+
+#### DESCRIPTION
+
+The `rtcGetBufferDataDevice` function returns a pointer to the buffer data
+of the specified buffer object (`buffer` argument) which can be used for
+accessing the data on the device. If Embree has no SYCL support or the SYCL
+device has host unified memory, the pointer is equal to the pointer returned
+by `rtcGetBufferData`.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcGetBufferData], [rtcNewBuffer]

doc/src/api/rtcGetDeviceError.md

@@ -39,16 +39,29 @@ Possible error codes returned by `rtcGetDeviceError` are:
 + `RTC_ERROR_CANCELLED`: The operation got canceled by a memory
   monitor callback or progress monitor callback function.
 
++ `RTC_ERROR_LEVEL_ZERO_RAYTRACING_SUPPORT_MISSING`: This error can occur when
+  creating an Embree device with SYCL support using `rtcNewSYCLDevice` fails.
+  This error probably means that the GPU driver is to old or not installed
+  properly. Install a new GPU driver and on Linux make sure that the package
+  `intel-level-zero-gpu-raytracing` is installed. For general driver installation
+  information for Linux refer to [https://dgpu-docs.intel.com](https://dgpu-docs.intel.com).
+
 When the device construction fails, `rtcNewDevice` returns `NULL` as
 device. To detect the error code of a such a failed device
 construction, pass `NULL` as device to the `rtcGetDeviceError`
 function. For all other invocations of `rtcGetDeviceError`, a proper
 device pointer must be specified.
 
+The API function `rtcGetDeviceLastErrorMessage` can be used to get more details
+about the last `RTCError` a `RTCDevice` encountered.
+
+For convenient reporting of a `RTCError`, the API function `rtcGetErrorString`
+can be used, which returns a string representation of a given `RTCError`.
+
 #### EXIT STATUS
 
 Returns the error code for the device.
 
 #### SEE ALSO
 
-[rtcSetDeviceErrorFunction]
+[rtcSetDeviceErrorFunction], [rtcGetDeviceLastErrorMessage], [rtcGetErrorString]

doc/src/api/rtcGetDeviceLastErrorMessage.md

@@ -0,0 +1,33 @@
+% rtcGetDeviceLastErrorMessage(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetDeviceLastErrorMessage - returns a message corresponding
+      to the last error code
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    const char* rtcGetDeviceLastErrorMessage(RTCDevice device);
+
+#### DESCRIPTION
+
+This function can be used to get a message corresponding to the last error code
+(returned by `rtcGetDeviceError`) which often provides details about the error
+that happened. The message is the same as the message that will written to
+console when verbosity is > 0 or which is passed as the `str` argument of the
+`RTCErrorFunction` (see [rtcSetDeviceErrorFunction]). However, when device
+construction fails this is the only way to get additional information about
+the error. In this case, `rtcNewDevice` returns `NULL` as device. To query the
+error message for such a failed device construction, pass `NULL` as device to the
+`rtcGetDeviceLastErrorMessage` function. For all other invocations of
+`rtcGetDeviceLastErrorMessage`, a proper device pointer must be specified.
+
+#### EXIT STATUS
+
+Returns a message corresponding to the last error code.
+
+#### SEE ALSO
+
+[rtcGetDeviceError], [rtcSetDeviceErrorFunction]

doc/src/api/rtcGetErrorString.md

@@ -0,0 +1,31 @@
+% rtcGetErrorString(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetErrorString - returns a string representation
+      of a given RTCError
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    const char* rtcGetErrorString(RTCError code);
+
+#### DESCRIPTION
+
+Returns a string representation for a `RTCError` error code. For example, for the
+`RTCError` RTC_ERROR_UNKNOWN this function will return the string "Unknown Error".
+This is purely a convenience function for printing error information on the user
+side.
+
+The returned strings should not be used for comparing different `RTCError` error
+codes or make other decisions based on the type of error that occurred. For such
+things only the `RTCError` enum values should be used.
+
+#### EXIT STATUS
+
+Returns a string representation of a given `RTCError` error code.
+
+#### SEE ALSO
+
+[rtcGetDeviceError]

doc/src/api/rtcGetGeometryBufferData.md

@@ -22,6 +22,10 @@ element of the buffer view attached to the specified buffer type and
 slot (`type` and `slot` argument) of the geometry (`geometry`
 argument).
 
+If the device associated with `geometry` is a SYCL device,
+the returned pointer only accessible on the host. The function
+`rtcGetGeometryBufferDataDevice` can be used to get a pointer
+accessible on the SYCL device.
 #### EXIT STATUS
 
 On failure an error code is set that can be queried using
@@ -29,4 +33,5 @@ On failure an error code is set that can be queried using
 
 #### SEE ALSO
 
-[rtcSetGeometryBuffer], [rtcSetSharedGeometryBuffer], [rtcSetNewGeometryBuffer]
+[rtcGetGeometryBufferDataDevice], [rtcSetGeometryBuffer],
+[rtcSetSharedGeometryBuffer], [rtcSetNewGeometryBuffer]

doc/src/api/rtcGetGeometryBufferDataDevice.md

@@ -0,0 +1,36 @@
+% rtcGetGeometryBufferDataDevice(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetGeometryBufferDataDevice - gets pointer to
+      the first buffer view element on the device.
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void* rtcGetGeometryBufferDataDevice(
+      RTCGeometry geometry,
+      enum RTCBufferType type,
+      unsigned int slot
+    );
+
+#### DESCRIPTION
+
+The `rtcGetGeometryBufferDataDevice` function returns a pointer to the first
+element of the buffer view attached to the specified buffer type and
+slot (`type` and `slot` argument) of the geometry (`geometry`
+argument) on the device.
+
+If the device associated with `geometry` is no SYCL device, or if Embree
+is executed on a system with host unified memory (e.g., on an iGPU),
+the returned pointer is the same as the one returned by `rtcGetGeometryBufferData`.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcGetGeometryBufferData], [rtcSetGeometryBuffer], [rtcSetSharedGeometryBuffer], [rtcSetNewGeometryBuffer]

doc/src/api/rtcGetGeometryTransformEx.md

@@ -0,0 +1,49 @@
+% rtcGetGeometryTransformEx(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetGeometryTransformEx - returns the interpolated instance
+    transformation for an instance of an instance array geometry for the
+    specified time.
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcGetGeometryTransformEx(
+      RTCGeometry geometry,
+      unsigned int instPrimID,
+      float time,
+      enum RTCFormat format,
+      void* xfm
+    );
+
+#### DESCRIPTION
+
+The `rtcGetGeometryTransformEx` function returns the interpolated local
+to world transformation (`xfm` parameter) of the `instPrimID`-th instance of an
+instance array geometry (`geometry` parameter) for a particular time (`time`
+parameter in range $[0,1]$) in the specified format (`format` parameter).
+The function can also be used when `geometry` refers to a regular instance, but
+then the `instPrimID` has to be $0$.
+
+Possible formats for the returned matrix are:
+
++ `RTC_FORMAT_FLOAT3X4_ROW_MAJOR`: The 3×4 float matrix is laid out
+  in row-major form.
+
++ `RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR`: The 3×4 float matrix is laid out
+  in column-major form.
+
++ `RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR`: The 3×4 float matrix is laid out
+  in column-major form as a 4×4 homogeneous matrix with last row equal
+  to (0, 0, 0, 1).
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[RTC_GEOMETRY_TYPE_INSTANCE_ARRAY]

doc/src/api/rtcGetGeometryTransformFromTraversable.md

@@ -0,0 +1,35 @@
+% rtcGetGeometryTransformFromTraversable(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetGeometryTransformFromTraversable - returns the interpolated instance
+      transformation for the specified time
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcGetGeometryTransformFromTraversable(
+      RTCTraversable traversable,
+      unsigned int geomID,
+      float time,
+      enum RTCFormat format,
+      void* xfm
+    );
+
+#### DESCRIPTION
+
+The `rtcGetGeometryTransformFromTraversable` function is equivalent to
+`rtcGetGeometryTransformFromScene` but takes traversable object (`traversable` argument)
+instead of a scene object. Using this method is optional on CPU but it is required for SYCL.
+
+For more details, refer to the documentation of `rtcGetGeometryTransformFromScene`.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcGetGeometryTransformFromScene], [rtcGetSceneTraversable]

doc/src/api/rtcGetGeometryUserDataFromTraversable.md

@@ -0,0 +1,30 @@
+% rtcGetGeometryUserDataFromTraversable(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetGeometryUserDataFromTraversable - returns the user data pointer
+      of the geometry through the scene object
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void* rtcGetGeometryUserDataFromTraversable(
+        RTCTraversable traversable,
+        unsigned int geomID);
+
+#### DESCRIPTION
+
+The `rtcGetGeometryUserDataFromTraversable` function is equivalent to
+`rtcGetGeometryUserDataFromScene` but takes traversable object
+(`traversable` argument) instead of a scene object.
+Using this method is optional on CPU but it is required for SYCL.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcGetGeometryUserDataFromScene], [rtcGetSceneTraversable]

doc/src/api/rtcGetSceneTraversable.md

@@ -0,0 +1,31 @@
+% rtcGetSceneTraversable(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcGetSceneTraversable - gets a read-only hand to a RTCScene object
+      that is also valid on a SYCL device and can be used for ray queries.
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    RTCTraversable rtcGetSceneTraversable(RTCScene scene);
+
+#### DESCRIPTION
+
+The `rtcGetSceneTraversable` function returns a read-only handle
+to a RTCScene object (`scene` argument). Traversable objects are
+used for ray queries on a SYCL device.
+
+Traversable objects can also be used in host/CPU code and Embree provides
+other API calls such as the `rtcTraversablePointQuery` to help
+write portable code compatible with CPU and SYCL device execution.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcTraversableIntersect1]

doc/src/api/rtcNewBuffer.md

@@ -26,6 +26,13 @@ and deallocated when the buffer is destroyed.
 ``` {include=src/api/inc/buffer_padding.md}
 ```
 
+If the `device` is a Embree SYCL device, the buffer will be allocated
+using SYCL USM shared memory, i.e. the buffer can be accessed on the host
+and device (GPU) and the SYCL runtime will handle buffer transfers automatically.
+
+For precise control over when memory is copied from host to device, 
+a buffer can also be created using `rtcNewBufferHostDevice`.
+
 #### EXIT STATUS
 
 On failure `NULL` is returned and an error code is set that can be
@@ -33,4 +40,4 @@ queried using `rtcGetDeviceError`.
 
 #### SEE ALSO
 
-[rtcRetainBuffer], [rtcReleaseBuffer]
+[rtcNewBufferHostDevice], [rtcRetainBuffer], [rtcReleaseBuffer]

doc/src/api/rtcNewBufferHostDevice.md

@@ -0,0 +1,43 @@
+% rtcNewBuffer(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcNewBufferHostDevice - creates a new data buffer with 
+      explicitly managed host and device memory allocations
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    RTCBuffer rtcNewBufferHostDevice(
+      RTCDevice device,
+      size_t byteSize
+    );
+
+#### DESCRIPTION
+
+The `rtcNewBufferHostDevice` function creates a new data buffer object of
+specified size in bytes (`byteSize` argument) that is bound to the
+specified device (`device` argument). The buffer object is reference
+counted with an initial reference count of 1. The returned buffer
+object can be released using the `rtcReleaseBuffer` API call. If Embree has SYCL
+support enabled and the SYCL device has no host unifed memory (e.g, a discrete GPU), 
+the buffer allocates memory on the host and device explicitly.
+After the buffer is modified on the host `rtcCommitBuffer` can be used to synchronize
+host and device memory by copying the buffer content from the host to device. If the Embree
+version has no SYCL support or the SYCL device has host unified memory, the buffer will behave 
+the same as a buffer created using `rtcNewBuffer`. The
+specified number of bytes are allocated at buffer construction time
+and deallocated when the buffer is destroyed.
+
+``` {include=src/api/inc/buffer_padding.md}
+```
+
+#### EXIT STATUS
+
+On failure `NULL` is returned and an error code is set that can be
+queried using `rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcCommitBuffer], [rtcNewBuffer], [rtcRetainBuffer], [rtcReleaseBuffer]

doc/src/api/rtcNewGeometry.md

@@ -33,7 +33,8 @@
      RTC_GEOMETRY_TYPE_DISC_POINT,
      RTC_GEOMETRY_TYPE_ORIENTED_DISC_POINT,
      RTC_GEOMETRY_TYPE_USER,
-     RTC_GEOMETRY_TYPE_INSTANCE
+     RTC_GEOMETRY_TYPE_INSTANCE,
+     RTC_GEOMETRY_TYPE_INSTANCE_ARRAY,
     };
 
     RTCGeometry rtcNewGeometry(
@@ -64,8 +65,9 @@ bases (`RTC_GEOMETRY_TYPE_FLAT_LINEAR_CURVE`, `RTC_GEOMETRY_TYPE_FLAT_BEZIER_CUR
 grid meshes (`RTC_GEOMETRY_TYPE_GRID`), point geometries
 (`RTC_GEOMETRY_TYPE_SPHERE_POINT`, `RTC_GEOMETRY_TYPE_DISC_POINT`,
 `RTC_TYPE_ORIENTED_DISC_POINT`),
-user-defined geometries (`RTC_GEOMETRY_TYPE_USER`), and instances
-(`RTC_GEOMETRY_TYPE_INSTANCE`).
+user-defined geometries (`RTC_GEOMETRY_TYPE_USER`), instances
+(`RTC_GEOMETRY_TYPE_INSTANCE`), and instance arrays
+(`RTC_GEOMETRY_TYPE_INSTANCE_ARRAY`).
 
 The types `RTC_GEOMETRY_TYPE_ROUND_BEZIER_CURVE`,
 `RTC_GEOMETRY_TYPE_ROUND_BSPLINE_CURVE`, and
@@ -123,4 +125,4 @@ queried using `rtcGetDeviceError`.
 [RTC_GEOMETRY_TYPE_TRIANGLE], [RTC_GEOMETRY_TYPE_QUAD],
 [RTC_GEOMETRY_TYPE_SUBDIVISION], [RTC_GEOMETRY_TYPE_CURVE],
 [RTC_GEOMETRY_TYPE_GRID], [RTC_GEOMETRY_TYPE_POINT],
-[RTC_GEOMETRY_TYPE_USER], [RTC_GEOMETRY_TYPE_INSTANCE]
+[RTC_GEOMETRY_TYPE_USER], [RTC_GEOMETRY_TYPE_INSTANCE], [RTC_GEOMETRY_TYPE_INSTANCE_ARRAY]

doc/src/api/rtcNewSharedBuffer.md

@@ -42,4 +42,4 @@ queried using `rtcGetDeviceError`.
 
 #### SEE ALSO
 
-[rtcRetainBuffer], [rtcReleaseBuffer]
+[rtcNewSharedBufferHostDevice], [rtcRetainBuffer], [rtcReleaseBuffer]

doc/src/api/rtcNewSharedBufferHostDevice.md

@@ -0,0 +1,50 @@
+% rtcNewSharedBufferHostDevice(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcNewSharedBufferHostDevice - creates a new shared data buffer
+      with a managed copy of the buffer in device memory
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    RTCBuffer rtcNewSharedBufferHostDevice(
+      RTCDevice device,
+      void* ptr,
+      size_t byteSize
+    );
+
+#### DESCRIPTION
+
+The `rtcNewSharedBufferHostDevice` function creates a new shared data buffer
+object bound to the specified device (`device` argument). The buffer
+object is reference counted with an initial reference count of 1. The
+buffer can be released using the `rtcReleaseBuffer` function.
+
+At construction time, the pointer to the user-managed buffer data
+(`ptr` argument) including its size in bytes (`byteSize` argument) is
+provided to create the buffer. At buffer construction time no buffer
+data is allocated on the host, but the buffer data provided by the application is
+used. A equal sized memory allocation is created on the device. 
+The buffer data must remain valid for as long as the buffer may
+be used, and the user is responsible to free the buffer data when no
+longer required.
+
+``` {include=src/api/inc/buffer_padding.md}
+```
+
+The data pointer (`ptr` argument) must be aligned to 4 bytes; otherwise
+the `rtcNewSharedBufferHostDevice` function will fail.
+
+The function `rtcCommitBuffer` can be used to synchronize the buffer
+content with the device.
+
+#### EXIT STATUS
+
+On failure `NULL` is returned and an error code is set that can be
+queried using `rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcNewSharedBuffer], [rtcCommitBuffer], [rtcRetainBuffer], [rtcReleaseBuffer]

doc/src/api/rtcSetGeometryInstancedScenes.md

@@ -0,0 +1,37 @@
+% rtcSetGeometryInstancedScenes(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcSetGeometryInstancedScenes - sets an array of scenes that can be
+    instanced by an instance array geometry
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcSetGeometryInstancedScenes(
+      RTCGeometry geometry,
+      RTCScene* scene,
+      size_t numScenes
+    );
+
+#### DESCRIPTION
+
+The `rtcSetGeometryInstancedScenes` function sets an array of type `RTCScene`
+with `numScenes` elements that the specified instance geometry (`geometry`
+argument) can instance. This call also requires setting an index buffer using
+either `rtcSetSharedGeometryBuffer` or `rtcSetNewGeometryBuffer` (similar to
+index buffers for triangle meshes), that specifies which instance of the
+instance array instances which scene in the scene array. If only one scene
+should be instanced the call `rtcSetGeometryInstancedScene` should be
+preferred.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[RTC_GEOMETRY_TYPE_INSTANCE_ARRAY], [rtcSetNewGeometryBuffer],
+[rtcSetSharedGeometryBuffer], [rtcSetGeometryInstancedScene]

doc/src/api/rtcSetNewGeometryBufferHostDevice.md

@@ -0,0 +1,62 @@
+% rtcSetNewGeometryBufferHostDevice(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcSetNewGeometryBufferHostDevice - creates and assigns a new host/device
+      data buffer pair to the geometry
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcSetNewGeometryBufferHostDevice(
+      RTCGeometry geometry,
+      enum RTCBufferType type,
+      unsigned int slot,
+      enum RTCFormat format,
+      size_t byteStride,
+      size_t itemCount,
+      void** ptr,
+      void** dptr
+    );
+
+#### DESCRIPTION
+
+The `rtcSetNewGeometryBufferHostDevice` function creates a new host/device data buffer pair of
+specified format (`format` argument), byte stride (`byteStride`
+argument), and number of items (`itemCount` argument), and assigns it
+to a geometry buffer slot (`type` and `slot` argument) of the
+specified geometry (`geometry` argument). The buffer data is managed
+internally and automatically freed when the geometry is destroyed.
+
+The byte stride (`byteStride` argument) must be aligned to 4 bytes;
+otherwise the `rtcSetNewGeometryBufferHostDevice` function will fail.
+
+The allocated buffer will be automatically over-allocated slightly
+when used as a vertex buffer, where a requirement is that each buffer
+element should be readable using 16-byte SSE load instructions.
+
+If `ptr` is not null, it will be used to return the address of the
+host data buffer. If `dptr` is not null, it will be used to return
+the address of the device data buffer. Either `ptr` or `dptr` or both
+can be null. In this case `rtcGetGeometryBufferData` and
+`rtcGetGeometryBufferDataDevice` can be used to get the addresses of the
+host and device data buffers.
+
+The application is responsible of keeping the host and device memory
+in sync. The host memory has to be updated before calls of
+`rtcCommitScene` involving the associated geometry.
+
+If Embree has no SYCL support or the associated Embree device is no
+SYCL device the `dptr` argument will return the same address as `ptr` if on null.
+In this case the function `rtcSetNewGeometryBufferHostDevice` will behave like
+`rtcSetSharedGeometryBuffer`.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+[rtcSetNewGeometryBuffer], [rtcSetGeometryBuffer], [rtcSetSharedGeometryBuffer],
+[rtcGetGeometryBufferData], [rtcGetGeometryBufferDataDevice]

doc/src/api/rtcSetSharedGeometryBuffer.md

@@ -43,6 +43,11 @@ The buffer data must remain valid for as long as the buffer may be
 used, and the user is responsible for freeing the buffer data when no
 longer required.
 
+If using a SYCL Embree device, `ptr` must be allocated using SYCL USM shared
+memory to be accessible on the device. Alternatively the function
+`rtcSetSharedGeometryBufferHostDevice` can be used to pass an explicitly
+managed host/device buffer pair.
+
 Sharing buffers can significantly reduce the memory required by the
 application, thus we recommend using this feature. When enabling the
 `RTC_SCENE_FLAG_COMPACT` scene flag, the spatial index structures index
@@ -55,4 +60,4 @@ On failure an error code is set that can be queried using
 
 #### SEE ALSO
 
-[rtcSetGeometryBuffer], [rtcSetNewGeometryBuffer]
+[rtcSetSharedGeometryBufferHostDevice], [rtcSetGeometryBuffer], [rtcSetNewGeometryBuffer]

doc/src/api/rtcSetSharedGeometryBufferHostDevice.md

@@ -0,0 +1,70 @@
+% rtcSetSharedGeometryBufferHostDevice(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcSetSharedGeometryBufferHostDevice - assigns views of shared data buffers
+      to a geometry using explicit host and device memory
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcSetSharedGeometryBuffer(
+      RTCGeometry geometry,
+      enum RTCBufferType type,
+      unsigned int slot,
+      enum RTCFormat format,
+      const void* ptr,
+      const void* dptr,
+      size_t byteOffset,
+      size_t byteStride,
+      size_t itemCount
+    );
+
+#### DESCRIPTION
+
+The `rtcSetSharedGeometryBufferHostDevice` function binds views of a
+shared user-managed data buffers to a geometry buffer
+type and slot (`type` and `slot` argument) of the specified geometry
+(`geometry` argument). The user-managed data buffers are passed using the
+`ptr` argument for the host side allocated memory and the `dptr` parameter
+for the memory allocated on the device.
+
+One can specify the start of the first buffer element in bytes
+(`byteOffset` argument), the byte stride between individual buffer
+elements (`byteStride` argument), the format of the buffer elements
+(`format` argument), and the number of elements to bind (`itemCount`).
+
+The start address (`byteOffset` argument) and stride (`byteStride`
+argument) must be both aligned to 4 bytes; otherwise the
+`rtcSetSharedGeometryBufferHostDevice` function will fail.
+
+``` {include=src/api/inc/buffer_padding.md}
+```
+
+The buffer data must remain valid for as long as the buffer may be
+used, and the user is responsible for freeing the buffer data when no
+longer required.
+
+The application is responsible of keeping the host and device memory
+in sync. The host memory has to be updated before calls of
+`rtcCommitScene` involving the associated geometry.
+
+If Embree has no SYCL support or the associated Embree device is no
+SYCL device the `dptr` argument must be a null pointer. In this case
+the function `rtcSetSharedGeometryBufferHostDevice` will behave like
+`rtcSetSharedGeometryBuffer`.
+
+Sharing buffers can significantly reduce the memory required by the
+application, thus we recommend using this feature. When enabling the
+`RTC_SCENE_FLAG_COMPACT` scene flag, the spatial index structures index
+into the vertex buffer, resulting in even higher memory savings.
+
+#### EXIT STATUS
+
+On failure an error code is set that can be queried using
+`rtcGetDeviceError`.
+
+#### SEE ALSO
+
+[rtcSetSharedGeometryBuffer], [rtcSetGeometryBuffer], [rtcSetNewGeometryBuffer]

doc/src/api/rtcTraversableForwardIntersect1.md

@@ -0,0 +1,43 @@
+% rtcTraversableForwardIntersect1/Ex(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTraversableForwardIntersect1/Ex - forwards a single ray to new scene
+      from user geometry callback
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableForwardIntersect1(
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay* ray,
+      unsigned int instID
+    );
+
+    void rtcTraversableForwardIntersect1Ex(
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay* ray,
+      unsigned int instID,
+      unsigned int instPrimID,
+    );
+
+#### DESCRIPTION
+
+The `rtcTraversalbeForwardIntersect1` and `rtcTraversableForwardIntersect1Ex` functions are
+equivalent to `rtcForwardIntersect1` and `rtcForwardIntersect1Ex` respectively
+but take a traversable object (`traversable` argument) instead of a scene
+object.
+
+For more details, refer to the documentation of `rtcForwardIntersect1/Ex`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcForwardIntersect1], [rtcGetSceneTraversable]

doc/src/api/rtcTraversableForwardIntersect4.md

@@ -0,0 +1,81 @@
+% rtcTraversableForwardIntersect4/8/16/Ex(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTraversableForwardIntersect4/8/16/Ex - forwards a ray packet to new scene
+      from user geometry callback
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableForwardIntersect4(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay4* ray,
+      unsigned int instID
+    );
+
+    void rtcTraversableForwardIntersect8(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay8* ray,
+      unsigned int instID
+    );
+
+    void rtcTraversableForwardIntersect16(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay16* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcTraversableForwardIntersect4Ex(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay4* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcTraversableForwardIntersect8Ex(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay8* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcTraversableForwardIntersect16Ex(
+      void int* valid,
+      const struct RTCIntersectFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay16* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+#### DESCRIPTION
+
+The `rtcTraversableForwardIntersect4/8/16/Ex` functions
+are equivalent to `rtcForwardIntersect4/8/16/Ex`
+but take a traversable object (`traversable` argument)
+instead of a scene object.
+
+For more details, refer to the documentation of `rtcForwardIntersect4/8/16/Ex`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcForwardIntersect4/8/16/Ex], [rtcGetSceneTraversable]
+

doc/src/api/rtcTraversableForwardOccluded1.md

@@ -0,0 +1,43 @@
+% rtcTraversableForwardOccluded1/Ex(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTraversableForwardOccluded1/Ex - forwards a single ray to new scene
+      from user geometry callback
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableForwardOccluded1(
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay* ray,
+      unsigned int instID
+    );
+
+    void rtcTraversableForwardOccluded1Ex(
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay* ray,
+      unsigned int instID,
+      unsigned int instPrimID,
+    );
+
+#### DESCRIPTION
+
+The `rtcTraversalbeForwardOccluded1` and `rtcTraversableForwardOccluded1Ex` functions are
+equivalent to `rtcForwardOccluded1` and `rtcForwardOccluded1Ex` respectively
+but take a traversable object (`traversable` argument) instead of a scene
+object.
+
+For more details, refer to the documentation of `rtcForwardOccluded1/Ex`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcForwardOccluded1], [rtcGetSceneTraversable]

doc/src/api/rtcTraversableForwardOccluded4.md

@@ -0,0 +1,79 @@
+% rtcTraversableForwardOccluded4/8/16/Ex(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTraversableForwardOccluded4/8/16/Ex - forwards a ray packet to new scene
+      from user geometry callback
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableForwardOccluded4(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay4* ray,
+      unsigned int instID
+    );
+
+    void rtcTraversableForwardOccluded8(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay8* ray,
+      unsigned int instID
+    );
+
+    void rtcTraversableForwardOccluded16(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay16* ray,
+      unsigned int instID
+    );
+
+    void rtcTraversableForwardOccluded4Ex(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay4* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcTraversableForwardOccluded8Ex(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay8* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+    void rtcTraversableForwardOccluded16Ex(
+      void int* valid,
+      const struct RTCOccludedFunctionNArguments* args,
+      RTCTraversable traversable,
+      struct RTCRay16* ray,
+      unsigned int instID,
+      unsigned int instPrimID
+    );
+
+#### DESCRIPTION
+
+The `rtcTraversableForwardOccluded4/8/16/Ex` functions
+are equivalent to `rtcForwardOccluded4/8/16/Ex`
+but take a traversable object (`traversable` argument)
+instead of a scene object.
+
+For more details, refer to the documentation of `rtcForwardOccluded4/8/16/Ex`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcForwardOccluded4/8/16/Ex], [rtcGetSceneTraversable]

doc/src/api/rtcTraversableIntersect1.md

@@ -0,0 +1,32 @@
+% rtcTraversableIntersect1(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTraversableIntersect1 - finds the closest hit for a single ray
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableIntersect1(
+      RTCTraversable traversable,
+      struct RTCRayHit* rayhit
+      struct RTCIntersectArguments* args = NULL
+    );
+
+#### DESCRIPTION
+
+The `rtcTraversableIntersect` function is equivalent to `rtcIntersect1`
+but takes traversable object (`traversable` argument) instead of a scene
+object. Using this method is optional on CPU but it is required for SYCL.
+
+For more details, refer to the documentation of `rtcIntersect1`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcIntersect1], [rtcGetSceneTraversable]

doc/src/api/rtcTraversableIntersect4.md

@@ -0,0 +1,48 @@
+% rtcTraversableIntersect4/8/16(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTravesableIntersect4/8/16 - finds the closest hits for a ray packet
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableIntersect4(
+      const int* valid,
+      RTCTraversable traversable,
+      struct RTCRayHit4* rayhit,
+      struct RTCIntersectArguments* args = NULL
+    );
+
+    void rtcTraversableIntersect8(
+      const int* valid,
+      RTCTraversable traversable,
+      struct RTCRayHit8* rayhit,
+      struct RTCIntersectArguments* args = NULL
+    );
+
+    void rtcTraversableIntersect16(
+      const int* valid,
+      RTCTraversable traversable,
+      struct RTCRayHit16* rayhit,
+      struct RTCIntersectArguments* args = NULL
+    );
+
+#### DESCRIPTION
+
+
+The `rtcTraversableIntersect4/8/16` functions are equivalent to `rtcIntersect4/8/16`
+but take a traversable object (`traversable` argument) instead of a scene
+object.
+
+For more details, refer to the documentation of `rtcIntersect4/8/16`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcIntersect4/8/16], [rtcGetSceneTraversable]

doc/src/api/rtcTraversableOccluded1.md

@@ -0,0 +1,32 @@
+% rtcTraversableOccluded1(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTraversableOccluded1 - finds any hit for a single ray
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableOccluded1(
+      RTCTraversable traversable,
+      struct RTCRay* ray,
+      struct RTCOccludedArguments* args = NULL
+    );
+
+#### DESCRIPTION
+
+The `rtcTraversableOccluded1` function is equivalent to `rtcOccluded1`
+but takes traversable object (`traversable` argument) instead of a scene
+object. Using this method is optional on CPU but it is required for SYCL.
+
+For more details, refer to the documentation of `rtcOccluded1`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcOccluded1], [rtcGetSceneTraversable]

doc/src/api/rtcTraversableOccluded4.md

@@ -0,0 +1,47 @@
+% rtcTraversableOccluded4/8/16(3) | Embree Ray Tracing Kernels 4
+
+#### NAME
+
+    rtcTraversableOccluded4/8/16 - finds any hits for a ray packet
+
+#### SYNOPSIS
+
+    #include <embree4/rtcore.h>
+
+    void rtcTraversableOccluded4(
+      const int* valid,
+      RTCTraversable traversable,
+      struct RTCRay4* ray,
+      struct RTCOccludedArguments* args = NULL
+    );
+
+    void rtcTraversableOccluded8(
+      const int* valid,
+      RTCTraversable traversable,
+      struct RTCRay8* ray,
+      struct RTCOccludedArguments* args = NULL
+    );
+
+    void rtcTraversableOccluded16(
+      const int* valid,
+      RTCTraversable traversable,
+      struct RTCRay16* ray,
+      struct RTCOccludedArguments* args = NULL
+    );
+
+#### DESCRIPTION
+
+The `rtcTraversableOccluded4/8/16` functions are equivalent to `rtcOccluded4/8/16`
+but take a traversable object (`traversable` argument) instead of a scene
+object.
+
+For more details, refer to the documentation of `rtcOccluded4/8/16`.
+
+#### EXIT STATUS
+
+For performance reasons this function does not do any error checks,
+thus will not set any error flags on failure.
+
+#### SEE ALSO
+
+[rtcOccluded4/8/16], [rtcGetSceneTraversable]