[Clang] Allow bitcode linking when the input is LLVM-IR

[jhuber6]

Summary

Clang provides the -mlink-bitcode-file and -mlink-builtin-bitcode
options to insert LLVM-IR into the current TU. These are usefuly
primarily for including LLVM-IR files that require special handling to
be correct and cannot be linked normally, such as GPU vendor libraries
like libdevice.10.bc. Currently these options can only be used if the
source input goes through the AST consumer path. This patch makes the
changes necessary to also support this when the input is LLVM-IR. This
will allow the following operation:

clang in.bc -Xclang -mlink-builtin-bitcode -Xclang libdevice.10.bc

[JonChesterfield]

That's cleaner than I expected, thanks. Might be reasonable to factor out the method as an initial NFC then insert the call to it along with the new test case as the functional change.

[tra]

clang in.bc -Xclang -mlink-builtin-bitcode -Xclang libdevice.10.bc

If that's something we intend to expose to the user, should we consider promoting it to a top-level driver option?

[jhuber6]

In D152391#4404106, @tra wrote:

clang in.bc -Xclang -mlink-builtin-bitcode -Xclang libdevice.10.bc

If that's something we intend to expose to the user, should we consider promoting it to a top-level driver option?

I'm not sure, this probably wouldn't be relevant for non-compiler developers so I'm fine keeping it as an -Xclang workaround.

[jhuber6]

Just realized this is probably going to be a bit more painful. the attribute propagation pass requires a CodeGenModule which isn't built without an ASTContext so it's not available here. Nothing those functions do explicitly requires the full CGM, they only use the options and the target info. So it's possible to make these functions more generic using only the CompilerInstance instead, but it'll be a lot nosier.

[jhuber6]

Scratch that, ASTContext::getFunctionFeatureMap is used. I wonder if we could forgo that.

Actually, that's only used if there's a GlobalDecl which we don't have when we call this. So that could be ignored here.

[jhuber6]

Updating, in order to do this starting with bitcode I had to expose a helper
that performs this operation using the options directly rather than through the
CodeGenModule. This should keep the existing interfaces intact by shuttling
them through a new internal function while also letting me write this standalone
function that can be used without a CodeGenModule.

[jhuber6]

Add a better test to show that attributes are replaced and default attributes are added.

[jhuber6]

ping

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2055–2109

can we reorder these functions to minimize the diffs? Also some comments about the difference among these functions may help.

clang/lib/CodeGen/CodeGenAction.cpp

266

Removing the default argument may make the code clearer

[jhuber6]

Addressing comments.

[jhuber6]

clang/lib/CodeGen/CGCall.cpp

2055–2109

I added some comments. These can't be reordered because they're new functions, I simply added a layer of indirection to go through a common helper so we can call the main bulk of the code without needing the full CGM.

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2105–2106

should this be removed? seems redundant

[jhuber6]

clang/lib/CodeGen/CGCall.cpp

2105–2106

I wasn't entirely sure, there's a lot of work done in that function but pretty much all of it goes away since we just use GlobalDecl(). So for the version that doesn't have access to CGM I copies out the only two relevant lines and put them in line 2063. We could probably make this use the same code but I wasn't sure if there was something else we might need it for.

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2105–2106

I tend to think they are equivalent. If GetCPUAndFeaturesAttributes did something else, it may be overwritten by the counterparts in mergeDefaultFunctionDefinitionAttributes.

[jhuber6]

Removing

[jhuber6]

Format

[JonChesterfield]

clang/lib/CodeGen/CGCall.cpp

2000–2001

I'm used to this sort of copy-some-args-and-not-others showing up in bug reports. Could this patch be re-ordered to make it apparent what functional changes are happening relative to the current head?

[jhuber6]

clang/lib/CodeGen/CGCall.cpp

2000–2001

It's the same that will get added in in the GetCPUAndFeaturesAttributes function if you pass in a null declaration.

[jhuber6]

ping
@ronlieb

[yaxunl]

LGTM. Thanks.

[MaskRay]

LGTM. CC1 -mlink-bitcode-file is an interesting option from 2011: f1d76db466b2a50781c0754b86ac994dd07b5041
I wonder what the original use case is...

clang/include/clang/CodeGen/CodeGenAction.h

private:
 
   /// Bitcode modules to link in to our module.
   SmallVector<LinkModule, 4> LinkModules;
   llvm::LLVMContext *VMContext;
   bool OwnsVMContext;
 
   std::unique_ptr<llvm::Module> loadModule(llvm::MemoryBufferRef MBRef);
 
+  /// Load bitcode modules to link into our module from the options.
+  bool loadLinkModules(CompilerInstance &CI);
+
 protected:
   /// Create a new code generation action.  If the optional \p _VMContext
   /// parameter is supplied, the action uses it without taking ownership,
   /// otherwise it creates a fresh LLVM context and takes ownership.
   CodeGenAction(unsigned _Act, llvm::LLVMContext *_VMContext = nullptr);
 
   bool hasIRSupport() const override;
 

clang/lib/CodeGen/CGCall.h

 namespace llvm {
 class Type;
 class Value;
 } // namespace llvm
 
 namespace clang {
 class Decl;
 class FunctionDecl;
+class TargetOptions;
 class VarDecl;
 
 namespace CodeGen {
 
 /// Abstract information about a function or function prototype.
 class CGCalleeInfo {
   /// The function prototype of the callee.
   const FunctionProtoType *CalleeProtoTy;
public:
 
   bool isNull() const { return !Addr.isValid(); }
   bool isVolatile() const { return IsVolatile; }
   Address getValue() const { return Addr; }
   bool isUnused() const { return IsUnused; }
   bool isExternallyDestructed() const { return IsExternallyDestructed; }
 };
 
+/// Helper to add attributes to \p F according to the CodeGenOptions and
+/// LangOptions without requiring a CodeGenModule to be constructed.
+void mergeDefaultFunctionDefinitionAttributes(llvm::Function &F,
+                                              const CodeGenOptions CodeGenOpts,
+                                              const LangOptions &LangOpts,
+                                              const TargetOptions &TargetOpts,
+                                              bool WillInternalize);
+
 } // end namespace CodeGen
 } // end namespace clang
 
 #endif

clang/lib/CodeGen/CGCall.cpp

 /// attributes in the linked library.
 static void
 addMergableDefaultFunctionAttributes(const CodeGenOptions &CodeGenOpts,
                                      llvm::AttrBuilder &FuncAttrs) {
   addDenormalModeAttrs(CodeGenOpts.FPDenormalMode, CodeGenOpts.FP32DenormalMode,
                        FuncAttrs);
 }
 
-void CodeGenModule::getTrivialDefaultFunctionAttributes(
-    StringRef Name, bool HasOptnone, bool AttrOnCallSite,
+static void getTrivialDefaultFunctionAttributes(
+    StringRef Name, bool HasOptnone, const CodeGenOptions &CodeGenOpts,
+    const LangOptions &LangOpts, bool AttrOnCallSite,
     llvm::AttrBuilder &FuncAttrs) {
   // OptimizeNoneAttr takes precedence over -Os or -Oz. No warning needed.
   if (!HasOptnone) {
     if (CodeGenOpts.OptimizeSize)
       FuncAttrs.addAttribute(llvm::Attribute::OptimizeForSize);
     if (CodeGenOpts.OptimizeSize == 2)
       FuncAttrs.addAttribute(llvm::Attribute::MinSize);
   }
case llvm::ZeroCallUsedRegs::ZeroCallUsedRegsKind::AllArg:
       FuncAttrs.addAttribute("zero-call-used-regs", "all-arg");
       break;
     case llvm::ZeroCallUsedRegs::ZeroCallUsedRegsKind::All:
       FuncAttrs.addAttribute("zero-call-used-regs", "all");
       break;
     }
   }
 
-  if (getLangOpts().assumeFunctionsAreConvergent()) {
+  if (LangOpts.assumeFunctionsAreConvergent()) {
     // Conservatively, mark all functions and calls in CUDA and OpenCL as
     // convergent (meaning, they may call an intrinsically convergent op, such
     // as __syncthreads() / barrier(), and so can't have certain optimizations
     // applied around them).  LLVM will remove this attribute where it safely
     // can.
     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
   }
 
   // TODO: NoUnwind attribute should be added for other GPU modes HIP,
   // OpenMP offload. AFAIK, neither of them support exceptions in device code.
-  if ((getLangOpts().CUDA && getLangOpts().CUDAIsDevice) ||
-      getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) {
+  if ((LangOpts.CUDA && LangOpts.CUDAIsDevice) || LangOpts.OpenCL ||
+      LangOpts.SYCLIsDevice) {
     FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
   }
 
   for (StringRef Attr : CodeGenOpts.DefaultFunctionAttrs) {
     StringRef Var, Value;
     std::tie(Var, Value) = Attr.split('=');
     FuncAttrs.addAttribute(Var, Value);
   }
 }
 
-void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
-                                                 bool HasOptnone,
-                                                 bool AttrOnCallSite,
-                                                 llvm::AttrBuilder &FuncAttrs) {
-  getTrivialDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite,
-                                      FuncAttrs);
-  if (!AttrOnCallSite) {
-    // If we're just getting the default, get the default values for mergeable
-    // attributes.
-    addMergableDefaultFunctionAttributes(CodeGenOpts, FuncAttrs);
-  }
-}
-
-void CodeGenModule::addDefaultFunctionDefinitionAttributes(llvm::Function &F) {
-  llvm::AttrBuilder FuncAttrs(F.getContext());
-  getDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
-                               /* AttrOnCallSite = */ false, FuncAttrs);
-  // TODO: call GetCPUAndFeaturesAttributes?
-  F.addFnAttrs(FuncAttrs);
-}
-
-/// Apply default attributes to \p F, accounting for merge semantics of
-/// attributes that should not overwrite existing attributes.
-void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
-    llvm::Function &F, bool WillInternalize) {
-  llvm::AttrBuilder FuncAttrs(F.getContext());
-  getTrivialDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
-                                      /*AttrOnCallSite=*/false, FuncAttrs);
-  GetCPUAndFeaturesAttributes(GlobalDecl(), FuncAttrs,
-                              /*AddTargetFeatures=*/false);
+/// Adds attributes to \p F according to our \p CodeGenOpts and \p LangOpts, as
+/// though we had emitted it ourselves. We remove any attributes on F that
+/// conflict with the attributes we add here.
+static void mergeDefaultFunctionDefinitionAttributes(
+    llvm::Function &F, const CodeGenOptions CodeGenOpts,
+    const LangOptions &LangOpts, bool WillInternalize,
+  ::getTrivialDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
+                                        CodeGenOpts, LangOpts,
+                                        /*AttrOnCallSite=*/false, FuncAttrs);
 
   if (!WillInternalize && F.isInterposable()) {
     // Do not promote "dynamic" denormal-fp-math to this translation unit's
     // setting for weak functions that won't be internalized. The user has no
     // real control for how builtin bitcode is linked, so we shouldn't assume
     // later copies will use a consistent mode.
     F.addFnAttrs(FuncAttrs);
     return;
FuncAttrs.addAttribute("denormal-fp-math-f32",
                            CodeGenOpts.FP32DenormalMode.str());
   }
 
   F.removeFnAttrs(AttrsToRemove);
   addDenormalModeAttrs(Merged, MergedF32, FuncAttrs);
   F.addFnAttrs(FuncAttrs);
 }
 
+void clang::CodeGen::mergeDefaultFunctionDefinitionAttributes(
+    llvm::Function &F, const CodeGenOptions CodeGenOpts,
+    const LangOptions &LangOpts, const TargetOptions &TargetOpts,
+    bool WillInternalize) {
+
+  // Here we only extract the options that are relevant compared to the version
+  // from GetCPUAndFeaturesAttributes below.
+  llvm::AttrBuilder FuncAttrs(F.getContext());
+  if (!TargetOpts.CPU.empty())
+    FuncAttrs.addAttribute("target-cpu", TargetOpts.CPU);
+  if (!TargetOpts.TuneCPU.empty())
+    FuncAttrs.addAttribute("tune-cpu", TargetOpts.TuneCPU);
+
+  ::mergeDefaultFunctionDefinitionAttributes(F, CodeGenOpts, LangOpts,
+                                             WillInternalize, FuncAttrs);
+}
+
+void CodeGenModule::getTrivialDefaultFunctionAttributes(
+    StringRef Name, bool HasOptnone, bool AttrOnCallSite,
+    llvm::AttrBuilder &FuncAttrs) {
+  ::getTrivialDefaultFunctionAttributes(Name, HasOptnone, getCodeGenOpts(),
+                                        getLangOpts(), AttrOnCallSite,
+                                        FuncAttrs);
+}
+
+void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
+                                                 bool HasOptnone,
+                                                 bool AttrOnCallSite,
+                                                 llvm::AttrBuilder &FuncAttrs) {
+  getTrivialDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite,
+                                      FuncAttrs);
+  // If we're just getting the default, get the default values for mergeable
+  // attributes.
+  if (!AttrOnCallSite)
+    addMergableDefaultFunctionAttributes(CodeGenOpts, FuncAttrs);
+}
+
+void CodeGenModule::addDefaultFunctionDefinitionAttributes(llvm::Function &F) {
+  llvm::AttrBuilder FuncAttrs(F.getContext());
+  getDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
+                               /* AttrOnCallSite = */ false, FuncAttrs);
+  // TODO: call GetCPUAndFeaturesAttributes?
+  F.addFnAttrs(FuncAttrs);
+}
+
+/// Apply default attributes to \p F, accounting for merge semantics of
+/// attributes that should not overwrite existing attributes.
+void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
+    llvm::Function &F, bool WillInternalize) {
+  llvm::AttrBuilder FuncAttrs(F.getContext());
+  GetCPUAndFeaturesAttributes(GlobalDecl(), FuncAttrs,
+                              /*AddTargetFeatures=*/false);
+
+  ::mergeDefaultFunctionDefinitionAttributes(F, getCodeGenOpts(), getLangOpts(),
+                                             WillInternalize, FuncAttrs);
+}
+
 void CodeGenModule::addDefaultFunctionDefinitionAttributes(
     llvm::AttrBuilder &attrs) {
   getDefaultFunctionAttributes(/*function name*/ "", /*optnone*/ false,
                                /*for call*/ false, attrs);
   GetCPUAndFeaturesAttributes(GlobalDecl(), attrs);
 }
 
 static void addNoBuiltinAttributes(llvm::AttrBuilder &FuncAttrs,

clang/lib/CodeGen/CodeGenAction.cpp

 //===--- CodeGenAction.cpp - LLVM Code Generation Frontend Action ---------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "clang/CodeGen/CodeGenAction.h"
+#include "CGCall.h"
 #include "CodeGenModule.h"
 #include "CoverageMappingGen.h"
 #include "MacroPPCallbacks.h"
 #include "clang/AST/ASTConsumer.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclGroup.h"
 #include "clang/Basic/DiagnosticFrontend.h"
public:
 
     void HandleInterestingDecl(DeclGroupRef D) override {
       // Ignore interesting decls from the AST reader after IRGen is finished.
       if (!IRGenFinished)
         HandleTopLevelDecl(D);
     }
 
     // Links each entry in LinkModules into our module.  Returns true on error.
-    bool LinkInModules() {
+    bool LinkInModules(llvm::Module *M) {
       for (auto &LM : LinkModules) {
         assert(LM.Module && "LinkModule does not actually have a module");
         if (LM.PropagateAttrs)
           for (Function &F : *LM.Module) {
             // Skip intrinsics. Keep consistent with how intrinsics are created
             // in LLVM IR.
             if (F.isIntrinsic())
               continue;
-            Gen->CGM().mergeDefaultFunctionDefinitionAttributes(F,
-                                                                LM.Internalize);
+            CodeGen::mergeDefaultFunctionDefinitionAttributes(
+                F, CodeGenOpts, LangOpts, TargetOpts, LM.Internalize);
           }
 
         CurLinkModule = LM.Module.get();
 
         bool Err;
         if (LM.Internalize) {
           Err = Linker::linkModules(
-              *getModule(), std::move(LM.Module), LM.LinkFlags,
+              *M, std::move(LM.Module), LM.LinkFlags,
               [](llvm::Module &M, const llvm::StringSet<> &GVS) {
                 internalizeModule(M, [&GVS](const llvm::GlobalValue &GV) {
                   return !GV.hasName() || (GVS.count(GV.getName()) == 0);
                 });
               });
         } else {
-          Err = Linker::linkModules(*getModule(), std::move(LM.Module),
-                                    LM.LinkFlags);
+          Err = Linker::linkModules(*M, std::move(LM.Module), LM.LinkFlags);
         }
 
         if (Err)
           return true;
       }
       LinkModules.clear();
       return false; // success
     }
void HandleTranslationUnit(ASTContext &C) override {
       }
 
       if (CodeGenOpts.DiagnosticsMisExpectTolerance) {
         Ctx.setDiagnosticsMisExpectTolerance(
             CodeGenOpts.DiagnosticsMisExpectTolerance);
       }
 
       // Link each LinkModule into our module.
-      if (LinkInModules())
+      if (LinkInModules(getModule()))
         return;
 
       for (auto &F : getModule()->functions()) {
         if (const Decl *FD = Gen->GetDeclForMangledName(F.getName())) {
           auto Loc = FD->getASTContext().getFullLoc(FD->getLocation());
           // TODO: use a fast content hash when available.
           auto NameHash = llvm::hash_value(F.getName());
           ManglingFullSourceLocs.push_back(std::make_pair(NameHash, Loc));
: Act(_Act), VMContext(_VMContext ? _VMContext : new LLVMContext),
       OwnsVMContext(!_VMContext) {}
 
 CodeGenAction::~CodeGenAction() {
   TheModule.reset();
   if (OwnsVMContext)
     delete VMContext;
 }
 
+bool CodeGenAction::loadLinkModules(CompilerInstance &CI) {
+  if (!LinkModules.empty())
+    return false;
+
+  for (const CodeGenOptions::BitcodeFileToLink &F :
+       CI.getCodeGenOpts().LinkBitcodeFiles) {
+    auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename);
+    if (!BCBuf) {
+      CI.getDiagnostics().Report(diag::err_cannot_open_file)
+          << F.Filename << BCBuf.getError().message();
+      LinkModules.clear();
+      return true;
+    }
+
+    Expected<std::unique_ptr<llvm::Module>> ModuleOrErr =
+        getOwningLazyBitcodeModule(std::move(*BCBuf), *VMContext);
+    if (!ModuleOrErr) {
+      handleAllErrors(ModuleOrErr.takeError(), [&](ErrorInfoBase &EIB) {
+        CI.getDiagnostics().Report(diag::err_cannot_open_file)
+            << F.Filename << EIB.message();
+      });
+      LinkModules.clear();
+      return true;
+    }
+    LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs,
+                           F.Internalize, F.LinkFlags});
+  }
+  return false;
+}
+
 bool CodeGenAction::hasIRSupport() const { return true; }
 
 void CodeGenAction::EndSourceFileAction() {
   // If the consumer creation failed, do nothing.
   if (!getCompilerInstance().hasASTConsumer())
     return;
 
   // Steal the module from the consumer.
CodeGenAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) {
   std::unique_ptr<raw_pwrite_stream> OS = CI.takeOutputStream();
   if (!OS)
     OS = GetOutputStream(CI, InFile, BA);
 
   if (BA != Backend_EmitNothing && !OS)
     return nullptr;
 
   // Load bitcode modules to link with, if we need to.
-  if (LinkModules.empty())
-    for (const CodeGenOptions::BitcodeFileToLink &F :
-         CI.getCodeGenOpts().LinkBitcodeFiles) {
-      auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename);
-      if (!BCBuf) {
-        CI.getDiagnostics().Report(diag::err_cannot_open_file)
-            << F.Filename << BCBuf.getError().message();
-        LinkModules.clear();
-        return nullptr;
-      }
-
-      Expected<std::unique_ptr<llvm::Module>> ModuleOrErr =
-          getOwningLazyBitcodeModule(std::move(*BCBuf), *VMContext);
-      if (!ModuleOrErr) {
-        handleAllErrors(ModuleOrErr.takeError(), [&](ErrorInfoBase &EIB) {
-          CI.getDiagnostics().Report(diag::err_cannot_open_file)
-              << F.Filename << EIB.message();
-        });
-        LinkModules.clear();
-        return nullptr;
-      }
-      LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs,
-                             F.Internalize, F.LinkFlags});
-    }
+  if (loadLinkModules(CI))
 
   CoverageSourceInfo *CoverageInfo = nullptr;
   // Add the preprocessor callback only when the coverage mapping is generated.
   if (CI.getCodeGenOpts().CoverageMapping)
     CoverageInfo = CodeGen::CoverageMappingModuleGen::setUpCoverageCallbacks(
         CI.getPreprocessor());
 
   std::unique_ptr<BackendConsumer> Result(new BackendConsumer(
if (!CI.getCodeGenOpts().ThinLTOIndexFile.empty()) {
     }
     Expected<std::unique_ptr<llvm::Module>> MOrErr =
         Bm->parseModule(*VMContext);
     if (!MOrErr)
       return DiagErrors(MOrErr.takeError());
     return std::move(*MOrErr);
   }
 
+  // Load bitcode modules to link with, if we need to.
+  if (loadLinkModules(CI))
+    return nullptr;
+
   llvm::SMDiagnostic Err;
   if (std::unique_ptr<llvm::Module> M = parseIR(MBRef, Err, *VMContext))
     return M;
 
   // Translate from the diagnostic info to the SourceManager location if
   // available.
   // TODO: Unify this with ConvertBackendLocation()
   SourceLocation Loc;
void CodeGenAction::ExecuteAction() {
 
   // Set clang diagnostic handler. To do this we need to create a fake
   // BackendConsumer.
   BackendConsumer Result(BA, CI.getDiagnostics(), &CI.getVirtualFileSystem(),
                          CI.getHeaderSearchOpts(), CI.getPreprocessorOpts(),
                          CI.getCodeGenOpts(), CI.getTargetOpts(),
                          CI.getLangOpts(), TheModule.get(),
                          std::move(LinkModules), *VMContext, nullptr);
+
+  // Link in each pending link module.
+  if (Result.LinkInModules(&*TheModule))
+    return;
+
   // PR44896: Force DiscardValueNames as false. DiscardValueNames cannot be
   // true here because the valued names are needed for reading textual IR.
   Ctx.setDiscardValueNames(false);
   Ctx.setDiagnosticHandler(
       std::make_unique<ClangDiagnosticHandler>(CodeGenOpts, &Result));
 
   Expected<std::unique_ptr<llvm::ToolOutputFile>> OptRecordFileOrErr =
       setupLLVMOptimizationRemarks(

clang/test/CodeGen/link-bitcode-file.c

 // RUN: %clang_cc1 -triple i386-pc-linux-gnu -DBITCODE -emit-llvm-bc -o %t.bc %s
 // RUN: %clang_cc1 -triple i386-pc-linux-gnu -DBITCODE2 -emit-llvm-bc -o %t-2.bc %s
 // RUN: %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file %t.bc \
 // RUN:     -O3 -emit-llvm -o - %s | FileCheck -check-prefix=CHECK-NO-BC %s
 // RUN: %clang_cc1 -triple i386-pc-linux-gnu -O3 -emit-llvm -o - \
 // RUN:     -mlink-bitcode-file %t.bc -mlink-bitcode-file %t-2.bc %s \
 // RUN:     | FileCheck -check-prefix=CHECK-NO-BC -check-prefix=CHECK-NO-BC2 %s
 // RUN: not %clang_cc1 -triple i386-pc-linux-gnu -DBITCODE -O3 -emit-llvm -o - \
 // RUN:     -mlink-bitcode-file %t.bc %s 2>&1 | FileCheck -check-prefix=CHECK-BC %s
 // Make sure we deal with failure to load the file.
 // RUN: not %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file no-such-file.bc \
 // RUN:    -emit-llvm -o - %s 2>&1 | FileCheck -check-prefix=CHECK-NO-FILE %s
 
+// Make sure we can perform the same options if the input is LLVM-IR
+// RUN: %clang_cc1 -triple i386-pc-linux-gnu -emit-llvm-bc -o %t-in.bc %s
+// RUN: %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file %t.bc \
+// RUN:     -O3 -emit-llvm -o - %t-in.bc | FileCheck -check-prefix=CHECK-NO-BC %s
+// RUN: %clang_cc1 -triple i386-pc-linux-gnu -O3 -emit-llvm -o - \
+// RUN:     -mlink-bitcode-file %t.bc -mlink-bitcode-file %t-2.bc %t-in.bc \
+// RUN:     | FileCheck -check-prefix=CHECK-NO-BC -check-prefix=CHECK-NO-BC2 %s
+
 int f(void);
 
 #ifdef BITCODE
 
 extern int f2(void);
 // CHECK-BC: fatal error: cannot link module {{.*}}'f': symbol multiply defined
 int f(void) {
   f2();

clang/test/CodeGen/link-builtin-bitcode.c

+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-attributes --check-globals --include-generated-funcs --version 2
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx803 -DBITCODE -emit-llvm-bc -o %t-lib.bc %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm-bc -o %t.bc %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm \
+// RUN:   -mlink-builtin-bitcode %t-lib.bc -o - %t.bc | FileCheck %s
+
+#ifdef BITCODE
+int foo(void) { return 42; }
+int x = 12;
+#endif
+
+extern int foo(void);
+extern int x;
+
+int bar() { return foo() + x; }
+//.
+// CHECK: @x = internal addrspace(1) global i32 12, align 4
+//.
+// CHECK: Function Attrs: noinline nounwind optnone
+// CHECK-LABEL: define dso_local i32 @bar
+// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// CHECK-NEXT:    [[CALL:%.*]] = call i32 @foo()
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @x to ptr), align 4
+// CHECK-NEXT:    [[ADD:%.*]] = add nsw i32 [[CALL]], [[TMP0]]
+// CHECK-NEXT:    ret i32 [[ADD]]
+//
+//
+// CHECK: Function Attrs: convergent noinline nounwind optnone
+// CHECK-LABEL: define internal i32 @foo
+// CHECK-SAME: () #[[ATTR1:[0-9]+]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// CHECK-NEXT:    ret i32 42
+//
+//.
+// CHECK: attributes #0 = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
+// CHECK: attributes #1 = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
+//.