diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 632e37e3cac8fe..f95b0f8cb317c7 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -595,6 +595,7 @@ Error linkBitcodeFiles(SmallVectorImpl &InputFiles, StringRef Arch = Args.getLastArgValue(OPT_arch_EQ); SmallVector BitcodeInputFiles; + DenseSet StrongResolutions; DenseSet UsedInRegularObj; DenseSet UsedInSharedLib; BumpPtrAllocator Alloc; @@ -608,6 +609,18 @@ Error linkBitcodeFiles(SmallVectorImpl &InputFiles, file_magic Type = identify_magic(Buffer.getBuffer()); switch (Type) { case file_magic::bitcode: { + Expected IRSymtabOrErr = readIRSymtab(Buffer); + if (!IRSymtabOrErr) + return IRSymtabOrErr.takeError(); + + // Check for any strong resolutions we need to preserve. + for (unsigned I = 0; I != IRSymtabOrErr->Mods.size(); ++I) { + for (const auto &Sym : IRSymtabOrErr->TheReader.module_symbols(I)) { + if (!Sym.isFormatSpecific() && Sym.isGlobal() && !Sym.isWeak() && + !Sym.isUndefined()) + StrongResolutions.insert(Saver.save(Sym.Name)); + } + } BitcodeInputFiles.emplace_back(std::move(File)); continue; } @@ -696,6 +709,7 @@ Error linkBitcodeFiles(SmallVectorImpl &InputFiles, // it is undefined or another definition has already been used. Res.Prevailing = !Sym.isUndefined() && + !(Sym.isWeak() && StrongResolutions.contains(Sym.getName())) && PrevailingSymbols.insert(Saver.save(Sym.getName())).second; // We need LTO to preseve the following global symbols: diff --git a/openmp/libomptarget/test/offloading/weak.c b/openmp/libomptarget/test/offloading/weak.c new file mode 100644 index 00000000000000..ca81db958356b2 --- /dev/null +++ b/openmp/libomptarget/test/offloading/weak.c @@ -0,0 +1,33 @@ +// RUN: %libomptarget-compile-generic -DA -c -o %t-a.o +// RUN: %libomptarget-compile-generic -DB -c -o %t-b.o +// RUN: %libomptarget-compile-generic %t-a.o %t-b.o && \ +// RUN: %libomptarget-run-generic | %fcheck-generic + +#if defined(A) +__attribute__((weak)) int x = 999; +#pragma omp declare target to(x) +#elif defined(B) +int x = 42; +#pragma omp declare target to(x) +__attribute__((weak)) int y = 42; +#pragma omp declare target to(y) +#else + +#include + +extern int x; +#pragma omp declare target to(x) +extern int y; +#pragma omp declare target to(y) + +int main() { + x = 0; + +#pragma omp target update from(x) +#pragma omp target update from(y) + + // CHECK: PASS + if (x == 42 && y == 42) + printf("PASS\n"); +} +#endif