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/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 &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..53242448954ab 100644 --- a/clang/lib/Sema/SemaLambda.cpp +++ b/clang/lib/Sema/SemaLambda.cpp @@ -356,7 +356,14 @@ Sema::getCurrentMangleNumberContext(const DeclContext *DC) { return std::make_tuple(&Context.getManglingNumberContext(DC), nullptr); } - return std::make_tuple(nullptr, nullptr); + if (ManglingContextDecl) { + // Lambdas defined in the initializer of a local variable are mangled + // in the enclosing function context. + if (isa(ManglingContextDecl) && + !cast(ManglingContextDecl)->hasGlobalStorage()) + ManglingContextDecl = nullptr; + } + return std::make_tuple(nullptr, ManglingContextDecl); } case NonInlineInModulePurview: 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 + auto dg_template = [] { return N; }; +} + +using namespace QL; +template +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@1@A" = internal global %class.anon undef, align 1 +// MSVC: @"?dg2@QL@@3V@1@A" = internal global %class.anon.0 undef, align 1 +// MSVC: @"?dg_inline1@QL@@3V@01@A" = linkonce_odr dso_local global %class.anon.2 undef, comdat, align 1 +// MSVC: @"??$dg_template@$02@QL@@3V@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@dg_inline1@QL@@@@YAXV@dg_inline1@QL@@@Z" +// MSVC: call noundef i32 @"??R@dg_inline1@QL@@QEBA?A?@@XZ" +// MSVC: define linkonce_odr dso_local void @"??$f@V@?$dg_template@$02@QL@@@@YAXV@?$dg_template@$02@QL@@@Z" +// MSVC: call noundef i32 @"??R@?$dg_template@$02@QL@@QEBA?A?@@XZ" +// MSVC: define internal noundef i32 @"??R@QL@@QEBA?A?@@XZ" +// MSVC: define internal noundef i32 @"??R@QL@@QEBA?A?@@XZ" +// MSVC: define linkonce_odr dso_local noundef i32 @"??R@dg_inline1@QL@@QEBA?A?@@XZ" +// MSVC: define linkonce_odr dso_local noundef i32 @"??R@?$dg_template@$02@QL@@QEBA?A?@@XZ" diff --git a/clang/test/CodeGenCUDA/anon-ns.cu b/clang/test/CodeGenCUDA/anon-ns.cu index d931f31d0207c..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_1UlvE_EEvT_\.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_1UlvE_EEvT___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 diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu index 6849df5a184ba..8b7e0510da9e0 100644 --- a/clang/test/CodeGenCUDA/unnamed-types.cu +++ b/clang/test/CodeGenCUDA/unnamed-types.cu @@ -1,14 +1,22 @@ -// 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" // 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 +// 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); @@ -20,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 __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 @@ -53,8 +61,58 @@ 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: 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 -// MSVC: __hipRegisterFunction{{.*}}@"??$k0@V@?0???R1?0??f1@@YAXPEAM@Z@QEBA@0@Z@@@YAXPEAMV@?0???R0?0??f1@@YAX0@Z@QEBA@0@Z@@Z{{.*}}@0 +// 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@?0???R1?0??f1@@YAXPEAM@Z@QEBA?A?@@0@Z@@@YAXPEAMV@?0???R0?0??f1@@YAX0@Z@QEBA?A?@@0@Z@@Z{{.*}}@0 // MSVC: __hipRegisterFunction{{.*}}@"??$k1@V@?0??f1@@YAXPEAM@Z@V@?0??2@YAX0@Z@V@?0??2@YAX0@Z@@@YAXPEAMV@?0??f1@@YAX0@Z@V@?0??1@YAX0@Z@V@?0??1@YAX0@Z@@Z{{.*}}@1 +// MSVC: __hipRegisterFunction{{.*}}@"??$f@V@dg1@QL@@@@YAXV@dg1@QL@@@Z"{{.*}}@2 +// MSVC: __hipRegisterFunction{{.*}}@"??$f@V@dg2@QL@@@@YAXV@dg2@QL@@@Z"{{.*}}@3 +// MSVC: __hipRegisterFunction{{.*}}@"??$f@V@dg_inline1@QL@@@@YAXV@dg_inline1@QL@@@Z"{{.*}}@4 +// MSVC: __hipRegisterFunction{{.*}}@"??$f@V@?$dg_template@$02@QL@@@@YAXV@?$dg_template@$02@QL@@@Z"{{.*}}@5 + + +// 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 +// 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 +auto dg_template = [] { return N; }; +} +using namespace QL; +template +__global__ void f(T 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/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 14366a092a1fe..c681a3a6c814d 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 @@ -47,7 +47,24 @@ 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: 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 + // 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 @_ZZ4mainENKUlvE0_clEv + // DEVICE: define internal spir_func void @_ZZ4mainENKUlvE0_clEv } diff --git a/clang/test/CodeGenSYCL/unnamed-types.cpp b/clang/test/CodeGenSYCL/unnamed-types.cpp new file mode 100644 index 0000000000000..9a702c48dfdc9 --- /dev/null +++ b/clang/test/CodeGenSYCL/unnamed-types.cpp @@ -0,0 +1,74 @@ +// 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 -triple spirv64-unknown-unknown \ +// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST + +// 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; }; + inline auto dg_inline1 = [] { return 1; }; +} + +namespace QL { + auto dg2 = [] { return 2; }; + template + auto dg_template = [] { return N; }; +} + +using namespace QL; +template +[[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 +// 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 +// 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 +// 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 +// 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