[SYCL-MLIR][clang] Attach SYCLKernelObjFunc to kernel body functions (

#12345)

`SemaSYCL` introduces kernels calling user-provided kernel body
functions to avoid OpenCL limitations on kernel arguments. Mark the
original kernel body functions with a new `SYCLKernelObjFunc` attribute
to handle them in codegen.

Signed-off-by: Victor Perez <victor.perez@codeplay.com>

clang/include/clang/Basic/Attr.td

@@ -1499,6 +1499,16 @@ def SYCLAccessorReadonly : Attr {
   let Documentation = [InternalOnly];
 }
 
+// Used to mark SYCL kernel object functions implementing SYCL kernels.
+// The Kernels arguments corresponds to the kernels this function implements.
+def SYCLKernelObjFunc : Attr {
+  // This attribute has no spellings as it is only ever created implicitly.
+  let Spellings = [];
+  let Args = [VariadicStringArgument<"Kernels">];
+  let SemaHandler = 0;
+  let Documentation = [InternalOnly];
+}
+
 // The attribute denotes that it is a function written in a scalar fashion, which
 // is used in ESIMD context and needs to be vectorized by a vector backend compiler.
 // For now, this attribute will be used only in internal implementation of

clang/lib/Sema/SemaSYCL.cpp

@@ -846,6 +846,26 @@ class SingleDeviceFunctionTracker {
     llvm::SmallVector<FunctionDecl *> CallStack;
     VisitCallNode(KernelNode, GetFDFromNode(KernelNode), CallStack);
 
+    // Mark kernel body function as such. Note a single function might have
+    // several kernel callers, e.g., when a range rounding kernel is added.
+    // Note the calling kernel name is already mangled.
+    if (KernelBody) {
+      llvm::SmallVector<llvm::StringRef, 2> KernelNames;
+      if (auto *Attr = KernelBody->getAttr<SYCLKernelObjFuncAttr>()) {
+        KernelNames.assign(Attr->kernels_begin(), Attr->kernels_end());
+      }
+      std::string KernelName = SYCLKernel->getNameAsString();
+      // Sorted insert to keep duplicates
+      auto *Pos = llvm::upper_bound(KernelNames, KernelName);
+      if (Pos == KernelNames.end() || *Pos != KernelName) {
+        KernelNames.insert(Pos, KernelName);
+        KernelBody->dropAttr<SYCLKernelObjFuncAttr>();
+        KernelBody->addAttr(SYCLKernelObjFuncAttr::CreateImplicit(
+            KernelBody->getASTContext(), KernelNames.data(),
+            KernelNames.size()));
+      }
+    }
+
     // Always inline the KernelBody in the kernel entry point. For ESIMD
     // inlining is handled later down the pipeline.
     if (KernelBody &&

clang/test/SemaSYCL/basic-kernel-wrapper.cpp

@@ -18,6 +18,12 @@ int main() {
   });
 }
 
+// Check kernel body function has SYCLKernelObjFuncAttr attribute attached
+
+// CHECK: CXXMethodDecl {{.*}} used constexpr operator() 'void () const'
+// CHECK: SYCLKernelObjFuncAttr
+// CHECK-SAME: Implicit _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E14kernel_wrapper
+
 // Check declaration of the kernel
 
 // CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'