Merge from 'sycl' to 'sycl-web'

Author: iclsrc

.github/CODEOWNERS

@@ -99,10 +99,10 @@ clang/tools/clang-offload-deps/ @sndmitriev @mlychkov @AlexeySachkov
 clang/tools/clang-offload-extract/ @sndmitriev @mlychkov @AlexeySachkov
 
 # Explicit SIMD
-SYCLLowerIR/ @kbobrovs @DenisBakhvalov
-esimd/ @kbobrovs @DenisBakhvalov
-sycl/include/sycl/ext/intel/experimental/esimd.hpp @kbobrovs @DenisBakhvalov
-sycl/doc/extensions/ExplicitSIMD/ @kbobrovs
+SYCLLowerIR/ @kbobrovs @DenisBakhvalov @kychendev
+esimd/ @kbobrovs @DenisBakhvalov @kychendev
+sycl/include/sycl/ext/intel/experimental/esimd.hpp @kbobrovs @DenisBakhvalov @kychendev
+sycl/doc/extensions/ExplicitSIMD/ @kbobrovs @DenisBakhvalov @kychendev
 
 # ITT annotations
 llvm/lib/Transforms/Instrumentation/SPIRITTAnnotations.cpp @MrSidims @vzakhari

.github/workflows/gh_pages.yml

@@ -16,7 +16,7 @@ jobs:
     - name: Install deps
       run: |
         sudo apt-get install -y doxygen graphviz ssh ninja-build
-        sudo pip3 install 'sphinx==3.0.0' 'recommonmark==0.6.0' 'sphinx_markdown_tables==0.0.12'
+        sudo pip3 install 'sphinx==4.1.2' 'myst-parser==0.15.1'
     - name: Build Docs
       run: |
         mkdir -p $GITHUB_WORKSPACE/build

libclc/ptx-nvidiacl/libspirv/group/collectives.cl

@@ -211,17 +211,17 @@ __CLC_SUBGROUP_COLLECTIVE(FAdd, __CLC_ADD, half, 0)
 __CLC_SUBGROUP_COLLECTIVE(FAdd, __CLC_ADD, float, 0)
 __CLC_SUBGROUP_COLLECTIVE(FAdd, __CLC_ADD, double, 0)
 
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, char, 0)
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, uchar, 0)
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, short, 0)
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, ushort, 0)
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, int, 0)
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, uint, 0)
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, long, 0)
-__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, ulong, 0)
-__CLC_SUBGROUP_COLLECTIVE(FMul, __CLC_MUL, half, 0)
-__CLC_SUBGROUP_COLLECTIVE(FMul, __CLC_MUL, float, 0)
-__CLC_SUBGROUP_COLLECTIVE(FMul, __CLC_MUL, double, 0)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, char, 1)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, uchar, 1)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, short, 1)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, ushort, 1)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, int, 1)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, uint, 1)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, long, 1)
+__CLC_SUBGROUP_COLLECTIVE(IMul, __CLC_MUL, ulong, 1)
+__CLC_SUBGROUP_COLLECTIVE(FMul, __CLC_MUL, half, 1)
+__CLC_SUBGROUP_COLLECTIVE(FMul, __CLC_MUL, float, 1)
+__CLC_SUBGROUP_COLLECTIVE(FMul, __CLC_MUL, double, 1)
 
 __CLC_SUBGROUP_COLLECTIVE(SMin, __CLC_MIN, char, CHAR_MAX)
 __CLC_SUBGROUP_COLLECTIVE(UMin, __CLC_MIN, uchar, UCHAR_MAX)
@@ -332,17 +332,17 @@ __CLC_GROUP_COLLECTIVE(FAdd, __CLC_ADD, float, 0)
 __CLC_GROUP_COLLECTIVE(FAdd, __CLC_ADD, double, 0)
 
 // There is no Mul group op in SPIR-V, use non-uniform variant instead.
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, char, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, uchar, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, short, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, ushort, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, int, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, uint, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, long, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, ulong, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformFMul, FMul, __CLC_MUL, half, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformFMul, FMul, __CLC_MUL, float, 0)
-__CLC_GROUP_COLLECTIVE(NonUniformFMul, FMul, __CLC_MUL, double, 0)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, char, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, uchar, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, short, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, ushort, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, int, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, uint, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, long, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformIMul, IMul, __CLC_MUL, ulong, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformFMul, FMul, __CLC_MUL, half, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformFMul, FMul, __CLC_MUL, float, 1)
+__CLC_GROUP_COLLECTIVE(NonUniformFMul, FMul, __CLC_MUL, double, 1)
 
 __CLC_GROUP_COLLECTIVE(SMin, __CLC_MIN, char, CHAR_MAX)
 __CLC_GROUP_COLLECTIVE(UMin, __CLC_MIN, uchar, UCHAR_MAX)

sycl/doc/PreprocessorMacros.md

@@ -2,38 +2,38 @@
 
 This file describes macros that have effect on SYCL compiler and run-time.
 
-### `RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR`
+- **RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR**
 
-The spec assumes that the SYCL implementation does address space deduction.
-However, for our implementation, the deduction is performed in the middle end,
-where it's hard to provide user friendly diagnositcs.
-Due to these problems writing to raw pointers obtained from `constant_ptr` is
-not diagnosed now.
-The user can enable diagnostics upon writing to such pointers via enabling the
-`RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR` macro.
-This allows `constant_ptr` to use constant pointers as underlying
-pointer types. Thus, conversions from `constant_ptr` to raw pointers will return
-constant pointers and writing to const pointers will be diagnosed by the
-front-end.
-This behavior is not following the SYCL spec since `constant_ptr` conversions to
-the underlying pointer types return pointers without any additional qualifiers
-so it's disabled by default.
+  The spec assumes that the SYCL implementation does address space deduction.
+  However, for our implementation, the deduction is performed in the middle end,
+  where it's hard to provide user friendly diagnositcs.
+  Due to these problems writing to raw pointers obtained from `constant_ptr` is
+  not diagnosed now.
+  The user can enable diagnostics upon writing to such pointers via enabling the
+  `RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR` macro.
+  This allows `constant_ptr` to use constant pointers as underlying
+  pointer types. Thus, conversions from `constant_ptr` to raw pointers will return
+  constant pointers and writing to const pointers will be diagnosed by the
+  front-end.
+  This behavior is not following the SYCL spec since `constant_ptr` conversions to
+  the underlying pointer types return pointers without any additional qualifiers
+  so it's disabled by default.
 
-### `DISABLE_SYCL_INSTRUMENTATION_METADATA`
+- **DISABLE_SYCL_INSTRUMENTATION_METADATA**
 
-This macro is used to disable passing of code location information to public
-methods.
+  This macro is used to disable passing of code location information to public
+  methods.
 
-### `SYCL2020_DISABLE_DEPRECATION_WARNINGS`
+- **SYCL2020_DISABLE_DEPRECATION_WARNINGS**
 
-Disables warnings coming from usage of SYCL 1.2.1 APIs, that are deprecated in
-SYCL 2020.
+  Disables warnings coming from usage of SYCL 1.2.1 APIs, that are deprecated in
+  SYCL 2020.
 
-### `SYCL_DISABLE_DEPRECATION_WARNINGS`
+- **SYCL_DISABLE_DEPRECATION_WARNINGS**
 
-Disables all deprecation warnings in SYCL runtime headers, including SYCL 1.2.1 deprecations.
+  Disables all deprecation warnings in SYCL runtime headers, including SYCL 1.2.1 deprecations.
 
-### Version macros
+## Version macros
 
 - `__LIBSYCL_MAJOR_VERSION` is set to SYCL runtime library major version.
 - `__LIBSYCL_MINOR_VERSION` is set to SYCL runtime library minor version.

sycl/doc/conf.py

@@ -32,8 +32,7 @@
 # extensions coming with Sphinx (named 'sphinx.ext.*') or your custom
 # ones.
 extensions = [
-    'recommonmark',
-    'sphinx_markdown_tables'
+    'myst_parser'
 ]
 
 # The name of the Pygments (syntax highlighting) style to use.
@@ -49,13 +48,11 @@
 # Extensions are mostly in asciidoc which has poor support in Sphinx
 exclude_patterns = ['extensions/*']
 
-source_parsers = {'.md': 'recommonmark.parser.CommonMarkParser'}
-
 suppress_warnings = [ 'misc.highlighting_failure' ]
 
 def on_missing_reference(app, env, node, contnode):
     if node['reftype'] == 'any':
-        contnode['refuri'] = "https://github.com/intel/llvm/tree/sycl/sycl/doc/" + contnode['refuri']
+        contnode['refuri'] = "https://github.com/intel/llvm/tree/sycl/sycl/doc/" + node['reftarget']
         return contnode
     else:
         return None

sycl/include/CL/sycl/buffer.hpp

@@ -263,7 +263,10 @@ class buffer {
   size_t get_count() const { return size(); }
   size_t size() const noexcept { return Range.size(); }
 
-  size_t get_size() const { return size() * sizeof(T); }
+  __SYCL2020_DEPRECATED(
+      "get_size() is deprecated, please use byte_size() instead")
+  size_t get_size() const { return byte_size(); }
+  size_t byte_size() const noexcept { return size() * sizeof(T); }
 
   AllocatorT get_allocator() const {
     return impl->template get_allocator<AllocatorT>();
@@ -343,7 +346,7 @@ class buffer {
   template <typename ReinterpretT, int ReinterpretDim>
   buffer<ReinterpretT, ReinterpretDim, AllocatorT>
   reinterpret(range<ReinterpretDim> reinterpretRange) const {
-    if (sizeof(ReinterpretT) * reinterpretRange.size() != get_size())
+    if (sizeof(ReinterpretT) * reinterpretRange.size() != byte_size())
       throw cl::sycl::invalid_object_error(
           "Total size in bytes represented by the type and range of the "
           "reinterpreted SYCL buffer does not equal the total size in bytes "
@@ -369,7 +372,7 @@ class buffer {
                                 (sizeof(ReinterpretT) != sizeof(T))),
       buffer<ReinterpretT, ReinterpretDim, AllocatorT>>::type
   reinterpret() const {
-    long sz = get_size(); // TODO: switch to byte_size() once implemented
+    long sz = byte_size();
     if (sz % sizeof(ReinterpretT) != 0)
       throw cl::sycl::invalid_object_error(
           "Total byte size of buffer is not evenly divisible by the size of "

sycl/include/CL/sycl/group_algorithm.hpp

@@ -624,14 +624,15 @@ joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
                      const ptrdiff_t &divisor) -> ptrdiff_t {
     return ((v + divisor - 1) / divisor) * divisor;
   };
-  typename InPtr::element_type x;
-  typename OutPtr::element_type carry = init;
+  typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
+      x;
+  typename detail::remove_pointer<OutPtr>::type carry = init;
   for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
     ptrdiff_t i = chunk + offset;
     if (i < N) {
       x = first[i];
     }
-    typename OutPtr::element_type out =
+    typename detail::remove_pointer<OutPtr>::type out =
         exclusive_scan_over_group(g, x, carry, binary_op);
     if (i < N) {
       result[i] = out;
@@ -664,13 +665,15 @@ joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
   // FIXME: Do not special-case for half precision
   static_assert(
       std::is_same<decltype(binary_op(*first, *first)),
-                   typename OutPtr::element_type>::value ||
-          (std::is_same<typename OutPtr::element_type, half>::value &&
+                   typename detail::remove_pointer<OutPtr>::type>::value ||
+          (std::is_same<typename detail::remove_pointer<OutPtr>::type,
+                        half>::value &&
            std::is_same<decltype(binary_op(*first, *first)), float>::value),
       "Result type of binary_op must match scan accumulation type.");
   return joint_exclusive_scan(
       g, first, last, result,
-      sycl::known_identity_v<BinaryOperation, typename OutPtr::element_type>,
+      sycl::known_identity_v<BinaryOperation,
+                             typename detail::remove_pointer<OutPtr>::type>,
       binary_op);
 }
 
@@ -791,14 +794,15 @@ joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                      const ptrdiff_t &divisor) -> ptrdiff_t {
     return ((v + divisor - 1) / divisor) * divisor;
   };
-  typename InPtr::element_type x;
-  typename OutPtr::element_type carry = init;
+  typename std::remove_const<typename detail::remove_pointer<InPtr>::type>::type
+      x;
+  typename detail::remove_pointer<OutPtr>::type carry = init;
   for (ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
     ptrdiff_t i = chunk + offset;
     if (i < N) {
       x = first[i];
     }
-    typename OutPtr::element_type out =
+    typename detail::remove_pointer<OutPtr>::type out =
         inclusive_scan_over_group(g, x, binary_op, carry);
     if (i < N) {
       result[i] = out;
@@ -830,13 +834,15 @@ joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
   // FIXME: Do not special-case for half precision
   static_assert(
       std::is_same<decltype(binary_op(*first, *first)),
-                   typename OutPtr::element_type>::value ||
-          (std::is_same<typename OutPtr::element_type, half>::value &&
+                   typename detail::remove_pointer<OutPtr>::type>::value ||
+          (std::is_same<typename detail::remove_pointer<OutPtr>::type,
+                        half>::value &&
            std::is_same<decltype(binary_op(*first, *first)), float>::value),
       "Result type of binary_op must match scan accumulation type.");
   return joint_inclusive_scan(
       g, first, last, result, binary_op,
-      sycl::known_identity_v<BinaryOperation, typename OutPtr::element_type>);
+      sycl::known_identity_v<BinaryOperation,
+                             typename detail::remove_pointer<OutPtr>::type>);
 }
 
 namespace detail {

sycl/include/CL/sycl/types.hpp

@@ -810,7 +810,10 @@ template <typename Type, int NumElements> class vec {
   __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
   static constexpr size_t get_count() { return size(); }
   static constexpr size_t size() noexcept { return NumElements; }
-  static constexpr size_t get_size() { return sizeof(m_Data); }
+  __SYCL2020_DEPRECATED(
+      "get_size() is deprecated, please use byte_size() instead")
+  static constexpr size_t get_size() { return byte_size(); }
+  static constexpr size_t byte_size() { return sizeof(m_Data); }
 
   template <typename convertT,
             rounding_mode roundingMode = rounding_mode::automatic>
@@ -1408,7 +1411,15 @@ class SwizzleOp {
   __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
   size_t get_count() const { return size(); }
   size_t size() const noexcept { return getNumElements(); }
-  template <int Num = getNumElements()> size_t get_size() const {
+
+  template <int Num = getNumElements()>
+  __SYCL2020_DEPRECATED(
+      "get_size() is deprecated, please use byte_size() instead")
+  size_t get_size() const {
+    return byte_size<Num>();
+  }
+
+  template <int Num = getNumElements()> size_t byte_size() const noexcept {
     return sizeof(DataT) * (Num == 3 ? 4 : Num);
   }
 

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

@@ -57,8 +57,7 @@ template <typename T, int SizeY, int StrideY, int SizeX, int StrideX>
 using region2d_t = region_base<true, T, SizeY, StrideY, SizeX, StrideX>;
 
 // A region with a single element.
-template <bool Is2D, typename T, int StrideY, int StrideX>
-using region_base_1 = region_base<Is2D, T, 1, StrideY, 1, StrideX>;
+template <typename T> using region_base_1 = region_base<false, T, 1, 0, 1, 0>;
 
 // simd_view forward declaration.
 template <typename BaseTy, typename RegionTy> class simd_view;

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

@@ -59,9 +59,9 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
       : M_base(Base), M_region(Region) {}
 
 public:
-  // Disallow copy and move constructors.
-  simd_view_impl(const simd_view_impl &Other) = delete;
-  simd_view_impl(simd_view_impl &&Other) = delete;
+  // Default copy and move constructors.
+  simd_view_impl(const simd_view_impl &Other) = default;
+  simd_view_impl(simd_view_impl &&Other) = default;
   /// @}
 
   /// Conversion to simd type.
@@ -80,6 +80,11 @@ template <typename BaseTy, typename RegionTy> class simd_view_impl {
   simd_view_impl &operator=(const value_type &Val) { return write(Val); }
   /// @}
 
+  /// Move assignment operator.
+  simd_view_impl &operator=(simd_view_impl &&Other) {
+    return write(Other.read());
+  }
+
   /// @{
   /// Region accessors.
   static constexpr bool is1D() { return !ShapeTy::Is_2D; }

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

@@ -139,7 +139,7 @@ template <typename Ty, int N> class simd {
   }
 
   /// View this simd object in a different element type.
-  template <typename EltTy> auto bit_cast_view() & {
+  template <typename EltTy> auto bit_cast_view() &[[clang::lifetimebound]] {
     using TopRegionTy = detail::compute_format_type_t<simd, EltTy>;
     using RetTy = simd_view<simd, TopRegionTy>;
     TopRegionTy R(0);
@@ -153,7 +153,8 @@ template <typename Ty, int N> class simd {
   }
 
   /// View as a 2-dimensional simd_view.
-  template <typename EltTy, int Height, int Width> auto bit_cast_view() & {
+  template <typename EltTy, int Height, int Width>
+  auto bit_cast_view() &[[clang::lifetimebound]] {
     using TopRegionTy =
         detail::compute_format_type_2d_t<simd, EltTy, Height, Width>;
     using RetTy = simd_view<simd, TopRegionTy>;
@@ -174,7 +175,8 @@ template <typename Ty, int N> class simd {
   /// \param Offset is the starting element offset.
   /// \return the representing region object.
   template <int Size, int Stride>
-  simd_view<simd, region1d_t<Ty, Size, Stride>> select(uint16_t Offset = 0) & {
+  simd_view<simd, region1d_t<Ty, Size, Stride>> select(uint16_t Offset = 0) &[
+      [clang::lifetimebound]] {
     region1d_t<Ty, Size, Stride> Reg(Offset);
     return {*this, Reg};
   }
@@ -200,7 +202,8 @@ template <typename Ty, int N> class simd {
   Ty operator()(int i) const { return data()[i]; }
 
   /// Return writable view of a single element.
-  simd_view<simd, region1d_t<Ty, 1, 0>> operator[](int i) {
+  simd_view<simd, region1d_t<Ty, 1, 0>> operator[](int i)
+      [[clang::lifetimebound]] {
     return select<1, 0>(i);
   }