Skip to content

Commit 8784b6a

Browse files
jhuber6jhuber6
committed
[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
1 parent 87da4b6 commit 8784b6a

File tree

6 files changed

+178
-63
lines changed

6 files changed

+178
-63
lines changed

‎clang/include/clang/CodeGen/CodeGenAction.h

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -53,6 +53,9 @@ class CodeGenAction : public ASTFrontendAction {
5353

5454
std::unique_ptr<llvm::Module> loadModule(llvm::MemoryBufferRef MBRef);
5555

56+
/// Load bitcode modules to link into our module from the options.
57+
bool loadLinkModules(CompilerInstance &CI);
58+
5659
protected:
5760
/// Create a new code generation action. If the optional \p _VMContext
5861
/// parameter is supplied, the action uses it without taking ownership,

‎clang/lib/CodeGen/CGCall.cpp

Lines changed: 68 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -1845,8 +1845,9 @@ addMergableDefaultFunctionAttributes(const CodeGenOptions &CodeGenOpts,
18451845
FuncAttrs);
18461846
}
18471847

1848-
void CodeGenModule::getTrivialDefaultFunctionAttributes(
1849-
StringRef Name, bool HasOptnone, bool AttrOnCallSite,
1848+
static void getTrivialDefaultFunctionAttributes(
1849+
StringRef Name, bool HasOptnone, const CodeGenOptions &CodeGenOpts,
1850+
const LangOptions &LangOpts, bool AttrOnCallSite,
18501851
llvm::AttrBuilder &FuncAttrs) {
18511852
// OptimizeNoneAttr takes precedence over -Os or -Oz. No warning needed.
18521853
if (!HasOptnone) {
@@ -1967,7 +1968,7 @@ void CodeGenModule::getTrivialDefaultFunctionAttributes(
19671968
}
19681969
}
19691970

1970-
if (getLangOpts().assumeFunctionsAreConvergent()) {
1971+
if (LangOpts.assumeFunctionsAreConvergent()) {
19711972
// Conservatively, mark all functions and calls in CUDA and OpenCL as
19721973
// convergent (meaning, they may call an intrinsically convergent op, such
19731974
// as __syncthreads() / barrier(), and so can't have certain optimizations
@@ -1978,8 +1979,8 @@ void CodeGenModule::getTrivialDefaultFunctionAttributes(
19781979

19791980
// TODO: NoUnwind attribute should be added for other GPU modes HIP,
19801981
// OpenMP offload. AFAIK, neither of them support exceptions in device code.
1981-
if ((getLangOpts().CUDA && getLangOpts().CUDAIsDevice) ||
1982-
getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) {
1982+
if ((LangOpts.CUDA && LangOpts.CUDAIsDevice) || LangOpts.OpenCL ||
1983+
LangOpts.SYCLIsDevice) {
19831984
FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
19841985
}
19851986

@@ -1990,36 +1991,25 @@ void CodeGenModule::getTrivialDefaultFunctionAttributes(
19901991
}
19911992
}
19921993

1993-
void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
1994-
bool HasOptnone,
1995-
bool AttrOnCallSite,
1996-
llvm::AttrBuilder &FuncAttrs) {
1997-
getTrivialDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite,
1998-
FuncAttrs);
1999-
if (!AttrOnCallSite) {
2000-
// If we're just getting the default, get the default values for mergeable
2001-
// attributes.
2002-
addMergableDefaultFunctionAttributes(CodeGenOpts, FuncAttrs);
2003-
}
2004-
}
1994+
/// Adds attributes to \p F according to our \p CodeGenOpts and \p LangOpts, as
1995+
/// though we had emitted it ourselves. We remove any attributes on F that
1996+
/// conflict with the attributes we add here.
1997+
static void mergeDefaultFunctionDefinitionAttributes(
1998+
llvm::Function &F, const CodeGenOptions CodeGenOpts,
1999+
const LangOptions &LangOpts, const TargetOptions &TargetOpts,
2000+
bool WillInternalize) {
20052001

2006-
void CodeGenModule::addDefaultFunctionDefinitionAttributes(llvm::Function &F) {
20072002
llvm::AttrBuilder FuncAttrs(F.getContext());
2008-
getDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
2009-
/* AttrOnCallSite = */ false, FuncAttrs);
2010-
// TODO: call GetCPUAndFeaturesAttributes?
2011-
F.addFnAttrs(FuncAttrs);
2012-
}
2003+
// Here we only extract the options that are relevant compared to the version
2004+
// from GetCPUAndFeaturesAttributes.
2005+
if (!TargetOpts.CPU.empty())
2006+
FuncAttrs.addAttribute("target-cpu", TargetOpts.CPU);
2007+
if (!TargetOpts.TuneCPU.empty())
2008+
FuncAttrs.addAttribute("tune-cpu", TargetOpts.TuneCPU);
20132009

2014-
/// Apply default attributes to \p F, accounting for merge semantics of
2015-
/// attributes that should not overwrite existing attributes.
2016-
void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
2017-
llvm::Function &F, bool WillInternalize) {
2018-
llvm::AttrBuilder FuncAttrs(F.getContext());
2019-
getTrivialDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
2020-
/*AttrOnCallSite=*/false, FuncAttrs);
2021-
GetCPUAndFeaturesAttributes(GlobalDecl(), FuncAttrs,
2022-
/*AddTargetFeatures=*/false);
2010+
::getTrivialDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
2011+
CodeGenOpts, LangOpts,
2012+
/*AttrOnCallSite=*/false, FuncAttrs);
20232013

20242014
if (!WillInternalize && F.isInterposable()) {
20252015
// Do not promote "dynamic" denormal-fp-math to this translation unit's
@@ -2064,6 +2054,52 @@ void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
20642054
F.addFnAttrs(FuncAttrs);
20652055
}
20662056

2057+
void clang::CodeGen::mergeDefaultFunctionDefinitionAttributes(
2058+
llvm::Function &F, const CodeGenOptions CodeGenOpts,
2059+
const LangOptions &LangOpts, const TargetOptions &TargetOpts,
2060+
bool WillInternalize) {
2061+
2062+
::mergeDefaultFunctionDefinitionAttributes(F, CodeGenOpts, LangOpts,
2063+
TargetOpts, WillInternalize);
2064+
}
2065+
2066+
void CodeGenModule::getTrivialDefaultFunctionAttributes(
2067+
StringRef Name, bool HasOptnone, bool AttrOnCallSite,
2068+
llvm::AttrBuilder &FuncAttrs) {
2069+
::getTrivialDefaultFunctionAttributes(Name, HasOptnone, getCodeGenOpts(),
2070+
getLangOpts(), AttrOnCallSite,
2071+
FuncAttrs);
2072+
}
2073+
2074+
void CodeGenModule::getDefaultFunctionAttributes(StringRef Name,
2075+
bool HasOptnone,
2076+
bool AttrOnCallSite,
2077+
llvm::AttrBuilder &FuncAttrs) {
2078+
getTrivialDefaultFunctionAttributes(Name, HasOptnone, AttrOnCallSite,
2079+
FuncAttrs);
2080+
// If we're just getting the default, get the default values for mergeable
2081+
// attributes.
2082+
if (!AttrOnCallSite)
2083+
addMergableDefaultFunctionAttributes(CodeGenOpts, FuncAttrs);
2084+
}
2085+
2086+
void CodeGenModule::addDefaultFunctionDefinitionAttributes(llvm::Function &F) {
2087+
llvm::AttrBuilder FuncAttrs(F.getContext());
2088+
getDefaultFunctionAttributes(F.getName(), F.hasOptNone(),
2089+
/* AttrOnCallSite = */ false, FuncAttrs);
2090+
// TODO: call GetCPUAndFeaturesAttributes?
2091+
F.addFnAttrs(FuncAttrs);
2092+
}
2093+
2094+
/// Apply default attributes to \p F, accounting for merge semantics of
2095+
/// attributes that should not overwrite existing attributes.
2096+
void CodeGenModule::mergeDefaultFunctionDefinitionAttributes(
2097+
llvm::Function &F, bool WillInternalize) {
2098+
::mergeDefaultFunctionDefinitionAttributes(F, getCodeGenOpts(), getLangOpts(),
2099+
getTarget().getTargetOpts(),
2100+
WillInternalize);
2101+
}
2102+
20672103
void CodeGenModule::addDefaultFunctionDefinitionAttributes(
20682104
llvm::AttrBuilder &attrs) {
20692105
getDefaultFunctionAttributes(/*function name*/ "", /*optnone*/ false,

‎clang/lib/CodeGen/CGCall.h

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,7 @@ class Value;
3030
namespace clang {
3131
class Decl;
3232
class FunctionDecl;
33+
class TargetOptions;
3334
class VarDecl;
3435

3536
namespace CodeGen {
@@ -377,6 +378,14 @@ class ReturnValueSlot {
377378
bool isExternallyDestructed() const { return IsExternallyDestructed; }
378379
};
379380

381+
/// Helper to add attributes to \p F according to the CodeGenOptions and
382+
/// LangOptions without requiring a CodeGenModule to be constructed.
383+
void mergeDefaultFunctionDefinitionAttributes(llvm::Function &F,
384+
const CodeGenOptions CodeGenOpts,
385+
const LangOptions &LangOpts,
386+
const TargetOptions &TargetOpts,
387+
bool WillInternalize);
388+
380389
} // end namespace CodeGen
381390
} // end namespace clang
382391

‎clang/lib/CodeGen/CodeGenAction.cpp

Lines changed: 48 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,7 @@
77
//===----------------------------------------------------------------------===//
88

99
#include "clang/CodeGen/CodeGenAction.h"
10+
#include "CGCall.h"
1011
#include "CodeGenModule.h"
1112
#include "CoverageMappingGen.h"
1213
#include "MacroPPCallbacks.h"
@@ -262,7 +263,7 @@ namespace clang {
262263
}
263264

264265
// Links each entry in LinkModules into our module. Returns true on error.
265-
bool LinkInModules() {
266+
bool LinkInModules(llvm::Module *M) {
266267
for (auto &LM : LinkModules) {
267268
assert(LM.Module && "LinkModule does not actually have a module");
268269
if (LM.PropagateAttrs)
@@ -271,24 +272,23 @@ namespace clang {
271272
// in LLVM IR.
272273
if (F.isIntrinsic())
273274
continue;
274-
Gen->CGM().mergeDefaultFunctionDefinitionAttributes(F,
275-
LM.Internalize);
275+
CodeGen::mergeDefaultFunctionDefinitionAttributes(
276+
F, CodeGenOpts, LangOpts, TargetOpts, LM.Internalize);
276277
}
277278

278279
CurLinkModule = LM.Module.get();
279280

280281
bool Err;
281282
if (LM.Internalize) {
282283
Err = Linker::linkModules(
283-
*getModule(), std::move(LM.Module), LM.LinkFlags,
284+
*M, std::move(LM.Module), LM.LinkFlags,
284285
[](llvm::Module &M, const llvm::StringSet<> &GVS) {
285286
internalizeModule(M, [&GVS](const llvm::GlobalValue &GV) {
286287
return !GV.hasName() || (GVS.count(GV.getName()) == 0);
287288
});
288289
});
289290
} else {
290-
Err = Linker::linkModules(*getModule(), std::move(LM.Module),
291-
LM.LinkFlags);
291+
Err = Linker::linkModules(*M, std::move(LM.Module), LM.LinkFlags);
292292
}
293293

294294
if (Err)
@@ -357,7 +357,7 @@ namespace clang {
357357
}
358358

359359
// Link each LinkModule into our module.
360-
if (LinkInModules())
360+
if (LinkInModules(getModule()))
361361
return;
362362

363363
for (auto &F : getModule()->functions()) {
@@ -993,6 +993,36 @@ CodeGenAction::~CodeGenAction() {
993993
delete VMContext;
994994
}
995995

996+
bool CodeGenAction::loadLinkModules(CompilerInstance &CI) {
997+
if (!LinkModules.empty())
998+
return false;
999+
1000+
for (const CodeGenOptions::BitcodeFileToLink &F :
1001+
CI.getCodeGenOpts().LinkBitcodeFiles) {
1002+
auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename);
1003+
if (!BCBuf) {
1004+
CI.getDiagnostics().Report(diag::err_cannot_open_file)
1005+
<< F.Filename << BCBuf.getError().message();
1006+
LinkModules.clear();
1007+
return true;
1008+
}
1009+
1010+
Expected<std::unique_ptr<llvm::Module>> ModuleOrErr =
1011+
getOwningLazyBitcodeModule(std::move(*BCBuf), *VMContext);
1012+
if (!ModuleOrErr) {
1013+
handleAllErrors(ModuleOrErr.takeError(), [&](ErrorInfoBase &EIB) {
1014+
CI.getDiagnostics().Report(diag::err_cannot_open_file)
1015+
<< F.Filename << EIB.message();
1016+
});
1017+
LinkModules.clear();
1018+
return true;
1019+
}
1020+
LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs,
1021+
F.Internalize, F.LinkFlags});
1022+
}
1023+
return false;
1024+
}
1025+
9961026
bool CodeGenAction::hasIRSupport() const { return true; }
9971027

9981028
void CodeGenAction::EndSourceFileAction() {
@@ -1048,30 +1078,8 @@ CodeGenAction::CreateASTConsumer(CompilerInstance &CI, StringRef InFile) {
10481078
return nullptr;
10491079

10501080
// Load bitcode modules to link with, if we need to.
1051-
if (LinkModules.empty())
1052-
for (const CodeGenOptions::BitcodeFileToLink &F :
1053-
CI.getCodeGenOpts().LinkBitcodeFiles) {
1054-
auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename);
1055-
if (!BCBuf) {
1056-
CI.getDiagnostics().Report(diag::err_cannot_open_file)
1057-
<< F.Filename << BCBuf.getError().message();
1058-
LinkModules.clear();
1059-
return nullptr;
1060-
}
1061-
1062-
Expected<std::unique_ptr<llvm::Module>> ModuleOrErr =
1063-
getOwningLazyBitcodeModule(std::move(*BCBuf), *VMContext);
1064-
if (!ModuleOrErr) {
1065-
handleAllErrors(ModuleOrErr.takeError(), [&](ErrorInfoBase &EIB) {
1066-
CI.getDiagnostics().Report(diag::err_cannot_open_file)
1067-
<< F.Filename << EIB.message();
1068-
});
1069-
LinkModules.clear();
1070-
return nullptr;
1071-
}
1072-
LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs,
1073-
F.Internalize, F.LinkFlags});
1074-
}
1081+
if (loadLinkModules(CI))
1082+
return nullptr;
10751083

10761084
CoverageSourceInfo *CoverageInfo = nullptr;
10771085
// Add the preprocessor callback only when the coverage mapping is generated.
@@ -1139,6 +1147,10 @@ CodeGenAction::loadModule(MemoryBufferRef MBRef) {
11391147
return std::move(*MOrErr);
11401148
}
11411149

1150+
// Load bitcode modules to link with, if we need to.
1151+
if (loadLinkModules(CI))
1152+
return nullptr;
1153+
11421154
llvm::SMDiagnostic Err;
11431155
if (std::unique_ptr<llvm::Module> M = parseIR(MBRef, Err, *VMContext))
11441156
return M;
@@ -1218,6 +1230,11 @@ void CodeGenAction::ExecuteAction() {
12181230
CI.getCodeGenOpts(), CI.getTargetOpts(),
12191231
CI.getLangOpts(), TheModule.get(),
12201232
std::move(LinkModules), *VMContext, nullptr);
1233+
1234+
// Link in each pending link module.
1235+
if (Result.LinkInModules(&*TheModule))
1236+
return;
1237+
12211238
// PR44896: Force DiscardValueNames as false. DiscardValueNames cannot be
12221239
// true here because the valued names are needed for reading textual IR.
12231240
Ctx.setDiscardValueNames(false);

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

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,14 @@
1111
// RUN: not %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file no-such-file.bc \
1212
// RUN: -emit-llvm -o - %s 2>&1 | FileCheck -check-prefix=CHECK-NO-FILE %s
1313

14+
// Make sure we can perform the same options if the input is LLVM-IR
15+
// RUN: %clang_cc1 -triple i386-pc-linux-gnu -emit-llvm-bc -o %t-in.bc %s
16+
// RUN: %clang_cc1 -triple i386-pc-linux-gnu -mlink-bitcode-file %t.bc \
17+
// RUN: -O3 -emit-llvm -o - %t-in.bc | FileCheck -check-prefix=CHECK-NO-BC %s
18+
// RUN: %clang_cc1 -triple i386-pc-linux-gnu -O3 -emit-llvm -o - \
19+
// RUN: -mlink-bitcode-file %t.bc -mlink-bitcode-file %t-2.bc %t-in.bc \
20+
// RUN: | FileCheck -check-prefix=CHECK-NO-BC -check-prefix=CHECK-NO-BC2 %s
21+
1422
int f(void);
1523

1624
#ifdef BITCODE

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

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-attributes --check-globals --include-generated-funcs --version 2
2+
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx803 -DBITCODE -emit-llvm-bc -o %t-lib.bc %s
3+
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm-bc -o %t.bc %s
4+
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm \
5+
// RUN: -mlink-builtin-bitcode %t-lib.bc -o - %t.bc | FileCheck %s
6+
7+
#ifdef BITCODE
8+
int foo(void) { return 42; }
9+
int x = 12;
10+
#endif
11+
12+
extern int foo(void);
13+
extern int x;
14+
15+
int bar() { return foo() + x; }
16+
//.
17+
// CHECK: @x = internal addrspace(1) global i32 12, align 4
18+
//.
19+
// CHECK: Function Attrs: noinline nounwind optnone
20+
// CHECK-LABEL: define dso_local i32 @bar
21+
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
22+
// CHECK-NEXT: entry:
23+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
24+
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
25+
// CHECK-NEXT: [[CALL:%.*]] = call i32 @foo()
26+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @x to ptr), align 4
27+
// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[CALL]], [[TMP0]]
28+
// CHECK-NEXT: ret i32 [[ADD]]
29+
//
30+
//
31+
// CHECK: Function Attrs: convergent noinline nounwind optnone
32+
// CHECK-LABEL: define internal i32 @foo
33+
// CHECK-SAME: () #[[ATTR1:[0-9]+]] {
34+
// CHECK-NEXT: entry:
35+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
36+
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
37+
// CHECK-NEXT: ret i32 42
38+
//
39+
//.
40+
// 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" }
41+
// 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" }
42+
//.

0 commit comments

Comments
 (0)