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,21 +662,24 @@ /// 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); - } + } CharUnits getVarSize() const { return VarSize; } void setVarSize(CharUnits Size) { VarSize = Size; } @@ -684,17 +688,20 @@ 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, - CharUnits VarSize, + registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName, + llvm::Constant *Addr, CharUnits VarSize, OMPTargetGlobalVarEntryKind Flags, llvm::GlobalValue::LinkageTypes Linkage); /// Checks if the variable with the given name has been registered already. @@ -707,6 +714,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 @@ -726,6 +735,8 @@ typedef llvm::StringMap OffloadEntriesDeviceGlobalVarTy; OffloadEntriesDeviceGlobalVarTy OffloadEntriesDeviceGlobalVar; + /// indexed by original name + llvm::StringMap OffloadEntriesDeviceGlobalVarNameMap; }; OffloadEntriesInfoManagerTy OffloadEntriesInfoManager; @@ -1924,6 +1935,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,20 +3032,38 @@ Action(D.first, F.first, P.first(), L.first, L.second); } +void CGOpenMPRuntime::OffloadEntriesInfoManagerTy:: + enterDeviceGlobalVarMangledName(StringRef OrigName, StringRef MangledName) { + if (!OrigName.equals(MangledName)) { + OffloadEntriesDeviceGlobalVarNameMap.try_emplace(OrigName, + MangledName.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, - CharUnits VarSize, + registerDeviceGlobalVarEntryInfo(StringRef VarName, StringRef OrigName, + llvm::Constant *Addr, CharUnits VarSize, OMPTargetGlobalVarEntryKind Flags, llvm::GlobalValue::LinkageTypes Linkage) { if (CGM.getLangOpts().OpenMPIsDevice) { @@ -3063,6 +3081,7 @@ Entry.setVarSize(VarSize); Entry.setLinkage(Linkage); Entry.setAddress(Addr); + Entry.setOrigName(OrigName); } else { if (hasDeviceGlobalVarEntryInfo(VarName)) { auto &Entry = OffloadEntriesDeviceGlobalVar[VarName]; @@ -3075,7 +3094,7 @@ return; } OffloadEntriesDeviceGlobalVar.try_emplace( - VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage); + VarName, OffloadingEntriesNum, Addr, VarSize, Flags, Linkage, OrigName); ++OffloadingEntriesNum; } } @@ -3090,8 +3109,9 @@ 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 +3204,12 @@ // - 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())}; + llvm::Metadata *Ops[] = {GetMDInt(E.getKind()), + GetMDString(MangledName), + 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 +3240,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))) { @@ -3260,15 +3283,18 @@ break; } - // Hidden or internal symbols on the device are not externally visible. We - // should not attempt to register them by creating an offloading entry. + // Hidden symbols on the device are not externally visible and constants + // don't need to be modified. 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() || + dyn_cast(GV)->isConstant()) 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 +3364,21 @@ /*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 +10785,39 @@ 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. + auto *GV = dyn_cast(Addr); + if (!CGM.getLangOpts().OpenMPIsDevice && !VD->isExternallyVisible() && + !GV->hasHiddenVisibility() && + !dyn_cast(GV)->isConstant()) { + VarName = + OffloadEntriesInfoManager.getOffloadEntryHostMangledName(OrigName); + if (VarName.empty()) { + unsigned DeviceID; + unsigned FileID; + unsigned Line; + SourceLocation Loc = VD->getCanonicalDecl()->getBeginLoc(); + getTargetEntryUniqueInfo(CGM.getContext(), Loc, DeviceID, FileID, Line); + { + Out << VD->getName() << "__static__" << llvm::format("%x", DeviceID) + << llvm::format("_%x_", FileID) << "l" << Line; + } + VarName = Buffer; + } + } else { + VarName = CGM.getMangledName(VD); + } + if (VD->hasDefinition(CGM.getContext()) != VarDecl::DeclarationOnly) { VarSize = CGM.getContext().getTypeSizeInChars(VD->getType()); assert(!VarSize.isZero() && "Expected non-zero size of the variable"); @@ -10801,7 +10864,7 @@ } OffloadEntriesInfoManager.registerDeviceGlobalVarEntryInfo( - VarName, Addr, VarSize, Flags, Linkage); + VarName, OrigName, Addr, VarSize, Flags, Linkage); } bool CGOpenMPRuntime::emitTargetGlobal(GlobalDecl GD) { @@ -11187,7 +11250,6 @@ isa(D) || isa(D)) && "Expecting either target enter, exit data, or update directives."); - CodeGenFunction::OMPTargetDataInfo InputInfo; llvm::Value *MapTypesArray = nullptr; llvm::Value *MapNamesArray = nullptr; 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 @@ -1120,9 +1120,10 @@ } void CGOpenMPRuntimeGPU::createOffloadEntry(llvm::Constant *ID, - llvm::Constant *Addr, - uint64_t Size, int32_t, - llvm::GlobalValue::LinkageTypes) { + llvm::Constant *Addr, uint64_t Size, + int32_t, + 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 @@ -1519,6 +1519,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 @@ -4274,13 +4291,13 @@ // Handle things which are present even on external declarations. if (D) { - if (LangOpts.OpenMP && !LangOpts.OpenMPSimd) - getOpenMPRuntime().registerTargetGlobalVariable(D, GV); - // FIXME: This code is overly simple and should be merged with other global // handling. GV->setConstant(isTypeConstant(D->getType(), false)); + if (LangOpts.OpenMP && !LangOpts.OpenMPSimd) + getOpenMPRuntime().registerTargetGlobalVariable(D, GV); + GV->setAlignment(getContext().getDeclAlign(D).getAsAlign()); setLinkageForGV(GV, D); @@ -4862,7 +4879,20 @@ !D->hasAttr()) Linkage = llvm::GlobalValue::InternalLinkage; - GV->setLinkage(Linkage); + // Make sure any variable with OpenMP declare target is visible to the runtime + // except for constants and those with hidden visibility + Optional DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(D); + if (DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_Any) && + getLangOpts().OpenMPIsDevice && D && !GV->hasHiddenVisibility() && + !GV->isConstant() && + !getOpenMPRuntime().hasRequiresUnifiedSharedMemory()) { + GV->setLinkage(llvm::GlobalValue::ExternalLinkage); + GV->setDSOLocal(false); + } else { + GV->setLinkage(Linkage); + } + if (D->hasAttr()) GV->setDLLStorageClass(llvm::GlobalVariable::DLLImportStorageClass); else if (D->hasAttr()) @@ -6968,6 +6998,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 @@ -7294,6 +7294,7 @@ const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const { if (GV->isDeclaration()) return; + const VarDecl *VD = dyn_cast_or_null(D); if (VD) { if (M.getLangOpts().CUDA) { 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 @@ -12959,20 +12959,19 @@ 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__.+]] = 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_only_one_side_compilation.cpp b/clang/test/OpenMP/declare_target_only_one_side_compilation.cpp --- a/clang/test/OpenMP/declare_target_only_one_side_compilation.cpp +++ b/clang/test/OpenMP/declare_target_only_one_side_compilation.cpp @@ -58,7 +58,7 @@ // TODO: It is odd, probably wrong, that we don't mangle all variables. // DEVICE-DAG: @G1 = {{.*}}global i32 0, align 4 -// DEVICE-DAG: @_ZL2G2 = internal {{.*}}global i32 0, align 4 +// DEVICE-DAG: @_ZL2G2 = {{.*}}global i32 0, align 4 // DEVICE-DAG: @G3 = {{.*}}global i32 0, align 4 // DEVICE-DAG: @_ZL2G4 = internal {{.*}}global i32 0, align 4 // DEVICE-DAG: @G5 = {{.*}}global i32 0, align 4 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]+_l[0-9]+}} +// HOST: @.omp_offloading.entry.x__static__{{[0-9a-z]+_[0-9a-z]+_l[0-9]+}} // 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* @_ZN2ns1aE1, 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) diff --git a/openmp/libomptarget/test/mapping/declare_target_static_var.c b/openmp/libomptarget/test/mapping/declare_target_static_var.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/declare_target_static_var.c @@ -0,0 +1,21 @@ +// RUN: %libomptarget-compile-run-and-check-generic + +#include + +#pragma omp declare target +static int y; +#pragma omp end declare target + +int main(void) { + y = 2; +#pragma omp target update to(y) + +#pragma omp target + { y += 3; } + +#pragma omp target update from(y) + + // CHECK: Declare target var update successful + printf("Declare target var update %s\n", (y == 5) ? "successful" : "failed"); + return 0; +}