[SYCL] Properly take care of unnamed lambdas (#827)

When two lambdas with the same signature but different content are
submitted to queue, compiler complains about duplicate names and
finishes with error. This patch introduces a more correct way to mangle lambdas.

Signed-off-by: Alexander Batashev <alexander.batashev@intel.com>

Author: Alexander Batashev

clang/lib/Sema/SemaSYCL.cpp

@@ -1104,8 +1104,7 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
 /// \param H           the integration header object
 /// \param Name        kernel name
 /// \param NameType    type representing kernel name (first template argument
-/// of
-///                      single_task, parallel_for, etc)
+/// of single_task, parallel_for, etc)
 /// \param KernelObjTy kernel object type
 static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
                               QualType NameType, CXXRecordDecl *KernelObjTy) {
@@ -1260,7 +1259,15 @@ 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, MC);
+
+  std::string Name;
+  // TODO SYCLIntegrationHeader also computes a unique stable name. It should
+  // probably lose this responsibility and only use the name provided here.
+  if (getLangOpts().SYCLUnnamedLambda)
+    Name = PredefinedExpr::ComputeName(
+        getASTContext(), PredefinedExpr::UniqueStableNameExpr, KernelNameType);
+  else
+    Name = constructKernelName(KernelNameType, MC);
 
   // TODO Maybe don't emit integration header inside the Sema?
   populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE);

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

@@ -8,5 +8,5 @@ int main() {
   return 0;
 }
 
-// CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_
-// CHECK: _ZTSZZ4mainENK3$_1clERN2cl4sycl7handlerEEUlvE_
+// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE6->12clES2_EUlvE6->54{{.*}}
+// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE7->12clES2_EUlvE7->54{{.*}}

sycl/test/regression/same_unnamed_kernels.cpp

@@ -0,0 +1,42 @@
+// RUN: %clangxx -fsycl %s -o %t.out -fsycl-unnamed-lambda
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+
+//==----- same_unnamed_kernels.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>
+
+template <typename F, typename B>
+void run(cl::sycl::queue &q, B &buf, const F &func) {
+  auto e = q.submit([&](cl::sycl::handler &cgh) {
+    auto acc = buf.template get_access<cl::sycl::access::mode::write>(cgh);
+    cgh.single_task([=]() { func(acc); });
+  });
+  e.wait();
+}
+
+int main() {
+  cl::sycl::queue q;
+
+  int A[1] = {1};
+  int B[1] = {1};
+  cl::sycl::buffer<int, 1> bufA(A, 1);
+  cl::sycl::buffer<int, 1> bufB(B, 1);
+
+  run(q, bufA,
+      [&](const cl::sycl::accessor<int, 1, cl::sycl::access::mode::write>
+              &acc) { acc[0] = 0; });
+  run(q, bufB,
+      [&](const cl::sycl::accessor<int, 1, cl::sycl::access::mode::write>
+              &acc) { acc[0] *= 2; });
+
+  if (A[0] != 0 || B[0] != 2)
+    return -1;
+
+  return 0;
+}