Merge branch 'sycl' into Realisation_of_sycl_link_build_compile

Author: HabKaffee

sycl/ReleaseNotes.md

@@ -800,7 +800,7 @@ Release notes for commit range 5d7e0925..5eebd1e4bfce
 ### SYCL Compiler
  - Allow for multiple build options for opencl-aot [5e5703f58449]
 ### SYCL Library
- - Implement [`SYCL_INTEL_mem_channel_property`](doc/extensions/MemChannel/SYCL_INTEL_mem_channel_property.asciidoc)
+ - Implement [`SYCL_INTEL_mem_channel_property`](doc/extensions/supported/SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY.asciidoc)
    extension [2f1f3167b7c6]
  - Add `marray` class as defined by SYCL 2020 provisional [5eebd1e4bfce]
  - Implement dynamic batch size adjusting when using Level-Zero plugin
@@ -1186,7 +1186,7 @@ Release notes for commit range 5976ff0..1fc0e4f
 
 ### Documentation
   - Added documentation for [`SPV_INTEL_usm_storage_classes`](doc/extensions/SPIRV/SPV_INTEL_usm_storage_classes.asciidoc)
-    and [SYCL_INTEL_usm_address_spaces](doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc) [781fbfc]
+    and [SYCL_INTEL_usm_address_spaces](doc/extensions/supported/SYCL_EXT_INTEL_USM_ADDRESS_SPACES.asciidoc) [781fbfc]
   - Fixed SPIR-V format name spelling [6e9bf3b]
   - Added extension [LocalMemory](doc/extensions/supported/SYCL_EXT_ONEAPI_LOCAL_MEMORY.asciidoc) draft specification [4b5308a]
   - Added extension [free functions queries](doc/extensions/experimental/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc) draft specification [8953bfd]

sycl/doc/extensions/MemChannel/MemChannel.asciidoc renamed to sycl/doc/extensions/supported/SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY.asciidoc

@@ -168,7 +168,8 @@ namespace sycl {
 
 template <typename T, access::address_space addressSpace =
                           access::address_space::global_space>
-class atomic {
+class __SYCL2020_DEPRECATED(
+    "sycl::atomic is deprecated since SYCL 2020") atomic {
   friend class atomic<T, access::address_space::global_space>;
   static_assert(detail::IsValidAtomicType<T>::value,
                 "Invalid SYCL atomic type. Valid types are: int, "

sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc renamed to sycl/doc/extensions/supported/SYCL_EXT_INTEL_USM_ADDRESS_SPACES.asciidoc

@@ -49,7 +49,7 @@ class simd_mask_impl
   /// Construct from an array. To allow e.g. simd_mask<N> m({1,0,0,1,...}).
   template <int N1, class = std::enable_if_t<N1 == N>>
   simd_mask_impl(const raw_element_type (&&Arr)[N1]) {
-    base_type::init_from_array(std::move(Arr));
+    base_type::template init_from_array<false>(std::move(Arr));
   }
 
   /// Implicit conversion from simd.

sycl/include/CL/sycl/atomic.hpp

@@ -153,13 +153,21 @@ class simd_obj_impl {
   static constexpr int length = N;
 
 protected:
+  template <bool UseSet = true>
   void init_from_array(const Ty (&&Arr)[N]) noexcept {
+    raw_vector_type tmp;
+
     if constexpr (is_wrapper_elem_type_v<Ty>) {
       for (auto I = 0; I < N; ++I) {
-        M_data[I] = bitcast_to_raw_type(Arr[I]);
+        tmp[I] = bitcast_to_raw_type(Arr[I]);
       }
     } else {
-      M_data = make_vector(std::move(Arr));
+      tmp = make_vector(std::move(Arr));
+    }
+    if constexpr (UseSet) {
+      set(std::move(tmp));
+    } else {
+      M_data = std::move(tmp);
     }
   }
 
@@ -216,7 +224,13 @@ class simd_obj_impl {
   template <int N1, class = std::enable_if_t<N1 == N>>
   simd_obj_impl(const Ty (&&Arr)[N1]) noexcept {
     __esimd_dbg_print(simd_obj_impl(const Ty(&&Arr)[N1]));
-    init_from_array(std::move(Arr));
+    init_from_array<false /*init M_data w/o using set(...)*/>(std::move(Arr));
+    // It is OK not to mark a write to M_data with __esimd_vstore (via 'set')
+    // here because:
+    // - __esimd_vstore/vload are need only to mark ESIMD_PRIVATE variable
+    //   access for the VC BE to generate proper code for them.
+    // - initializers are not allowed for ESIMD_PRIVATE vars, so only the
+    //   default ctor can be used for them
   }
 
   /// Load constructor.
@@ -239,34 +253,15 @@ class simd_obj_impl {
     copy_from(acc, offset, Flags{});
   }
 
-  // Load the object's value from array.
-  template <int N1>
-  std::enable_if_t<N1 == N> copy_from(const RawTy (&&Arr)[N1]) {
-    __esimd_dbg_print(copy_from(const RawTy(&&Arr)[N1]));
-    raw_vector_type Tmp;
-    for (auto I = 0; I < N; ++I) {
-      Tmp[I] = Arr[I];
-    }
-    set(Tmp);
-  }
-
-  // Store the object's value to array.
-  template <int N1> std::enable_if_t<N1 == N> copy_to(RawTy (&&Arr)[N1]) const {
-    __esimd_dbg_print(copy_to(RawTy(&&Arr)[N1]));
-    for (auto I = 0; I < N; ++I) {
-      Arr[I] = data()[I];
-    }
+  // Load the object's value from an rvalue array.
+  template <int N1> std::enable_if_t<N1 == N> copy_from(const Ty (&&Arr)[N1]) {
+    __esimd_dbg_print(copy_from(const Ty(&&Arr)[N1]));
+    init_from_array(std::move(Arr));
   }
 
-  /// @{
-  /// Conversion operators.
-  explicit operator const raw_vector_type &() const & {
-    __esimd_dbg_print(explicit operator const raw_vector_type &() const &);
-    return M_data;
-  }
-  explicit operator raw_vector_type &() & {
-    __esimd_dbg_print(explicit operator raw_vector_type &() &);
-    return M_data;
+  explicit operator raw_vector_type() const {
+    __esimd_dbg_print(explicit operator raw_vector_type());
+    return data();
   }
 
   /// Type conversion into a scalar:
@@ -277,7 +272,6 @@ class simd_obj_impl {
     __esimd_dbg_print(operator Ty());
     return bitcast_to_wrapper_type<Ty>(data()[0]);
   }
-  /// @}
 
   raw_vector_type data() const {
     __esimd_dbg_print(raw_vector_type data());

sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_mask_impl.hpp

@@ -68,8 +68,6 @@ class simd_view_impl {
 protected:
   simd_view_impl(BaseTy &Base, RegionTy Region)
       : M_base(Base), M_region(Region) {}
-  simd_view_impl(BaseTy &&Base, RegionTy Region)
-      : M_base(Base), M_region(Region) {}
 
   simd_view_impl(BaseTy &Base) : M_base(Base), M_region(RegionTy(0)) {}
 

sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_obj_impl.hpp

@@ -1,7 +1,4 @@
 // RUN: %clangxx %fsycl-host-only -fsyntax-only -sycl-std=2020 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -o %t.out
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -o %t.out
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -sycl-std=2017 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -o %t.out
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -sycl-std=1.2.1 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -o %t.out
 
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/online_compiler.hpp>
@@ -176,6 +173,11 @@ int main() {
 
   // expected-warning@+1{{'barrier' is deprecated: use 'ext_oneapi_barrier' instead}}
   Queue.submit([&](sycl::handler &CGH) { CGH.barrier(); });
-  
+
+  cl::sycl::multi_ptr<int, cl::sycl::access::address_space::global_space> a(
+      nullptr);
+  // expected-warning@+1 {{'atomic<int, sycl::access::address_space::global_space>' is deprecated: sycl::atomic is deprecated since SYCL 2020}}
+  cl::sycl::atomic<int> b(a);
+
   return 0;
 }