Ensure launched kernels are fully inlined

Author: joeatodd

sycl/include/syclcompat/launch_policy.hpp

@@ -192,6 +192,18 @@ launch_policy(dim3, dim3, Ts...) -> launch_policy<
     detail::has_type<local_mem_size, std::tuple<Ts...>>::value>;
 
 namespace detail {
+// Custom std::apply helpers to enable inlining
+template <class F, class Tuple, size_t... Is>
+__syclcompat_inline__ constexpr void
+apply_expand(F f, Tuple t, std::index_sequence<Is...>) {
+    [[clang::always_inline]] f(get<Is>(t)...);
+}
+
+template <class F, class Tuple>
+__syclcompat_inline__ constexpr void
+apply_helper(F f, Tuple t) {
+    apply_expand(f, t, std::make_index_sequence<std::tuple_size<Tuple>{}>{});
+}
 
 template <auto F, typename Range, typename KProps, bool HasLocalMem,
           typename... Args>
@@ -212,11 +224,11 @@ struct KernelFunctor {
     if constexpr (HasLocalMem) {
       char *local_mem_ptr = static_cast<char *>(
           _local_acc.template get_multi_ptr<sycl::access::decorated::no>().get());
-      std::apply(
-          [lmem_ptr = local_mem_ptr](auto &&...args) { F(args..., lmem_ptr); },
-          _argument_tuple);
+      apply_helper(
+           [lmem_ptr = local_mem_ptr](auto &&...args) { [[clang::always_inline]] F(args..., lmem_ptr); },
+           _argument_tuple);
     } else {
-      std::apply([](auto &&...args) { F(args...); }, _argument_tuple);
+      apply_helper([](auto &&...args) { [[clang::always_inline]] F(args...); }, _argument_tuple);
     }
   }