Merge from 'sycl' to 'sycl-web' (12 commits)

  CONFLICT (content): Merge conflict in sycl/CMakeLists.txt

Author: jsji

clang/include/clang/Driver/Options.td

@@ -4209,6 +4209,10 @@ def fsycl_remove_unused_external_funcs : Flag<["-"], "fsycl-remove-unused-extern
   Group<sycl_Group>, HelpText<"Allow removal of unused `SYCL_EXTERNAL` functions (default)">;
 def fno_sycl_remove_unused_external_funcs : Flag<["-"], "fno-sycl-remove-unused-external-funcs">,
   Group<sycl_Group>, HelpText<"Prevent removal of unused `SYCL_EXTERNAL` functions">;
+def fsycl_allow_device_dependencies : Flag<["-"], "fsycl-allow-device-dependencies">,
+  Group<sycl_Group>, HelpText<"Allow dependencies between device code images">;
+def fno_sycl_allow_device_dependencies : Flag<["-"], "fno-sycl-allow-device-dependencies">,
+  Group<sycl_Group>, HelpText<"Do not allow dependencies between device code images (default)">;
 
 def fsave_optimization_record : Flag<["-"], "fsave-optimization-record">,
   Visibility<[ClangOption, FlangOption]>,

clang/lib/Driver/ToolChains/Clang.cpp

@@ -10762,6 +10762,14 @@ static void addArgs(ArgStringList &DstArgs, const llvm::opt::ArgList &Alloc,
   }
 }
 
+static bool supportDynamicLinking(const llvm::opt::ArgList &TCArgs) {
+  if (TCArgs.hasFlag(options::OPT_fsycl_allow_device_dependencies,
+                     options::OPT_fno_sycl_allow_device_dependencies,
+                     false))
+    return true;
+  return false;
+}
+
 static void getNonTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
                                               const JobAction &JA,
                                               const llvm::opt::ArgList &TCArgs,
@@ -10788,6 +10796,9 @@ static void getNonTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
   if (TCArgs.hasFlag(options::OPT_fno_sycl_esimd_force_stateless_mem,
                      options::OPT_fsycl_esimd_force_stateless_mem, false))
     addArgs(PostLinkArgs, TCArgs, {"-lower-esimd-force-stateless-mem=false"});
+
+  if (supportDynamicLinking(TCArgs))
+    addArgs(PostLinkArgs, TCArgs, {"-support-dynamic-linking"});
 }
 
 // Add any sycl-post-link options that rely on a specific Triple in addition
@@ -10835,6 +10846,8 @@ static void getTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
                        options::OPT_fsycl_remove_unused_external_funcs,
                        false) &&
        !isSYCLNativeCPU(TC)) &&
+      // When supporting dynamic linking, non-kernels in a device image can be called
+      !supportDynamicLinking(TCArgs) &&
       !Triple.isNVPTX() && !Triple.isAMDGPU())
     addArgs(PostLinkArgs, TCArgs, {"-emit-only-kernels-as-entry-points"});
 

clang/test/Driver/sycl-offload-old-model.cpp

@@ -174,9 +174,12 @@
 // RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_OPT_PASS %s
 // CHECK_SYCL_POST_LINK_OPT_PASS: sycl-post-link{{.*}}emit-only-kernels-as-entry-points
 // RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fno-sycl-remove-unused-external-funcs %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_OPT_NO_PASS %s
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_OPT_NO_PASS %s
 // CHECK_SYCL_POST_LINK_OPT_NO_PASS-NOT: sycl-post-link{{.*}}emit-only-kernels-as-entry-points
 
 /// Check selective passing of -support-dynamic-linking to sycl-post-link tool
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_fpga -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SHARED_PASS %s
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SHARED_PASS %s
 // TODO: Enable when SYCL RT supports dynamic linking
 // RUNx: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_fpga -shared %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SHARED_PASS %s
 // RUNx: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -shared %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SHARED_PASS %s

llvm/lib/SYCLLowerIR/ModuleSplitter.cpp

@@ -182,11 +182,8 @@ class DependencyGraph {
       FuncTypeToFuncsMap[F.getFunctionType()].insert(&F);
     }
 
-    // We add every function into the graph except if
-    // SupportDynamicLinking is true
     for (const auto &F : M.functions()) {
-
-      if (SupportDynamicLinking && canBeImportedFunction(F))
+      if (canBeImportedFunction(F))
         continue;
 
       // case (1), see comment above the class definition
@@ -1312,8 +1309,26 @@ splitSYCLModule(std::unique_ptr<Module> M, ModuleSplitterSettings Settings) {
 }
 
 bool canBeImportedFunction(const Function &F) {
+  // It may be theoretically possible to determine what is importable
+  // based solely on function F, but the "SYCL/imported symbols"
+  // property list MUST NOT have any imported symbols that are not supplied
+  // the exported symbols from another device image.  This will lead to a
+  // runtime crash "No device image found for external symbol". Generating
+  // precise "SYCL/imported symbols" can be difficult because there exist
+  // functions that may look like they can be imported, but are supplied outside
+  // of user device code (e.g. _Z38__spirv_JointMatrixWorkItemLength...) In
+  // order to be safe and not require perfect name analysis just start with this
+  // simple check.
+  if (!SupportDynamicLinking)
+    return false;
+
+  // SYCL_EXTERNAL property is not recorded for a declaration
+  // in a header file.  Thus SYCL IR that is a declaration
+  // will be considered as SYCL_EXTERNAL for the purposes of
+  // this function.
   if (F.isIntrinsic() || F.getName().starts_with("__") ||
-      !llvm::sycl::utils::isSYCLExternalFunction(&F))
+      isSpirvSyclBuiltin(F.getName()) || isESIMDBuiltin(F.getName()) ||
+      (!F.isDeclaration() && !llvm::sycl::utils::isSYCLExternalFunction(&F)))
     return false;
 
   bool ReturnValue = true;

llvm/test/tools/sycl-post-link/emit_imported_symbols.ll

@@ -1,12 +1,12 @@
 ; This test checks that the -emit-imported-symbols option generates a list of imported symbols
 ; Function names were chosen so that no function with a 'inside' in their function name is imported
-;
+; Note that -emit-imported-symbols will not emit any imported symbols without -support-dynamic-linking.
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ; Test with -split=kernel
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 
-; RUN: sycl-post-link -properties -symbols -emit-imported-symbols -split=kernel -S < %s -o %t_kernel.table
+; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -emit-imported-symbols -split=kernel -S < %s -o %t_kernel.table
 
 ; RUN: FileCheck %s -input-file=%t_kernel_0.sym --check-prefixes CHECK-KERNEL-SYM-0
 ; RUN: FileCheck %s -input-file=%t_kernel_1.sym --check-prefixes CHECK-KERNEL-SYM-1
@@ -23,29 +23,29 @@
 
 ; CHECK-KERNEL-SYM-1: foo
 ; CHECK-KERNEL-IMPORTED-SYM-1: [SYCL/imported symbols]
+; CHECK-KERNEL-IMPORTED-SYM-1-NEXT: middle
 ; CHECK-KERNEL-IMPORTED-SYM-1-NEXT: childA
 ; CHECK-KERNEL-IMPORTED-SYM-1-NEXT: childC
-; CHECK-KERNEL-IMPORTED-SYM-1-NEXT: childD
 ; CHECK-KERNEL-IMPORTED-SYM-1-EMPTY:
 
 
 ; CHECK-KERNEL-SYM-2: bar
 ; CHECK-KERNEL-IMPORTED-SYM-2: [SYCL/imported symbols]
+; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: middle
 ; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: childB
 ; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: childC
-; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: childD
 ; CHECK-KERNEL-IMPORTED-SYM-2-NEXT: _Z7outsidev
 ; CHECK-KERNEL-IMPORTED-SYM-2-EMPTY:
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ; Test with -split=source
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 
-; RUN: sycl-post-link -properties -symbols -emit-imported-symbols -split=source -S < %s -o %t_source.table
+; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -emit-imported-symbols -split=source -S < %s -o %t_source.table
 ; RUN: FileCheck %s -input-file=%t_source_0.sym --check-prefixes CHECK-SOURCE-SYM-0
 ; RUN: FileCheck %s -input-file=%t_source_0.prop --check-prefixes CHECK-SOURCE-IMPORTED-SYM-0
 
-; RUN: sycl-post-link -properties -symbols -emit-imported-symbols -split=source -S < %s -o %t_source.table -O0
+; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -emit-imported-symbols -split=source -S < %s -o %t_source.table -O0
 ; RUN: FileCheck %s -input-file=%t_source_0.sym --check-prefixes CHECK-SOURCE-SYM-0
 ; RUN: FileCheck %s -input-file=%t_source_0.prop --check-prefixes CHECK-SOURCE-IMPORTED-SYM-0
 
@@ -73,7 +73,7 @@ define weak_odr spir_kernel void @foo() #0 {
 }
 
 define weak_odr spir_kernel void @bar() #0 {
-  ;; Functions that are not SYCL External (i.e. they have no sycl-module-id) cannot be imported
+  ;; Functions whose name start with '__' cannot be imported
   call spir_func void @__itt_offload_wi_start_wrapper()
 
   call void @childB()

llvm/test/tools/sycl-post-link/internalize_functions.ll

@@ -1,5 +1,5 @@
 ; Test that when -support-dynamic-linking is used
-; non SYCL-EXTERNAL functions are internalized.
+; non SYCL_EXTERNAL functions are internalized.
 ; Variables must not be internalized.
 
 ; RUN: sycl-post-link -symbols -support-dynamic-linking -split=kernel -S < %s -o %t.table
@@ -8,8 +8,8 @@
 
 ; CHECK-SYM-0: foo0
 
-; Non SYCL-EXTERNAL Functions are internalized
-; foo0 is a SYCL-EXTERNAL function
+; Non SYCL_EXTERNAL Functions are internalized
+; foo0 is a SYCL_EXTERNAL function
 ; CHECK-LL-0-DAG: define weak_odr spir_kernel void @foo0() #0 {
 ; Internalize does not change available_externally
 ; CHECK-LL-0-DAG: define available_externally spir_func void @internalA() {

sycl/CMakeLists.txt

@@ -42,6 +42,45 @@ list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules")
 include(AddSYCLExecutable)
 include(AddSYCL)
 include(SYCLUtils)
+
+if(MSVC)
+  set_property(GLOBAL PROPERTY USE_FOLDERS ON)
+  # Skip asynchronous C++ exceptions catching and assume "extern C" functions
+  # never throw C++ exceptions.
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /EHsc")
+
+  # Add PDB debug information
+  list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}")
+  include(CheckLinkerFlag)
+  check_linker_flag(CXX "/DEBUG" LINKER_SUPPORTS_DEBUG)
+  if(LINKER_SUPPORTS_DEBUG)
+    # sccache is not compatible with /Zi flag
+    if (CMAKE_CXX_COMPILER_LAUNCHER STREQUAL "sccache")
+      # CMake may put /Zi by default
+      if(CMAKE_BUILD_TYPE STREQUAL "Debug")
+        string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}")
+        string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG}")
+      elseif(CMAKE_BUILD_TYPE STREQUAL "Release")
+        string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}")
+        string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE}")
+      elseif(CMAKE_BUILD_TYPE STREQUAL "RelWithDebInfo")
+        string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_RELWITHDEBINFO "${CMAKE_CXX_FLAGS_RELWITHDEBINFO}")
+        string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_RELWITHDEBINFO "${CMAKE_C_FLAGS_RELWITHDEBINFO}")
+      endif()
+      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /Z7")
+    else()
+      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /Zi")
+    endif()
+    add_link_options("/DEBUG")
+
+    # Enable unreferenced removal and ICF in Release mode.
+    check_linker_flag(CXX "/OPT:REF /OPT:ICF" LINKER_SUPPORTS_OPTS)
+    if (LINKER_SUPPORTS_OPTS AND uppercase_CMAKE_BUILD_TYPE STREQUAL "RELEASE")
+      add_link_options("/OPT:REF" "/OPT:ICF")
+    endif()
+  endif()
+endif()
+
 include(FetchUnifiedRuntime)
 
 # The change in SYCL_MAJOR_VERSION must be accompanied with the same update in
@@ -86,44 +125,6 @@ endif()
 # of the SYCL runtime and expect enabling
 option(SYCL_ENABLE_XPTI_TRACING "Enable tracing of SYCL constructs" OFF)
 
-if(MSVC)
-  set_property(GLOBAL PROPERTY USE_FOLDERS ON)
-  # Skip asynchronous C++ exceptions catching and assume "extern C" functions
-  # never throw C++ exceptions.
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /EHsc")
-
-  # Add PDB debug information
-  list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}")
-  include(CheckLinkerFlag)
-  check_linker_flag(CXX "/DEBUG" LINKER_SUPPORTS_DEBUG)
-  if(LINKER_SUPPORTS_DEBUG)
-    # sccache is not compatible with /Zi flag
-    if (CMAKE_CXX_COMPILER_LAUNCHER STREQUAL "sccache")
-      # CMake may put /Zi by default
-      if(CMAKE_BUILD_TYPE STREQUAL "Debug")
-        string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}")
-        string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG}")
-      elseif(CMAKE_BUILD_TYPE STREQUAL "Release")
-        string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}")
-        string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE}")
-      elseif(CMAKE_BUILD_TYPE STREQUAL "RelWithDebInfo")
-        string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_RELWITHDEBINFO "${CMAKE_CXX_FLAGS_RELWITHDEBINFO}")
-        string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_RELWITHDEBINFO "${CMAKE_C_FLAGS_RELWITHDEBINFO}")
-      endif()
-      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /Z7")
-    else()
-      set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /Zi")
-    endif()
-    add_link_options("/DEBUG")
-
-    # Enable unreferenced removal and ICF in Release mode.
-    check_linker_flag(CXX "/OPT:REF /OPT:ICF" LINKER_SUPPORTS_OPTS)
-    if (LINKER_SUPPORTS_OPTS AND uppercase_CMAKE_BUILD_TYPE STREQUAL "RELEASE")
-      add_link_options("/OPT:REF" "/OPT:ICF")
-    endif()
-  endif()
-endif()
-
 # Get clang's version
 include(VersionFromVCS)
 set(PACKAGE_VERSION "${LLVM_PACKAGE_VERSION}")

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -115,14 +115,14 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
       CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
   endfunction()
 
-set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 3e762e00bcf13d158fb58e8e8c2eabcfc8934b4e
-  # Merge: c805a71a a2a053de
+  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
+  # commit c5d2175b5823d5b74de1e7e0d6081ab6d885bc34
+  # Merge: 99489ad4 c86beb60
   # Author: Omar Ahmed <[email protected]>
-  # Date:   Wed Jul 31 12:26:34 2024 +0100
-  #     Merge pull request #1884 from callumfare/callum/fix_printtrace
-  #     Enable PrintTrace when SYCL UR tracing is enabled
-  set(UNIFIED_RUNTIME_TAG 3e762e00bcf13d158fb58e8e8c2eabcfc8934b4e)
+  # Date:   Wed Jul 31 14:52:26 2024 +0100
+  #     Merge pull request #1882 from przemektmalon/przemek/interop-map-memory
+  #    [Bindless][Exp] Add interop memory mapping to USM.
+  set(UNIFIED_RUNTIME_TAG c5d2175b5823d5b74de1e7e0d6081ab6d885bc34)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -1712,19 +1712,35 @@ image_mem_handle map_external_image_memory(
     external_mem externalMemHandle,
     const image_descriptor &imageDescriptor,
     const sycl::queue &syclQueue);
+
+void *map_external_linear_memory(
+    external_mem externalMemHandle,
+    uint64_t size, uint64_t offset,
+    const sycl::device &syclDevice,
+    const sycl::context &syclContext);
+void *map_external_linear_memory(
+    external_mem externalMemHandle,
+    uint64_t size, uint64_t offset,
+    const sycl::queue &syclQueue);
 }
 ```
 
 The resulting `external_mem` can then be mapped, where the resulting type 
-is an `image_mem_handle`. This can be used to construct images in the same way 
-as memory allocated through `alloc_image_mem`. The `ext_oneapi_copy` operations 
-also work with imported memory mapped to `image_mem_handle` types.
-
-When calling `create_image` with an `image_mem_handle` mapped from an external 
-memory object, the user must ensure that the image descriptor they pass to 
-`create_image` has members that match or map to those of the external API. 
-A mismatch between any of the `width`, `height`, `depth`, `image_channel_type`,
-or `num_channels` members will result in undefined behavior.
+is an `image_mem_handle` or a `void *`. This can be used to construct images in
+the same way as memory allocated through `alloc_image_mem`, 
+`pitched_alloc_device`, or another USM allocation method. The `ext_oneapi_copy` 
+operations also work with imported memory mapped to `image_mem_handle` and 
+`void *` types.
+
+When calling `create_image` with an `image_mem_handle` or `void *` mapped from 
+an external memory object, the user must ensure that the image descriptor they 
+pass to `create_image` has members that match or map to those of the external 
+API. A mismatch between any of the `width`, `height`, `depth`, 
+`image_channel_type`, or `num_channels` members will result in undefined 
+behavior. Likewise, if the image is mapped to a linear USM (`void *`) region, 
+the pitch value passed to `create_image` needs to match the pitch of the image 
+as defined by the external API. Note that when external memory is mapped to a 
+linear USM region, this is device-side USM, and not accessible on the host.
 
 Additionally, the `image_type` describing the image must match to the image of 
 the external API. The current supported importable image types are `standard` 
@@ -2866,4 +2882,6 @@ These features still need to be handled:
                   - Removed `handle` keyword from `interop_xxx_handle` to 
                     clear up possible confusion between 3rd party interop 
                     handles and the imported `interop_xxx_handle`.
+|5.17|2024-07-30| - Add support for mapping external memory to linear USM using
+                    `map_external_linear_memory`.
 |======================

sycl/doc/extensions/proposed/sycl_ext_intel_cache_controls.asciidoc

@@ -8,7 +8,6 @@
 
 #pragma once
 
-#include <sycl/detail/boolean.hpp>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/generic_type_traits.hpp>
 #include <sycl/ext/intel/esimd/detail/math_intrin.hpp>

sycl/include/sycl/builtins_esimd.hpp

@@ -10,7 +10,6 @@
 
 #include <sycl/access/access.hpp>              // for address_space, decorated
 #include <sycl/aliases.hpp>                    // for half
-#include <sycl/detail/boolean.hpp>             // for Boolean
 #include <sycl/detail/defines_elementary.hpp>  // for __SYCL_ALWAYS_INLINE
 #include <sycl/detail/generic_type_traits.hpp> // for is_svgenfloat, is_sge...
 #include <sycl/detail/type_list.hpp>           // for is_contained, type_list