[SYCL] Free function traits and APIs (#13885)

This PR implements the traits and other APIs for free function support.

The free function spec is here:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc.

Later PRs will add the remaining parts of the spec. This is the current
status:

1. Free functions are supported at file scope only.
2. The SYCL_EXTERNAL markup is needed for free functions in addition to
the SYCL_EXT_ONEAPI_FUNCTION_PROPERTY property.
3. The compiler does not yet diagnose an error if the application
violates any of the restrictions listed in the specification under the
section "Defining a free function kernel".
4. Device code generation is supported for scalars and USM pointers
only. It is not supported for complex kernel argument types requiring
decomposition like accessor, local_accessor, or stream.
5. The implementation has not been tested to handle the case when a
kernel argument is optimized away. The switch
-fno-sycl-dead-args-optimization could be used to disable this
optimization, if needed
6. The kernel information descriptor info::kernel::num_args cannot yet
be used to query the number of arguments in a free function kernel.

Author: rdeodhar

clang/lib/Sema/SemaSYCL.cpp

@@ -1109,6 +1109,21 @@ static bool isFreeFunction(SemaSYCL &SemaSYCLRef, const FunctionDecl *FD) {
   return false;
 }
 
+static int getFreeFunctionRangeDim(SemaSYCL &SemaSYCLRef,
+                                   const FunctionDecl *FD) {
+  for (auto *IRAttr : FD->specific_attrs<SYCLAddIRAttributesFunctionAttr>()) {
+    SmallVector<std::pair<std::string, std::string>, 4> NameValuePairs =
+        IRAttr->getAttributeNameValuePairs(SemaSYCLRef.getASTContext());
+    for (const auto &NameValuePair : NameValuePairs) {
+      if (NameValuePair.first == "sycl-nd-range-kernel")
+        return std::stoi(NameValuePair.second);
+      if (NameValuePair.first == "sycl-single-task-kernel")
+        return 0;
+    }
+  }
+  return false;
+}
+
 // Creates a name for the free function kernel function.
 // Consider a free function named "MyFunction". The normal device function will
 // be given its mangled name, say "_Z10MyFunctionIiEvPT_S0_". The corresponding
@@ -2568,7 +2583,6 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler {
 
 // A type to Create and own the FunctionDecl for the kernel.
 class SyclKernelDeclCreator : public SyclKernelFieldHandler {
-  bool IsFreeFunction = false;
   FunctionDecl *KernelDecl = nullptr;
   llvm::SmallVector<ParmVarDecl *, 8> Params;
   Sema::ContextRAII FuncContext;
@@ -2788,9 +2802,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 public:
   static constexpr const bool VisitInsideSimpleContainers = false;
   SyclKernelDeclCreator(SemaSYCL &S, SourceLocation Loc, bool IsInline,
-                        bool IsSIMDKernel, bool IsFreeFunction,
-                        FunctionDecl *SYCLKernel)
-      : SyclKernelFieldHandler(S), IsFreeFunction(IsFreeFunction),
+                        bool IsSIMDKernel, FunctionDecl *SYCLKernel)
+      : SyclKernelFieldHandler(S),
         KernelDecl(
             createKernelDecl(S.getASTContext(), Loc, IsInline, IsSIMDKernel)),
         FuncContext(SemaSYCLRef.SemaRef, KernelDecl) {
@@ -5110,7 +5123,7 @@ void SemaSYCL::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
 
   SyclKernelDeclCreator kernel_decl(*this, KernelObj->getLocation(),
                                     KernelCallerFunc->isInlined(), IsSIMDKernel,
-                                    false /*IsFreeFunction*/, KernelCallerFunc);
+                                    KernelCallerFunc);
   SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj,
                                     KernelCallerFunc, IsSIMDKernel,
                                     CallOperator);
@@ -5152,7 +5165,7 @@ void ConstructFreeFunctionKernel(SemaSYCL &SemaSYCLRef, FunctionDecl *FD) {
                                             false /*IsSIMDKernel*/);
   SyclKernelDeclCreator kernel_decl(SemaSYCLRef, FD->getLocation(),
                                     FD->isInlined(), false /*IsSIMDKernel */,
-                                    true /*IsFreeFunction*/, FD);
+                                    FD);
 
   FreeFunctionKernelBodyCreator kernel_body(SemaSYCLRef, kernel_decl, FD);
 
@@ -6052,6 +6065,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
 
   O << "#include <sycl/detail/defines_elementary.hpp>\n";
   O << "#include <sycl/detail/kernel_desc.hpp>\n";
+  O << "#include <sycl/ext/oneapi/experimental/free_function_traits.hpp>\n";
 
   O << "\n";
 
@@ -6301,7 +6315,102 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "} // namespace detail\n";
   O << "} // namespace _V1\n";
   O << "} // namespace sycl\n";
-  O << "\n";
+
+  unsigned ShimCounter = 1;
+  int FreeFunctionCount = 0;
+  for (const KernelDesc &K : KernelDescs) {
+    if (!isFreeFunction(S, K.SyclKernel))
+      continue;
+
+    ++FreeFunctionCount;
+    // Generate forward declaration for free function.
+    O << "\n// Definition of " << K.Name << " as a free function kernel\n";
+    if (K.SyclKernel->getLanguageLinkage() == CLanguageLinkage)
+      O << "extern \"C\" ";
+    std::string ParmList;
+    bool FirstParam = true;
+    for (ParmVarDecl *Param : K.SyclKernel->parameters()) {
+      if (FirstParam)
+        FirstParam = false;
+      else
+        ParmList += ", ";
+      ParmList += Param->getType().getCanonicalType().getAsString();
+    }
+    FunctionTemplateDecl *FTD = K.SyclKernel->getPrimaryTemplate();
+    Policy.SuppressDefinition = true;
+    Policy.PolishForDeclaration = true;
+    if (FTD) {
+      FTD->print(O, Policy);
+    } else {
+      K.SyclKernel->print(O, Policy);
+    }
+    O << ";\n";
+
+    // Generate a shim function that returns the address of the free function.
+    O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n";
+    O << "  return (void (*)(" << ParmList << "))"
+      << K.SyclKernel->getIdentifier()->getName().data();
+    if (FTD) {
+      const TemplateArgumentList *TAL =
+          K.SyclKernel->getTemplateSpecializationArgs();
+      ArrayRef<TemplateArgument> A = TAL->asArray();
+      bool FirstParam = true;
+      O << "<";
+      for (auto X : A) {
+        if (FirstParam)
+          FirstParam = false;
+        else
+          O << ", ";
+        X.print(Policy, O, true);
+      }
+      O << ">";
+    }
+    O << ";\n";
+    O << "}\n";
+
+    // Generate is_kernel, is_single_task_kernel and nd_range_kernel functions.
+    O << "namespace sycl {\n";
+    O << "template <>\n";
+    O << "struct ext::oneapi::experimental::is_kernel<__sycl_shim"
+      << ShimCounter << "()";
+    O << "> {\n";
+    O << "  static constexpr bool value = true;\n";
+    O << "};\n";
+    int Dim = getFreeFunctionRangeDim(S, K.SyclKernel);
+    O << "template <>\n";
+    if (Dim > 0)
+      O << "struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim"
+        << ShimCounter << "(), " << Dim;
+    else
+      O << "struct "
+           "ext::oneapi::experimental::is_single_task_kernel<__sycl_shim"
+        << ShimCounter << "()";
+    O << "> {\n";
+    O << "  static constexpr bool value = true;\n";
+    O << "};\n";
+    O << "}\n";
+    ++ShimCounter;
+  }
+
+  if (FreeFunctionCount > 0) {
+    O << "\n#include <sycl/kernel_bundle.hpp>\n";
+  }
+  ShimCounter = 1;
+  for (const KernelDesc &K : KernelDescs) {
+    if (!isFreeFunction(S, K.SyclKernel))
+      continue;
+
+    O << "\n// Definition of kernel_id of " << K.Name << "\n";
+    O << "namespace sycl {\n";
+    O << "template <>\n";
+    O << "kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim"
+      << ShimCounter << "()>() {\n";
+    O << "  return sycl::detail::get_kernel_id_impl(std::string_view{\""
+      << K.Name << "\"});\n";
+    O << "}\n";
+    O << "}\n";
+    ++ShimCounter;
+  }
 }
 
 bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) {

clang/test/CodeGenSYCL/free_function_int_header.cpp

@@ -82,4 +82,126 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 8, 16 },
 
 // CHECK:        { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
-// CHECK-NEXT: };
+// CHECK-NEXT: };
+
+// CHECK: Definition of _Z18__sycl_kernel_ff_2Piii as a free function kernel
+// CHECK-NEXT: void ff_2(int *ptr, int start, int end);
+// CHECK-NEXT: static constexpr auto __sycl_shim1() {
+// CHECK-NEXT:   return (void (*)(int *, int, int))ff_2;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim1()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim1()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT: }
+
+// CHECK: Definition of _Z18__sycl_kernel_ff_2Piiii as a free function kernel
+// CHECK-NEXT: void ff_2(int *ptr, int start, int end, int value);
+// CHECK-NEXT: static constexpr auto __sycl_shim2() {
+// CHECK-NEXT:   return (void (*)(int *, int, int, int))ff_2;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim2()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim2()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT: }
+
+// CHECK: Definition of _Z18__sycl_kernel_ff_3IiEvPT_S0_S0_ as a free function kernel
+// CHECK-NEXT: template <typename T> void ff_3(T *ptr, T start, T end);
+// CHECK-NEXT: static constexpr auto __sycl_shim3() {
+// CHECK-NEXT:   return (void (*)(int *, int, int))ff_3<int>;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim3()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim3(), 2> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT: }
+ 
+// CHECK: Definition of _Z18__sycl_kernel_ff_3IfEvPT_S0_S0_ as a free function kernel
+// CHECK-NEXT: template <typename T> void ff_3(T *ptr, T start, T end);
+// CHECK-NEXT: static constexpr auto __sycl_shim4() {
+// CHECK-NEXT:   return (void (*)(float *, float, float))ff_3<float>;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim4()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim4(), 2> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT: }
+
+// CHECK: Definition of _Z18__sycl_kernel_ff_3IdEvPT_S0_S0_ as a free function kernel
+// CHECK-NEXT: template <typename T> void ff_3(T *ptr, T start, T end);
+// CHECK-NEXT: static constexpr auto __sycl_shim5() {
+// CHECK-NEXT:   return (void (*)(double *, double, double))ff_3<double>;
+// CHECK-NEXT: }
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_kernel<__sycl_shim5()> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim5(), 2> {
+// CHECK-NEXT:     static constexpr bool value = true;
+// CHECK-NEXT:   };
+// CHECK-NEXT: }
+
+// CHECK: #include <sycl/kernel_bundle.hpp>
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim1()>() {
+// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_2Piii"});
+// CHECK-NEXT:   }
+// CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piiii
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim2()>() {
+// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_2Piiii"});
+// CHECK-NEXT:   }
+// CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_3IiEvPT_S0_S0_
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim3()>() {
+// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IiEvPT_S0_S0_"});
+// CHECK-NEXT:   }
+// CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_3IfEvPT_S0_S0_
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim4()>() {
+// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IfEvPT_S0_S0_"});
+// CHECK-NEXT:   }
+// CHECK-NEXT: }
+
+// CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_3IdEvPT_S0_S0_
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT:   template <>
+// CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim5()>() {
+// CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IdEvPT_S0_S0_"});
+// CHECK-NEXT:   }
+// CHECK-NEXT: }

clang/test/SemaSYCL/Inputs/sycl/ext/oneapi/experimental/free_function_traits.hpp

@@ -0,0 +1,40 @@
+//==-------- free_function_traits.hpp - SYCL free function queries --------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+template <auto *Func, int Dims> struct is_nd_range_kernel {
+  static constexpr bool value = false;
+};
+
+template <auto *Func> struct is_single_task_kernel {
+  static constexpr bool value = false;
+};
+
+template <auto *Func, int Dims>
+inline constexpr bool is_nd_range_kernel_v =
+    is_nd_range_kernel<Func, Dims>::value;
+
+template <auto *Func>
+inline constexpr bool is_single_task_kernel_v =
+    is_single_task_kernel<Func>::value;
+
+template <auto *Func> struct is_kernel {
+  static constexpr bool value = false;
+};
+
+template <auto *Func>
+inline constexpr bool is_kernel_v = is_kernel<Func>::value;
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

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

@@ -0,0 +1,40 @@
+//==-------- free_function_traits.hpp - SYCL free function queries --------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+template <auto *Func, int Dims> struct is_nd_range_kernel {
+  static constexpr bool value = false;
+};
+
+template <auto *Func> struct is_single_task_kernel {
+  static constexpr bool value = false;
+};
+
+template <auto *Func, int Dims>
+inline constexpr bool is_nd_range_kernel_v =
+    is_nd_range_kernel<Func, Dims>::value;
+
+template <auto *Func>
+inline constexpr bool is_single_task_kernel_v =
+    is_single_task_kernel<Func>::value;
+
+template <auto *Func> struct is_kernel {
+  static constexpr bool value = false;
+};
+
+template <auto *Func>
+inline constexpr bool is_kernel_v = is_kernel<Func>::value;
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl