[SYCL] Fix kernel name mangling for unnamed lambdas

The problem was that new mangling context for every kernel function was
created. So if the code contains several unnamed lambdas with the same
signature it caused equal mangled names for them.

Signed-off-by: Viktoria Maksimova <viktoria.maksimova@intel.com>

Author: vmaksimo

clang/include/clang/Sema/Sema.h

@@ -11512,7 +11512,7 @@ class Sema {
     KernelCallVariadicFunction
  };
   DeviceDiagBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
-  void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc);
+  void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC);
   void MarkDevice(void);
   bool CheckSYCLCall(SourceLocation Loc, FunctionDecl *Callee);
 };

clang/lib/Sema/SemaSYCL.cpp

@@ -1176,14 +1176,11 @@ static std::string eraseAnonNamespace(std::string S) {
 
 // Creates a mangled kernel name for given kernel name type
 static std::string constructKernelName(QualType KernelNameType,
-                                       ASTContext &AC) {
-  std::unique_ptr<MangleContext> MC(AC.createMangleContext());
-
+                                       MangleContext &MC) {
   SmallString<256> Result;
   llvm::raw_svector_ostream Out(Result);
 
-  MC->mangleTypeName(KernelNameType, Out);
-
+  MC.mangleTypeName(KernelNameType, Out);
   return Out.str();
 }
 
@@ -1209,7 +1206,8 @@ static std::string constructKernelName(QualType KernelNameType,
 //   }
 //
 //
-void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc) {
+void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
+                                 MangleContext &MC) {
   CXXRecordDecl *LE = getKernelObjectType(KernelCallerFunc);
   assert(LE && "invalid kernel caller");
 
@@ -1223,7 +1221,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc) {
   assert(TemplateArgs && "No template argument info");
   QualType KernelNameType = TypeName::getFullyQualifiedType(
       TemplateArgs->get(0).getAsType(), getASTContext(), true);
-  std::string Name = constructKernelName(KernelNameType, getASTContext());
+  std::string Name = constructKernelName(KernelNameType, MC);
 
   // TODO Maybe don't emit integration header inside the Sema?
   populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE);

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -17,6 +17,7 @@
 #include "clang/AST/DependentDiagnostic.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
+#include "clang/AST/Mangle.h"
 #include "clang/AST/PrettyDeclStackTrace.h"
 #include "clang/AST/TypeLoc.h"
 #include "clang/Sema/Initialization.h"
@@ -5520,6 +5521,8 @@ NamedDecl *Sema::FindInstantiatedDecl(SourceLocation Loc, NamedDecl *D,
 /// Performs template instantiation for all implicit template
 /// instantiations we have seen until this point.
 void Sema::PerformPendingInstantiations(bool LocalOnly) {
+  std::unique_ptr<MangleContext> MangleCtx(
+      getASTContext().createMangleContext());
   while (!PendingLocalImplicitInstantiations.empty() ||
          (!LocalOnly && !PendingInstantiations.empty())) {
     PendingImplicitInstantiation Inst;
@@ -5538,7 +5541,8 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) {
                                 TSK_ExplicitInstantiationDefinition;
       if (Function->isMultiVersion()) {
         getASTContext().forEachMultiversionedFunctionVersion(
-            Function, [this, Inst, DefinitionRequired](FunctionDecl *CurFD) {
+            Function, [this, Inst, DefinitionRequired,
+                       MangleCtx = move(MangleCtx)](FunctionDecl *CurFD) {
               InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, CurFD, true,
                                             DefinitionRequired, true);
               if (CurFD->isDefined()) {
@@ -5547,7 +5551,7 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) {
                 // so we are checking for SYCL kernel attribute after instantination.
                 if (getLangOpts().SYCLIsDevice &&
                         CurFD->hasAttr<SYCLKernelAttr>()) {
-                  ConstructOpenCLKernel(CurFD);
+                  ConstructOpenCLKernel(CurFD, *MangleCtx);
                 }
                 CurFD->setInstantiationIsPending(false);
               }
@@ -5561,7 +5565,7 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) {
           // so we are checking for SYCL kernel attribute after instantination.
           if (getLangOpts().SYCLIsDevice &&
                   Function->hasAttr<SYCLKernelAttr>()) {
-              ConstructOpenCLKernel(Function);
+            ConstructOpenCLKernel(Function, *MangleCtx);
           }
           Function->setInstantiationIsPending(false);
         }

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -161,6 +161,22 @@ class sampler {
   void use(void) const {}
 };
 
+class event {};
+class queue {
+public:
+  template <typename T>
+  event submit(T cgf) { return event{}; }
+};
+#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
+class handler {
+public:
+  template <typename KernelName, typename KernelType>
+  ATTR_SYCL_KERNEL void single_task(KernelType kernelFunc) {}
+
+  template <typename KernelType>
+  ATTR_SYCL_KERNEL void single_task(KernelType kernelFunc) {}
+};
+
 } // namespace sycl
 } // namespace cl
 

clang/test/SemaSYCL/mangle-unnamed-kernel.cpp

@@ -0,0 +1,12 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -I %S/Inputs -fsycl-is-device -fsycl-unnamed-lambda -ast-dump %s | FileCheck %s
+#include <sycl.hpp>
+
+int main() {
+  cl::sycl::queue q;
+  q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
+  q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
+  return 0;
+}
+
+// CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_
+// CHECK: _ZTSZZ4mainENK3$_1clERN2cl4sycl7handlerEEUlvE_