[SYCL][InvokeSimd] Add error for callee with struct return or argument

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{{.*}}
+}