[clang] 8784b6a - [Clang] Allow bitcode linking when the input is LLVM-IR
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Tue Jun 20 06:03:07 PDT 2023
Author: Joseph Huber
Date: 2023-06-20T08:02:58-05:00
New Revision: 8784b6a8540f4a333690e7233587859f1d82f620
URL: https://github.com/llvm/llvm-project/commit/8784b6a8540f4a333690e7233587859f1d82f620
DIFF: https://github.com/llvm/llvm-project/commit/8784b6a8540f4a333690e7233587859f1d82f620.diff
LOG: [Clang] Allow bitcode linking when the input is LLVM-IR
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
```
Reviewed By: yaxunl
Differential Revision: https://reviews.llvm.org/D152391
Added:
clang/test/CodeGen/link-builtin-bitcode.c
Modified:
clang/include/clang/CodeGen/CodeGenAction.h
clang/lib/CodeGen/CGCall.cpp
clang/lib/CodeGen/CGCall.h
clang/lib/CodeGen/CodeGenAction.cpp
clang/test/CodeGen/link-bitcode-file.c
Removed:
################################################################################
diff --git a/clang/include/clang/CodeGen/CodeGenAction.h b/clang/include/clang/CodeGen/CodeGenAction.h
index b5721344046d0..821e80919fc84 100644
--- a/clang/include/clang/CodeGen/CodeGenAction.h
+++ b/clang/include/clang/CodeGen/CodeGenAction.h
@@ -53,6 +53,9 @@ class CodeGenAction : public ASTFrontendAction {
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,
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index fc77566c790c1..61029c9226d7b 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -1845,8 +1845,9 @@ addMergableDefaultFunctionAttributes(const CodeGenOptions &CodeGenOpts,
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) {
@@ -1967,7 +1968,7 @@ void CodeGenModule::getTrivialDefaultFunctionAttributes(
}
}
- 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
@@ -1978,8 +1979,8 @@ void CodeGenModule::getTrivialDefaultFunctionAttributes(
// 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);
}
@@ -1990,36 +1991,25 @@ void CodeGenModule::getTrivialDefaultFunctionAttributes(
}
}
-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);
- }
-}
+/// 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, const TargetOptions &TargetOpts,
+ bool WillInternalize) {
-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);
-}
+ // Here we only extract the options that are relevant compared to the version
+ // from GetCPUAndFeaturesAttributes.
+ if (!TargetOpts.CPU.empty())
+ FuncAttrs.addAttribute("target-cpu", TargetOpts.CPU);
+ if (!TargetOpts.TuneCPU.empty())
+ FuncAttrs.addAttribute("tune-cpu", TargetOpts.TuneCPU);
-/// 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);
+ ::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
@@ -2064,6 +2054,52 @@ void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
F.addFnAttrs(FuncAttrs);
}
+void clang::CodeGen::mergeDefaultFunctionDefinitionAttributes(
+ llvm::Function &F, const CodeGenOptions CodeGenOpts,
+ const LangOptions &LangOpts, const TargetOptions &TargetOpts,
+ bool WillInternalize) {
+
+ ::mergeDefaultFunctionDefinitionAttributes(F, CodeGenOpts, LangOpts,
+ TargetOpts, WillInternalize);
+}
+
+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) {
+ ::mergeDefaultFunctionDefinitionAttributes(F, getCodeGenOpts(), getLangOpts(),
+ getTarget().getTargetOpts(),
+ WillInternalize);
+}
+
void CodeGenModule::addDefaultFunctionDefinitionAttributes(
llvm::AttrBuilder &attrs) {
getDefaultFunctionAttributes(/*function name*/ "", /*optnone*/ false,
diff --git a/clang/lib/CodeGen/CGCall.h b/clang/lib/CodeGen/CGCall.h
index 59c3f304f59b9..824f0a9a88299 100644
--- a/clang/lib/CodeGen/CGCall.h
+++ b/clang/lib/CodeGen/CGCall.h
@@ -30,6 +30,7 @@ class Value;
namespace clang {
class Decl;
class FunctionDecl;
+class TargetOptions;
class VarDecl;
namespace CodeGen {
@@ -377,6 +378,14 @@ class ReturnValueSlot {
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
diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp
index 4aa51e956655f..4879bcd6a42a5 100644
--- a/clang/lib/CodeGen/CodeGenAction.cpp
+++ b/clang/lib/CodeGen/CodeGenAction.cpp
@@ -7,6 +7,7 @@
//===----------------------------------------------------------------------===//
#include "clang/CodeGen/CodeGenAction.h"
+#include "CGCall.h"
#include "CodeGenModule.h"
#include "CoverageMappingGen.h"
#include "MacroPPCallbacks.h"
@@ -262,7 +263,7 @@ namespace clang {
}
// 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)
@@ -271,8 +272,8 @@ namespace clang {
// 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();
@@ -280,15 +281,14 @@ namespace clang {
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)
@@ -357,7 +357,7 @@ namespace clang {
}
// Link each LinkModule into our module.
- if (LinkInModules())
+ if (LinkInModules(getModule()))
return;
for (auto &F : getModule()->functions()) {
@@ -993,6 +993,36 @@ CodeGenAction::~CodeGenAction() {
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() {
@@ -1048,30 +1078,8 @@ CodeGenAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) {
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))
+ return nullptr;
CoverageSourceInfo *CoverageInfo = nullptr;
// Add the preprocessor callback only when the coverage mapping is generated.
@@ -1139,6 +1147,10 @@ CodeGenAction::loadModule(MemoryBufferRef MBRef) {
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;
@@ -1218,6 +1230,11 @@ void CodeGenAction::ExecuteAction() {
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);
diff --git a/clang/test/CodeGen/link-bitcode-file.c b/clang/test/CodeGen/link-bitcode-file.c
index df04ec2bec0d9..58fee64a95138 100644
--- a/clang/test/CodeGen/link-bitcode-file.c
+++ b/clang/test/CodeGen/link-bitcode-file.c
@@ -11,6 +11,14 @@
// 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
diff --git a/clang/test/CodeGen/link-builtin-bitcode.c b/clang/test/CodeGen/link-builtin-bitcode.c
new file mode 100644
index 0000000000000..b3b54badf3f82
--- /dev/null
+++ b/clang/test/CodeGen/link-builtin-bitcode.c
@@ -0,0 +1,42 @@
+// 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" }
+//.
More information about the cfe-commits
mailing list