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.ammarguel...@intel.com>
Date: Tue, 16 Sep 2025 08:47:40 -0700
Subject: [PATCH 1/2] 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@V<lambda_1>@?0??main@@9@V1?0??2@9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32
%kernelFunc.coerce) #{{[0-9]+}} {
+// CHECK-HOST-WINDOWS: define internal void
@"??$kernel_single_task@V<lambda_1>@lambda@?0??main@@9@V12?0??3@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.ammarguel...@intel.com>
Date: Wed, 17 Sep 2025 07:59:38 -0700
Subject: [PATCH 2/2] 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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
