From 4c4117089599cb5b6c6fa5635c28462ffd1bddf4 Mon Sep 17 00:00:00 2001 From: Saiyedul Islam Date: Tue, 28 Sep 2021 14:26:54 +0000 Subject: [PATCH] [Clang][OpenMP] Add partial support for Static Device Libraries An archive containing device code object files can be passed to clang command line for linking. For each given offload target it creates a device specific archives which is either passed to llvm-link if the target is amdgpu, or to clang-nvlink-wrapper if the target is nvptx. -L/-l flags are used to specify these fat archives on the command line. E.g. clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp -L. -lmylib It currently doesn't support linking an archive directly, like: clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp libmylib.a Linking with x86 offload also does not work. Reviewed By: ye-luo Differential Revision: https://reviews.llvm.org/D105191 --- clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp | 4 + clang/lib/Driver/ToolChains/Clang.cpp | 37 ++- clang/lib/Driver/ToolChains/CommonArgs.cpp | 287 +++++++++++++++++++++ clang/lib/Driver/ToolChains/CommonArgs.h | 33 +++ clang/lib/Driver/ToolChains/Cuda.cpp | 7 +- .../openmp_static_device_link/libFatArchive.a | Bin 0 -> 41104 bytes clang/test/Driver/fat_archive_amdgpu.cpp | 81 ++++++ clang/test/Driver/fat_archive_nvptx.cpp | 81 ++++++ .../clang-offload-bundler/ClangOffloadBundler.cpp | 26 +- 9 files changed, 551 insertions(+), 5 deletions(-) create mode 100644 clang/test/Driver/Inputs/openmp_static_device_link/libFatArchive.a create mode 100644 clang/test/Driver/fat_archive_amdgpu.cpp create mode 100644 clang/test/Driver/fat_archive_nvptx.cpp diff --git a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp index 135e369..5400e26 100644 --- a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp +++ b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp @@ -114,6 +114,10 @@ const char *AMDGCN::OpenMPLinker::constructLLVMLinkCommand( } } + AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, CmdArgs, "amdgcn", + SubArchName, + /* bitcode SDL?*/ true, + /* PostClang Link? */ false); // Add an intermediate output file. CmdArgs.push_back("-o"); const char *OutputFileName = diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 369c12a..65dfe0a 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -7734,12 +7734,28 @@ void OffloadBundler::ConstructJob(Compilation &C, const JobAction &JA, Triples += Action::GetOffloadKindName(CurKind); Triples += '-'; Triples += CurTC->getTriple().normalize(); - if ((CurKind == Action::OFK_HIP || CurKind == Action::OFK_OpenMP || - CurKind == Action::OFK_Cuda) && + if ((CurKind == Action::OFK_HIP || CurKind == Action::OFK_Cuda) && CurDep->getOffloadingArch()) { Triples += '-'; Triples += CurDep->getOffloadingArch(); } + + // TODO: Replace parsing of -march flag. Can be done by storing GPUArch + // with each toolchain. + StringRef GPUArchName; + if (CurKind == Action::OFK_OpenMP) { + // Extract GPUArch from -march argument in TC argument list. + for (unsigned ArgIndex = 0; ArgIndex < TCArgs.size(); ArgIndex++) { + auto ArchStr = StringRef(TCArgs.getArgString(ArgIndex)); + auto Arch = ArchStr.startswith_insensitive("-march="); + if (Arch) { + GPUArchName = ArchStr.substr(7); + Triples += "-"; + break; + } + } + Triples += GPUArchName.str(); + } } CmdArgs.push_back(TCArgs.MakeArgString(Triples)); @@ -7813,12 +7829,27 @@ void OffloadBundler::ConstructJobMultipleOutputs( Triples += '-'; Triples += Dep.DependentToolChain->getTriple().normalize(); if ((Dep.DependentOffloadKind == Action::OFK_HIP || - Dep.DependentOffloadKind == Action::OFK_OpenMP || Dep.DependentOffloadKind == Action::OFK_Cuda) && !Dep.DependentBoundArch.empty()) { Triples += '-'; Triples += Dep.DependentBoundArch; } + // TODO: Replace parsing of -march flag. Can be done by storing GPUArch + // with each toolchain. + StringRef GPUArchName; + if (Dep.DependentOffloadKind == Action::OFK_OpenMP) { + // Extract GPUArch from -march argument in TC argument list. + for (uint ArgIndex = 0; ArgIndex < TCArgs.size(); ArgIndex++) { + StringRef ArchStr = StringRef(TCArgs.getArgString(ArgIndex)); + auto Arch = ArchStr.startswith_insensitive("-march="); + if (Arch) { + GPUArchName = ArchStr.substr(7); + Triples += "-"; + break; + } + } + Triples += GPUArchName.str(); + } } CmdArgs.push_back(TCArgs.MakeArgString(Triples)); diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp index 9f18954..c3abdf4 100644 --- a/clang/lib/Driver/ToolChains/CommonArgs.cpp +++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp @@ -34,6 +34,7 @@ #include "clang/Driver/Util.h" #include "clang/Driver/XRayArgs.h" #include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/SmallSet.h" #include "llvm/ADT/SmallString.h" #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/StringSwitch.h" @@ -1587,6 +1588,292 @@ void tools::addX86AlignBranchArgs(const Driver &D, const ArgList &Args, } } +/// SDLSearch: Search for Static Device Library +/// The search for SDL bitcode files is consistent with how static host +/// libraries are discovered. That is, the -l option triggers a search for +/// files in a set of directories called the LINKPATH. The host library search +/// procedure looks for a specific filename in the LINKPATH. The filename for +/// a host library is lib.a or lib.so. For SDLs, there is an +/// ordered-set of filenames that are searched. We call this ordered-set of +/// filenames as SEARCH-ORDER. Since an SDL can either be device-type specific, +/// architecture specific, or generic across all architectures, a naming +/// convention and search order is used where the file name embeds the +/// architecture name (nvptx or amdgcn) and the GPU device type +/// such as sm_30 and gfx906. is absent in case of +/// device-independent SDLs. To reduce congestion in host library directories, +/// the search first looks for files in the “libdevice” subdirectory. SDLs that +/// are bc files begin with the prefix “lib”. +/// +/// Machine-code SDLs can also be managed as an archive (*.a file). The +/// convention has been to use the prefix “lib”. To avoid confusion with host +/// archive libraries, we use prefix "libbc-" for the bitcode SDL archives. +/// +bool tools::SDLSearch(const Driver &D, const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args, + SmallVector LibraryPaths, std::string Lib, + StringRef Arch, StringRef Target, bool isBitCodeSDL, + bool postClangLink) { + SmallVector SDLs; + + std::string LibDeviceLoc = "/libdevice"; + std::string LibBcPrefix = "/libbc-"; + std::string LibPrefix = "/lib"; + + if (isBitCodeSDL) { + // SEARCH-ORDER for Bitcode SDLs: + // libdevice/libbc---.a + // libbc---.a + // libdevice/libbc--.a + // libbc--.a + // libdevice/libbc-.a + // libbc-.a + // libdevice/lib--.bc + // lib--.bc + // libdevice/lib-.bc + // lib-.bc + // libdevice/lib.bc + // lib.bc + + for (StringRef Base : {LibBcPrefix, LibPrefix}) { + const auto *Ext = Base.contains(LibBcPrefix) ? ".a" : ".bc"; + + for (auto Suffix : {Twine(Lib + "-" + Arch + "-" + Target).str(), + Twine(Lib + "-" + Arch).str(), Twine(Lib).str()}) { + SDLs.push_back(Twine(LibDeviceLoc + Base + Suffix + Ext).str()); + SDLs.push_back(Twine(Base + Suffix + Ext).str()); + } + } + } else { + // SEARCH-ORDER for Machine-code SDLs: + // libdevice/lib--.a + // lib--.a + // libdevice/lib-.a + // lib-.a + + const auto *Ext = ".a"; + + for (auto Suffix : {Twine(Lib + "-" + Arch + "-" + Target).str(), + Twine(Lib + "-" + Arch).str()}) { + SDLs.push_back(Twine(LibDeviceLoc + LibPrefix + Suffix + Ext).str()); + SDLs.push_back(Twine(LibPrefix + Suffix + Ext).str()); + } + } + + // The CUDA toolchain does not use a global device llvm-link before the LLVM + // backend generates ptx. So currently, the use of bitcode SDL for nvptx is + // only possible with post-clang-cc1 linking. Clang cc1 has a feature that + // will link libraries after clang compilation while the LLVM IR is still in + // memory. This utilizes a clang cc1 option called “-mlink-builtin-bitcode”. + // This is a clang -cc1 option that is generated by the clang driver. The + // option value must a full path to an existing file. + bool FoundSDL = false; + for (auto LPath : LibraryPaths) { + for (auto SDL : SDLs) { + auto FullName = Twine(LPath + SDL).str(); + if (llvm::sys::fs::exists(FullName)) { + if (postClangLink) + CC1Args.push_back("-mlink-builtin-bitcode"); + CC1Args.push_back(DriverArgs.MakeArgString(FullName)); + FoundSDL = true; + break; + } + } + if (FoundSDL) + break; + } + return FoundSDL; +} + +/// Search if a user provided archive file lib.a exists in any of +/// the library paths. If so, add a new command to clang-offload-bundler to +/// unbundle this archive and create a temporary device specific archive. Name +/// of this SDL is passed to the llvm-link (for amdgcn) or to the +/// clang-nvlink-wrapper (for nvptx) commands by the driver. +bool tools::GetSDLFromOffloadArchive( + Compilation &C, const Driver &D, const Tool &T, const JobAction &JA, + const InputInfoList &Inputs, const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args, SmallVector LibraryPaths, + StringRef Lib, StringRef Arch, StringRef Target, bool isBitCodeSDL, + bool postClangLink) { + + // We don't support bitcode archive bundles for nvptx + if (isBitCodeSDL && Arch.contains("nvptx")) + return false; + + bool FoundAOB = false; + SmallVector AOBFileNames; + std::string ArchiveOfBundles; + for (auto LPath : LibraryPaths) { + ArchiveOfBundles.clear(); + + AOBFileNames.push_back(Twine(LPath + "/libdevice/lib" + Lib + ".a").str()); + AOBFileNames.push_back(Twine(LPath + "/lib" + Lib + ".a").str()); + + for (auto AOB : AOBFileNames) { + if (llvm::sys::fs::exists(AOB)) { + ArchiveOfBundles = AOB; + FoundAOB = true; + break; + } + } + + if (!FoundAOB) + continue; + + StringRef Prefix = isBitCodeSDL ? "libbc-" : "lib"; + std::string OutputLib = D.GetTemporaryPath( + Twine(Prefix + Lib + "-" + Arch + "-" + Target).str(), "a"); + + C.addTempFile(C.getArgs().MakeArgString(OutputLib.c_str())); + + ArgStringList CmdArgs; + SmallString<128> DeviceTriple; + DeviceTriple += Action::GetOffloadKindName(JA.getOffloadingDeviceKind()); + DeviceTriple += '-'; + std::string NormalizedTriple = T.getToolChain().getTriple().normalize(); + DeviceTriple += NormalizedTriple; + if (!Target.empty()) { + DeviceTriple += '-'; + DeviceTriple += Target; + } + + std::string UnbundleArg("-unbundle"); + std::string TypeArg("-type=a"); + std::string InputArg("-inputs=" + ArchiveOfBundles); + std::string OffloadArg("-targets=" + std::string(DeviceTriple)); + std::string OutputArg("-outputs=" + OutputLib); + + const char *UBProgram = DriverArgs.MakeArgString( + T.getToolChain().GetProgramPath("clang-offload-bundler")); + + ArgStringList UBArgs; + UBArgs.push_back(C.getArgs().MakeArgString(UnbundleArg.c_str())); + UBArgs.push_back(C.getArgs().MakeArgString(TypeArg.c_str())); + UBArgs.push_back(C.getArgs().MakeArgString(InputArg.c_str())); + UBArgs.push_back(C.getArgs().MakeArgString(OffloadArg.c_str())); + UBArgs.push_back(C.getArgs().MakeArgString(OutputArg.c_str())); + + // Add this flag to not exit from clang-offload-bundler if no compatible + // code object is found in heterogenous archive library. + std::string AdditionalArgs("-allow-missing-bundles"); + UBArgs.push_back(C.getArgs().MakeArgString(AdditionalArgs.c_str())); + + C.addCommand(std::make_unique( + JA, T, ResponseFileSupport::AtFileCurCP(), UBProgram, UBArgs, Inputs, + InputInfo(&JA, C.getArgs().MakeArgString(OutputLib.c_str())))); + if (postClangLink) + CC1Args.push_back("-mlink-builtin-bitcode"); + + CC1Args.push_back(DriverArgs.MakeArgString(OutputLib)); + break; + } + + return FoundAOB; +} + +// Wrapper function used by driver for adding SDLs during link phase. +void tools::AddStaticDeviceLibsLinking(Compilation &C, const Tool &T, + const JobAction &JA, + const InputInfoList &Inputs, + const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args, + StringRef Arch, StringRef Target, + bool isBitCodeSDL, bool postClangLink) { + AddStaticDeviceLibs(&C, &T, &JA, &Inputs, C.getDriver(), DriverArgs, CC1Args, + Arch, Target, isBitCodeSDL, postClangLink); +} + +// Wrapper function used for post clang linking of bitcode SDLS for nvptx by +// the CUDA toolchain. +void tools::AddStaticDeviceLibsPostLinking(const Driver &D, + const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args, + StringRef Arch, StringRef Target, + bool isBitCodeSDL, bool postClangLink) { + AddStaticDeviceLibs(nullptr, nullptr, nullptr, nullptr, D, DriverArgs, + CC1Args, Arch, Target, isBitCodeSDL, postClangLink); +} + +// User defined Static Device Libraries(SDLs) can be passed to clang for +// offloading GPU compilers. Like static host libraries, the use of a SDL is +// specified with the -l command line option. The primary difference between +// host and SDLs is the filenames for SDLs (refer SEARCH-ORDER for Bitcode SDLs +// and SEARCH-ORDER for Machine-code SDLs for the naming convention). +// SDLs are of following types: +// +// * Bitcode SDLs: They can either be a *.bc file or an archive of *.bc files. +// For NVPTX, these libraries are post-clang linked following each +// compilation. For AMDGPU, these libraries are linked one time +// during the application link phase. +// +// * Machine-code SDLs: They are archive files. For NVPTX, the archive members +// contain cubin for Nvidia GPUs and are linked one time during the +// link phase by the CUDA SDK linker called nvlink. For AMDGPU, the +// process for machine code SDLs is still in development. But they +// will be linked by the LLVM tool lld. +// +// * Bundled objects that contain both host and device codes: Bundled objects +// may also contain library code compiled from source. For NVPTX, the +// bundle contains cubin. For AMDGPU, the bundle contains bitcode. +// +// For Bitcode and Machine-code SDLs, current compiler toolchains hardcode the +// inclusion of specific SDLs such as math libraries and the OpenMP device +// library libomptarget. +void tools::AddStaticDeviceLibs(Compilation *C, const Tool *T, + const JobAction *JA, + const InputInfoList *Inputs, const Driver &D, + const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args, + StringRef Arch, StringRef Target, + bool isBitCodeSDL, bool postClangLink) { + + SmallVector LibraryPaths; + // Add search directories from LIBRARY_PATH env variable + llvm::Optional LibPath = + llvm::sys::Process::GetEnv("LIBRARY_PATH"); + if (LibPath) { + SmallVector Frags; + const char EnvPathSeparatorStr[] = {llvm::sys::EnvPathSeparator, '\0'}; + llvm::SplitString(*LibPath, Frags, EnvPathSeparatorStr); + for (StringRef Path : Frags) + LibraryPaths.emplace_back(Path.trim()); + } + + // Add directories from user-specified -L options + for (std::string Search_Dir : DriverArgs.getAllArgValues(options::OPT_L)) + LibraryPaths.emplace_back(Search_Dir); + + // Add path to lib-debug folders + SmallString<256> DefaultLibPath = llvm::sys::path::parent_path(D.Dir); + llvm::sys::path::append(DefaultLibPath, Twine("lib") + CLANG_LIBDIR_SUFFIX); + LibraryPaths.emplace_back(DefaultLibPath.c_str()); + + // Build list of Static Device Libraries SDLs specified by -l option + llvm::SmallSet SDLNames; + static const StringRef HostOnlyArchives[] = { + "omp", "cudart", "m", "gcc", "gcc_s", "pthread", "hip_hcc"}; + for (auto SDLName : DriverArgs.getAllArgValues(options::OPT_l)) { + if (!HostOnlyArchives->contains(SDLName)) { + SDLNames.insert(SDLName); + } + } + + // The search stops as soon as an SDL file is found. The driver then provides + // the full filename of the SDL to the llvm-link or clang-nvlink-wrapper + // command. If no SDL is found after searching each LINKPATH with + // SEARCH-ORDER, it is possible that an archive file lib.a exists + // and may contain bundled object files. + for (auto SDLName : SDLNames) { + // This is the only call to SDLSearch + if (!SDLSearch(D, DriverArgs, CC1Args, LibraryPaths, SDLName, Arch, Target, + isBitCodeSDL, postClangLink)) { + GetSDLFromOffloadArchive(*C, D, *T, *JA, *Inputs, DriverArgs, CC1Args, + LibraryPaths, SDLName, Arch, Target, + isBitCodeSDL, postClangLink); + } + } +} + static llvm::opt::Arg * getAMDGPUCodeObjectArgument(const Driver &D, const llvm::opt::ArgList &Args) { // The last of -mcode-object-v3, -mno-code-object-v3 and diff --git a/clang/lib/Driver/ToolChains/CommonArgs.h b/clang/lib/Driver/ToolChains/CommonArgs.h index 8e48f3e..00291a3 100644 --- a/clang/lib/Driver/ToolChains/CommonArgs.h +++ b/clang/lib/Driver/ToolChains/CommonArgs.h @@ -49,6 +49,39 @@ void AddRunTimeLibs(const ToolChain &TC, const Driver &D, llvm::opt::ArgStringList &CmdArgs, const llvm::opt::ArgList &Args); +void AddStaticDeviceLibsLinking(Compilation &C, const Tool &T, + const JobAction &JA, + const InputInfoList &Inputs, + const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CmdArgs, + StringRef Arch, StringRef Target, + bool isBitCodeSDL, bool postClangLink); +void AddStaticDeviceLibsPostLinking(const Driver &D, + const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CmdArgs, + StringRef Arch, StringRef Target, + bool isBitCodeSDL, bool postClangLink); +void AddStaticDeviceLibs(Compilation *C, const Tool *T, const JobAction *JA, + const InputInfoList *Inputs, const Driver &D, + const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CmdArgs, StringRef Arch, + StringRef Target, bool isBitCodeSDL, + bool postClangLink); + +bool SDLSearch(const Driver &D, const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CmdArgs, + SmallVector LibraryPaths, std::string Lib, + StringRef Arch, StringRef Target, bool isBitCodeSDL, + bool postClangLink); + +bool GetSDLFromOffloadArchive(Compilation &C, const Driver &D, const Tool &T, + const JobAction &JA, const InputInfoList &Inputs, + const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args, + SmallVector LibraryPaths, + StringRef Lib, StringRef Arch, StringRef Target, + bool isBitCodeSDL, bool postClangLink); + const char *SplitDebugName(const JobAction &JA, const llvm::opt::ArgList &Args, const InputInfo &Input, const InputInfo &Output); diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp index e4a6fb8..18351da 100644 --- a/clang/lib/Driver/ToolChains/Cuda.cpp +++ b/clang/lib/Driver/ToolChains/Cuda.cpp @@ -610,8 +610,11 @@ void NVPTX::OpenMPLinker::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back(CubinF); } + AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, CmdArgs, "nvptx", GPUArch, + false, false); + const char *Exec = - Args.MakeArgString(getToolChain().GetProgramPath("nvlink")); + Args.MakeArgString(getToolChain().GetProgramPath("clang-nvlink-wrapper")); C.addCommand(std::make_unique( JA, *this, ResponseFileSupport{ResponseFileSupport::RF_Full, llvm::sys::WEM_UTF8, @@ -741,6 +744,8 @@ void CudaToolChain::addClangTargetOptions( addOpenMPDeviceRTL(getDriver(), DriverArgs, CC1Args, BitcodeSuffix, getTriple()); + AddStaticDeviceLibsPostLinking(getDriver(), DriverArgs, CC1Args, "nvptx", GpuArch, + /* bitcode SDL?*/ true, /* PostClang Link? */ true); } } diff --git a/clang/test/Driver/Inputs/openmp_static_device_link/libFatArchive.a b/clang/test/Driver/Inputs/openmp_static_device_link/libFatArchive.a new file mode 100644 index 0000000000000000000000000000000000000000..ebd7e55898b3a3de38ec596fd844f4de06d708cb GIT binary patch literal 41104 zcmeG_4_s7bw%?sQ%pGyKb5Y7M0q+PFn$m!XgQCs=>R6~`RAedjypeCIpg`Tkw*xmTZRa}<=!3%!jW91H(|tyo%RW6y(_*MU_PrP7CQN()K+V4`?n<=AjNyhtGLe_bgO67mvk1qsE;$^Q+?xHdC#TX4q* z^!sqhM7i5}Yp9^fguIG%c4t*`QeI_w-b_RM>f2=a%HpaNLoxuaWlPDD4J9NcB#SuK z(1#XeE>utyn3NbOGLKx%$xuSN60W9ALefNT1Wfb2=6mk4RG)AGY;paC#r0R}Kd@|3 zK8+}r`rOzO{K>MsnGtJL3Ls~!-BRz2&9>BM#F{Mi-;b@rA1h`puM1Rf?z#fOQorvQ zX0`YFwBZ@CoMj8Ju+%GKfmIw9K+I!|ywyVDksC0|1jLfeEt#=df$UzPp-5o3$>*FX zl>8pCUqiU1{?AM5`RFY8^^Gl$HkC&)$^flUeF=%(>?K zQA76!sK|RanuGfUms$UY`y}Y451OA@|4#jiSY^Ei`%cz}TGYp^LHgiO9~#t$TB`5D z?%(1!)Y;(n8}xZ4zK{pJg;lr!2wy{YAtuK@lloc7*5+WXqH z_r0wRSN;Kb{efGO(x$E50#avEzKr?T^)#Ic>NGn&G1?S9YyEZgGD_J8lQW!R3(4tX4|snQllyUOkY96y!Anc@6MWE>7{U zeMRmq**Z8FJNNq&u3(F_o2J5eKAsLgguIT!;W%7m&f1Dq6|2`*48(>xhBZceAL&!D20y2-|9RgYv2E7;!{9eZ@IY9G`5(Q9vZDnQa#GmPQ$La@)3Y_s!$QLFjB-@gclT!+ljK-vt z#N=5s@@EzlCmRb5#l`uBwwZ+~2{VgRlBY<*2-6Lusg27CRmD*gG>qB{x;FqUQo=ci zaT?2|u{tx$FJhZ1xAIZCQK9H!c{8gMN>KHMaAhGbgd4E`%mPgs%crq%X{^D_8Z+1| zH`Acn67s0Bk>0FW&t?~})-<-n%-R>TSygOy8f%5b8aZ( z!Vry}DAohAWv~VmZL^}BH5Ra@G}dBfvlg=kdOIbH3iQB{V7o6N0|>ZL`di^$XpfSd z0m~d3=wU3W1>{n|zdi|F&t^uEQ z4frTvI2@m;*MQ?nJRJTU(7QgM$9cfRuvCnM60~K5l)qn!v1c@JUBW#F1djwC4Kd

mcq+tUaLM?>MhEiPe}Vje3waF4pM^`te;I)gg7IG!!2%pjs87njFU8153zy92 zcBvy|J`x(qmEeevdXn*LgYacN)Ca(^UfF(s7QiPwfR8@_Kd`+W0q{URA76vdrvdoL zM-FEL@W%j9mj64*6oTAdotPzM5gpPl96<$X_0D2pzKdn@Sk-WV0O!Lx3dCL|q%v@&9$VX-4LPyu6atmCor^#^k)@r0LS?J-w{7 zVr|v*l@)7&YE`+tAg_3}Bd^kFbCwq5l~$BG^ODvQI8T<@;+;iR&Y;sEtas&jIK4`` zD$iMDE3X8IbEPv6&!LshA_pYatSyC;tjKEcnKTd2q@_?Maj~zhEXi})94m{Qc~D?4 zSiQEwnOAPJ+lw3$U%3>kl1r`%84uz+oNMheLuk>eyaHQU8Hul0?JSC4wDj)jfKsq3 z06}hYF|?V);~Ji4D=dV?e&vuOwA>bKourzglDuMvt-NRmu5jGWs}PE=V`~y{fEN#q z@dX50<#5{aA>L3f#6qVyRu|fwHV{LI5ho!&AsHGryalYb7gdzor`yU4R~A%YgjY$W zZMt~)9A5tGx?M0>5{cnjCzx4LVSIdoq%gsp2LPu-{l^8}%DAFDuoz~7JMi|w)J3n!$meN4s}PXlLuTv6Deem;UvUs%V#MJ49f>ZZm` z^(K73CHFCYR|oK|=Lwk#apI)9QGEVFqQ5n^NvEF-KjUxFY})R%<`?GjvBfn#nxks3 zS;;rW5PBM2`f!}>Y}C;wT^;-{AL6HG@DD!oD{n`s?wKdQGcBZzZ{y;WhPo_`@>yR( zol6-*DwWL!+SH)m`is(z@fjp^M%?I($3`iCq6#N%${X&B)n$Y$HGV^N;@G$yPntR9 zrz(Xiyf*()zDc3W;ObHoC0wWyk%}0<*NkIyhU!Y^-w!;L{F( zc!xlC4u(Jjt?royW}TYZSmL(xjKk}3GKLKQhmm<@3ryOr5!@vMtOnSsx`IM2Ffj}uEp4vf_uBP>{C2W1U)O(YKyi)msR zznS#7!<2gR&gQvOzJ}g>2o8;XG+|xyNvuopmCL|vPy<^}vdgm`)Y9y6co|>Rf!3(E zDSn?GIybC)TjXwiWcRkP69YZc6TPQ8@)#9Gj6}7*$^%;6XkkzuXbVvzq#4*URW5{7 zg0!Fq6XSx8hz7PttE1a&k$>72)?Q)E3 z3dr)zx`7KL4jy}I1N2%GQ>FH-5jv~Vz9STOG7Sssc3icrPzR>ToS(0RvaBocOnnE6gamrH>SqyY^9LuZg!Zy^1 zHqfN#Vc?7>*Z_Bw*YBa{C>>X5Uhks62dP)O5_f>c5j%wpxcS=)w3h?htpM4B4s|a_V zW6#>zBPMORf$huE`cj(cd55XUZps8qI{YbH>A5?qQ?{lXyVB=&q^ES!EvhwErkrPw z1BoEdd4smBMRmx_24)ZY250-&^Bgk>0U-PkhJs5YdNegI^&VF@NeGc~%qC^2 zsmyy5olWKw3uUBTtcYX3Wj9XaTHj6Nm&6AK^8-_pJGC}_&erq^K0g`&bWjHkz1+hd zA>2t&X_x9m7W)aup6F6t$YKwfw5NG(xm|nQ#GZGtU{~%rucl9^OP{-)2HM#>{i&oD zA<+A^bgB;_DLc$_YSDg7+Ot_&U<;Z7b$td#$ZJoVK!@!4KK86&S4XYv5#Vpo9&ZNB zp1^)@HTu%$l25E`zn6g-keIc-y!MkU?2DtINt5<-p*wsxwO2=hb!`VzF{3?lDgQ-FDJ*J7vh>L6B5X(bjYrYdG{kjXTm)KxW_ri5>LZ?Lr5M9LR*NpI{#{ z)-ZAxiEY*<8@q%)y*deWF%#;J2{@*GRtqNPlQE`p{+EQ`6sDO>wv(sK9B00Xo6zv7 zfacSZQaYbPvk^}ni(%czv%?e}N~*z@iRp z*AT{G!cqreoGwJR1LrQ)DQJbr-!K}h=0ewYss0`2pC0yjGqb+s89NPA{YJIthOjv; z?zubM*y%9cdC{nhadY_$i+X(G^1`|;{NRt=Yn)P-0rMbZ0|O8LA9C+N{WZ)k>LwAv-|20Z>om;Q!F z3&$%oaEvwqfGADVNcyjhZ{5rbj1WXz*<|yIUr#Dm8zLbGn@o@f1*`hSK9@*^Sb|`0 zjQ)_Kyh0X#P_U`ve@!cRq!I!O_M;iywZsou;gC*oqh0N)>b9T3*8V;opm@AbjE0?d zc&4{u)Xj@_Dbp1TNr{eX&R4nEL&T-XnFMij4HpsjNNu=ll5$Gx>ND=dxazbx&!k&x z6t&GoN>|62IJ25O590%G^~_MaBHMc-izBGHb?L@he?yv!*-*t)Hm%VdCEAm)P`KC` znzZP))+pGX6dOVWG&f8Suzbyhm9Vp_0k+TIgAJ_3V~2gTonBYvvGbnso#E^WFWYO- zRv5I0`d~3+XK?LN2VnE1k7~0;7aQv5yz0i;dIy~f%PCA7oqZAOtjzaHm|4O;aTe?@ zVSjC9V4t|IKLpsWcX{$)Cjwhe*nT!uPdnAep6X#g_pv7obbbf!sl}ajv=d}1WZcOGSzpCeHQkwc-m66oL_Z+x5cRkxeUQSve5XFD&!$?j$!+ZP&uK}Y3ktc^mU=jS zP90?Z4R3_v7HmDp10yVsaWqxu9)hLJ#~#UI&lzG`_n6^$F|NBe>Hr8}883}Ln7JEZ zbKy_j*_r}-mc4Xpr)9?)*p6A5HK4$3&-jmbu_vtBlRdbC!4?3PELa;WIQ9_i0$@kb zgX{m{EVkbebMS->k3rqN<4)+K;KB9((NljAkyb zFf}1D?NP7(=g%5`Lp5yEEvaGs=a4S&5AUmr92G@n>kCS3aC4?~W081{mrnDAMe9ln zin11`^O;5Kipu_3f2k;~aIP+|D7C{al?o@1qTN+mZZ9k1({H_%Uy)`uLmPfM6I%ke zZEJXIj5;m3hA)Xxw`6Z>-8 z_U}J%tnJvI1N)CI9QC_~?rwnXJJ9iI_uILh-EF&%F4XVX^4W4*!3(P@7suRM`SS^r zg*$b>g84-c?HnEIh=XyUWtdjLq>c&W;l`vxR`!!D z+!P(g!6s~tT3pq548d8T$^#aRV<6l9l-~&Z{OxYo;VV|bgc7;GH*!B8dDw=($AlAE z*S5$G(Qv8vVLHWI;!I~M>_DD^!4G>m*yq3s<)`dfwXiFPiO6qu;Lh+fJA1Z|{Tw$B zHlC1eFhP%1E}n)nGTu^+-Akr#GNt`l)wNs~8$nZ_Pc(Jn%5ug4b1B4=7Ik>ycbe~JQ=aUCPxZgLBpUqE>Us<*~-&U6AEG>+ejpjDlY=ZIQ6NW;9c;GWJ z&Rdnq2KNo=6IxowKgaFo!|u_n?omC^qAziWQ8$(89yhimYw9AiCv1id-YR#Ws=URP zO{V}5rU|@@yHX0hEsvYCDwmjWhM?K%p>oNuis?%xIHD)Jsm!Zs*RA?Y?ndanM4ZK0HRE!3zdD_gB!xkt}C z@kVo~lhe#fYSLxSDqL!!xy=)55{Uc08e4N(mD|s$rq<3d#@U{zj#F){HM1#++`(U) z9#!Y>)ayE|zc=$W_toS#r*Y~Bb@66eSCOqqY$`G~Sew&IG|zNc>vyZGxGZ18=%c3D zF@01*nMv73o_fq@_eXeJr}=HlSsUyhGXAjY$huaiKT_Rlr|)D64euCu^FeCIub8fz zCz^ixAURoA-Rk0X8OO3h+pi}5SCeg5lhvDP+(*0x*M$)3PNlWO{$*Wbe_dnt9=+kc z-(IY1HNNIgeXYk(y)0y2elcUvym&I^=9@*#4_`*i{<=4k5GQFthpE4A%S@kP?@kfE zX{P4Iy)TOC#ghTP_O*+32}?{7lrAy%RT18(d-tP$pCRM5h!oc{xHFl1ms0)1`3qLW z^w(X^T=IG5M}NucoF5fVX&!IXEj==)^U!E^Z2yWL`f+0UVWVz*tU-G^X5PV=2M=Y= z%l~d;Ptu1KUqBZAd$KOS?oWL^eaYxOod3$j_Dz}Vy_a8aOKrPC!ds@bpZ(;;Ww{-P zV~!T8m(J*~`=i{<*Ol`eihuK`&fAlead_O)!_02imvvV%+RcOx|Itql*Sz!2bmQ!k zxBsE%tYh?}OLIOt3?jw+XxYbNb2_Mqj2R1>%nKCf`nSvwHQ8TRKG$au6opN0daEbv zhsRfd#^z1BJeprJk00eq`m%1#V#T?>3#SI!lKq%3Rcr|U<`;qH*Pe`7*_;B^ElJOK zus7qu9!-;|!z0KI`v`;=x!u2D%nbuI_SY?#x#w_(GCf1}{>8etu%%Mf`S*mUPYpJy6U%9rMKGQK(d@hx_YFQy%O2K&oeX`Bo9zdE)4R8h@xaEISo*&zPhq ziq)MjT&$bEfL{LLOEP^&xW&@DjCp^r#3Fy@=8TY+8g+YQd$^-#{$qaQ*o6MNr`48% zSzAvV8gE+?p2$hsxO;KU2P*=axzgiH{rqIiy(ZO_|1dQ6jaz#DjqVOWdJf|p}vK# zWna=f?|QZM3G1`93gq)0(*al~K^|ONMsFr@o{)6X;8!X*MPVC9cGc>I5^*iM z1#reH&B9uQN0l8B3vr1hI3d?q4L^ucwZn1O`hze|G0DAV3b&9}wskb;hITww&HY&E zSjhRqT1~1C!W7|U91905VNJ~9-i@#&yM^)7*gblZamf*|)C5mP2&hk0Y1oT~n+Upw zEYyX%!?RhAs!a##FngiRC<6u?4?n_sxk!+ahAhhH&7vdl&&{|-B5Z;xL-Smidwf%M z2&q-OIh$ry-v>bPvjdy{8{xSkAAj8oe`yl0WpelItUpShpwpmd#MbpKwyhKlrsy!+JH zjxe)kmhx#`XQC#vf53Xnl;KbZ8!Ns z6q;F@2eMTDuvPrlR%d#Ts%=_zLRPpcty(pTf0yyJdyzV#jwc*4>}dW}f-5R+E1CBN(!ym-l~Q zPWt>*3?p=m>PKprDpF6zJn;+SZ@Fk$cCl{p6wQm>^V&WZd*#JbG09trzbEFaQ!y8c zrd0#N4U=)&IB>D<$=Jv3oWFK%ZTJ+Qp?lsl=jP>4eKQ7!)Q!Ki+*PZo1>^*n(fd1o zd$D8u_O|&yyKu_zo-#v~a2e<9yxF_`%ilYc@xvzmyH5uUufuxF!`nx#x>9yIrW(eX zQ&C^|rP=z?y!;iR>C<;M_Px^Bq%v1J!6hhkpZYFW_}^Sx6s zFTJ2%dKXUNmh$^E>5Rt;zO*!6touy_>Ue8h8 zQ3pT#>XVEU%FH_Z?ES#}jB_RY@$E!|QNTH5=`y!^YNU9U9u{l2m3iAi@-r(=F_ zGUfu9LDt0%=&Uob^m3H5BP>anCi~SOXUb&Fi~SwDFVmA#;LTf|o6MC&;Mt9;dvsIO|d9J$C`i_sO9*=Z5FX3Y&qeD&V&G%$m)ikAC z)6&eDoR*f-y)R_=YsG|HQ$ zfjLVRQDf~|yDKT*nxR*#?q9F;hnYiEFKDVcX(mX5nLwD>aHn~ume*yk>>~Q?MBR@S zwj>k(6V2Pda7B(+uUX8mR3_}_G=+x$&R&&IK{A2DJ>dj$wR<-Ls zg{|F{1>3}zQo8Isxd~*t#mjG6mXv-vW*tm_Xe8N>z;@|m%%2K`>=~Wn+_8E2v&Fej z9FH$a6J@TeCIS}Kuv0NlRHCKL%QrHDZ6B@FSndWJp0D8+kjhcXN&$S3Wk|j^ZHeS4oG*$`9+Xl+Z+y^)O#m6JWrigZd>qw5u z)en&%2j)lD;iJ%ESi_5JQzFYvs#P90I{w^5M%y4#CVRn6D|xq7D@rjG1%d{H6X@2zg(<1 zpXvaP^D4#&S5G741xRB%M|BinXcuxU<6{GmuN=#89l&v*$g#{P zL4wP%jE|2vVL!{UjE|2v;gjBSEaR^MI6nR-$1;4A1eaqO{v^OrPC1tO1Bghs9fTO+ zUg#)<;hjr4MtCjcu@CT(9yvxh-je(-z)*fUMz~TQ^b$Tq!r$8iAHk9Pl!88$9~HVE z=sbW6HWQCFBb{i!i8DuDQB_evUis?6A~;4%559H~5{0A~#12HQjL(OsOiGI!fxz>{)4cs&cN=OgRb`avE`o}^VlM+~*a3FsXE|e*_emgRS z5GqCkzrpoifRf?rSL4bKL73Pe#`+O2xc(o2P*^`s=7C|ORR6d@Af#pe$?<29z_irN z0qkapzGMI*CW7TZ3dF<7&r9-KBzlw|{!lVR{&Q0O!SZv`3l15x)#^nox;2jk-62mkHC`tO$L zxgY{rewl9>ln>Yb7ODMK2@&;=eJ|6??KcBGoq1(76uICU2%I#3An(8sBn~+MfMD4X z|CltO`y~R5P^94cX&{zRTu?j^J`nztLP-Cl0g#Z8a0h^xACCT=1QBI8`~hM}NIrVP z8>Jor8VmLojy7{((I0-jg!lEVrqA;bHm{HsCm5nnv?@f-2Q z1E0J=U%bE1_s%%Phwt0Jy62zci+2IW=7=xe@VT@kM|WP!H*+O^Uc??YH&Zv1<(m%zQbHsi`FVVq1 zq2JDa{+jdP`Tsl4gFS&xWdA(pLBn9@!9MmpI0pr|C>{n#)({A|7X{(`gPah-NprA| zq~v}G{tV!9AUYv}_oD2p2+p41oG8gg%J?x9xRO1dNJfKi=@hb(bV&|n&6~)L;7zui z>;@N9KCCW`yGzF5lo?)CyXsd7u8=<372%pej!8IY4~m|u&w;lx{7HRG7gGy9Cr0#r zqf%c57eadOR(C2ve-&_ncpn{C%fW@ij_#A-Z#CdQ_=T5+N8gWuW1|6lTG7KaeRd0Y z&w)z`Z9)hIUN=S`Ej-7LzIDbfdNu*y)z5)#f)<=l54Eu1Vg^sJgRh0;>d3R7TUl^0 z1ph?v!T1U=qlt2EtvmH8^wOHpK&S4d1+UKygc;<)3Z9G{@IeQdbL2U5WVadIBxeJO zAdlz?Xxtu{rNEIepp;5^m7d!k=qD&ZI)*?xhd{c9K=uxSfIL?v5jgK33;{)f6JisM z&YL^TMlX$yK;RX#qZRx*(L)AGpK_=4(9wZTkqAj21OYt{YG&)dmbc@Guj7cXqu^5b zzu(vK3phI6$HQ#}e9maZ*Kx$xam3fL!a3sWIO6L#;_E2BF)J@GK4HYyam3eg#Mg1e z*AcufMtmJdd>u!89Y=f}M|>Sed>w@+v(t zh;i^U(cjkB@m5&z@R`mLU&q0E9`SV)^E~9-zfdY0C<@LG`p_VFFdEQ}Q3_&O{2lo2U_$cDWTE_<`#P%Q2C63} zAvxmfIO6L#;_Eo#>-c}&*YU4>ME2^G__tHv*4MGXP?%t}&6-h|T$Gq>m}$!|GNdFJ z3-j}fXBHcaXUr%zq!cC0OqwD|GUDqvyzfwON9~5i^m@KKxDEEI!+m!U-#3DA*_Y^l z>brx}0~+}M>AQn?h8X$op!hN_Gy#0S2cFa9@8^wtcM!D#Izh!izz84Kll3FNrwHMA z&XD2qSr_-*fxe6Kb~qSr0z3@Ca1koPM|>UO3xe>jGlK6n8~N@ae3J3+{O;hjw@48KDIAaQw@fub(^o5YuR`}I7P!2WDLGPU$qQ60v2Xyuj`@bG>21e0! znhcJ>>$Dc=slUgrr9K(PJg)V6Nx_3$9>~|(uD^3AeB#6d~C@2?1a z0|aiKAS{534ulI(@kOq*@cLZ|AVv7jzv+DV08nB>ngK!!yS;%TA@Mx?hMC65_g)Dv z7576tiIZMbhcM|ShLeFIe6>@;x*F{QIGp^h0;J7kj{ukc0uQAo3Bv$q6H z4!(0EM8c8ZzkxUdX~4m?7MD}xBd@!NgK3lq>%wNBDiB9%l2DhNrVOwjo(ko0O7bk? zMJovp@c=wjExcRMMCk}UUPqMU#J^+eP9TSNOY9)j}Xu=3$<_Z(yqF3E<8p}aHGGDFu2A@Nr+#QQ3;JhV^H@2Gzl_-jLF zf&67Tj9P5xNiF}ic`w=*&V2oSD*6_PEk=g^1nfV~cmeyT9lWzui$aKt%M4-e4S{~k zR4N%z-`-Qb*gohVS$>u376DJL=LFF6?eLXjW|lxV`pP2UUy?goh4lMR3HAs5uO(DJ z+P6GDLdU)-==m~^-%D7oK6AC49MTexFdqJv{Z9mX3h0&6_!s%}-&+?9LJT2g)J9b5`DilR`zj7>Io#f6n<0#V6zV3gMx>#OorHgt(Sr z$3gqblP@z?u>Xg2&^}m7V0^=O%V2@)3sHUoAI~H?;b7k;=pQ=t6R`T7xsw3?>_a3+ z%ffh~3v589H*D2Bg8bx%}lL%$%d@RnTogdK4K zy&huZbAf5Y{?qZ=aS9!0$cf=WeliB|W+@Qw472%?(7%wcgS--Bh6nmtuvEw8e2(L* zePSff!nk;dTSn_=!G{ujxq6uIkp6S#=1ao-33;R!;6O(Uc^>#Te{BN$QR)|wX3CdK z3w+2@tx}ZZEMq;0^9Ri`iYCgCqXhd!JzzT^rf+E36wxr-T^s&oH|L{R(iP(+Y4tN6;tZ%e^4-SL-Ex66hb)n-*eh z59t3A+9v`2Is)wy>T?=E7<|srwQ~U+m{yYfB`$z zih9QR#|QSSWkq}CAr|zAc_t3}5A3sBELVfw@_I(m40>Uu{>A?Dok#tHUH~8ANY4u~ zFVeS9RI^||sDB>D9omVOhkm5UQjky3pC0pI|5|>WARnWJJgg(a{15umYO_Ee;h8MU zp&}!-Qn?`CQ#ii?KCD;cpj|kwSvFk2yDW|`HG5?_>Nis_*RLO_U(i2~d8l6x?L+-A zFrHj=3G^?23HK**y*z6X=9L(T#r_BTg!x6wngnBEYBjySF6XJ zFqbRXGviA+DfK@?aD0{IFu1Q#77z6Q1!fZJU#(Pu-eCR=f!Ig?IY%YUr)QvjCK%t0 zFrV(rQA_jZC}IBm81xVHAb+Svq#q5jPqjZM?8-P{{e%5MH0g$uB(6uqyIqUx3h_S6 zs)hKu%S4?2=};lxa=r6EcEYy_xYrMUPd@T9;SYqnU^NL0h#wg6{p-Wg|4^bE6I2*X zkALq7szPA@EbgKpfDH=)aqtuhn`0r!=hG=-)-HhuqBvA11g{tk$cTSq2@EU6%n7E=8w;f{oc{6t52C1ozSdyELoo)49_5Fmibq&LXoj>%9{_}; zN%~fTtT?0xlZfC4qmdVYcvyPrR*y--M|uz`gqQMigdc$#j@~8F%O62OdeC!hFg+X; z1)XAfV37Y_fxH7la2~()f~vk2Js(I6>0tVx;=e2V5ikv|9~sHtt7!=`j$j%7zkjc$ zLYxDJJOlkrzgKet6pA_xQZs3mCX14wUX{nb491fjqMBrY6o1q2)yQK2$8unhzgL6z zzygC@zibEFpl~?(=Spi3A4n|W27-T_ O87HC)hd)3JiT?toP9~KA literal 0 HcmV?d00001 diff --git a/clang/test/Driver/fat_archive_amdgpu.cpp b/clang/test/Driver/fat_archive_amdgpu.cpp new file mode 100644 index 0000000..b64ba8b --- /dev/null +++ b/clang/test/Driver/fat_archive_amdgpu.cpp @@ -0,0 +1,81 @@ +// REQUIRES: clang-driver +// REQUIRES: x86-registered-target +// REQUIRES: amdgpu-registered-target + +// See the steps to create a fat archive are given at the end of the file. + +// Given a FatArchive, clang-offload-bundler should be called to create a +// device specific archive, which should be passed to llvm-link. +// RUN: %clang -O2 -### -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906 %s -L%S/Inputs/openmp_static_device_link -lFatArchive 2>&1 | FileCheck %s +// CHECK: clang{{.*}}"-cc1"{{.*}}"-triple" "amdgcn-amd-amdhsa"{{.*}}"-emit-llvm-bc"{{.*}}"-target-cpu" "[[GPU:gfx[0-9]+]]"{{.*}}"-o" "[[HOSTBC:.*.bc]]" "-x" "c++"{{.*}}.cpp +// CHECK: clang-offload-bundler" "-unbundle" "-type=a" "-inputs={{.*}}/Inputs/openmp_static_device_link/libFatArchive.a" "-targets=openmp-amdgcn-amd-amdhsa-[[GPU]]" "-outputs=[[DEVICESPECIFICARCHIVE:.*.a]]" "-allow-missing-bundles" +// CHECK: llvm-link{{.*}}"[[HOSTBC]]" "[[DEVICESPECIFICARCHIVE]]" "-o" "{{.*}}-[[GPU]]-linked-{{.*}}.bc" +// CHECK: ld"{{.*}}" "-L{{.*}}/Inputs/openmp_static_device_link" "{{.*}} "-lFatArchive" "{{.*}}" "-lomp{{.*}}-lomptarget" +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +#define N 10 + +#pragma omp declare target +// Functions defined in Fat Archive. +extern "C" void func_present(float *, float *, unsigned); + +#ifdef MISSING +// Function not defined in the fat archive. +extern "C" void func_missing(float *, float *, unsigned); +#endif + +#pragma omp end declare target + +int main() { + float in[N], out[N], sum = 0; + unsigned i; + +#pragma omp parallel for + for (i = 0; i < N; ++i) { + in[i] = i; + } + + func_present(in, out, N); // Returns out[i] = a[i] * 0 + +#ifdef MISSING + func_missing(in, out, N); // Should throw an error here +#endif + +#pragma omp parallel for reduction(+ \ + : sum) + for (i = 0; i < N; ++i) + sum += out[i]; + + if (!sum) + return 0; + return sum; +} + +#endif + +/*********************************************** + Steps to create Fat Archive (libFatArchive.a) +************************************************ +***************** File: func_1.c *************** +void func_present(float* in, float* out, unsigned n){ + unsigned i; + #pragma omp target teams distribute parallel for map(to: in[0:n]) map(from: out[0:n]) + for(i=0; i&1 | FileCheck %s +// CHECK: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64"{{.*}}"-target-cpu" "[[GPU:sm_[0-9]+]]"{{.*}}"-o" "[[HOSTBC:.*.s]]" "-x" "c++"{{.*}}.cpp +// CHECK: clang-offload-bundler" "-unbundle" "-type=a" "-inputs={{.*}}/Inputs/openmp_static_device_link/libFatArchive.a" "-targets=openmp-nvptx64-[[GPU]]" "-outputs=[[DEVICESPECIFICARCHIVE:.*.a]]" "-allow-missing-bundles" +// CHECK: clang-nvlink-wrapper{{.*}}"-o" "{{.*}}.out" "-arch" "[[GPU]]" "{{.*}}[[DEVICESPECIFICARCHIVE]]" +// CHECK: ld"{{.*}}" "-L{{.*}}/Inputs/openmp_static_device_link" "{{.*}} "-lFatArchive" "{{.*}}" "-lomp{{.*}}-lomptarget" +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +#define N 10 + +#pragma omp declare target +// Functions defined in Fat Archive. +extern "C" void func_present(float *, float *, unsigned); + +#ifdef MISSING +// Function not defined in the fat archive. +extern "C" void func_missing(float *, float *, unsigned); +#endif + +#pragma omp end declare target + +int main() { + float in[N], out[N], sum = 0; + unsigned i; + +#pragma omp parallel for + for (i = 0; i < N; ++i) { + in[i] = i; + } + + func_present(in, out, N); // Returns out[i] = a[i] * 0 + +#ifdef MISSING + func_missing(in, out, N); // Should throw an error here +#endif + +#pragma omp parallel for reduction(+ \ + : sum) + for (i = 0; i < N; ++i) + sum += out[i]; + + if (!sum) + return 0; + return sum; +} + +#endif + +/*********************************************** + Steps to create Fat Archive (libFatArchive.a) +************************************************ +***************** File: func_1.c *************** +void func_present(float* in, float* out, unsigned n){ + unsigned i; + #pragma omp target teams distribute parallel for map(to: in[0:n]) map(from: out[0:n]) + for(i=0; i