[SYCL][InvokeSIMD] Add error for invalid uniform arguments (#8916)

The spec requires uniform arguments to contain a trivially copyable
type. Add an error if it isn't.

---------

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

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

@@ -187,11 +187,13 @@ template <class... SpmdArgs> struct all_uniform_types {
 // - the case when there is nothing to unwrap
 template <typename T> struct unwrap_uniform {
   static auto impl(T val) { return val; }
+  using type = T;
 };
 
 // - the real unwrapping case
 template <typename T> struct unwrap_uniform<uniform<T>> {
   static T impl(uniform<T> val) { return val; }
+  using type = T;
 };
 
 // Verify the callee return type matches the subgroup size as is required by the
@@ -361,6 +363,20 @@ template <typename T>
 using strip_regcall_from_function_ptr_t =
     typename strip_regcall_from_function_ptr<T>::type;
 
+template <typename T> struct is_non_trivially_copyable_uniform {
+  static constexpr bool value =
+      is_uniform_type<T>::value &&
+      !std::is_trivially_copyable_v<typename unwrap_uniform<T>::type>;
+};
+
+template <> struct is_non_trivially_copyable_uniform<void> {
+  static constexpr bool value = false;
+};
+
+template <typename T>
+inline constexpr bool is_non_trivially_copyable_uniform_v =
+    is_non_trivially_copyable_uniform<T>::value;
+
 template <typename Ret, typename... Args>
 constexpr bool has_ref_arg(Ret (*)(Args...)) {
   return (... || std::is_reference_v<Args>);
@@ -371,7 +387,12 @@ constexpr bool has_ref_ret(Ret (*)(Args...)) {
   return std::is_reference_v<Ret>;
 }
 
-template <class Callable> constexpr void verify_no_ref() {
+template <typename Ret, typename... Args>
+constexpr bool has_non_trivially_copyable_uniform_ret(Ret (*)(Args...)) {
+  return is_non_trivially_copyable_uniform_v<Ret>;
+}
+
+template <class Callable> constexpr void verify_callable() {
   if constexpr (is_function_ptr_or_ref_v<Callable>) {
     using RemoveRef =
         remove_ref_from_func_ptr_ref_type_t<std::remove_reference_t<Callable>>;
@@ -388,9 +409,33 @@ template <class Callable> constexpr void verify_no_ref() {
     static_assert(
         !callable_has_ref_arg,
         "invoke_simd does not support callables with reference arguments");
+#ifdef __SYCL_DEVICE_ONLY__
+    constexpr bool callable_has_uniform_non_trivially_copyable_ret =
+        has_non_trivially_copyable_uniform_ret(obj);
+    static_assert(!callable_has_uniform_non_trivially_copyable_ret,
+                  "invoke_simd does not support callables returning uniforms "
+                  "that are not trivially copyable");
+#endif
   }
 }
 
+template <class... Ts>
+constexpr void verify_no_uniform_non_trivially_copyable_args() {
+#ifdef __SYCL_DEVICE_ONLY__
+  constexpr bool has_non_trivially_copyable_uniform_arg =
+      (... || is_non_trivially_copyable_uniform_v<Ts>);
+  static_assert(!has_non_trivially_copyable_uniform_arg,
+                "Uniform arguments must be trivially copyable");
+#endif
+}
+
+template <class Callable, class... Ts>
+constexpr void verify_valid_args_and_ret() {
+  verify_no_uniform_non_trivially_copyable_args<Ts...>();
+
+  verify_callable<Callable>();
+}
+
 } // namespace detail
 
 // --- The main API
@@ -420,7 +465,7 @@ __attribute__((always_inline)) auto invoke_simd(sycl::sub_group sg,
   // what the subgroup size is and arguments don't need widening and return
   // value does not need shrinking by this library or SPMD compiler, so 0
   // is fine in this case.
-  detail::verify_no_ref<Callable>();
+  detail::verify_valid_args_and_ret<Callable, T...>();
   constexpr int N = detail::get_sg_size<Callable, T...>();
   using RetSpmd = detail::SpmdRetType<N, Callable, T...>;
   detail::verify_return_type_matches_sg_size<

sycl/test/invoke_simd/not-trivially-copyable-uniform.cpp

@@ -0,0 +1,46 @@
+// RUN: not %clangxx -fsycl -fsycl-device-only -Xclang -fsycl-allow-func-ptr -S %s -o /dev/null 2>&1 | FileCheck -check-prefix CHECK-ARG %s
+// RUN: not %clangxx -fsycl -fsycl-device-only -Xclang -fsycl-allow-func-ptr -DRET -S %s -o /dev/null 2>&1 | FileCheck -check-prefix CHECK-RET %s
+
+#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
+#include <sycl/sycl.hpp>
+
+using namespace sycl::ext::oneapi::experimental;
+using namespace sycl;
+namespace esimd = sycl::ext::intel::esimd;
+struct B {
+  virtual ~B() {}
+};
+struct D : public B {
+  ~D() override {}
+};
+
+#ifdef RET
+[[intel::device_indirectly_callable]] uniform<D> callee() {}
+#else
+[[intel::device_indirectly_callable]] void callee(D d) {}
+#endif
+
+void foo() {
+  constexpr unsigned Size = 1024;
+  constexpr unsigned GroupSize = 64;
+  sycl::range<1> GlobalRange{Size};
+  sycl::range<1> LocalRange{GroupSize};
+  sycl::nd_range<1> Range(GlobalRange, LocalRange);
+  queue q;
+  auto e = q.submit([&](handler &cgh) {
+    cgh.parallel_for(Range, [=](nd_item<1> ndi) {
+#ifdef RET
+      invoke_simd(ndi.get_sub_group(), callee);
+#else
+      D d;
+      invoke_simd(ndi.get_sub_group(), callee, uniform{d});
+#endif
+    });
+  });
+}
+
+int main() {
+  foo();
+  // CHECK-ARG: {{.*}}error:{{.*}}static assertion failed due to requirement '!has_non_trivially_copyable_uniform_arg': Uniform arguments must be trivially copyable
+  // CHECK-RET: {{.*}}error:{{.*}}static assertion failed due to requirement '!callable_has_uniform_non_trivially_copyable_ret': invoke_simd does not support callables returning uniforms that are not trivially copyable
+}