[SYCL] Forbid non-trivially-copyable type in annotated_ptr/ref class (i…

…ntel#11798)

1. Explicitly state that use non-trivially-copyable type in
annotated_ptr and annotated_ref is not supported and error out if
annotated_ptr is constructed with non-trivially-copyable type

2. `void` is allowed on annotated_ptr as an exception

3. Use pass-by-value in annotated_ref operator= now, and updated the
spec
this fix volatile type not supported in the following code
```
void func(const T& a) { ... }
volatile T b;
func(b); // this will error out
```

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/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/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);
 }