[SYCL][InvokeSimd] Add error for callee with struct return or argument (#10121)

It doesn't work at runtime even if we make it through the compiler
because of an IGC issue.

Once the IGC issue is fixed, we will need make some changes to the
compiler and then lift this restriction.

---------

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

Author: sarnex

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

@@ -364,6 +364,16 @@ constexpr bool has_ref_ret(Ret (*)(Args...)) {
   return std::is_reference_v<Ret>;
 }
 
+template <typename Ret, typename... Args>
+constexpr bool has_struct_arg(Ret (*)(Args...)) {
+  return (... || (std::is_class_v<Args> && !is_simd_or_mask_type<Args>::value));
+}
+
+template <typename Ret, typename... Args>
+constexpr bool has_struct_ret(Ret (*)(Args...)) {
+  return std::is_class_v<Ret> && !is_simd_or_mask_type<Ret>::value;
+}
+
 template <typename Ret, typename... Args>
 constexpr bool has_non_trivially_copyable_uniform_ret(Ret (*)(Args...)) {
   return is_non_trivially_copyable_uniform_v<Ret>;
@@ -386,6 +396,16 @@ template <class Callable> constexpr void verify_callable() {
     static_assert(
         !callable_has_ref_arg,
         "invoke_simd does not support callables with reference arguments");
+#ifndef __INVOKE_SIMD_ENABLE_STRUCTS
+    constexpr bool callable_has_struct_ret = has_struct_ret(obj);
+    static_assert(
+        !callable_has_struct_ret,
+        "invoke_simd does not support callables returning structures");
+    constexpr bool callable_has_struct_arg = has_struct_arg(obj);
+    static_assert(
+        !callable_has_struct_arg,
+        "invoke_simd does not support callables with structure arguments");
+#endif
 #ifdef __SYCL_DEVICE_ONLY__
     constexpr bool callable_has_uniform_non_trivially_copyable_ret =
         has_non_trivially_copyable_uniform_ret(obj);

sycl/test/invoke_simd/return-type-struct.cpp

@@ -0,0 +1,31 @@
+// RUN: not %clangxx -fsycl -fsycl-device-only -S %s -o /dev/null 2>&1 | FileCheck %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;
+
+[[intel::device_indirectly_callable]] std::tuple<simd<int, 4>, simd<int, 4>>
+callee(simd<int, 8>) {
+  return std::make_tuple(simd<int, 4>(), simd<int, 4>());
+}
+
+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) {
+      auto x = invoke_simd(ndi.get_sub_group(), callee, 0);
+    });
+  });
+}
+
+int main() {
+  foo();
+  // CHECK: {{.*}}error:{{.*}}static assertion failed due to requirement '!callable_has_struct_ret': invoke_simd does not support callables returning structures{{.*}}
+}

sycl/test/invoke_simd/struct-arg.cpp

@@ -0,0 +1,32 @@
+// RUN: not %clangxx -fsycl -fsycl-device-only -S %s -o /dev/null 2>&1 | FileCheck %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;
+
+[[intel::device_indirectly_callable]] void
+callee(simd<int, 8>, std::tuple<simd<int, 4>, simd<int, 4>>) {
+  return std::make_tuple(simd<int, 4>(), simd<int, 4>());
+}
+
+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) {
+      std::tuple<simd<int, 4>, simd<int, 4>> arg;
+      invoke_simd(ndi.get_sub_group(), callee, 0, arg);
+    });
+  });
+}
+
+int main() {
+  foo();
+  // CHECK: {{.*}}error:{{.*}}static assertion failed due to requirement '!callable_has_struct_arg': invoke_simd does not support callables with structure arguments{{.*}}
+}