[SYCL] Aligned set_arg behaviour with SYCL specification (#2159)

According to SYCL 1.2.1 specification, rev. 7, paragraph 4.8.5:

template <typename T>
void set_arg(int argIndex, T &&arg)
...
The argument can be either a SYCL accessor, a SYCL sampler or
a trivially copyable and standard-layout C++ type.

Signed-off-by: Ivan Karachun <[email protected]>

Author: Ivan Karachun

sycl/include/CL/sycl/handler.hpp

@@ -796,13 +796,42 @@ class __SYCL_EXPORT handler {
     }
   }
 
+  template <typename T>
+  using remove_cv_ref_t =
+      typename std::remove_cv<detail::remove_reference_t<T>>::type;
+
+  template <typename U, typename T>
+  using is_same_type = std::is_same<remove_cv_ref_t<U>, remove_cv_ref_t<T>>;
+
+  template <typename T> struct ShouldEnableSetArg {
+    static constexpr bool value =
+        std::is_trivially_copyable<T>::value
+#if CL_SYCL_LANGUAGE_VERSION && CL_SYCL_LANGUAGE_VERSION <= 121
+            && std::is_standard_layout<T>::value
+#endif
+        || is_same_type<sampler, T>::value // Sampler
+        || (!is_same_type<cl_mem, T>::value &&
+            std::is_pointer<remove_cv_ref_t<T>>::value) // USM
+        || is_same_type<cl_mem, T>::value;              // Interop
+  };
+
   /// Sets argument for OpenCL interoperability kernels.
   ///
   /// Registers Arg passed as argument # ArgIndex.
   ///
   /// \param ArgIndex is a positional number of argument to be set.
   /// \param Arg is an argument value to be set.
-  template <typename T> void set_arg(int ArgIndex, T &&Arg) {
+  template <typename T>
+  typename std::enable_if<ShouldEnableSetArg<T>::value, void>::type
+  set_arg(int ArgIndex, T &&Arg) {
+    setArgHelper(ArgIndex, std::move(Arg));
+  }
+
+  template <typename DataT, int Dims, access::mode AccessMode,
+            access::target AccessTarget, access::placeholder IsPlaceholder>
+  void
+  set_arg(int ArgIndex,
+          accessor<DataT, Dims, AccessMode, AccessTarget, IsPlaceholder> Arg) {
     setArgHelper(ArgIndex, std::move(Arg));
   }
 

sycl/test/basic_tests/set_arg_error.cpp

@@ -0,0 +1,47 @@
+// RUN: %clangxx -fsycl -Xclang -verify %s -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning -fsyntax-only
+
+#include <CL/sycl.hpp>
+
+struct TriviallyCopyable {
+  int a;
+  int b;
+};
+
+struct NonTriviallyCopyable {
+  NonTriviallyCopyable() = default;
+  NonTriviallyCopyable(NonTriviallyCopyable const &) {}
+  int a;
+  int b;
+};
+
+struct NonStdLayout {
+  int a;
+
+private:
+  int b;
+};
+
+int main() {
+  constexpr size_t size = 1;
+  cl::sycl::buffer<int> buf(size);
+  cl::sycl::queue q;
+  q.submit([&](cl::sycl::handler &h) {
+    auto global_acc = buf.get_access<cl::sycl::access::mode::write>(h);
+    cl::sycl::sampler samp(cl::sycl::coordinate_normalization_mode::normalized,
+                           cl::sycl::addressing_mode::clamp,
+                           cl::sycl::filtering_mode::nearest);
+    cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                       cl::sycl::access::target::local>
+        local_acc({size}, h);
+    h.set_arg(0, local_acc);
+    h.set_arg(1, global_acc);
+    h.set_arg(2, samp);
+    h.set_arg(3, TriviallyCopyable{});
+    h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}}
+        4, NonTriviallyCopyable{});
+#if CL_SYCL_LANGUAGE_VERSION && CL_SYCL_LANGUAGE_VERSION <= 121
+    h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}}
+        5, NonStdLayout{});
+#endif
+  });
+}