diff --git a/.github/workflows/continuous.yml b/.github/workflows/continuous.yml index aa464dcd17..d0369cb061 100644 --- a/.github/workflows/continuous.yml +++ b/.github/workflows/continuous.yml @@ -194,7 +194,7 @@ jobs: with: runs-on: '[ "Windows", "dg2" ]' msvc-version: 2022 - env-from-files: ./.github/workflows/dpcpp-sycl-nightly.env ./.github/workflows/gfx-windows-public.env + env-from-files: ./.github/workflows/gfx-windows-public.env artifact-in: windows-JIT-build cmd: | $EmbreeDir = .\\scripts\\prepare-test-files.ps1 -TestType ctest @@ -207,7 +207,7 @@ jobs: with: runs-on: '[ "Windows", "bmg" ]' msvc-version: 2022 - env-from-files: ./.github/workflows/dpcpp-sycl-nightly.env ./.github/workflows/gfx-windows-public.env + env-from-files: ./.github/workflows/gfx-windows-public.env artifact-in: windows-JIT-build cmd: | $EmbreeDir = .\\scripts\\prepare-test-files.ps1 -TestType ctest @@ -220,7 +220,7 @@ jobs: with: runs-on: '[ "Windows", "lnl" ]' msvc-version: 2022 - env-from-files: ./.github/workflows/dpcpp-sycl-nightly.env ./.github/workflows/gfx-windows-public.env + env-from-files: ./.github/workflows/gfx-windows-public.env artifact-in: windows-JIT-build cmd: | $EmbreeDir = .\\scripts\\prepare-test-files.ps1 -TestType ctest @@ -233,7 +233,7 @@ jobs: with: runs-on: '[ "Windows", "dg2" ]' msvc-version: 2022 - env-from-files: ./.github/workflows/dpcpp-sycl-nightly.env ./.github/workflows/gfx-windows-internal.env + env-from-files: ./.github/workflows/gfx-windows-internal.env artifact-in: windows-JIT-build cmd: | $EmbreeDir = .\\scripts\\prepare-test-files.ps1 -TestType ctest diff --git a/.github/workflows/dpcpp-sycl-nightly.env b/.github/workflows/dpcpp-sycl-nightly.env index 674ff1ce36..d04958443f 100644 --- a/.github/workflows/dpcpp-sycl-nightly.env +++ b/.github/workflows/dpcpp-sycl-nightly.env @@ -1 +1 @@ -DPCPP_VERSION=intel-llvm/v6.2.0 \ No newline at end of file +DPCPP_VERSION=intel-llvm/v7.2.0 \ No newline at end of file diff --git a/.github/workflows/nightly.yml b/.github/workflows/nightly.yml index efa11b1543..bc4ed2c7cd 100644 --- a/.github/workflows/nightly.yml +++ b/.github/workflows/nightly.yml @@ -95,7 +95,7 @@ jobs: options: --device=/dev/dri:/dev/dri runs-on: '[ "Linux", "docker", "dg2" ]' env-from-files: ./.github/workflows/gfx-ubuntu24-public.env - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-in: nightly-linux-DG2-JIT-PUBLIC-build cmd: | EMBREE_PATH=$(./scripts/prepare-test-files.sh --type ctest) @@ -140,7 +140,7 @@ jobs: options: --device=/dev/dri:/dev/dri runs-on: '[ "Linux", "docker", "pvc" ]' env-from-files: ./.github/workflows/gfx-ubuntu24-public.env - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-in: nightly-linux-PVC-JIT-build cmd: | export NEOReadDebugKeys=1 @@ -159,7 +159,7 @@ jobs: options: --device=/dev/dri:/dev/dri runs-on: '[ "Linux", "docker", "pvc" ]' env-from-files: ./.github/workflows/gfx-ubuntu24-public.env - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-in: nightly-linux-PVC-JIT-build cmd: | export NEOReadDebugKeys=1 @@ -210,7 +210,7 @@ jobs: runs-on: '[ "Windows", "dg2" ]' env-from-files: ./.github/workflows/gfx-windows-public.env msvc-version: 2022 - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-out: nightly-windows-DG2-JIT-INTELLLVM-bulid artifact-path: ./build/*.zip cmd: | @@ -242,7 +242,7 @@ jobs: runs-on: '[ "Windows", "dg2" ]' env-from-files: ./.github/workflows/gfx-windows-public.env msvc-version: 2022 - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-in: nightly-windows-DG2-JIT-INTELLLVM-bulid cmd: | $EmbreeDir = .\scripts\prepare-test-files.ps1 -TestType ctest @@ -255,7 +255,7 @@ jobs: runs-on: '[ "Windows", "dg2" ]' env-from-files: ./.github/workflows/gfx-windows-public.env msvc-version: 2022 - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-out: nightly-windows-DG2-JIT-PUBLIC-build artifact-path: ./build/*.zip cmd: | @@ -287,7 +287,7 @@ jobs: runs-on: '[ "Windows", "dg2" ]' env-from-files: ./.github/workflows/gfx-windows-public.env msvc-version: 2022 - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-in: nightly-windows-DG2-JIT-PUBLIC-build cmd: | $EmbreeDir = .\scripts\prepare-test-files.ps1 -TestType ctest diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index c2e63edbb5..6d4b414ff4 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -23,7 +23,7 @@ jobs: artifact-out: build-windows artifact-path: ./build/embree-*.zip msvc-version: 2019 - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 cmd: | $opts = "-DCMAKE_BUILD_TYPE=Release" $opts += " -DBUILD_TESTING=ON" @@ -167,7 +167,7 @@ jobs: runs-on: '[ "Linux", "build" ]' artifact-out: build-linux artifact-path: ./build/*.tar.gz - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 pre-cmd: | dnf -y install epel-release dnf -y config-manager --set-enabled powertools @@ -215,7 +215,7 @@ jobs: with: runs-on: '[ "Windows", "build" ]' msvc-version: 2019 - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-out: build-windows-sycl artifact-path: ./build/*.zip cmd: | @@ -243,7 +243,7 @@ jobs: with: image: rockylinux:8.7 runs-on: '[ "Linux", "build" ]' - dpcpp-version: intel-llvm/v6.2.0 + dpcpp-version: intel-llvm/v7.2.0 artifact-out: build-linux-sycl artifact-path: ./build/*.tar.gz pre-cmd: | diff --git a/README.md b/README.md index fac9b1e8e5..4165e4e49e 100644 --- a/README.md +++ b/README.md @@ -286,7 +286,7 @@ C++11. Embree is tested with the following compilers: Linux - Intel® oneAPI DPC++/C++ Compiler 2025.2 - - oneAPI DPC++/C++ Compiler 6.2.0 + - oneAPI DPC++/C++ Compiler 7.2.0 - Clang 20 - GCC 15 (Fedora 42) AVX512 support - Intel® Implicit SPMD Program Compiler 1.22.0 @@ -397,7 +397,7 @@ The "oneAPI DPC++ Compiler" is more up-to-date than the "Intel(R) oneAPI DPC++/C++ Compiler" but less stable. The current tested version of the "oneAPI DPC++ compiler is - - [oneAPI DPC++ Compiler 6.2.0](https://github.com/intel/llvm/releases/tag/v6.2.0) + - [oneAPI DPC++ Compiler 7.2.0](https://github.com/intel/llvm/releases/tag/v7.2.0) The compiler can be downloaded and simply extracted. The oneAPI DPC++ compiler can be set up executing the following commands in a Linux (bash) shell: @@ -481,7 +481,7 @@ Windows Embree is tested using the following compilers under Windows: - Intel® oneAPI DPC++/C++ Compiler 2025.2 - - oneAPI DPC++/C++ Compiler 6.2.0 + - oneAPI DPC++/C++ Compiler 7.2.0 - Visual Studio 2022 - Visual Studio 2019 - Visual Studio 2017 @@ -602,7 +602,7 @@ The "oneAPI DPC++ Compiler" is more up-to-date than the "Intel(R) oneAPI DPC++/C++ Compiler" but less stable. The current tested version of the oneAPI DPC++ compiler is - - [oneAPI DPC++ Compiler 6.2.0](https://github.com/intel/llvm/releases/tag/v6.2.0) + - [oneAPI DPC++ Compiler 7.2.0](https://github.com/intel/llvm/releases/tag/v7.2.0) Download and unpack the archive and open the "x64 Native Tools Command Prompt" of Visual Studio and execute the following lines to properly configure the diff --git a/common/cmake/package.cmake b/common/cmake/package.cmake index daa4f49fa9..baefb11db5 100644 --- a/common/cmake/package.cmake +++ b/common/cmake/package.cmake @@ -45,6 +45,9 @@ IF (EMBREE_SYCL_SUPPORT AND EMBREE_INSTALL_DEPENDENCIES) 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}/ur_adapter_level_zero_V2.dll") + INSTALL(FILES "${DPCPP_COMPILER_DIR}/ur_adapter_level_zero_V2.dll" DESTINATION "${CMAKE_INSTALL_BINDIR}" COMPONENT lib) + ENDIF() # copy unified runtime loader dll IF(EXISTS "${DPCPP_COMPILER_DIR}/ur_loader.dll") @@ -84,6 +87,7 @@ IF (EMBREE_SYCL_SUPPORT AND EMBREE_INSTALL_DEPENDENCIES) "${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_adapter_level_zero_V2.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) diff --git a/common/sys/sycl.h b/common/sys/sycl.h index 6bef829fc2..a0e6d0262c 100644 --- a/common/sys/sycl.h +++ b/common/sys/sycl.h @@ -32,158 +32,17 @@ using sycl::ushort3; using sycl::ushort2; #ifdef __SYCL_DEVICE_ONLY__ -#define GLOBAL __attribute__((opencl_global)) -#define LOCAL __attribute__((opencl_local)) - -SYCL_EXTERNAL extern int work_group_reduce_add(int x); -SYCL_EXTERNAL extern float work_group_reduce_min(float x); -SYCL_EXTERNAL extern float work_group_reduce_max(float x); - -SYCL_EXTERNAL extern float atomic_min(volatile GLOBAL float *p, float val); -SYCL_EXTERNAL extern float atomic_min(volatile LOCAL float *p, float val); -SYCL_EXTERNAL extern float atomic_max(volatile GLOBAL float *p, float val); -SYCL_EXTERNAL extern float atomic_max(volatile LOCAL float *p, float val); - SYCL_EXTERNAL extern "C" unsigned int intel_sub_group_ballot(bool valid); - -SYCL_EXTERNAL extern "C" void __builtin_IB_assume_uniform(void *p); - -// Load message caching control - - enum LSC_LDCC { - LSC_LDCC_DEFAULT, - LSC_LDCC_L1UC_L3UC, // Override to L1 uncached and L3 uncached - LSC_LDCC_L1UC_L3C, // Override to L1 uncached and L3 cached - LSC_LDCC_L1C_L3UC, // Override to L1 cached and L3 uncached - LSC_LDCC_L1C_L3C, // Override to L1 cached and L3 cached - LSC_LDCC_L1S_L3UC, // Override to L1 streaming load and L3 uncached - LSC_LDCC_L1S_L3C, // Override to L1 streaming load and L3 cached - LSC_LDCC_L1IAR_L3C, // Override to L1 invalidate-after-read, and L3 cached - }; - - - -// Store message caching control (also used for atomics) - - enum LSC_STCC { - LSC_STCC_DEFAULT, - LSC_STCC_L1UC_L3UC, // Override to L1 uncached and L3 uncached - LSC_STCC_L1UC_L3WB, // Override to L1 uncached and L3 written back - LSC_STCC_L1WT_L3UC, // Override to L1 written through and L3 uncached - LSC_STCC_L1WT_L3WB, // Override to L1 written through and L3 written back - LSC_STCC_L1S_L3UC, // Override to L1 streaming and L3 uncached - LSC_STCC_L1S_L3WB, // Override to L1 streaming and L3 written back - LSC_STCC_L1WB_L3WB, // Override to L1 written through and L3 written back - }; - - - -/////////////////////////////////////////////////////////////////////// - -// LSC Loads - -/////////////////////////////////////////////////////////////////////// - -SYCL_EXTERNAL /* extern "C" */ uint32_t __builtin_IB_lsc_load_global_uchar_to_uint (const GLOBAL uint8_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D8U32 -SYCL_EXTERNAL /* extern "C" */ uint32_t __builtin_IB_lsc_load_global_ushort_to_uint(const GLOBAL uint16_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D16U32 -SYCL_EXTERNAL /* extern "C" */ uint32_t __builtin_IB_lsc_load_global_uint (const GLOBAL uint32_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V1 -SYCL_EXTERNAL /* extern "C" */ sycl::uint2 __builtin_IB_lsc_load_global_uint2 (const GLOBAL sycl::uint2 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V2 -SYCL_EXTERNAL /* extern "C" */ sycl::uint3 __builtin_IB_lsc_load_global_uint3 (const GLOBAL sycl::uint3 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V3 -SYCL_EXTERNAL /* extern "C" */ sycl::uint4 __builtin_IB_lsc_load_global_uint4 (const GLOBAL sycl::uint4 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V4 -SYCL_EXTERNAL /* extern "C" */ sycl::uint8 __builtin_IB_lsc_load_global_uint8 (const GLOBAL sycl::uint8 *base, int elemOff, enum LSC_LDCC cacheOpt); //D32V8 -SYCL_EXTERNAL /* extern "C" */ uint64_t __builtin_IB_lsc_load_global_ulong (const GLOBAL uint64_t *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V1 -SYCL_EXTERNAL /* extern "C" */ sycl::ulong2 __builtin_IB_lsc_load_global_ulong2 (const GLOBAL sycl::ulong2 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V2 -SYCL_EXTERNAL /* extern "C" */ sycl::ulong3 __builtin_IB_lsc_load_global_ulong3 (const GLOBAL sycl::ulong3 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V3 -SYCL_EXTERNAL /* extern "C" */ sycl::ulong4 __builtin_IB_lsc_load_global_ulong4 (const GLOBAL sycl::ulong4 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V4 -SYCL_EXTERNAL /* extern "C" */ sycl::ulong8 __builtin_IB_lsc_load_global_ulong8 (const GLOBAL sycl::ulong8 *base, int elemOff, enum LSC_LDCC cacheOpt); //D64V8 - -// global address space -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uchar_from_uint (GLOBAL uint8_t *base, int immElemOff, uint32_t val, enum LSC_STCC cacheOpt); //D8U32 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ushort_from_uint(GLOBAL uint16_t *base, int immElemOff, uint32_t val, enum LSC_STCC cacheOpt); //D16U32 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint (GLOBAL uint32_t *base, int immElemOff, uint32_t val, enum LSC_STCC cacheOpt); //D32V1 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint2 (GLOBAL sycl::uint2 *base, int immElemOff, sycl::uint2 val, enum LSC_STCC cacheOpt); //D32V2 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint3 (GLOBAL sycl::uint3 *base, int immElemOff, sycl::uint3 val, enum LSC_STCC cacheOpt); //D32V3 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint4 (GLOBAL sycl::uint4 *base, int immElemOff, sycl::uint4 val, enum LSC_STCC cacheOpt); //D32V4 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_uint8 (GLOBAL sycl::uint8 *base, int immElemOff, sycl::uint8 val, enum LSC_STCC cacheOpt); //D32V8 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong (GLOBAL uint64_t *base, int immElemOff, uint64_t val, enum LSC_STCC cacheOpt); //D64V1 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong2 (GLOBAL sycl::ulong2 *base, int immElemOff, sycl::ulong2 val, enum LSC_STCC cacheOpt); //D64V2 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong3 (GLOBAL sycl::ulong3 *base, int immElemOff, sycl::ulong3 val, enum LSC_STCC cacheOpt); //D64V3 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong4 (GLOBAL sycl::ulong4 *base, int immElemOff, sycl::ulong4 val, enum LSC_STCC cacheOpt); //D64V4 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_store_global_ulong8 (GLOBAL sycl::ulong8 *base, int immElemOff, sycl::ulong8 val, enum LSC_STCC cacheOpt); //D64V8 - -/////////////////////////////////////////////////////////////////////// -// prefetching -/////////////////////////////////////////////////////////////////////// -// -// LSC Pre-Fetch Load functions with CacheControls -// global address space -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uchar (const GLOBAL uint8_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D8U32 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ushort(const GLOBAL uint16_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D16U32 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint (const GLOBAL uint32_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V1 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint2 (const GLOBAL sycl::uint2 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V2 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint3 (const GLOBAL sycl::uint3 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V3 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint4 (const GLOBAL sycl::uint4 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V4 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_uint8 (const GLOBAL sycl::uint8 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D32V8 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong (const GLOBAL uint64_t *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V1 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong2(const GLOBAL sycl::ulong2 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V2 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong3(const GLOBAL sycl::ulong3 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V3 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong4(const GLOBAL sycl::ulong4 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V4 -SYCL_EXTERNAL extern "C" void __builtin_IB_lsc_prefetch_global_ulong8(const GLOBAL sycl::ulong8 *base, int immElemOff, enum LSC_LDCC cacheOpt); //D64V8 - #else - -#define GLOBAL -#define LOCAL - -/* dummy functions for host */ -inline int work_group_reduce_add(int x) { return x; } -inline float work_group_reduce_min(float x) { return x; } -inline float work_group_reduce_max(float x) { return x; } - -inline float atomic_min(volatile float *p, float val) { return val; }; -inline float atomic_max(volatile float *p, float val) { return val; }; - inline uint32_t intel_sub_group_ballot(bool valid) { return 0; } - #endif -/* creates a temporary that is enforced to be uniform */ -#define SYCL_UNIFORM_VAR(Ty,tmp,k) \ - Ty tmp##_data; \ - Ty* p##tmp##_data = (Ty*) sub_group_broadcast((uint64_t)&tmp##_data,k); \ - Ty& tmp = *p##tmp##_data; - #if !defined(__forceinline) #define __forceinline inline __attribute__((always_inline)) #endif -#if __SYCL_COMPILER_VERSION < 20210801 -#define all_of_group all_of -#define any_of_group any_of -#define none_of_group none_of -#define group_broadcast broadcast -#define reduce_over_group reduce -#define exclusive_scan_over_group exclusive_scan -#define inclusive_scan_over_group inclusive_scan -#endif - namespace embree { - template - __forceinline T cselect(const bool mask, const T &a, const T &b) - { - return sycl::select(b,a,(int)mask); - } - - template - __forceinline T cselect(const M &mask, const T &a, const T &b) - { - return sycl::select(b,a,mask); - } - -#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(); @@ -191,112 +50,8 @@ namespace embree return sycl::ext::oneapi::experimental::this_sub_group(); #endif } - + __forceinline const uint32_t get_sub_group_local_id() { return this_sub_group().get_local_id()[0]; } - - __forceinline const uint32_t get_sub_group_size() { - return this_sub_group().get_max_local_range().size(); - } - - __forceinline const uint32_t get_sub_group_id() { - return this_sub_group().get_group_id()[0]; - } - - __forceinline const uint32_t get_num_sub_groups() { - return this_sub_group().get_group_range().size(); - } - - __forceinline uint32_t sub_group_ballot(bool pred) { - return intel_sub_group_ballot(pred); - } - - __forceinline bool sub_group_all_of(bool pred) { - return sycl::all_of_group(this_sub_group(),pred); - } - - __forceinline bool sub_group_any_of(bool pred) { - return sycl::any_of_group(this_sub_group(),pred); - } - - __forceinline bool sub_group_none_of(bool pred) { - return sycl::none_of_group(this_sub_group(),pred); - } - - template __forceinline T sub_group_broadcast(T x, sycl::id<1> local_id) { - return sycl::group_broadcast(this_sub_group(),x,local_id); - } - - template __forceinline T sub_group_make_uniform(T x) { - return sub_group_broadcast(x,sycl::ctz(intel_sub_group_ballot(true))); - } - - __forceinline void assume_uniform_array(void* ptr) { -#ifdef __SYCL_DEVICE_ONLY__ - __builtin_IB_assume_uniform(ptr); -#endif - } - - template __forceinline T sub_group_reduce(T x, BinaryOperation binary_op) { - return sycl::reduce_over_group(this_sub_group(),x,binary_op); - } - - template __forceinline T sub_group_reduce(T x, T init, BinaryOperation binary_op) { - return sycl::reduce_over_group(this_sub_group(),x,init,binary_op); - } - - template __forceinline T sub_group_reduce_min(T x, T init) { - return sub_group_reduce(x, init, sycl::ext::oneapi::minimum()); - } - - template __forceinline T sub_group_reduce_min(T x) { - return sub_group_reduce(x, sycl::ext::oneapi::minimum()); - } - - template __forceinline T sub_group_reduce_max(T x) { - return sub_group_reduce(x, sycl::ext::oneapi::maximum()); - } - - template __forceinline T sub_group_reduce_add(T x) { - return sub_group_reduce(x, sycl::ext::oneapi::plus()); - } - - template __forceinline T sub_group_exclusive_scan(T x, BinaryOperation binary_op) { - return sycl::exclusive_scan_over_group(this_sub_group(),x,binary_op); - } - - template __forceinline T sub_group_exclusive_scan_min(T x) { - return sub_group_exclusive_scan(x,sycl::ext::oneapi::minimum()); - } - - template __forceinline T sub_group_exclusive_scan(T x, T init, BinaryOperation binary_op) { - return sycl::exclusive_scan_over_group(this_sub_group(),x,init,binary_op); - } - - template __forceinline T sub_group_inclusive_scan(T x, BinaryOperation binary_op) { - return sycl::inclusive_scan_over_group(this_sub_group(),x,binary_op); - } - - template __forceinline T sub_group_inclusive_scan(T x, BinaryOperation binary_op, T init) { - return sycl::inclusive_scan_over_group(this_sub_group(),x,binary_op,init); - } - - template __forceinline T sub_group_load(const void* src) { - return this_sub_group().load(sycl::multi_ptr((T*)src)); - } - - template __forceinline void sub_group_store(void* dst, const T& x) { - this_sub_group().store(sycl::multi_ptr((T*)dst),x); - } } - -#if __SYCL_COMPILER_VERSION < 20210801 -#undef all_of_group -#undef any_of_group -#undef none_of_group -#undef group_broadcast -#undef reduce_over_group -#undef exclusive_scan_over_group -#undef inclusive_scan_over_group -#endif diff --git a/doc/src/compilation.md b/doc/src/compilation.md index 0e883c2a60..29f3ab1dc3 100644 --- a/doc/src/compilation.md +++ b/doc/src/compilation.md @@ -18,7 +18,7 @@ C++11. Embree is tested with the following compilers: Linux - Intel® oneAPI DPC++/C++ Compiler 2025.2 - - oneAPI DPC++/C++ Compiler 6.2.0 + - oneAPI DPC++/C++ Compiler 7.2.0 - Clang 20 - GCC 15 (Fedora 42) AVX512 support - Intel® Implicit SPMD Program Compiler 1.22.0 @@ -129,7 +129,7 @@ The "oneAPI DPC++ Compiler" is more up-to-date than the "Intel(R) oneAPI DPC++/C++ Compiler" but less stable. The current tested version of the "oneAPI DPC++ compiler is - - [oneAPI DPC++ Compiler 6.2.0](https://github.com/intel/llvm/releases/tag/v6.2.0) + - [oneAPI DPC++ Compiler 7.2.0](https://github.com/intel/llvm/releases/tag/v7.2.0) The compiler can be downloaded and simply extracted. The oneAPI DPC++ compiler can be set up executing the following commands in a Linux (bash) shell: @@ -213,7 +213,7 @@ Windows Embree is tested using the following compilers under Windows: - Intel® oneAPI DPC++/C++ Compiler 2025.2 - - oneAPI DPC++/C++ Compiler 6.2.0 + - oneAPI DPC++/C++ Compiler 7.2.0 - Visual Studio 2022 - Visual Studio 2019 - Visual Studio 2017 @@ -334,7 +334,7 @@ The "oneAPI DPC++ Compiler" is more up-to-date than the "Intel(R) oneAPI DPC++/C++ Compiler" but less stable. The current tested version of the oneAPI DPC++ compiler is - - [oneAPI DPC++ Compiler 6.2.0](https://github.com/intel/llvm/releases/tag/v6.2.0) + - [oneAPI DPC++ Compiler 7.2.0](https://github.com/intel/llvm/releases/tag/v7.2.0) Download and unpack the archive and open the "x64 Native Tools Command Prompt" of Visual Studio and execute the following lines to properly configure the