Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/include/clang/Basic/DiagnosticGroups.td

@@ -780,9 +780,11 @@ def SyclIvdepAttribute : DiagGroup<"ivdep-compat">;
 def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">;
 def UnknownAttributes : DiagGroup<"unknown-attributes">;
 def IgnoredAttributes : DiagGroup<"ignored-attributes", [SyclIvdepAttribute]>;
+def IncorrectSubGroupSize: DiagGroup<"incorrect-sub-group-size">;
 def AcceptedAttributes : DiagGroup<"accepted-attributes">;
 def Attributes : DiagGroup<"attributes", [UnknownAttributes,
                                           IgnoredAttributes,
+                                          IncorrectSubGroupSize,
                                           AcceptedAttributes]>;
 def UnknownSanitizers : DiagGroup<"unknown-sanitizers">;
 def UnnamedTypeTemplateArgs : DiagGroup<"unnamed-type-template-args",

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -3543,7 +3543,7 @@ def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed"
 def warn_reqd_sub_group_attribute_n
     : Warning<"attribute argument %0 is invalid and will be ignored; %1 "
               "requires sub_group size %2">,
-  InGroup<IgnoredAttributes>;
+  InGroup<IncorrectSubGroupSize>;
 def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with"
   " exception specification; attribute ignored">,
   InGroup<IgnoredAttributes>;

clang/test/SemaSYCL/reqd-sub-group-size-silence-warnings.cpp

@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple nvptx -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s
+// RUN: %clang_cc1 -fsycl-is-device -triple nvptx -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s
+//
+// Sub group size of 8 is incompatible with both CUDA and HIP, expect it to be
+// silenced. Check both the dedicated switch '-Wno-incorrect-sub-group-size' and
+// the catch all '-Wno-attributes'.
+#include "sycl.hpp"
+
+
+// expected-no-diagnostics
+int main() {
+
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {});
+  });
+
+  return 0;
+}

sycl/doc/extensions/experimental/sycl_ext_oneapi_annotated_ptr.asciidoc

@@ -183,6 +183,31 @@ information provided as compile-time constant properties through all uses of the
 pointer unless noted otherwise.
 
 .Unsupported Usage Example
+[source,c++]
+----
+using sycl::ext::oneapi::experimental;
+struct MyType {
+  MyType(const MyType& otherM) {
+     ...
+  }
+};
+
+struct MyKernel {
+  annotated_ptr<MyType, properties<PropA>> a; //non-trivially-copyable type is not allowed
+  annotated_ptr<void, properties<PropB>> b;   //void is legal 
+  ...
+  void operator()() const {
+    ...
+    *b = ...; //deference is not allowed
+  }
+};
+----
+It is ill-formed to instantiate`annotated_ptr` and `annotated_ref` for a non-trivially-copyable
+type `T`. The only exception is that `annotated_ptr` is allowed to be instantiated with `void`.
+
+In the above example, encapsulating `annotated_ptr` within `MyType` is illegal. Encapsulating
+`annotated_ptr` with `void` is legal, but it can not be deferenced.
+
 [source,c++]
 ----
 using sycl::ext::oneapi::experimental;
@@ -524,10 +549,10 @@ class annotated_ref {
   public:
     annotated_ref(const annotated_ref&) = delete;
     operator T() const;
-    T operator=(const T &) const;
+    T operator=(T) const;
     T operator=(const annotated_ref&) const;
     // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
-    T operatorOP(const T &) const;
+    T operatorOP(T) const;
     T operator++() const;
     T operator++(int) const;
     T operator--() const;
@@ -556,7 +581,7 @@ annotations when the object is loaded from memory.
 a|
 [source,c++]
 ----
-T operator=(const T &) const;
+T operator=(T) const;
 ----
 |
 Writes an object of type `T` to the location referenced by this wrapper,
@@ -583,7 +608,7 @@ Does not rebind the reference!
 a|
 [source,c++]
 ----
-T operatorOP(const T &) const;
+T operatorOP(T) const;
 ----
 a|
 Where [code]#OP# is: [code]#pass:[+=]#, [code]#-=#,[code]#*=#, [code]#/=#, [code]#%=#, [code]#+<<=+#, [code]#>>=#, [code]#&=#, [code]#\|=#, [code]#^=#.

sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp

@@ -32,7 +32,7 @@ namespace experimental {
 
 namespace {
 #define PROPAGATE_OP(op)                                                       \
-  T operator op##=(const T &rhs) const {                                       \
+  T operator op##=(T rhs) const {                                              \
     T t = *this;                                                               \
     t op## = rhs;                                                              \
     *this = t;                                                                 \
@@ -78,6 +78,10 @@ template <typename T, typename... Props>
 class annotated_ref<T, detail::properties_t<Props...>> {
   using property_list_t = detail::properties_t<Props...>;
 
+  static_assert(
+      std::is_trivially_copyable_v<T>,
+      "annotated_ref can only encapsulate a trivially-copyable type!");
+
 private:
   T *m_Ptr;
   annotated_ref(T *Ptr) : m_Ptr(Ptr) {}
@@ -95,7 +99,7 @@ class annotated_ref<T, detail::properties_t<Props...>> {
 #endif
   }
 
-  T operator=(const T &Obj) const {
+  T operator=(T Obj) const {
 #ifdef __SYCL_DEVICE_ONLY__
     *__builtin_intel_sycl_ptr_annotation(
         m_Ptr, detail::PropertyMetaInfo<Props>::name...,
@@ -178,6 +182,11 @@ class __SYCL_SPECIAL_CLASS
 __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
   using property_list_t = detail::properties_t<Props...>;
 
+  static_assert(std::is_same_v<T, void> || std::is_trivially_copyable_v<T>,
+                "annotated_ptr can only encapsulate either "
+                "a trivially-copyable type "
+                "or void!");
+
   // buffer_location and alignment are allowed for annotated_ref
   // Cache controls are allowed for annotated_ptr
   using allowed_properties =
@@ -249,6 +258,7 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
   template <typename... PropertyValueTs>
   explicit annotated_ptr(T *Ptr, const PropertyValueTs &...props) noexcept
       : m_Ptr(Ptr) {
+
     static constexpr bool has_same_properties = std::is_same<
         property_list_t,
         detail::merged_properties_t<property_list_t,

sycl/include/sycl/types.hpp

@@ -1544,8 +1544,37 @@ template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 class SwizzleOp {
   using DataT = typename VecT::element_type;
-  using CommonDataT = std::common_type_t<typename OperationLeftT::DataT,
-                                         typename OperationRightT::DataT>;
+  // Certain operators return a vector with a different element type. Also, the
+  // left and right operand types may differ. CommonDataT selects a result type
+  // based on these types to ensure that the result value can be represented.
+  //
+  // Example 1:
+  //   sycl::vec<unsigned char, 4> vec{...};
+  //   auto result = 300u + vec.x();
+  //
+  // CommonDataT is std::common_type_t<OperationLeftT, OperationRightT> since
+  // it's larger than unsigned char.
+  //
+  // Example 2:
+  //   sycl::vec<bool, 1> vec{...};
+  //   auto result = vec.template swizzle<sycl::elem::s0>() && vec;
+  //
+  // CommonDataT is DataT since operator&& returns a vector with element type
+  // int8_t, which is larger than bool.
+  //
+  // Example 3:
+  //   sycl::vec<std::byte, 4> vec{...}; auto swlo = vec.lo();
+  //   auto result = swlo == swlo;
+  //
+  // CommonDataT is DataT since operator== returns a vector with element type
+  // int8_t, which is the same size as std::byte. std::common_type_t<DataT, ...>
+  // can't be used here since there's no type that int8_t and std::byte can both
+  // be implicitly converted to.
+  using OpLeftDataT = typename OperationLeftT::DataT;
+  using OpRightDataT = typename OperationRightT::DataT;
+  using CommonDataT = std::conditional_t<
+      sizeof(DataT) >= sizeof(std::common_type_t<OpLeftDataT, OpRightDataT>),
+      DataT, std::common_type_t<OpLeftDataT, OpRightDataT>>;
   static constexpr int getNumElements() { return sizeof...(Indexes); }
 
   using rel_t = detail::rel_t<DataT>;

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -57,11 +57,11 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 712d79120e0a9293daf5741f9c0b5269ddd862ce
+  # commit b3de31844ca45d1815c54d85f248f13ad81b6755
   # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Fri Nov 24 12:34:51 2023 +0000
-  #     Set version to v0.8.0
-  set(UNIFIED_RUNTIME_TAG v0.8.0)
+  # Date:   Tue Nov 28 16:21:43 2023 +0000
+  #     Set version to v0.8.1
+  set(UNIFIED_RUNTIME_TAG v0.8.1)
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")

sycl/test-e2e/Annotated_arg_ptr/annotated_ptr.cpp

@@ -28,6 +28,17 @@ int main() {
 
   auto *d = malloc_shared<int>(4, Q);
   auto d_ptr = annotated_ptr{d};
+
+  auto *e1 = malloc_shared<int>(1, Q);
+  *e1 = 0;
+  volatile int *e1_vol = e1;
+  auto e1_ptr = annotated_ptr{e1_vol};
+
+  auto *e2 = malloc_shared<int>(1, Q);
+  *e2 = 5;
+  const volatile int *e2_vol = e2;
+  auto e2_ptr = annotated_ptr{e2_vol};
+
   for (int i = 0; i < 4; i++)
     d_ptr[i] = i;
   Q.single_task([=]() {
@@ -60,6 +71,8 @@ int main() {
      };
 
      d_ptr[3] = func(d_ptr[0], d_ptr[1], d_ptr[2]);
+
+     e1_ptr[0] = e2_ptr[0];
    }).wait();
 
   assert(a_ptr[0] == -1 && "a_ptr[0] value does not match.");
@@ -78,11 +91,15 @@ int main() {
 
   assert(d_ptr[3] == -1 && "d_ptr[3] value does not match.");
 
+  assert(e1_ptr[0] == 5 && "e_ptr[0] value does not match.");
+
   free(a, Q);
   free(b, Q);
   free(c->b, Q);
   free(c, Q);
   free(d, Q);
+  free(e1, Q);
+  free(e2, Q);
 
   return 0;
 }

sycl/test-e2e/Regression/vec_rel_swizzle_ops.cpp

@@ -0,0 +1,68 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// RUN: %if preview-breaking-changes-supported %{ %clangxx -fsycl -fpreview-breaking-changes %s -o %t2.out %}
+// RUN: %if preview-breaking-changes-supported %{ %{run} %t2.out %}
+
+#include <cstdlib>
+#include <sycl/sycl.hpp>
+
+template <typename T, typename ResultT>
+bool testAndOperator(const std::string &typeName) {
+  constexpr int N = 5;
+  std::array<ResultT, N> results{};
+
+  sycl::queue q;
+  sycl::buffer<ResultT, 1> buffer{results.data(), N};
+  q.submit([&](sycl::handler &cgh) {
+     sycl::accessor acc{buffer, cgh, sycl::write_only};
+     cgh.parallel_for(sycl::range<1>{1}, [=](sycl::id<1> id) {
+       auto testVec1 = sycl::vec<T, 1>(static_cast<T>(1));
+       auto testVec2 = sycl::vec<T, 1>(static_cast<T>(2));
+       sycl::vec<ResultT, 1> resVec;
+
+       ResultT expected = static_cast<ResultT>(
+           -(static_cast<ResultT>(1) && static_cast<ResultT>(2)));
+       acc[0] = expected;
+
+       // LHS swizzle
+       resVec = testVec1.template swizzle<sycl::elem::s0>() && testVec2;
+       acc[1] = resVec[0];
+
+       // RHS swizzle
+       resVec = testVec1 && testVec2.template swizzle<sycl::elem::s0>();
+       acc[2] = resVec[0];
+
+       // No swizzle
+       resVec = testVec1 && testVec2;
+       acc[3] = resVec[0];
+
+       // Both swizzle
+       resVec = testVec1.template swizzle<sycl::elem::s0>() &&
+                testVec2.template swizzle<sycl::elem::s0>();
+       acc[4] = resVec[0];
+     });
+   }).wait();
+
+  bool passed = true;
+  ResultT expected = results[0];
+
+  std::cout << "Testing with T = " << typeName << std::endl;
+  std::cout << "Expected: " << (int)expected << std::endl;
+  for (int i = 1; i < N; i++) {
+    std::cout << "Test " << (i - 1) << ": " << ((int)results[i]) << std::endl;
+    passed &= expected == results[i];
+  }
+  std::cout << std::endl;
+  return passed;
+}
+
+int main() {
+  bool passed = true;
+  passed &= testAndOperator<bool, std::int8_t>("bool");
+  passed &= testAndOperator<std::int8_t, std::int8_t>("std::int8_t");
+  passed &= testAndOperator<float, std::int32_t>("float");
+  passed &= testAndOperator<int, std::int32_t>("int");
+  std::cout << (passed ? "Pass" : "Fail") << std::endl;
+  return (passed ? EXIT_SUCCESS : EXIT_FAILURE);
+}

sycl/test/extensions/annotated_ptr/annotated_ptr.cpp

@@ -157,6 +157,18 @@ void TestVectorAddWithAnnotatedMMHosts() {
     std::cout << raw[i] << std::endl;
   }
 
+  class test {
+    int n;
+
+  public:
+    test(int n_) : n(n_) {}
+    test(const test &t) { n = t.n; }
+  };
+  // expected-error-re@sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp:* {{static assertion failed due to requirement {{.+}}: annotated_ptr can only encapsulate either a trivially-copyable type or void!}}
+  annotated_ptr<test> non_trivially_copyable;
+
+  annotated_ptr<void> void_type;
+
   free(raw, q);
 }