[SYCL] Support unnamed lambda kernels

Add support for kernels without kernel name. Requires compiling
with -DUNNAMED_LAMBDA_EXT. Uses new __unique_stable_name
compiler built-in.

Signed-off-by: Roland Schulz <[email protected]>
Signed-off-by: Alexey Bader <[email protected]>

Author: rolandschulz

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9768,7 +9768,7 @@ def err_builtin_launder_invalid_arg : Error<
 // SYCL-specific diagnostics
 def err_sycl_attribute_address_space_invalid : Error<
   "address space is outside the valid range of values">;
-def err_sycl_kernel_name_class_not_top_level : Error<
+def warn_sycl_kernel_name_class_not_top_level : Warning<
  "kernel name class and its template argument classes' declarations can only "
  "nest in a namespace: %0">;
 def err_sycl_restrict : Error<

clang/include/clang/Sema/Sema.h

@@ -317,7 +317,8 @@ class SYCLIntegrationHeader {
 
   ///  Signals that subsequent parameter descriptor additions will go to
   ///  the kernel with given name. Starts new kernel invocation descriptor.
-  void startKernel(StringRef KernelName, QualType KernelNameType);
+  void startKernel(StringRef KernelName, QualType KernelNameType,
+                   StringRef KernelStableName);
 
   /// Adds a kernel parameter descriptor to current kernel invocation
   /// descriptor.
@@ -352,6 +353,9 @@ class SYCLIntegrationHeader {
     /// Kernel name type.
     QualType NameType;
 
+    /// Kernel name with stable lamba name mangling
+    std::string StableName;
+
     /// Descriptor of kernel actual parameters.
     SmallVector<KernelParamDesc, 8> Params;
 

clang/lib/Sema/SemaSYCL.cpp

@@ -797,7 +797,9 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
 
   ASTContext &Ctx = KernelObjTy->getASTContext();
   const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(KernelObjTy);
-  H.startKernel(Name, NameType);
+  const std::string StableName = PredefinedExpr::ComputeName(
+      Ctx, PredefinedExpr::UniqueStableNameExpr, NameType);
+  H.startKernel(Name, NameType, StableName);
 
   auto populateHeaderForAccessor = [&](const QualType &ArgTy, uint64_t Offset) {
     // The parameter is a SYCL accessor object.
@@ -1112,7 +1114,7 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
           // defined class constituting the kernel name is not globally
           // accessible - contradicts the spec
           Diag.Report(D->getSourceRange().getBegin(),
-                      diag::err_sycl_kernel_name_class_not_top_level);
+                      diag::warn_sycl_kernel_name_class_not_top_level);
         }
       }
       break;
@@ -1238,12 +1240,14 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";
 
   O << "\n";
+  O << "#ifndef UNNAMED_LAMBDA_EXT\n";
   O << "// Forward declarations of templated kernel function types:\n";
 
   llvm::SmallPtrSet<const void *, 4> Printed;
   for (const KernelDesc &K : KernelDescs) {
     emitForwardClassDecls(O, K.NameType, Printed);
   }
+  O << "#endif\n";
   O << "\n";
 
   O << "namespace cl {\n";
@@ -1305,19 +1309,21 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   }
   O << "};\n\n";
 
-  O << "// Specializations of this template class encompasses information\n";
-  O << "// about a kernel. The kernel is identified by the template\n";
-  O << "// parameter type.\n";
-  O << "template <class KernelNameType> struct KernelInfo;\n";
-  O << "\n";
-
   O << "// Specializations of KernelInfo for kernel function types:\n";
   CurStart = 0;
 
   for (const KernelDesc &K : KernelDescs) {
     const size_t N = K.Params.size();
+    O << "#ifdef UNNAMED_LAMBDA_EXT\n";
+    O << "template <> struct KernelInfoData<";
+    O << "'" << K.StableName.front();
+    for (char c : StringRef(K.StableName).substr(1))
+      O << "', '" << c;
+    O << "'> {\n";
+    O << "#else\n";
     O << "template <> struct KernelInfo<"
       << eraseAnonNamespace(K.NameType.getAsString()) << "> {\n";
+    O << "#endif\n";
     O << "  DLL_LOCAL\n";
     O << "  static constexpr const char* getName() { return \"" << K.Name
       << "\"; }\n";
@@ -1355,10 +1361,12 @@ bool SYCLIntegrationHeader::emit(const StringRef &IntHeaderName) {
 }
 
 void SYCLIntegrationHeader::startKernel(StringRef KernelName,
-                                        QualType KernelNameType) {
+                                        QualType KernelNameType,
+                                        StringRef KernelStableName) {
   KernelDescs.resize(KernelDescs.size() + 1);
   KernelDescs.back().Name = KernelName;
   KernelDescs.back().NameType = KernelNameType;
+  KernelDescs.back().StableName = KernelStableName;
 }
 
 void SYCLIntegrationHeader::addParamDesc(kernel_param_kind_t Kind, int Info,

clang/test/CodeGenSYCL/integration_header.cpp

@@ -47,7 +47,6 @@
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 //
-// CHECK: template <class KernelNameType> struct KernelInfo;
 // CHECK: template <> struct KernelInfo<class first_kernel> {
 // CHECK: template <> struct KernelInfo<::second_namespace::second_kernel<char>> {
 // CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point<X> >> {

clang/test/CodeGenSYCL/wrapped-accessor.cpp

@@ -27,8 +27,6 @@
 // CHECK-NEXT:  0 // _ZTSZ4mainE14wrapped_access
 // CHECK-NEXT: };
 
-// CHECK: template <class KernelNameType> struct KernelInfo;
-
 // CHECK: template <> struct KernelInfo<class wrapped_access> {
 
 #include <sycl.hpp>

clang/test/Misc/warning-flags.c

@@ -18,7 +18,7 @@ This test serves two purposes:
 
 The list of warnings below should NEVER grow.  It should gradually shrink to 0.
 
-CHECK: Warnings without flags (74):
+CHECK: Warnings without flags (75):
 CHECK-NEXT:   ext_excess_initializers
 CHECK-NEXT:   ext_excess_initializers_in_char_array_initializer
 CHECK-NEXT:   ext_expected_semi_decl_list
@@ -84,6 +84,7 @@ CHECK-NEXT:   warn_property_getter_owning_mismatch
 CHECK-NEXT:   warn_register_objc_catch_parm
 CHECK-NEXT:   warn_related_result_type_compatibility_class
 CHECK-NEXT:   warn_related_result_type_compatibility_protocol
+CHECK-NEXT:   warn_sycl_kernel_name_class_not_top_level
 CHECK-NEXT:   warn_template_export_unsupported
 CHECK-NEXT:   warn_template_spec_extra_headers
 CHECK-NEXT:   warn_tentative_incomplete_array

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -51,6 +51,7 @@ struct kernel_param_desc_t {
   int offset;
 };
 
+#ifndef UNNAMED_LAMBDA_EXT
 template <class KernelNameType> struct KernelInfo {
   static constexpr unsigned getNumParams() { return 0; }
   static const kernel_param_desc_t &getParamDesc(int Idx) {
@@ -59,6 +60,29 @@ template <class KernelNameType> struct KernelInfo {
   }
   static constexpr const char *getName() { return ""; }
 };
+#else
+template <char...> struct KernelInfoData; // Should this have dummy impl?
+
+// C++14 like index_sequence and make_index_sequence
+// not needed C++14 members (value_type, size) not implemented
+template <class T, T...> struct integer_sequence {};
+template <size_t... I> using index_sequence = integer_sequence<size_t, I...>;
+template <size_t N>
+using make_index_sequence = __make_integer_seq<integer_sequence, size_t, N>;
+
+template <typename T> struct KernelInfoImpl {
+private:
+  static constexpr auto n = __unique_stable_name(T);
+  template <std::size_t... I>
+  static KernelInfoData<n[I]...> impl(index_sequence<I...>) {
+    return {};
+  }
+
+public:
+  using type = decltype(impl(make_index_sequence<__builtin_strlen(n)>{}));
+};
+template <typename T> using KernelInfo = typename KernelInfoImpl<T>::type;
+#endif
 
 } // namespace detail
 } // namespace sycl

sycl/test/regression/kernel_name_class.cpp

@@ -3,6 +3,8 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
+// RUN: %clangxx -fsycl %s -o %t.ext.out -lOpenCL -DUNNAMED_LAMBDA_EXT
+// RUN: %CPU_RUN_PLACEHOLDER %t.ext.out
 
 //==-- kernel_name_class.cpp - SYCL kernel naming variants test ------------==//
 //

sycl/test/regression/kernel_unnamed.cpp

@@ -0,0 +1,67 @@
+// RUN: %clangxx -fsycl %s -o %t.out -lOpenCL -DUNNAMED_LAMBDA_EXT
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==-- kernel_unnamed.cpp - SYCL kernel naming variants test ------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+#define GOLD 10
+static int NumTestCases = 0;
+
+template <class F>
+void foo(cl::sycl::queue &deviceQueue, cl::sycl::buffer<int, 1> &buf, F f) {
+  deviceQueue.submit([&](cl::sycl::handler &cgh) {
+    auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+    cgh.single_task([=]() { acc[0] = f(acc[0], GOLD); });
+  });
+}
+
+namespace nm {
+struct Wrapper {
+
+  int test() {
+    int arr[] = {0};
+    {
+      // Simple test
+      cl::sycl::queue deviceQueue;
+      cl::sycl::buffer<int, 1> buf(arr, 1);
+      deviceQueue.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+        cgh.single_task([=]() { acc[0] += GOLD; });
+      });
+      ++NumTestCases;
+
+// Test lambdas with different ordinal because of macro expansion
+#ifdef __SYCL_DEVICE_ONLY__
+      [] {}();
+#endif
+      deviceQueue.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+        cgh.single_task([=]() { acc[0] += GOLD; });
+      });
+      ++NumTestCases;
+
+      // Test lambda passed to function
+      foo(deviceQueue, buf, [](int a, int b) { return a + b; });
+      ++NumTestCases;
+    }
+    return arr[0];
+  }
+};
+} // namespace nm
+
+int main() {
+  nm::Wrapper w;
+  int res = w.test();
+  assert (res == GOLD * NumTestCases && "Wrong result");
+}