review feedback 2

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

Author: sarnex

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

@@ -363,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>);
@@ -373,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>>;
@@ -390,23 +409,27 @@ template <class Callable> constexpr void verify_no_ref() {
     static_assert(
         !callable_has_ref_arg,
         "invoke_simd does not support callables with reference arguments");
+    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");
   }
 }
 
 template <class... Ts>
 constexpr void verify_no_uniform_non_trivially_copyable_args() {
   constexpr bool has_non_trivially_copyable_uniform_arg =
-      (... ||
-       (is_uniform_type<Ts>::value &&
-        !std::is_trivially_copyable_v<typename unwrap_uniform<Ts>::type>));
+      (... || is_non_trivially_copyable_uniform_v<Ts>);
   static_assert(!has_non_trivially_copyable_uniform_arg,
                 "Uniform arguments must be trivially copyable");
 }
 
-template <class Callable, class... Ts> constexpr void verify_valid_args() {
+template <class Callable, class... Ts>
+constexpr void verify_valid_args_and_ret() {
   verify_no_uniform_non_trivially_copyable_args<Ts...>();
 
-  verify_no_ref<Callable>();
+  verify_callable<Callable>();
 }
 
 } // namespace detail
@@ -438,7 +461,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_valid_args<Callable, T...>();
+  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-arg.cpp renamed to sycl/test/invoke_simd/not-trivially-copyable-uniform.cpp

@@ -1,4 +1,6 @@
-// RUN: not %clangxx -fsycl -fsycl-device-only -Xclang -fsycl-allow-func-ptr -S %s -o /dev/null 2>&1 | FileCheck %s
+// 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>
 
@@ -12,7 +14,11 @@ 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;
@@ -23,13 +29,18 @@ void foo() {
   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: {{.*}}error:{{.*}}static assertion failed due to requirement '!has_non_trivially_copyable_uniform_arg': Uniform arguments must be trivially copyable
+  // 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
 }