[clang] Fix Lambda Mangling in Namespace-Scope Variable Initializers. (PR #159115)
Zahira Ammarguellat via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 19 06:22:09 PDT 2025
https://github.com/zahiraam updated https://github.com/llvm/llvm-project/pull/159115
>From b197e92e844472ccef1999cac653a76109183ee4 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat at intel.com>
Date: Tue, 16 Sep 2025 08:47:40 -0700
Subject: [PATCH 1/6] Fix Lambda Mangling in Namespace-Scope Variable
Initializers.
---
clang/lib/AST/ASTContext.cpp | 1 +
clang/lib/Sema/SemaLambda.cpp | 2 +-
clang/test/CodeGenCUDA/anon-ns.cu | 8 ++--
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 2 +-
clang/test/CodeGenSYCL/unique_stable_name.cpp | 40 ++++++++--------
.../unique_stable_name_windows_diff.cpp | 22 +++++++--
clang/test/CodeGenSYCL/unnamed-types.cpp | 46 +++++++++++++++++++
7 files changed, 90 insertions(+), 31 deletions(-)
create mode 100644 clang/test/CodeGenSYCL/unnamed-types.cpp
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 5240054c2f36b..55dcf4fc3a335 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -13395,6 +13395,7 @@ MangleNumberingContext &
ASTContext::getManglingNumberContext(const DeclContext *DC) {
assert(LangOpts.CPlusPlus); // We don't need mangling numbers for plain C.
std::unique_ptr<MangleNumberingContext> &MCtx = MangleNumberingContexts[DC];
+ DC = DC->getPrimaryContext();
if (!MCtx)
MCtx = createMangleNumberingContext();
return *MCtx;
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index fbc2e7eb30676..f82069b0b4e2b 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -356,7 +356,7 @@ Sema::getCurrentMangleNumberContext(const DeclContext *DC) {
return std::make_tuple(&Context.getManglingNumberContext(DC), nullptr);
}
- return std::make_tuple(nullptr, nullptr);
+ return std::make_tuple(nullptr, ManglingContextDecl);
}
case NonInlineInModulePurview:
diff --git a/clang/test/CodeGenCUDA/anon-ns.cu b/clang/test/CodeGenCUDA/anon-ns.cu
index d931f31d0207c..83e863d0e396f 100644
--- a/clang/test/CodeGenCUDA/anon-ns.cu
+++ b/clang/test/CodeGenCUDA/anon-ns.cu
@@ -26,14 +26,14 @@
// HIP-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv\.intern\.b04fd23c98500190]](
// HIP-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT_\.intern\.b04fd23c98500190]](
-// HIP-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT_\.intern\.b04fd23c98500190]](
+
// HIP-DAG: @[[VM:_ZN12_GLOBAL__N_12vmE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
// HIP-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE\.static\.b04fd23c98500190]] = addrspace(4) externally_initialized constant
// HIP-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
// CUDA-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190]](
// CUDA-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT___intern__b04fd23c98500190]](
-// CUDA-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT___intern__b04fd23c98500190]](
+
// CUDA-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE__static__b04fd23c98500190]] = addrspace(4) externally_initialized constant
// CUDA-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE__static__b04fd23c98500190]] = addrspace(1) externally_initialized global
@@ -45,14 +45,14 @@
// COMMON-DAG: @[[KERNSTR:.*]] = {{.*}} c"[[KERN]]\00"
// COMMON-DAG: @[[KTXSTR:.*]] = {{.*}} c"[[KTX]]\00"
-// COMMON-DAG: @[[KTLSTR:.*]] = {{.*}} c"[[KTL]]\00"
+
// HIP-DAG: @[[VMSTR:.*]] = {{.*}} c"[[VM]]\00"
// COMMON-DAG: @[[VCSTR:.*]] = {{.*}} c"[[VC]]\00"
// COMMON-DAG: @[[VTSTR:.*]] = {{.*}} c"[[VT]]\00"
// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KERNSTR]]
// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTXSTR]]
-// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTLSTR]]
+
// HIP-DAG: call void @__{{.*}}RegisterManagedVar({{.*}}@[[VMSTR]]
// COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VCSTR]]
// COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VTSTR]]
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index cd1d4d801951d..5b3b73812d286 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -90,7 +90,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
-// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task at V<lambda_1>@?0??main@@9 at V1?0??2 at 9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %kernelFunc.coerce) #{{[0-9]+}} {
+// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task at V<lambda_1>@lambda@?0??main@@9 at V12?0??3 at 9@@@YAXV<lambda_1>@lambda@?0??main@@9@@Z"(i32 %kernelFunc.coerce) #{{[0-9]+}} {
// CHECK-HOST-WINDOWS-NEXT: entry:
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon, align 4
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %kernelFunc, i32 0, i32 0
diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp
index 3ab7e3b8f2e7a..1af2c874b74d9 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp
@@ -3,20 +3,20 @@
// CHECK: @[[INT1:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
// CHECK: @[[STRING:[^\w]+]] = private unnamed_addr constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00",
// CHECK: @[[INT2:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
-// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00"
-// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE0_\00"
-// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE1_\00"
-// CHECK: @{{.*}} = private unnamed_addr constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE2_\00", align 1
-// CHECK: @{{.*}} = private unnamed_addr constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE3_\00", align 1
-// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE4_\00"
-// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE5_\00"
+// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE_\00"
+// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE0_\00"
+// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE1_\00"
+// CHECK: @{{.*}} = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE2_\00"
+// CHECK: @{{.*}} = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE3_\00"
+// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE4_\00"
+// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE5_\00"
// CHECK: @[[INT3:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
-// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [[LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00"
+// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [30 x i8] c"_ZTSZZ4mainENKUlvE_clEvEUlvE_\00"
// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE_\00",
-// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_EvvEUlvE_\00",
// CHECK: @[[LAMBDA_NO_DEP:[^\w]+]] = private unnamed_addr constant [[NO_DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ13lambda_no_depIidEvT_T0_EUlidE_\00",
-// CHECK: @[[LAMBDA_TWO_DEP:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA1_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_EvvEUlvE_\00",
-// CHECK: @[[LAMBDA_TWO_DEP2:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA2_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_TWO_DEP:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA1_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_TWO_DEP2:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA2_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE_clEvEUldE_ZZ4mainENKS0_clEvEUliE_EvvEUlvE_\00",
extern "C" void puts(const char *) {}
@@ -101,7 +101,7 @@ int main() {
// CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
// CHECK: declare noundef ptr @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
// CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
- // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvT0_
+ // CHECK: _Z26unnamed_kernel_single_taskIZ4mainEUlvE_EvT_
unnamed_kernel_single_task(
[]() {
@@ -124,13 +124,13 @@ int main() {
// CHECK: call void @_Z14template_paramIiEvv
template_param<decltype(x)>();
- // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
+ // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE_clEvEUlvE_Evv
lambda_in_dependent_function<int>();
// CHECK: call void @_Z28lambda_in_dependent_functionIiEvv
lambda_in_dependent_function<decltype(x)>();
- // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
+ // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_Evv
lambda_no_dep<int, double>(3, 5.5);
// CHECK: call void @_Z13lambda_no_depIidEvT_T0_(i32 noundef 3, double noundef 5.500000e+00)
@@ -140,30 +140,30 @@ int main() {
auto y = [](int a) { return a; };
auto z = [](double b) { return b; };
lambda_two_dep<decltype(y), decltype(z)>();
- // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+ // CHECK: @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
lambda_two_dep<decltype(z), decltype(y)>();
- // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+ // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
});
}
// CHECK: define linkonce_odr void @_Z14template_paramIiEvv
// CHECK: call void @puts(ptr noundef @[[INT3]])
-// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
+// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE_clEvEUlvE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA]])
// CHECK: define linkonce_odr void @_Z28lambda_in_dependent_functionIiEvv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_IN_DEP_INT]])
-// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
+// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_IN_DEP_X]])
// CHECK: define linkonce_odr void @_Z13lambda_no_depIidEvT_T0_(i32 noundef %a, double noundef %b)
// CHECK: call void @puts(ptr noundef @[[LAMBDA_NO_DEP]])
-// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_TWO_DEP]])
-// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_TWO_DEP2]])
diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
index 14366a092a1fe..33c6d461aeb02 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) '
-// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE='
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
+// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
template<typename KN, typename Func>
@@ -47,7 +47,19 @@ int main() {
// Make sure the following 3 are the same between the host and device compile.
// Note that these are NOT the same value as each other, they differ by the
// signature.
- // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUlvE_\00"
- // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUliE_\00"
- // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUldE_\00"
+ // HOST: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUlvE_\00"
+ // HOST: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUliE_\00"
+ // HOST: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUldE_\00"
+
+ // DEVICE: define dso_local spir_kernel void @_ZTSZ4mainE2K1
+ // DEVICE: call spir_func void @_ZZ4mainENKUlvE_clEv
+ // DEVICE: define internal spir_func void @_ZZ4mainENKUlvE_clEv
+ // DEVICE: define dso_local spir_kernel void @_ZTSZ4mainE2K2
+ // DEVICE: call spir_func void @_ZZ4mainENKUliE_clEi
+ // DEVICE: define internal spir_func void @_ZZ4mainENKUliE_clEi
+ // DEVICE: define dso_local spir_kernel void @_ZTSZ4mainE2K3
+ // DEVICE: call spir_func void @_ZZ4mainENKUldE_clEd
+ // DEVICE: define internal spir_func void @_ZZ4mainENKUldE_clEd
+ // DEVICE: define dso_local spir_kernel void @_ZTSZ4mainE2K4
+ // DEVICE: call spir_func void @_ZZ4mainENKUlvE_clEv
}
diff --git a/clang/test/CodeGenSYCL/unnamed-types.cpp b/clang/test/CodeGenSYCL/unnamed-types.cpp
new file mode 100644
index 0000000000000..d87d9a57b752e
--- /dev/null
+++ b/clang/test/CodeGenSYCL/unnamed-types.cpp
@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -fsycl-is-device -O0 \
+// RUN: -triple spirv64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
+
+// RUN: %clang_cc1 -fsycl-is-host -O0 \
+// RUN: -triple spirv64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
+
+// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spir-unknown--unknown %s -o - | FileCheck %s --check-prefix=MSVC
+
+namespace QL {
+auto dg1 = [] { return 1; };
+}
+namespace QL {
+auto dg2 = [] { return 2; };
+}
+using namespace QL;
+template<typename T>
+[[clang::sycl_kernel_entry_point(T)]] void f(T t) {
+ t();
+}
+void g() {
+ f(dg1);
+ f(dg2);
+}
+
+// HOST: @_ZN2QL3dg1E = internal global %class.anon undef, align 1
+// HOST: @_ZN2QL3dg2E = internal global %class.anon.0 undef, align 1
+
+// DEVICE: define spir_kernel void @_ZTSN2QL3dg1MUlvE_E
+// DEVICE: call spir_func noundef i32 @_ZNK2QL3dg1MUlvE_clEv
+// DEVICE: define internal spir_func noundef i32 @_ZNK2QL3dg1MUlvE_clEv
+// DEVICE: define spir_kernel void @_ZTSN2QL3dg2MUlvE_E
+// DEVICE: call spir_func noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+// DEVICE: define internal spir_func noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+
+// HOST: define spir_func void @_Z1gv
+// HOST: call spir_func void @_Z1fIN2QL3dg1MUlvE_EEvT_
+// HOST: call spir_func void @_Z1fIN2QL3dg2MUlvE_EEvT_
+// HOST: define internal spir_func void @_Z1fIN2QL3dg1MUlvE_EEvT
+
+// MSVC: define dso_local spir_kernel void @_ZTSN2QL3dg1MUlvE_E
+// MSVC: call spir_func noundef i32 @_ZNK2QL3dg1MUlvE_clEv
+// MSVC: define internal spir_func noundef i32 @_ZNK2QL3dg1MUlvE_clEv
+// MSVC: define dso_local spir_kernel void @_ZTSN2QL3dg2MUlvE_E
+// MSVC: call spir_func noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+// MSVC: define internal spir_func noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+
>From a5292c2e3fa445a224dacff66a95129b25e7151b Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat at intel.com>
Date: Wed, 17 Sep 2025 07:59:38 -0700
Subject: [PATCH 2/6] Addressed review comments.
---
clang/test/CodeGenCUDA/anon-ns.cu | 8 ++--
clang/test/CodeGenCUDA/unnamed-types.cpp | 40 +++++++++++++++++++
clang/test/CodeGenSYCL/unique_stable_name.cpp | 8 ++--
clang/test/CodeGenSYCL/unnamed-types.cpp | 13 +++---
4 files changed, 56 insertions(+), 13 deletions(-)
create mode 100644 clang/test/CodeGenCUDA/unnamed-types.cpp
diff --git a/clang/test/CodeGenCUDA/anon-ns.cu b/clang/test/CodeGenCUDA/anon-ns.cu
index 83e863d0e396f..ae8e1abe52a9f 100644
--- a/clang/test/CodeGenCUDA/anon-ns.cu
+++ b/clang/test/CodeGenCUDA/anon-ns.cu
@@ -26,14 +26,14 @@
// HIP-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv\.intern\.b04fd23c98500190]](
// HIP-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT_\.intern\.b04fd23c98500190]](
-
+// HIP-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_16lambdaMUlvE_EEvT_.intern.b04fd23c98500190]](
// HIP-DAG: @[[VM:_ZN12_GLOBAL__N_12vmE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
// HIP-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE\.static\.b04fd23c98500190]] = addrspace(4) externally_initialized constant
// HIP-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
// CUDA-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190]](
// CUDA-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT___intern__b04fd23c98500190]](
-
+// CUDA-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_16lambdaMUlvE_EEvT___intern__b04fd23c98500190]](
// CUDA-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE__static__b04fd23c98500190]] = addrspace(4) externally_initialized constant
// CUDA-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE__static__b04fd23c98500190]] = addrspace(1) externally_initialized global
@@ -45,14 +45,14 @@
// COMMON-DAG: @[[KERNSTR:.*]] = {{.*}} c"[[KERN]]\00"
// COMMON-DAG: @[[KTXSTR:.*]] = {{.*}} c"[[KTX]]\00"
-
+// COMMON-DAG: @[[KTLSTR:.*]] = {{.*}} c"[[KTL]]\00"
// HIP-DAG: @[[VMSTR:.*]] = {{.*}} c"[[VM]]\00"
// COMMON-DAG: @[[VCSTR:.*]] = {{.*}} c"[[VC]]\00"
// COMMON-DAG: @[[VTSTR:.*]] = {{.*}} c"[[VT]]\00"
// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KERNSTR]]
// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTXSTR]]
-
+// COMMON-DAG: call i32 @__{{.*}}RegisterFunction({{.*}}@[[KTLSTR]]
// HIP-DAG: call void @__{{.*}}RegisterManagedVar({{.*}}@[[VMSTR]]
// COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VCSTR]]
// COMMON-DAG: call void @__{{.*}}RegisterVar({{.*}}@[[VTSTR]]
diff --git a/clang/test/CodeGenCUDA/unnamed-types.cpp b/clang/test/CodeGenCUDA/unnamed-types.cpp
new file mode 100644
index 0000000000000..eb36594cb0395
--- /dev/null
+++ b/clang/test/CodeGenCUDA/unnamed-types.cpp
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN: -aux-triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN: -x cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN: -aux-triple nvptx64-nvidia-cuda \
+// RUN: -x cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
+
+int cudaConfigureCall(int, int, decltype(sizeof(int)) = 0, void* = nullptr);
+namespace QL {
+auto dg1 = [] { return 1; };
+}
+namespace QL {
+auto dg2 = [] { return 2; };
+}
+using namespace QL;
+template<typename T>
+__attribute__((global)) void f(T t) {
+ t();
+}
+void g() {
+ f<<<1,1>>>(dg1);
+ f<<<1,1>>>(dg2);
+}
+
+// HOST: @_ZN2QL3dg1E = internal global %class.anon undef, align 1
+// HOST: @_ZN2QL3dg2E = internal global %class.anon.0 undef, align 1
+
+// DEVICE: define void @_Z1fIN2QL3dg1MUlvE_EEvT_
+// DEVICE: call noundef i32 @_ZNK2QL3dg1MUlvE_clEv
+// DEVICE: define internal noundef i32 @_ZNK2QL3dg1MUlvE_clEv
+// DEVICE: define void @_Z1fIN2QL3dg2MUlvE_EEvT_
+// DEVICE: call noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+// DEVICE: define internal noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+
+// HOST: define dso_local void @_Z1gv
+// HOST: call void @_Z16__device_stub__fIN2QL3dg1MUlvE_EEvT_
+// HOST: call void @_Z16__device_stub__fIN2QL3dg2MUlvE_EEvT_
+// HOST: define internal void @_Z16__device_stub__fIN2QL3dg1MUlvE_EEvT_
+// HOST: define internal void @_Z16__device_stub__fIN2QL3dg2MUlvE_EEvT_
diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp
index 1af2c874b74d9..9160249ca0fc4 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp
@@ -11,7 +11,7 @@
// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE4_\00"
// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE5_\00"
// CHECK: @[[INT3:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
-// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [30 x i8] c"_ZTSZZ4mainENKUlvE_clEvEUlvE_\00"
+// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE_\00"
// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE_\00",
// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_EvvEUlvE_\00",
// CHECK: @[[LAMBDA_NO_DEP:[^\w]+]] = private unnamed_addr constant [[NO_DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ13lambda_no_depIidEvT_T0_EUlidE_\00",
@@ -101,7 +101,7 @@ int main() {
// CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
// CHECK: declare noundef ptr @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
// CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
- // CHECK: _Z26unnamed_kernel_single_taskIZ4mainEUlvE_EvT_
+ // CHECK: define internal void @_Z26unnamed_kernel_single_taskIZ4mainEUlvE_EvT_
unnamed_kernel_single_task(
[]() {
@@ -130,7 +130,7 @@ int main() {
// CHECK: call void @_Z28lambda_in_dependent_functionIiEvv
lambda_in_dependent_function<decltype(x)>();
- // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_Evv
+ // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_Evv
lambda_no_dep<int, double>(3, 5.5);
// CHECK: call void @_Z13lambda_no_depIidEvT_T0_(i32 noundef 3, double noundef 5.500000e+00)
@@ -140,7 +140,7 @@ int main() {
auto y = [](int a) { return a; };
auto z = [](double b) { return b; };
lambda_two_dep<decltype(y), decltype(z)>();
- // CHECK: @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+ // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
lambda_two_dep<decltype(z), decltype(y)>();
// CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
diff --git a/clang/test/CodeGenSYCL/unnamed-types.cpp b/clang/test/CodeGenSYCL/unnamed-types.cpp
index d87d9a57b752e..64b8c166f8a16 100644
--- a/clang/test/CodeGenSYCL/unnamed-types.cpp
+++ b/clang/test/CodeGenSYCL/unnamed-types.cpp
@@ -1,10 +1,12 @@
-// RUN: %clang_cc1 -fsycl-is-device -O0 \
-// RUN: -triple spirv64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
+// RUN: %clang_cc1 -fsycl-is-device -O0 -triple spirv64-unknown-unknown \
+// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
-// RUN: %clang_cc1 -fsycl-is-host -O0 \
-// RUN: -triple spirv64-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
+// RUN: %clang_cc1 -fsycl-is-host -O0 -triple spirv64-unknown-unknown \
+// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
-// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-pc-windows-msvc -triple spir-unknown--unknown %s -o - | FileCheck %s --check-prefix=MSVC
+// RUN: %clang_cc1 -fsycl-is-device -emit-llvm \
+// RUN: -aux-triple x86_64-pc-windows-msvc -triple spir-unknown--unknown \
+// RUN: %s -o - | FileCheck %s --check-prefix=MSVC
namespace QL {
auto dg1 = [] { return 1; };
@@ -36,6 +38,7 @@ void g() {
// HOST: call spir_func void @_Z1fIN2QL3dg1MUlvE_EEvT_
// HOST: call spir_func void @_Z1fIN2QL3dg2MUlvE_EEvT_
// HOST: define internal spir_func void @_Z1fIN2QL3dg1MUlvE_EEvT
+// HOST: define internal spir_func void @_Z1fIN2QL3dg2MUlvE_EEvT_
// MSVC: define dso_local spir_kernel void @_ZTSN2QL3dg1MUlvE_E
// MSVC: call spir_func noundef i32 @_ZNK2QL3dg1MUlvE_clEv
>From b70972c7df0ea2fc8f6297185c5f5e3824f9d20c Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat at intel.com>
Date: Wed, 17 Sep 2025 13:07:07 -0700
Subject: [PATCH 3/6] Fix LIT tests again
---
clang/include/clang/Sema/Sema.h | 8 +++-
clang/test/CodeGenCUDA/unnamed-types.cpp | 40 -------------------
clang/test/CodeGenCUDA/unnamed-types.cu | 37 ++++++++++++++++-
clang/test/CodeGenSYCL/unique_stable_name.cpp | 1 +
4 files changed, 43 insertions(+), 43 deletions(-)
delete mode 100644 clang/test/CodeGenCUDA/unnamed-types.cpp
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 7e00085685b21..7097898324d17 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -9178,7 +9178,13 @@ class Sema final : public SemaBase {
};
/// Compute the mangling number context for a lambda expression or
- /// block literal. Also return the extra mangling decl if any.
+ /// block literal that appears in the specified declaration context in
+ /// consideration of the current expression evaluation and template
+ /// instantiation contexts. If the mangling context requires external linkage,
+ /// then a mangling number context is returned in the first tuple
+ /// element. If the mangling context is non-normal (specialized for
+ /// lambda and block types relative to other entities), the overriding
+ /// declaration is returned in the second tuple element.
///
/// \param DC - The DeclContext containing the lambda expression or
/// block literal.
diff --git a/clang/test/CodeGenCUDA/unnamed-types.cpp b/clang/test/CodeGenCUDA/unnamed-types.cpp
deleted file mode 100644
index eb36594cb0395..0000000000000
--- a/clang/test/CodeGenCUDA/unnamed-types.cpp
+++ /dev/null
@@ -1,40 +0,0 @@
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
-// RUN: -aux-triple nvptx64-nvidia-cuda -fcuda-is-device \
-// RUN: -x cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
-
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
-// RUN: -aux-triple nvptx64-nvidia-cuda \
-// RUN: -x cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
-
-int cudaConfigureCall(int, int, decltype(sizeof(int)) = 0, void* = nullptr);
-namespace QL {
-auto dg1 = [] { return 1; };
-}
-namespace QL {
-auto dg2 = [] { return 2; };
-}
-using namespace QL;
-template<typename T>
-__attribute__((global)) void f(T t) {
- t();
-}
-void g() {
- f<<<1,1>>>(dg1);
- f<<<1,1>>>(dg2);
-}
-
-// HOST: @_ZN2QL3dg1E = internal global %class.anon undef, align 1
-// HOST: @_ZN2QL3dg2E = internal global %class.anon.0 undef, align 1
-
-// DEVICE: define void @_Z1fIN2QL3dg1MUlvE_EEvT_
-// DEVICE: call noundef i32 @_ZNK2QL3dg1MUlvE_clEv
-// DEVICE: define internal noundef i32 @_ZNK2QL3dg1MUlvE_clEv
-// DEVICE: define void @_Z1fIN2QL3dg2MUlvE_EEvT_
-// DEVICE: call noundef i32 @_ZNK2QL3dg2MUlvE_clEv
-// DEVICE: define internal noundef i32 @_ZNK2QL3dg2MUlvE_clEv
-
-// HOST: define dso_local void @_Z1gv
-// HOST: call void @_Z16__device_stub__fIN2QL3dg1MUlvE_EEvT_
-// HOST: call void @_Z16__device_stub__fIN2QL3dg2MUlvE_EEvT_
-// HOST: define internal void @_Z16__device_stub__fIN2QL3dg1MUlvE_EEvT_
-// HOST: define internal void @_Z16__device_stub__fIN2QL3dg2MUlvE_EEvT_
diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu
index 6849df5a184ba..0ffe8359fd066 100644
--- a/clang/test/CodeGenCUDA/unnamed-types.cu
+++ b/clang/test/CodeGenCUDA/unnamed-types.cu
@@ -6,10 +6,13 @@
// HOST: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
// HOST: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
+// HOST: @2 = private unnamed_addr constant [25 x i8] c"_Z1fIN2QL3dg1MUlvE_EEvT_\00", align 1
+// HOST: @3 = private unnamed_addr constant [25 x i8] c"_Z1fIN2QL3dg2MUlvE_EEvT_\00", align 1
// Check that, on MSVC, the same device kernel mangling name is generated.
// MSVC: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
// MSVC: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
-
+// MSVC: @2 = private unnamed_addr constant [25 x i8] c"_Z1fIN2QL3dg1MUlvE_EEvT_\00", align 1
+// MSVC: @3 = private unnamed_addr constant [25 x i8] c"_Z1fIN2QL3dg2MUlvE_EEvT_\00", align 1
__device__ float d0(float x) {
return [](float x) { return x + 1.f; }(x);
}
@@ -42,7 +45,6 @@ void f0(float *p) {
// The inner/outer lambdas are required to be mangled following ODR but their
// linkages are still required to keep the original `internal` linkage.
-
// HOST: define internal void @_ZZ2f1PfENKUlS_E_clES_(
void f1(float *p) {
[](float *p) {
@@ -53,8 +55,39 @@ void f1(float *p) {
[] __device__ (float x, float y) { return x * y; },
[] __device__ (float x) { return x + 5.f; });
}
+// HOST: define internal void @_Z16__device_stub__fIN2QL3dg1MUlvE_EEvT_(
+// HOST: define internal void @_Z16__device_stub__fIN2QL3dg2MUlvE_EEvT_(
+
// HOST: @__hip_register_globals
// HOST: __hipRegisterFunction{{.*}}@_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_{{.*}}@0
// HOST: __hipRegisterFunction{{.*}}@_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_{{.*}}@1
+// HOST: __hipRegisterFunction{{.*}}@_Z1fIN2QL3dg1MUlvE_EEvT_{{.*}}@2
+// HOST: __hipRegisterFunction{{.*}}_Z1fIN2QL3dg2MUlvE_EEvT_{{.*}}@3
+
// MSVC: __hipRegisterFunction{{.*}}@"??$k0 at V<lambda_1>@?0???R1?0??f1@@YAXPEAM at Z@QEBA at 0@Z@@@YAXPEAMV<lambda_1>@?0???R0?0??f1@@YAX0 at Z@QEBA at 0@Z@@Z{{.*}}@0
// MSVC: __hipRegisterFunction{{.*}}@"??$k1 at V<lambda_2>@?0??f1@@YAXPEAM at Z@V<lambda_3>@?0??2 at YAX0@Z at V<lambda_4>@?0??2 at YAX0@Z@@@YAXPEAMV<lambda_2>@?0??f1@@YAX0 at Z@V<lambda_3>@?0??1 at YAX0@Z at V<lambda_4>@?0??1 at YAX0@Z@@Z{{.*}}@1
+// MSVC: __hipRegisterFunction{{.*}}@"??$f at V<lambda_1>@dg1 at QL@@@@YAXV<lambda_1>@dg1 at QL@@@Z"{{.*}}@2
+// MSVC: __hipRegisterFunction{{.*}}@"??$f at V<lambda_1>@dg2 at QL@@@@YAXV<lambda_1>@dg2 at QL@@@Z"{{.*}}@3
+
+// DEVICE: define amdgpu_kernel void @_Z1fIN2QL3dg1MUlvE_EEvT_(
+// DEVICE: call noundef i32 @_ZNK2QL3dg1MUlvE_clEv(
+// DEVICE: define internal noundef i32 @_ZNK2QL3dg1MUlvE_clEv(
+// DEVICE define amdgpu_kernel void @_Z1fIN2QL3dg2MUlvE_EEvT_(
+// DEVICE: call noundef i32 @_ZNK2QL3dg2MUlvE_clEv(
+// DEVICE: define internal noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+
+namespace QL {
+auto dg1 = [] { return 1; };
+}
+namespace QL {
+auto dg2 = [] { return 2; };
+}
+using namespace QL;
+template<typename T>
+__global__ void f(T t) {
+ t();
+}
+void g() {
+ f<<<1,1>>>(dg1);
+ f<<<1,1>>>(dg2);
+}
diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp
index 9160249ca0fc4..9d54696b0fc7a 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp
@@ -102,6 +102,7 @@ int main() {
// CHECK: declare noundef ptr @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
// CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
// CHECK: define internal void @_Z26unnamed_kernel_single_taskIZ4mainEUlvE_EvT_
+ // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE_S0_EvT0_
unnamed_kernel_single_task(
[]() {
>From ea29248e926638dcb178332a37316eed1305c740 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat at intel.com>
Date: Wed, 17 Sep 2025 13:13:52 -0700
Subject: [PATCH 4/6] Added new line
---
clang/test/CodeGenCUDA/unnamed-types.cu | 1 +
1 file changed, 1 insertion(+)
diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu
index 0ffe8359fd066..ac9b17088c180 100644
--- a/clang/test/CodeGenCUDA/unnamed-types.cu
+++ b/clang/test/CodeGenCUDA/unnamed-types.cu
@@ -45,6 +45,7 @@ void f0(float *p) {
// The inner/outer lambdas are required to be mangled following ODR but their
// linkages are still required to keep the original `internal` linkage.
+
// HOST: define internal void @_ZZ2f1PfENKUlS_E_clES_(
void f1(float *p) {
[](float *p) {
>From 1b279d853dd60311b63b803fe38280e31019dad3 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat at intel.com>
Date: Wed, 17 Sep 2025 13:15:41 -0700
Subject: [PATCH 5/6] Added new line
---
clang/test/CodeGenCUDA/unnamed-types.cu | 1 +
1 file changed, 1 insertion(+)
diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu
index ac9b17088c180..ee148e14a555e 100644
--- a/clang/test/CodeGenCUDA/unnamed-types.cu
+++ b/clang/test/CodeGenCUDA/unnamed-types.cu
@@ -13,6 +13,7 @@
// MSVC: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
// MSVC: @2 = private unnamed_addr constant [25 x i8] c"_Z1fIN2QL3dg1MUlvE_EEvT_\00", align 1
// MSVC: @3 = private unnamed_addr constant [25 x i8] c"_Z1fIN2QL3dg2MUlvE_EEvT_\00", align 1
+
__device__ float d0(float x) {
return [](float x) { return x + 1.f; }(x);
}
>From 80ca41fe372865ec19f9f09884b71097046cb800 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat at intel.com>
Date: Fri, 19 Sep 2025 06:21:55 -0700
Subject: [PATCH 6/6] Fix CodeGenSYCL LIT tests.
---
clang/lib/Sema/SemaLambda.cpp | 7 ++++
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 2 +-
clang/test/CodeGenSYCL/unique_stable_name.cpp | 41 +++++++++----------
.../unique_stable_name_windows_diff.cpp | 6 ++-
4 files changed, 33 insertions(+), 23 deletions(-)
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index f82069b0b4e2b..3c241a59b6a77 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -356,6 +356,13 @@ Sema::getCurrentMangleNumberContext(const DeclContext *DC) {
return std::make_tuple(&Context.getManglingNumberContext(DC), nullptr);
}
+ if (ManglingContextDecl) {
+ // Lambdas defined in the initializer of a local variable are mangled
+ // in the enclosing function context.
+ if (isa<VarDecl>(ManglingContextDecl) &&
+ !cast<VarDecl>(ManglingContextDecl)->hasGlobalStorage())
+ ManglingContextDecl = nullptr;
+ }
return std::make_tuple(nullptr, ManglingContextDecl);
}
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index 5b3b73812d286..cd1d4d801951d 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -90,7 +90,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
-// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task at V<lambda_1>@lambda@?0??main@@9 at V12?0??3 at 9@@@YAXV<lambda_1>@lambda@?0??main@@9@@Z"(i32 %kernelFunc.coerce) #{{[0-9]+}} {
+// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task at V<lambda_1>@?0??main@@9 at V1?0??2 at 9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %kernelFunc.coerce) #{{[0-9]+}} {
// CHECK-HOST-WINDOWS-NEXT: entry:
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon, align 4
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %kernelFunc, i32 0, i32 0
diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp
index 9d54696b0fc7a..3ab7e3b8f2e7a 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp
@@ -3,20 +3,20 @@
// CHECK: @[[INT1:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
// CHECK: @[[STRING:[^\w]+]] = private unnamed_addr constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00",
// CHECK: @[[INT2:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
-// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE_\00"
-// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE0_\00"
-// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE1_\00"
-// CHECK: @{{.*}} = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE2_\00"
-// CHECK: @{{.*}} = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE3_\00"
-// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE4_\00"
-// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE5_\00"
+// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00"
+// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE0_\00"
+// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE1_\00"
+// CHECK: @{{.*}} = private unnamed_addr constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE2_\00", align 1
+// CHECK: @{{.*}} = private unnamed_addr constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE3_\00", align 1
+// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE4_\00"
+// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE5_\00"
// CHECK: @[[INT3:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
-// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE]] c"_ZTSZZ4mainENKUlvE_clEvEUlvE_\00"
+// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [[LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00"
// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE_\00",
-// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_EvvEUlvE_\00",
// CHECK: @[[LAMBDA_NO_DEP:[^\w]+]] = private unnamed_addr constant [[NO_DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ13lambda_no_depIidEvT_T0_EUlidE_\00",
-// CHECK: @[[LAMBDA_TWO_DEP:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA1_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_EvvEUlvE_\00",
-// CHECK: @[[LAMBDA_TWO_DEP2:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA2_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE_clEvEUldE_ZZ4mainENKS0_clEvEUliE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_TWO_DEP:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA1_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_TWO_DEP2:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA2_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_EvvEUlvE_\00",
extern "C" void puts(const char *) {}
@@ -101,8 +101,7 @@ int main() {
// CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
// CHECK: declare noundef ptr @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
// CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
- // CHECK: define internal void @_Z26unnamed_kernel_single_taskIZ4mainEUlvE_EvT_
- // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE_S0_EvT0_
+ // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvT0_
unnamed_kernel_single_task(
[]() {
@@ -125,13 +124,13 @@ int main() {
// CHECK: call void @_Z14template_paramIiEvv
template_param<decltype(x)>();
- // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE_clEvEUlvE_Evv
+ // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
lambda_in_dependent_function<int>();
// CHECK: call void @_Z28lambda_in_dependent_functionIiEvv
lambda_in_dependent_function<decltype(x)>();
- // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_Evv
+ // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
lambda_no_dep<int, double>(3, 5.5);
// CHECK: call void @_Z13lambda_no_depIidEvT_T0_(i32 noundef 3, double noundef 5.500000e+00)
@@ -141,30 +140,30 @@ int main() {
auto y = [](int a) { return a; };
auto z = [](double b) { return b; };
lambda_two_dep<decltype(y), decltype(z)>();
- // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+ // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
lambda_two_dep<decltype(z), decltype(y)>();
- // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+ // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
});
}
// CHECK: define linkonce_odr void @_Z14template_paramIiEvv
// CHECK: call void @puts(ptr noundef @[[INT3]])
-// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE_clEvEUlvE_Evv
+// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA]])
// CHECK: define linkonce_odr void @_Z28lambda_in_dependent_functionIiEvv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_IN_DEP_INT]])
-// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE_clEvEUlvE_Evv
+// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_IN_DEP_X]])
// CHECK: define linkonce_odr void @_Z13lambda_no_depIidEvT_T0_(i32 noundef %a, double noundef %b)
// CHECK: call void @puts(ptr noundef @[[LAMBDA_NO_DEP]])
-// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_TWO_DEP]])
-// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
// CHECK: call void @puts(ptr noundef @[[LAMBDA_TWO_DEP2]])
diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
index 33c6d461aeb02..8e34176909080 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -51,6 +51,9 @@ int main() {
// HOST: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUliE_\00"
// HOST: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUldE_\00"
+ // DEVICE: private unnamed_addr addrspace(1) constant [17 x i8] c"_ZTSZ4mainEUlvE_\00"
+ // DEVICE: private unnamed_addr addrspace(1) constant [17 x i8] c"_ZTSZ4mainEUliE_\00"
+ // DEVICE: private unnamed_addr addrspace(1) constant [17 x i8] c"_ZTSZ4mainEUldE_\00"
// DEVICE: define dso_local spir_kernel void @_ZTSZ4mainE2K1
// DEVICE: call spir_func void @_ZZ4mainENKUlvE_clEv
// DEVICE: define internal spir_func void @_ZZ4mainENKUlvE_clEv
@@ -61,5 +64,6 @@ int main() {
// DEVICE: call spir_func void @_ZZ4mainENKUldE_clEd
// DEVICE: define internal spir_func void @_ZZ4mainENKUldE_clEd
// DEVICE: define dso_local spir_kernel void @_ZTSZ4mainE2K4
- // DEVICE: call spir_func void @_ZZ4mainENKUlvE_clEv
+ // DEVICE: call spir_func void @_ZZ4mainENKUlvE0_clEv
+ // DEVICE: define internal spir_func void @_ZZ4mainENKUlvE0_clEv
}
More information about the cfe-commits
mailing list