Merge branch 'sycl' into peter/3channel

# Conflicts:
#	sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

Author: ProGTX

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-97f373a",
-      "version": "97f373a",
-      "updated_at": "2024-08-25T22:46:47Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1852379596/zip",
+      "github_tag": "igc-dev-427a492",
+      "version": "427a492",
+      "updated_at": "2024-08-27T03:48:42Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1857739280/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -157,7 +157,7 @@ def : TargetInfo<"__TestDeprecatedAspectList",
 
 def : TargetInfo<"spir64", [], [], "", "", 1>;
 def : TargetInfo<"spir64_gen", [], [], "", "", 1>;
-def : TargetInfo<"spir64_x86_64", [AspectFp16, AspectFp64, AspectAtomic64], [4, 8, 16, 32, 64], "", "", 1>;
+def : TargetInfo<"spir64_x86_64", [AspectFp16, AspectFp64, AspectAtomic64, AspectExt_oneapi_ballot_group], [4, 8, 16, 32, 64], "", "", 1>;
 def : TargetInfo<"spir64_fpga", [], [], "", "", 1>;
 def : TargetInfo<"x86_64", [], [], "", "", 1>;
 // Examples of how to use a combination of explicitly specified values + predefined lists

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -11,9 +11,9 @@ endif()
 
 # Options to override the default behaviour of the FetchContent to include UR
 # source code.
-set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO
+set(SYCL_UR_OVERRIDE_FETCH_CONTENT_REPO
   "" CACHE STRING "Override the Unified Runtime FetchContent repository")
-set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_TAG
+set(SYCL_UR_OVERRIDE_FETCH_CONTENT_TAG
   "" CACHE STRING "Override the Unified Runtime FetchContent tag")
 
 # Options to disable use of FetchContent to include Unified Runtime source code
@@ -160,11 +160,11 @@ if(SYCL_UR_USE_FETCH_CONTENT)
     ${UNIFIED_RUNTIME_TAG}
   )
 
-  if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
-    set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")
+  if(SYCL_UR_OVERRIDE_FETCH_CONTENT_REPO)
+    set(UNIFIED_RUNTIME_REPO "${SYCL_UR_OVERRIDE_FETCH_CONTENT_REPO}")
   endif()
-  if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_TAG)
-    set(UNIFIED_RUNTIME_TAG "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_TAG}")
+  if(SYCL_UR_OVERRIDE_FETCH_CONTENT_TAG)
+    set(UNIFIED_RUNTIME_TAG "${SYCL_UR_OVERRIDE_FETCH_CONTENT_TAG}")
   endif()
 
   message(STATUS "Will fetch Unified Runtime from ${UNIFIED_RUNTIME_REPO}")

sycl/doc/GetStartedGuide.md

@@ -348,10 +348,10 @@ specific repository URL and revision tag used can be found in the file
 In order to enable developers, a number of CMake variables are available to
 control which revision of Unified Runtime should be used when building DPC++:
 
-* `SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO` is a variable which can be used to
+* `SYCL_UR_OVERRIDE_FETCH_CONTENT_REPO` is a variable which can be used to
   override the `UNIFIED_RUNTIME_REPO` variable used by `FetchContent` to attain
   the Unified Runtime source code.
-* `SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_TAG` is a variable which can be used to
+* `SYCL_UR_OVERRIDE_FETCH_CONTENT_TAG` is a variable which can be used to
   override the `UNIFIED_RUNTIME_TAG` variable used by `FetchContent` to attain
   the Unified Runtime source code.
 * `SYCL_UR_USE_FETCH_CONTENT` is an option to control if CMake should use

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -114,11 +114,13 @@ implementation supports.
 
 [frame="none",options="header"]
 |======================
-|Rev |Description
-|1 |Initial draft of the proposal
-|2 |Second revision of the proposal
-|3 |Third revision of the proposal
-|4 |Fourth revision of the proposal
+|Rev | Description
+|1   | Initial draft of the proposal
+|2   | Second revision of the proposal
+|3   | Third revision of the proposal
+|4   | Fourth revision of the proposal
+|5   | Fifth revision of the proposal
+|6   | Sixth revision of the proposal
 |======================
 
 See the revision history at the bottom of this document for features added in 
@@ -2877,7 +2879,7 @@ These features still need to be handled:
                     funcs.
 |5.14|2024-07-17| - Rename `destroy_external_semaphore` to 
                     `release_external_semaphore`.
-|5.15|2024-07-19 - Add missing device queries for image arrays, sampling USM 
+|5.15|2024-07-19| - Add missing device queries for image arrays, sampling USM 
                    images and unique addressing per dimension.
                  - Remove aspects for semaphore export, memory export and fetch 
                    3D USM images as they are not supported on any platform.
@@ -2892,5 +2894,7 @@ These features still need to be handled:
                     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`.
-|5.18|2024-08-27| - Allow 3-channel image formats on some backends.
+|6 |2024-08-29 | - Collated all changes since revision 5.
+                 - Bumped SYCL_EXT_ONEAPI_BINDLESS_IMAGES to number 6.
+                 - Allow 3-channel image formats on some backends.
 |======================

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -1284,7 +1284,7 @@ begin_recording(queue& recordingQueue,
 ----
 
 |Synchronously changes the state of `recordingQueue` to the
-`queue_state::recording` state. This operation is a no-op if `recordingQueue`
+`queue_state::recording` state. This operation is an error if `recordingQueue`
 is already in the `queue_state::recording` state.
 
 Parameters:
@@ -1299,7 +1299,7 @@ Parameters:
 Exceptions:
 
 * Throws synchronously with error code `invalid` if `recordingQueue` is
-  already recording to a different graph.
+  already recording to a graph.
 
 * Throws synchronously with error code `invalid` if `recordingQueue` is
   associated with a device or context that is different from the device
@@ -1314,7 +1314,7 @@ begin_recording(const std::vector<queue>& recordingQueues,
 ----
 
 |Synchronously changes the state of each queue in `recordingQueues` to the
-`queue_state::recording` state. This operation is a no-op for any queue in
+`queue_state::recording` state. This operation is an error for any queue in
 `recordingQueues` that is already in the `queue_state::recording` state.
 
 Parameters:
@@ -1328,8 +1328,8 @@ Parameters:
 
 Exceptions:
 
-* Throws synchronously with error code `invalid` if the any queue in
-  `recordingQueues` is already recording to a different graph.
+* Throws synchronously with error code `invalid` if any queue in
+  `recordingQueues` is already recording to a graph.
 
 * Throws synchronously with error code `invalid` if any of `recordingQueues`
   is associated with a device or context that is different from the device

sycl/include/sycl/ext/intel/esimd/simd.hpp

@@ -77,14 +77,19 @@ class simd : public detail::simd_obj_impl<
   }
 
   // Implicit conversion constructor from sycl::ext::oneapi::experimental::simd
-  template <
-      int N1 = N, class Ty1 = Ty,
-      class SFINAE = std::enable_if_t<
-          (N1 == N) && (N1 <= std::experimental::simd_abi::max_fixed_size<
-                                  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
+  template <int N1 = N, class Ty1 = Ty,
+            class SFINAE = std::enable_if_t<
+                (N1 == N) &&
+                (N1 <= std::experimental::simd_abi::max_fixed_size<Ty>) &&
+                !detail::is_wrapper_elem_type_v<Ty1>>>
   simd(const sycl::ext::oneapi::experimental::simd<Ty, N1> &v)
       : simd(static_cast<raw_vector_type>(v)) {}
 
+  // Implicit conversion constructor from 1D simd_view
+  template <typename BaseTy, int Stride>
+  simd(simd_view<BaseTy, region_base<false, Ty, 1, 1, N, Stride>> &v)
+      : simd(v.read()) {}
+
   /// Broadcast constructor with conversion. Converts given value to
   /// #element_type and replicates it in all elements.
   /// Available when \c T1 is a valid simd element type.
@@ -113,11 +118,11 @@ class simd : public detail::simd_obj_impl<
   /// object. Available when the number of elements does not exceed maximum
   /// fixed size of the oneapi's simd_abi and (TODO, temporary limitation) the
   /// element type is a primitive type (e.g. can't be sycl::half).
-  template <
-      int N1, class Ty1 = Ty,
-      class SFINAE = std::enable_if_t<
-          (N1 == N) && (N1 <= std::experimental::simd_abi::max_fixed_size<
-                                  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
+  template <int N1, class Ty1 = Ty,
+            class SFINAE = std::enable_if_t<
+                (N1 == N) &&
+                (N1 <= std::experimental::simd_abi::max_fixed_size<Ty>) &&
+                !detail::is_wrapper_elem_type_v<Ty1>>>
   operator sycl::ext::oneapi::experimental::simd<Ty, N1>() {
     return sycl::ext::oneapi::experimental::simd<Ty, N1>(base_type::data());
   }
@@ -210,7 +215,8 @@ template <int N> using simd_mask = detail::simd_mask_type<N>;
 template <typename Ty, int N>
 std::ostream &operator<<(std::ostream &OS, const __ESIMD_NS::simd<Ty, N> &V)
 #ifdef __SYCL_DEVICE_ONLY__
-    {}
+{
+}
 #else
 {
   __ESIMD_UNSUPPORTED_ON_HOST;