diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h --- a/clang/lib/CodeGen/CGOpenMPRuntime.h +++ b/clang/lib/CodeGen/CGOpenMPRuntime.h @@ -319,7 +319,8 @@ /// address \a Addr, size \a Size, and flags \a Flags. virtual void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr, uint64_t Size, int32_t Flags, - llvm::GlobalValue::LinkageTypes Linkage); + llvm::GlobalValue::LinkageTypes Linkage, + StringRef MangledName); /// Helper to emit outlined function for 'target' directive. /// \param D Directive to emit. @@ -661,19 +662,22 @@ /// Type of the global variable. CharUnits VarSize; llvm::GlobalValue::LinkageTypes Linkage; + StringRef OrigName; public: OffloadEntryInfoDeviceGlobalVar() : OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar) {} explicit OffloadEntryInfoDeviceGlobalVar(unsigned Order, - OMPTargetGlobalVarEntryKind Flags) - : OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar, Order, Flags) {} + OMPTargetGlobalVarEntryKind Flags, + StringRef OrigName) + : OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar, Order, Flags), + OrigName(OrigName) {} explicit OffloadEntryInfoDeviceGlobalVar( unsigned Order, llvm::Constant *Addr, CharUnits VarSize, OMPTargetGlobalVarEntryKind Flags, - llvm::GlobalValue::LinkageTypes Linkage) + llvm::GlobalValue::LinkageTypes Linkage, StringRef OrigName) : OffloadEntryInfo(OffloadingEntryInfoDeviceGlobalVar, Order, Flags), - VarSize(VarSize), Linkage(Linkage) { + VarSize(VarSize), Linkage(Linkage), OrigName(OrigName) { setAddress(Addr); } @@ -684,16 +688,21 @@ static bool classof(const OffloadEntryInfo *Info) { return Info->getKind() == OffloadingEntryInfoDeviceGlobalVar; } + StringRef getOrigName() const { return OrigName; } + void setOrigName( StringRef Name ) { OrigName = Name; } }; /// Initialize device global variable entry. void initializeDeviceGlobalVarEntryInfo(StringRef Name, OMPTargetGlobalVarEntryKind Flags, - unsigned Order); + unsigned Order, + StringRef OrigName); + void enterDeviceGlobalVarMangledName(StringRef OrigName, StringRef Name); /// Register device global variable entry. void - registerDeviceGlobalVarEntryInfo(StringRef VarName, llvm::Constant *Addr, + registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName, + llvm::Constant *Addr, CharUnits VarSize, OMPTargetGlobalVarEntryKind Flags, llvm::GlobalValue::LinkageTypes Linkage); @@ -707,7 +716,8 @@ OffloadDeviceGlobalVarEntryInfoActTy; void actOnDeviceGlobalVarEntriesInfo( const OffloadDeviceGlobalVarEntryInfoActTy &Action); - + /// Return host mangled name + StringRef getOffloadEntryHostMangledName(StringRef VarName); private: // Storage for target region entries kind. The storage is to be indexed by // file ID, device ID, parent function name and line number. @@ -726,6 +736,8 @@ typedef llvm::StringMap OffloadEntriesDeviceGlobalVarTy; OffloadEntriesDeviceGlobalVarTy OffloadEntriesDeviceGlobalVar; + /// indexed by original name + llvm::StringMap OffloadEntriesDeviceGlobalVarNameMap; }; OffloadEntriesInfoManagerTy OffloadEntriesInfoManager; @@ -1924,6 +1936,9 @@ /// Returns true if the variable is a local variable in untied task. bool isLocalVarInUntiedTask(CodeGenFunction &CGF, const VarDecl *VD) const; + + /// Returns the mangled name for declare target global + StringRef getHostMangledDeclareTargetGlobal( StringRef VarName ); }; /// Class supports emissionof SIMD-only code. diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -3032,19 +3032,35 @@ Action(D.first, F.first, P.first(), L.first, L.second); } +void CGOpenMPRuntime::OffloadEntriesInfoManagerTy:: + enterDeviceGlobalVarMangledName(StringRef OrigName, StringRef Name) { + OffloadEntriesDeviceGlobalVarNameMap.try_emplace(OrigName, Name.str()); +} + void CGOpenMPRuntime::OffloadEntriesInfoManagerTy:: initializeDeviceGlobalVarEntryInfo(StringRef Name, OMPTargetGlobalVarEntryKind Flags, - unsigned Order) { + unsigned Order, + StringRef OrigName) { assert(CGM.getLangOpts().OpenMPIsDevice && "Initialization of entries is " "only required for the device " "code generation."); - OffloadEntriesDeviceGlobalVar.try_emplace(Name, Order, Flags); + OffloadEntriesDeviceGlobalVar.try_emplace(Name, Order, Flags, OrigName); ++OffloadingEntriesNum; } +StringRef CGOpenMPRuntime::OffloadEntriesInfoManagerTy:: + getOffloadEntryHostMangledName(StringRef VarName) { + if ( OffloadEntriesDeviceGlobalVarNameMap.find(VarName) != + OffloadEntriesDeviceGlobalVarNameMap.end() ) { + return OffloadEntriesDeviceGlobalVarNameMap[VarName]; + } + return StringRef(); +} + void CGOpenMPRuntime::OffloadEntriesInfoManagerTy:: - registerDeviceGlobalVarEntryInfo(StringRef VarName, llvm::Constant *Addr, + registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName, + llvm::Constant *Addr, CharUnits VarSize, OMPTargetGlobalVarEntryKind Flags, llvm::GlobalValue::LinkageTypes Linkage) { @@ -3063,6 +3079,7 @@ Entry.setVarSize(VarSize); Entry.setLinkage(Linkage); Entry.setAddress(Addr); + Entry.setOrigName(OrigName); } else { if (hasDeviceGlobalVarEntryInfo(VarName)) { auto &Entry = OffloadEntriesDeviceGlobalVar[VarName]; @@ -3075,7 +3092,7 @@ return; } OffloadEntriesDeviceGlobalVar.try_emplace( - VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage); + VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage, OrigName); ++OffloadingEntriesNum; } } @@ -3090,8 +3107,10 @@ void CGOpenMPRuntime::createOffloadEntry( llvm::Constant *ID, llvm::Constant *Addr, uint64_t Size, int32_t Flags, - llvm::GlobalValue::LinkageTypes Linkage) { - OMPBuilder.emitOffloadingEntry(ID, Addr->getName(), Size, Flags); + llvm::GlobalValue::LinkageTypes Linkage, + StringRef MangledName) { + StringRef VarName = (MangledName.empty()) ? Addr->getName() : MangledName; + OMPBuilder.emitOffloadingEntry(ID, VarName, Size, Flags); } void CGOpenMPRuntime::createOffloadEntriesAndInfoMetadata() { @@ -3184,10 +3203,11 @@ // - Entry 1 -> Mangled name of the variable. // - Entry 2 -> Declare target kind. // - Entry 3 -> Order the entry was created. + // - Entry 4 -> Original name of the variable. // The first element of the metadata node is the kind. llvm::Metadata *Ops[] = { GetMDInt(E.getKind()), GetMDString(MangledName), - GetMDInt(E.getFlags()), GetMDInt(E.getOrder())}; + GetMDInt(E.getFlags()), GetMDInt(E.getOrder()), GetMDString(E.getOrigName())}; // Save this entry in the right position of the ordered entries array. OrderedEntries[E.getOrder()] = @@ -3218,7 +3238,8 @@ continue; } createOffloadEntry(CE->getID(), CE->getAddress(), /*Size=*/0, - CE->getFlags(), llvm::GlobalValue::WeakAnyLinkage); + CE->getFlags(), llvm::GlobalValue::WeakAnyLinkage, + /*MangledName*/StringRef()); } else if (const auto *CE = dyn_cast( std::get<0>(E))) { @@ -3263,12 +3284,13 @@ // Hidden or internal symbols on the device are not externally visible. We // should not attempt to register them by creating an offloading entry. if (auto *GV = dyn_cast(CE->getAddress())) - if (GV->hasLocalLinkage() || GV->hasHiddenVisibility()) + if (GV->hasHiddenVisibility()) continue; + StringRef MangledName = std::get<2>(E); createOffloadEntry(CE->getAddress(), CE->getAddress(), CE->getVarSize().getQuantity(), Flags, - CE->getLinkage()); + CE->getLinkage(), MangledName); } else { llvm_unreachable("Unsupported entry kind."); } @@ -3338,12 +3360,20 @@ /*MangledName=*/GetMDString(1), static_cast( /*Flags=*/GetMDInt(2)), - /*Order=*/GetMDInt(3)); + /*Order=*/GetMDInt(3), + /*OrigName=*/GetMDString(4)); + OffloadEntriesInfoManager.enterDeviceGlobalVarMangledName( + /*OrigName=*/GetMDString(4), + /*MangledName=*/GetMDString(1)); break; } } } +StringRef CGOpenMPRuntime::getHostMangledDeclareTargetGlobal( StringRef VarName ) { + return OffloadEntriesInfoManager.getOffloadEntryHostMangledName(VarName); +} + void CGOpenMPRuntime::emitKmpRoutineEntryT(QualType KmpInt32Ty) { if (!KmpRoutineEntryPtrTy) { // Build typedef kmp_int32 (* kmp_routine_entry_t)(kmp_int32, void *); type. @@ -10750,11 +10780,29 @@ StringRef VarName; CharUnits VarSize; llvm::GlobalValue::LinkageTypes Linkage; + StringRef OrigName = VD->getName(); + SmallString<256> Buffer; + llvm::raw_svector_ostream Out(Buffer); if (*Res == OMPDeclareTargetDeclAttr::MT_To && !HasRequiresUnifiedSharedMemory) { Flags = OffloadEntriesInfoManagerTy::OMPTargetGlobalVarEntryTo; - VarName = CGM.getMangledName(VD); + // We don't need to mangle the host side of declare target global variables but we + // need to create offload entry that matches the device side which gets mangled. + if (!CGM.getLangOpts().OpenMPIsDevice && + !VD->isExternallyVisible()) { + VarName = OffloadEntriesInfoManager.getOffloadEntryHostMangledName(OrigName); + if ( VarName.empty() ) { + Out<hasDefinition(CGM.getContext()) != VarDecl::DeclarationOnly) { VarSize = CGM.getContext().getTypeSizeInChars(VD->getType()); assert(!VarSize.isZero() && "Expected non-zero size of the variable"); @@ -10801,7 +10849,7 @@ } OffloadEntriesInfoManager.registerDeviceGlobalVarEntryInfo( - VarName, Addr, VarSize, Flags, Linkage); + VarName, OrigName, Addr, VarSize, Flags, Linkage); } bool CGOpenMPRuntime::emitTargetGlobal(GlobalDecl GD) { diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h @@ -70,7 +70,8 @@ /// address \a Addr, size \a Size, and flags \a Flags. void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr, uint64_t Size, int32_t Flags, - llvm::GlobalValue::LinkageTypes Linkage) override; + llvm::GlobalValue::LinkageTypes Linkage, + StringRef MangledName) override; /// Emit outlined function specialized for the Fork-Join /// programming model for applicable target directives on the NVPTX device. diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -1122,7 +1122,8 @@ void CGOpenMPRuntimeGPU::createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr, uint64_t Size, int32_t, - llvm::GlobalValue::LinkageTypes) { + llvm::GlobalValue::LinkageTypes, + StringRef) { // TODO: Add support for global variables on the device after declare target // support. llvm::Function *Fn = dyn_cast(Addr); diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1501,6 +1501,23 @@ const auto *ND = cast(GD.getDecl()); std::string MangledName = getMangledNameImpl(*this, GD, ND); + if ( getLangOpts().OpenMPIsDevice ) { + if ( isa(GD.getDecl()) ) { + const auto *VD = dyn_cast(GD.getDecl()); + llvm::Optional Res = + OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD); + + if ( Res && (*Res == OMPDeclareTargetDeclAttr::MT_To) && + !getOpenMPRuntime().hasRequiresUnifiedSharedMemory() && + !VD->isExternallyVisible() ) { + StringRef HostMangledName = + getOpenMPRuntime().getHostMangledDeclareTargetGlobal(VD->getName()); + if (!HostMangledName.empty()) + MangledName = HostMangledName.str(); + } + } + } + // Ensure either we have different ABIs between host and device compilations, // says host compilation following MSVC ABI but device compilation follows // Itanium C++ ABI or, if they follow the same ABI, kernel names after @@ -6886,6 +6903,7 @@ SM.getDiagnostics().Report(diag::err_cannot_open_file) << PLoc.getFilename() << EC.message(); } + OS << llvm::format("%x", ID.getFile()) << llvm::format("%x", ID.getDevice()) << "_" << llvm::utohexstr(Result.low(), /*LowerCase=*/true, /*Width=*/8); } else { diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -7289,8 +7289,20 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes( const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const { + // Make sure any variable with OpenMP declare target is visible to runtime + // except for those with hidden visibility + if ( D && isa(D) && D->hasAttr() && + isa(GV) ) { + if ( (GV->hasInternalLinkage() || + GV->hasPrivateLinkage()) && + !GV->hasHiddenVisibility() ) { + GV->setLinkage(llvm::GlobalVariable::ExternalLinkage); + GV->setDSOLocal(false); + } + } if (GV->isDeclaration()) return; + const VarDecl *VD = dyn_cast_or_null(D); if (VD) { if (M.getLangOpts().CUDA) { @@ -9424,6 +9436,18 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes( const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const { + // Make sure any variable with OpenMP declare target is visible to runtime + // except for those with hidden visibility + if ( D && isa(D) && D->hasAttr() && + isa(GV) ) { + if ( (GV->hasInternalLinkage() || + GV->hasPrivateLinkage()) && + !GV->hasHiddenVisibility() ) { + GV->setLinkage(llvm::GlobalVariable::ExternalLinkage); + GV->setDSOLocal(false); + } + } + if (requiresAMDGPUProtectedVisibility(D, GV)) { GV->setVisibility(llvm::GlobalValue::ProtectedVisibility); GV->setDSOLocal(true); diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -12961,20 +12961,18 @@ return hasClauses(Clauses, K) || hasClauses(Clauses, ClauseTypes...); } -/// Check if the variables in the mapping clause are externally visible. +/// Check if the variables in the mapping clause have hidden visibility attribute static bool isClauseMappable(ArrayRef Clauses) { for (const OMPClause *C : Clauses) { if (auto *TC = dyn_cast(C)) return llvm::all_of(TC->all_decls(), [](ValueDecl *VD) { return !VD || !VD->hasAttr() || - (VD->isExternallyVisible() && - VD->getVisibility() != HiddenVisibility); + (VD->getVisibility() != HiddenVisibility); }); else if (auto *FC = dyn_cast(C)) return llvm::all_of(FC->all_decls(), [](ValueDecl *VD) { return !VD || !VD->hasAttr() || - (VD->isExternallyVisible() && - VD->getVisibility() != HiddenVisibility); + (VD->getVisibility() != HiddenVisibility); }); } diff --git a/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp --- a/clang/test/OpenMP/declare_target_codegen.cpp +++ b/clang/test/OpenMP/declare_target_codegen.cpp @@ -43,7 +43,7 @@ // CHECK-DAG: @d ={{ protected | }}global i32 0, // CHECK-DAG: @c = external global i32, // CHECK-DAG: @globals ={{ protected | }}global %struct.S zeroinitializer, -// CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer, +// CHECK-DAG: [[STAT:@stat__static__.+]] = internal global %struct.S zeroinitializer, // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]] // CHECK-DAG: @out_decl_target ={{ protected | }}global i32 0, // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)], @@ -247,8 +247,8 @@ // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}} -// CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}} -// CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}} +// CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}, !"aaa"} +// CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}, !"ccc"} // CHECK-DAG: !{{{.+}}virtual_foo #ifdef OMP5 diff --git a/clang/test/OpenMP/declare_target_link_codegen.cpp b/clang/test/OpenMP/declare_target_link_codegen.cpp --- a/clang/test/OpenMP/declare_target_link_codegen.cpp +++ b/clang/test/OpenMP/declare_target_link_codegen.cpp @@ -85,5 +85,5 @@ // HOST: [[C:%.*]] = load i32, i32* @c, // HOST: store i32 [[C]], i32* % -// CHECK: !{i32 1, !"c_decl_tgt_ref_ptr", i32 1, i32 {{[0-9]+}}} +// CHECK: !{i32 1, !"c_decl_tgt_ref_ptr", i32 1, i32 {{[0-9]+}}, !"c"} #endif // HEADER diff --git a/clang/test/OpenMP/declare_target_visibility_codegen.cpp b/clang/test/OpenMP/declare_target_visibility_codegen.cpp --- a/clang/test/OpenMP/declare_target_visibility_codegen.cpp +++ b/clang/test/OpenMP/declare_target_visibility_codegen.cpp @@ -8,8 +8,8 @@ // HOST: @[[X:.+]] = internal global i32 0, align 4 // HOST: @y = hidden global i32 0 // HOST: @z = global i32 0 -// HOST-NOT: @.omp_offloading.entry.c -// HOST-NOT: @.omp_offloading.entry.x +// HOST: @.omp_offloading.entry.c__static__{{[0-9a-z]+_[0-9a-z]+}} +// HOST: @.omp_offloading.entry.x__static__{{[0-9a-z]+_[0-9a-z]+}} // HOST-NOT: @.omp_offloading.entry.y // HOST: @.omp_offloading.entry.z C() : x(0) {} diff --git a/clang/test/OpenMP/nvptx_allocate_codegen.cpp b/clang/test/OpenMP/nvptx_allocate_codegen.cpp --- a/clang/test/OpenMP/nvptx_allocate_codegen.cpp +++ b/clang/test/OpenMP/nvptx_allocate_codegen.cpp @@ -89,7 +89,7 @@ // CHECK1-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 // CHECK1-NEXT: [[B:%.*]] = alloca double, align 8 // CHECK1-NEXT: store i32 0, i32* [[RETVAL]], align 4 -// CHECK1-NEXT: store i32 2, i32* @_ZZ4mainE1a, align 4 +// CHECK1-NEXT: store i32 2, i32* @a1, align 4 // CHECK1-NEXT: store double 3.000000e+00, double* [[B]], align 8 // CHECK1-NEXT: [[CALL:%.*]] = call noundef i32 @_Z3fooIiET_v() #[[ATTR7:[0-9]+]] // CHECK1-NEXT: ret i32 [[CALL]] diff --git a/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp b/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp --- a/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp +++ b/clang/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp @@ -15,7 +15,7 @@ // SIMD-ONLY-NOT: {{__kmpc|__tgt}} -// DEVICE-DAG: [[C_ADDR:.+]] = internal global i32 0, +// DEVICE-DAG: [[C_ADDR:.+]] = global i32 0, // DEVICE-DAG: [[CD_ADDR:@.+]] ={{ protected | }}global %struct.S zeroinitializer, // HOST-DAG: @[[C_ADDR:.+]] = internal global i32 0, // HOST-DAG: @[[CD_ADDR:.+]] ={{( protected | dso_local)?}} global %struct.S zeroinitializer, @@ -72,6 +72,8 @@ // DEVICE-DAG: call void // DEVICE-DAG: ret void +// HOST-DAG: @.omp_offloading.entry_name = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_ENTRY_NAME:c__static__.+]]\00" +// HOST-DAG: @.omp_offloading.entry.[[C_ENTRY_NAME]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (i32* @[[C_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name, i32 0, i32 0), i64 4, i32 0, i32 0 }, section "omp_offloading_entries", align 1 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_ADDR]]\00" // HOST-DAG: @.omp_offloading.entry.[[CD_ADDR]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (%struct.S* @[[CD_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section "omp_offloading_entries", align 1 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_CTOR]]\00" @@ -97,8 +99,8 @@ // HOST: [[C:%.*]] = load i32, i32* @[[C_ADDR]], // HOST: store i32 [[C]], i32* % -// HOST-DAG: !{i32 1, !"[[CD_ADDR]]", i32 0, i32 {{[0-9]+}}} -// HOST-DAG: !{i32 1, !"[[C_ADDR]]", i32 0, i32 {{[0-9]+}}} +// HOST-DAG: !{i32 1, !"[[CD_ADDR]]", i32 0, i32 {{[0-9]+}}, !"cd"} +// HOST-DAG: !{i32 1, !"[[C_ENTRY_NAME]]", i32 0, i32 {{[0-9]+}}, !"c"} // DEVICE: !nvvm.annotations // DEVICE-DAG: !{void ()* [[C_CTOR]], !"kernel", i32 1} diff --git a/clang/test/OpenMP/target_update_messages.cpp b/clang/test/OpenMP/target_update_messages.cpp --- a/clang/test/OpenMP/target_update_messages.cpp +++ b/clang/test/OpenMP/target_update_messages.cpp @@ -14,13 +14,6 @@ argc = x; // expected-warning {{variable 'x' is uninitialized when used here}} } -static int y; -#pragma omp declare target(y) - -void yyy() { -#pragma omp target update to(y) // expected-error {{the host cannot update a declare target variable that is not externally visible.}} -} - int __attribute__((visibility("hidden"))) z; #pragma omp declare target(z)