From b197e92e844472ccef1999cac653a76109183ee4 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat@intel.com>
Date: Tue, 16 Sep 2025 08:47:40 -0700
Subject: [PATCH 01/10] 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.ammarguellat@intel.com>
Date: Wed, 17 Sep 2025 07:59:38 -0700
Subject: [PATCH 02/10] 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@intel.com>
Date: Wed, 17 Sep 2025 13:07:07 -0700
Subject: [PATCH 03/10] 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@V<lambda_1>@?0???R1?0??f1@@YAXPEAM@Z@QEBA@0@Z@@@YAXPEAMV<lambda_1>@?0???R0?0??f1@@YAX0@Z@QEBA@0@Z@@Z{{.*}}@0
 // MSVC: __hipRegisterFunction{{.*}}@"??$k1@V<lambda_2>@?0??f1@@YAXPEAM@Z@V<lambda_3>@?0??2@YAX0@Z@V<lambda_4>@?0??2@YAX0@Z@@@YAXPEAMV<lambda_2>@?0??f1@@YAX0@Z@V<lambda_3>@?0??1@YAX0@Z@V<lambda_4>@?0??1@YAX0@Z@@Z{{.*}}@1
+// MSVC: __hipRegisterFunction{{.*}}@"??$f@V<lambda_1>@dg1@QL@@@@YAXV<lambda_1>@dg1@QL@@@Z"{{.*}}@2
+// MSVC: __hipRegisterFunction{{.*}}@"??$f@V<lambda_1>@dg2@QL@@@@YAXV<lambda_1>@dg2@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@intel.com>
Date: Wed, 17 Sep 2025 13:13:52 -0700
Subject: [PATCH 04/10] 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@intel.com>
Date: Wed, 17 Sep 2025 13:15:41 -0700
Subject: [PATCH 05/10] 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@intel.com>
Date: Fri, 19 Sep 2025 06:21:55 -0700
Subject: [PATCH 06/10] 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@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:      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-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
 }

From b9a581ab101959f8fd9d45d5ff464ffcf7996ef1 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat@intel.com>
Date: Fri, 19 Sep 2025 06:25:17 -0700
Subject: [PATCH 07/10] Added a new line for clarity

---
 clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
index 8e34176909080..c681a3a6c814d 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -54,6 +54,7 @@ int main() {
   // 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

From 8fb0464710063047340003520d98badd24794a86 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat@intel.com>
Date: Fri, 19 Sep 2025 06:30:06 -0700
Subject: [PATCH 08/10] Fix format

---
 clang/lib/Sema/SemaLambda.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index 3c241a59b6a77..53242448954ab 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -361,7 +361,7 @@ Sema::getCurrentMangleNumberContext(const DeclContext *DC) {
       // in the enclosing function context.
       if (isa<VarDecl>(ManglingContextDecl) &&
           !cast<VarDecl>(ManglingContextDecl)->hasGlobalStorage())
-          ManglingContextDecl = nullptr;
+        ManglingContextDecl = nullptr;
     }
     return std::make_tuple(nullptr, ManglingContextDecl);
   }

From c96de5a64f57ab14ff8d0b6c223d144ad2bc0cda Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat@intel.com>
Date: Sat, 20 Sep 2025 05:55:57 -0700
Subject: [PATCH 09/10] Fix LIT tests (after Tom's solution)

---
 clang/test/CodeGenCUDA/unnamed-types.cu  | 35 ++++++++++++++++++++----
 clang/test/CodeGenSYCL/unnamed-types.cpp | 31 +++++++++++++++++++--
 2 files changed, 57 insertions(+), 9 deletions(-)

diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu
index ee148e14a555e..8b7e0510da9e0 100644
--- a/clang/test/CodeGenCUDA/unnamed-types.cu
+++ b/clang/test/CodeGenCUDA/unnamed-types.cu
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
-// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-pc-windows-msvc -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=MSVC
-// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
+// RUN: %clang_cc1 -std=c++17 -x hip -triple x86_64-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
+// RUN: %clang_cc1 -std=c++17 -x hip -triple x86_64-pc-windows-msvc -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=MSVC
+// RUN: %clang_cc1 -std=c++17 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE
 
 #include "Inputs/cuda.h"
 
@@ -8,11 +8,15 @@
 // 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
+// HOST: @4 = private unnamed_addr constant [33 x i8] c"_Z1fIN2QL10dg_inline1MUlvE_EEvT_\00", align 1
+// HOST: @5 = private unnamed_addr constant [40 x i8] c"_Z1fIN2QL11dg_templateILi3EEMUlvE_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
+// MSVC: @4 = private unnamed_addr constant [33 x i8] c"_Z1fIN2QL10dg_inline1MUlvE_EEvT_\00", align 1
+// MSVC: @5 = private unnamed_addr constant [40 x i8] c"_Z1fIN2QL11dg_templateILi3EEMUlvE_EEvT_\00", align 1
 
 __device__ float d0(float x) {
   return [](float x) { return x + 1.f; }(x);
@@ -24,13 +28,13 @@ __device__ float d1(float x) {
 
 // DEVICE: amdgpu_kernel void @_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_(
 // DEVICE: define internal noundef float @_ZZZ2f1PfENKUlS_E_clES_ENKUlfE_clEf(
+// DEVICE: define amdgpu_kernel void @_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_(
+// DEVICE: define internal noundef float @_ZZ2f1PfENKUlfE_clEf(
 template <typename F>
 __global__ void k0(float *p, F f) {
   p[0] = f(p[0]) + d0(p[1]) + d1(p[2]);
 }
 
-// DEVICE: amdgpu_kernel void @_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_(
-// DEVICE: define internal noundef float @_ZZ2f1PfENKUlfE_clEf(
 // DEVICE: define internal noundef float @_ZZ2f1PfENKUlffE_clEff(
 // DEVICE: define internal noundef float @_ZZ2f1PfENKUlfE0_clEf(
 template <typename F0, typename F1, typename F2>
@@ -59,17 +63,25 @@ void f1(float *p) {
 }
 // HOST: define internal void @_Z16__device_stub__fIN2QL3dg1MUlvE_EEvT_(
 // HOST: define internal void @_Z16__device_stub__fIN2QL3dg2MUlvE_EEvT_(
+// HOST: define linkonce_odr void @_Z16__device_stub__fIN2QL10dg_inline1MUlvE_EEvT_
+// HOST: define internal void @_Z17__device_stub__k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_
+// HOST: define linkonce_odr void @_ZN4dim3C2Ejjj
 
 // 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
+// HOST: __hipRegisterFunction{{.*}}_Z1fIN2QL10dg_inline1MUlvE_EEvT_{{.*}}@4
+// HOST: __hipRegisterFunction{{.*}}_Z1fIN2QL11dg_templateILi3EEMUlvE_EEvT_{{.*}}@5
 
-// MSVC: __hipRegisterFunction{{.*}}@"??$k0@V<lambda_1>@?0???R1?0??f1@@YAXPEAM@Z@QEBA@0@Z@@@YAXPEAMV<lambda_1>@?0???R0?0??f1@@YAX0@Z@QEBA@0@Z@@Z{{.*}}@0
+// MSVC: __hipRegisterFunction{{.*}}@"??$k0@V<lambda_1>@?0???R1?0??f1@@YAXPEAM@Z@QEBA?A?<auto>@@0@Z@@@YAXPEAMV<lambda_1>@?0???R0?0??f1@@YAX0@Z@QEBA?A?<auto>@@0@Z@@Z{{.*}}@0
 // MSVC: __hipRegisterFunction{{.*}}@"??$k1@V<lambda_2>@?0??f1@@YAXPEAM@Z@V<lambda_3>@?0??2@YAX0@Z@V<lambda_4>@?0??2@YAX0@Z@@@YAXPEAMV<lambda_2>@?0??f1@@YAX0@Z@V<lambda_3>@?0??1@YAX0@Z@V<lambda_4>@?0??1@YAX0@Z@@Z{{.*}}@1
 // MSVC: __hipRegisterFunction{{.*}}@"??$f@V<lambda_1>@dg1@QL@@@@YAXV<lambda_1>@dg1@QL@@@Z"{{.*}}@2
 // MSVC: __hipRegisterFunction{{.*}}@"??$f@V<lambda_1>@dg2@QL@@@@YAXV<lambda_1>@dg2@QL@@@Z"{{.*}}@3
+// MSVC: __hipRegisterFunction{{.*}}@"??$f@V<lambda_1>@dg_inline1@QL@@@@YAXV<lambda_1>@dg_inline1@QL@@@Z"{{.*}}@4
+// MSVC: __hipRegisterFunction{{.*}}@"??$f@V<lambda_1>@?$dg_template@$02@QL@@@@YAXV<lambda_1>@?$dg_template@$02@QL@@@Z"{{.*}}@5
+
 
 // DEVICE: define amdgpu_kernel void @_Z1fIN2QL3dg1MUlvE_EEvT_(
 // DEVICE: call noundef i32 @_ZNK2QL3dg1MUlvE_clEv(
@@ -77,12 +89,21 @@ void f1(float *p) {
 // DEVICE  define amdgpu_kernel void @_Z1fIN2QL3dg2MUlvE_EEvT_(
 // DEVICE: call noundef i32 @_ZNK2QL3dg2MUlvE_clEv(
 // DEVICE: define internal noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+// DEVICE: define amdgpu_kernel void @_Z1fIN2QL10dg_inline1MUlvE_EEvT_
+// DEVICE: call noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// DEVICE: define linkonce_odr noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// DEVICE: define amdgpu_kernel void @_Z1fIN2QL11dg_templateILi3EEMUlvE_EEvT_
+// DEVICE: call noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv
+// DEVICE: define linkonce_odr noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv
 
 namespace QL {
 auto dg1 = [] { return 1; };
+inline auto dg_inline1 = [] { return 1; };
 }
 namespace QL {
 auto dg2 = [] { return 2; };
+template<int N>
+auto dg_template = [] { return N; };
 }
 using namespace QL;
 template<typename T>
@@ -92,4 +113,6 @@ __global__ void f(T t) {
 void g() {
   f<<<1,1>>>(dg1);
   f<<<1,1>>>(dg2);
+  f<<<1,1>>>(dg_inline1);
+  f<<<1,1>>>(dg_template<3>);
 }
diff --git a/clang/test/CodeGenSYCL/unnamed-types.cpp b/clang/test/CodeGenSYCL/unnamed-types.cpp
index 64b8c166f8a16..9a702c48dfdc9 100644
--- a/clang/test/CodeGenSYCL/unnamed-types.cpp
+++ b/clang/test/CodeGenSYCL/unnamed-types.cpp
@@ -9,23 +9,33 @@
 // RUN: %s -o - | FileCheck %s --check-prefix=MSVC
 
 namespace QL {
-auto dg1 = [] { return 1; };
+  auto dg1 = [] { return 1; };
+  inline auto dg_inline1 = [] { return 1; };
 }
+
 namespace QL {
-auto dg2 = [] { return 2; };
+  auto dg2 = [] { return 2; };
+  template<int N>
+  auto dg_template = [] { return N; };
 }
+
 using namespace QL;
 template<typename T>
 [[clang::sycl_kernel_entry_point(T)]] void f(T t) {
   t();
 }
+
 void g() {
   f(dg1);
   f(dg2);
+  f(dg_inline1);
+  f(dg_template<3>);
 }
 
 // HOST: @_ZN2QL3dg1E = internal global %class.anon undef, align 1
 // HOST: @_ZN2QL3dg2E = internal global %class.anon.0 undef, align 1
+// HOST: @_ZN2QL10dg_inline1E = linkonce_odr global %class.anon.2 undef, comdat, align 1
+// HOST: @_ZN2QL11dg_templateILi3EEE = linkonce_odr global %class.anon.4 undef, comdat, align 1
 
 // DEVICE: define spir_kernel void @_ZTSN2QL3dg1MUlvE_E
 // DEVICE: call spir_func noundef i32 @_ZNK2QL3dg1MUlvE_clEv
@@ -33,12 +43,22 @@ void g() {
 // DEVICE: define spir_kernel void @_ZTSN2QL3dg2MUlvE_E
 // DEVICE: call spir_func noundef i32 @_ZNK2QL3dg2MUlvE_clEv
 // DEVICE: define internal spir_func noundef i32 @_ZNK2QL3dg2MUlvE_clEv
+// DEVICE: define spir_kernel void @_ZTSN2QL10dg_inline1MUlvE_E
+// DEVICE: call spir_func noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// DEVICE: define linkonce_odr spir_func noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// DEVICE: define spir_kernel void @_ZTSN2QL11dg_templateILi3EEMUlvE_E
+// DEVICE: call spir_func noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv
+// DEVICE: define linkonce_odr spir_func noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv
 
 // HOST: define spir_func void @_Z1gv
 // HOST: call spir_func void @_Z1fIN2QL3dg1MUlvE_EEvT_
 // HOST: call spir_func void @_Z1fIN2QL3dg2MUlvE_EEvT_
+// HOST: call spir_func void @_Z1fIN2QL10dg_inline1MUlvE_EEvT_
+// HOST: call spir_func void @_Z1fIN2QL11dg_templateILi3EEMUlvE_EEvT_
 // HOST: define internal spir_func void @_Z1fIN2QL3dg1MUlvE_EEvT
 // HOST: define internal spir_func void @_Z1fIN2QL3dg2MUlvE_EEvT_
+// HOST: define linkonce_odr spir_func void @_Z1fIN2QL10dg_inline1MUlvE_EEvT_
+// HOST: define linkonce_odr spir_func void @_Z1fIN2QL11dg_templateILi3EEMUlvE_EEvT_
 
 // MSVC: define dso_local spir_kernel void @_ZTSN2QL3dg1MUlvE_E
 // MSVC: call spir_func noundef i32 @_ZNK2QL3dg1MUlvE_clEv
@@ -46,4 +66,9 @@ void g() {
 // 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
-
+// MSVC: define dso_local spir_kernel void @_ZTSN2QL10dg_inline1MUlvE_E
+// MSVC: call spir_func noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// MSVC: define linkonce_odr spir_func noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// MSVC: define dso_local spir_kernel void @_ZTSN2QL11dg_templateILi3EEMUlvE_E
+// MSVC: call spir_func noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv
+// MSVC: define linkonce_odr spir_func noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv

From 7de57f29a23222535b0e47e51489814af46ba194 Mon Sep 17 00:00:00 2001
From: Zahira Ammarguellat <zahira.ammarguellat@intel.com>
Date: Mon, 22 Sep 2025 06:56:43 -0700
Subject: [PATCH 10/10] Add C++ LIT test

---
 clang/test/CodeGen/unnamed-types.cpp | 61 ++++++++++++++++++++++++++++
 1 file changed, 61 insertions(+)
 create mode 100644 clang/test/CodeGen/unnamed-types.cpp

diff --git a/clang/test/CodeGen/unnamed-types.cpp b/clang/test/CodeGen/unnamed-types.cpp
new file mode 100644
index 0000000000000..b47552278f479
--- /dev/null
+++ b/clang/test/CodeGen/unnamed-types.cpp
@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 -O0 -triple x86_64-unknown-unknown \
+// RUN: -emit-llvm  %s -o - | FileCheck %s
+
+// RUN: %clang_cc1 -O0 -triple x86_64-pc-windows-msvc \
+// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefix=MSVC
+
+namespace QL {
+  auto dg1 = [] { return 1; };
+  inline auto dg_inline1 = [] { return 1; };
+}
+
+namespace QL {
+  auto dg2 = [] { return 2; };
+  template<int N>
+  auto dg_template = [] { return N; };
+}
+
+using namespace QL;
+template<typename T>
+void f(T t) {
+  t();
+}
+
+void g() {
+  f(dg1);
+  f(dg2);
+  f(dg_inline1);
+  f(dg_template<3>);
+}
+
+// CHECK: @_ZN2QL3dg1E = internal global %class.anon undef, align 1
+// CHECK: @_ZN2QL3dg2E = internal global %class.anon.0 undef, align 1
+// CHECK: @_ZN2QL10dg_inline1E = linkonce_odr global %class.anon.2 undef, comdat, align 1
+// CHECK: @_ZN2QL11dg_templateILi3EEE = linkonce_odr global %class.anon.4 undef, comdat, align 1
+
+// MSVC: @"?dg1@QL@@3V<lambda_0>@1@A" = internal global %class.anon undef, align 1
+// MSVC: @"?dg2@QL@@3V<lambda_1>@1@A" = internal global %class.anon.0 undef, align 1
+// MSVC: @"?dg_inline1@QL@@3V<lambda_1>@01@A" = linkonce_odr dso_local global %class.anon.2 undef, comdat, align 1
+// MSVC: @"??$dg_template@$02@QL@@3V<lambda_1>@01@A" = linkonce_odr dso_local global %class.anon.4 undef, comdat, align 1
+
+
+// CHECK: define internal void @"_Z1fIN2QL3$_0EEvT_"
+// CHECK: call noundef i32 @"_ZNK2QL3$_0clEv"
+// CHECK: define internal void @"_Z1fIN2QL3$_1EEvT_"
+// CHECK: define linkonce_odr void @_Z1fIN2QL10dg_inline1MUlvE_EEvT_
+// CHECK: call noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// CHECK: define linkonce_odr void @_Z1fIN2QL11dg_templateILi3EEMUlvE_EEvT_
+// CHECK: call noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv
+// CHECK: define internal noundef i32 @"_ZNK2QL3$_0clEv"
+// CHECK: define internal noundef i32 @"_ZNK2QL3$_1clEv"
+// CHECK: define linkonce_odr noundef i32 @_ZNK2QL10dg_inline1MUlvE_clEv
+// CHECK: define linkonce_odr noundef i32 @_ZNK2QL11dg_templateILi3EEMUlvE_clEv
+
+// MSVC: define linkonce_odr dso_local void @"??$f@V<lambda_1>@dg_inline1@QL@@@@YAXV<lambda_1>@dg_inline1@QL@@@Z"
+// MSVC: call noundef i32 @"??R<lambda_1>@dg_inline1@QL@@QEBA?A?<auto>@@XZ"
+// MSVC: define linkonce_odr dso_local void @"??$f@V<lambda_1>@?$dg_template@$02@QL@@@@YAXV<lambda_1>@?$dg_template@$02@QL@@@Z"
+// MSVC: call noundef i32 @"??R<lambda_1>@?$dg_template@$02@QL@@QEBA?A?<auto>@@XZ"
+// MSVC: define internal noundef i32 @"??R<lambda_0>@QL@@QEBA?A?<auto>@@XZ"
+// MSVC: define internal noundef i32 @"??R<lambda_1>@QL@@QEBA?A?<auto>@@XZ"
+// MSVC: define linkonce_odr dso_local noundef i32 @"??R<lambda_1>@dg_inline1@QL@@QEBA?A?<auto>@@XZ"
+// MSVC: define linkonce_odr dso_local noundef i32 @"??R<lambda_1>@?$dg_template@$02@QL@@QEBA?A?<auto>@@XZ"