diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f7cffc4008317..fb2914b727ce2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1091,8 +1091,7 @@ static void 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) { @@ -1246,7 +1245,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); diff --git a/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp b/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp index f5225163e0779..5ebfe228a073d 100644 --- a/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp @@ -8,5 +8,5 @@ int main() { return 0; } -// CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_ -// CHECK: _ZTSZZ4mainENK3$_1clERN2cl4sycl7handlerEEUlvE_ \ No newline at end of file +// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE6->12clES2_EUlvE6->54{{.*}} +// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE7->12clES2_EUlvE7->54{{.*}} diff --git a/sycl/test/regression/same_unnamed_kernels.cpp b/sycl/test/regression/same_unnamed_kernels.cpp new file mode 100644 index 0000000000000..5d8fb11387152 --- /dev/null +++ b/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 + +template +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(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 bufA(A, 1); + cl::sycl::buffer bufB(B, 1); + + run(q, bufA, + [&](const cl::sycl::accessor + &acc) { acc[0] = 0; }); + run(q, bufB, + [&](const cl::sycl::accessor + &acc) { acc[0] *= 2; }); + + if (A[0] != 0 || B[0] != 2) + return -1; + + return 0; +}