Formalize is_function_v w/a, extend invoke_simd test.

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Author: kbobrovs

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

@@ -226,24 +226,42 @@ simd_call_helper(const void *obj_ptr,
   return f(simd_args...);
 };
 
+#ifdef _GLIBCXX_RELEASE
+#if _GLIBCXX_RELEASE < 10
+#define __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
+#endif // _GLIBCXX_RELEASE < 10
+#endif // _GLIBCXX_RELEASE
+
+#ifdef __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
 // TODO This is a workaround for libstdc++ version 9 buggy behavior which
 // returns false in the code below. Version 10 works fine. Once required
 // minimum libstdc++ version is bumped to 10, this w/a should be removed.
 //   template <class F> bool foo(F &&f) {
 //     return std::is_function_v<std::remove_reference_t<F>>;
 //   }
 // where F is a function type with __regcall.
-template <class F> struct is_regcall_function : std::false_type {};
+template <class F> struct is_regcall_function_ptr_or_ref_v : std::false_type {};
 
 template <class Ret, class... Args>
-struct is_regcall_function<Ret(__regcall *)(Args...)> : std::true_type {};
+struct is_regcall_function_ptr_or_ref_v<Ret(__regcall &)(Args...)>
+    : std::true_type {};
 
 template <class Ret, class... Args>
-struct is_regcall_function<Ret(__regcall &)(Args...)> : std::true_type {};
+struct is_regcall_function_ptr_or_ref_v<Ret(__regcall *)(Args...)>
+    : std::true_type {};
 
 template <class F>
-static constexpr bool is_regcall_function_v = is_regcall_function<F>::value;
-
+static constexpr bool is_regcall_function_ptr_or_ref_v =
+    is_regcall_function_ptr_or_ref_v<F>::value;
+#endif // __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
+
+template <class Callable>
+static constexpr bool is_function_ptr_or_ref_v =
+    std::is_function_v<std::remove_pointer_t<std::remove_reference_t<Callable>>>
+#ifdef __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
+    || is_regcall_function_ptr_or_ref_v<Callable>
+#endif // __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
+    ;
 } // namespace detail
 
 // --- The main API
@@ -274,11 +292,7 @@ __attribute__((always_inline)) auto invoke_simd(sycl::sub_group sg,
   // is fine in this case.
   constexpr int N = detail::get_sg_size<Callable, T...>();
   using RetSpmd = detail::SpmdRetType<N, Callable, T...>;
-
-  using CallableNoRef = std::remove_reference_t<Callable>;
-  using CallableNoRefNoPtr = std::remove_pointer_t<CallableNoRef>;
-  constexpr bool is_function = std::is_function_v<CallableNoRefNoPtr> ||
-                               detail::is_regcall_function_v<Callable>;
+  constexpr bool is_function = detail::is_function_ptr_or_ref_v<Callable>;
 
   if constexpr (is_function) {
     return __builtin_invoke_simd<is_function, RetSpmd>(

sycl/test/invoke_simd/invoke_simd.cpp

@@ -260,3 +260,51 @@ SYCL_EXTERNAL auto barx(sub_group sg, float a, char ch,
   auto x = invoke_simd(sg, f, 1.f, uniform{a});
   static_assert(std::is_same_v<decltype(x), uniform<char>>);
 }
+
+// Internal is_function_ref_v meta-API checks {
+template <class F> void assert_is_func(F &&f) {
+  static_assert(
+      sycl::ext::oneapi::experimental::detail::is_function_ptr_or_ref_v<F>);
+}
+
+template <class F> void assert_is_not_func(F &&f) {
+  static_assert(
+      !sycl::ext::oneapi::experimental::detail::is_function_ptr_or_ref_v<F>);
+}
+
+void ordinary_func();
+
+// clang-format off
+void check_f(
+  int(*func_ptr)(float*), int(__regcall* func_ptr_regcall)(float*),
+  int(&func_ref)(float*), int(__regcall& func_ref_regcall)(float*),
+  int(func)(float*), int(__regcall func_regcall)(float*)) {
+
+  assert_is_func(SIMD_CALLEE);
+  assert_is_func(ordinary_func);
+
+  assert_is_func(func_ptr);
+  assert_is_func(func_ptr_regcall);
+
+  assert_is_func(func_ref);
+  assert_is_func(func_ref_regcall);
+
+  assert_is_func(func);
+  assert_is_func(func_regcall);
+}
+// clang-format on
+
+void check_not_f(char ch) {
+  assert_is_not_func(SIMD_FUNCTOR{10});
+  const auto capt_lambda = [=] [[gnu::regcall]] (simd<float, 16>, float) {
+    // capturing lambda
+    return ch;
+  };
+  const auto non_capt_lambda = [](simd<float, 16>, float) {
+    // non-capturing lambda
+    return 10;
+  };
+  assert_is_not_func(capt_lambda);
+  assert_is_not_func(non_capt_lambda);
+}
+// }