diff options
| author | pvanhout <pierre.vanhoutryve@amd.com> | 2023-06-06 13:32:00 +0200 |
|---|---|---|
| committer | pvanhout <pierre.vanhoutryve@amd.com> | 2023-06-07 15:51:52 +0200 |
| commit | 23431b52460328c554ad244fd7b50ecb751cec31 (patch) | |
| tree | 1ba3db4336580769601403717bace5c9f08fbc7a /clang/lib/CodeGen/CodeGenModule.cpp | |
| parent | dcc8f9490f790e7ecbe08a655abec272422df163 (diff) | |
[clang][CodeGen] Fix GPU-specific attributes being dropped by bitcode linking
Device libs make use of patterns like this:
```
__attribute__((target("gfx11-insts")))
static unsigned do_intrin_stuff(void)
{
return __builtin_amdgcn_s_sendmsg_rtnl(0x0);
}
```
For functions that are assumed to be eliminated if the currennt GPU target doesn't support them.
At O0 such functions aren't eliminated by common optimizations but often by AMDGPURemoveIncompatibleFunctions instead, which sees the "+gfx11-insts" attribute on, say, GFX9 and knows it's not valid, so it removes the function.
D142907 accidentally made it so such attributes were dropped during bitcode linking, making it impossible for RemoveIncompatibleFunctions to catch the functions and causing ISel to catch fire eventually.
This fixes the issue and adds a new test to ensure we don't accidentally fall into this trap again.
Fixes SWDEV-403642
Reviewed By: arsenm, yaxunl
Differential Revision: https://reviews.llvm.org/D152251
Diffstat (limited to 'clang/lib/CodeGen/CodeGenModule.cpp')
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 5 |
1 files changed, 3 insertions, 2 deletions
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 49b670487749..41ceebfba404 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2226,7 +2226,8 @@ void CodeGenModule::SetCommonAttributes(GlobalDecl GD, llvm::GlobalValue *GV) { } bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD, - llvm::AttrBuilder &Attrs) { + llvm::AttrBuilder &Attrs, + bool SetTargetFeatures) { // Add target-cpu and target-features attributes to functions. If // we have a decl for the function and it has a target attribute then // parse that and add it to the feature set. @@ -2286,7 +2287,7 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD, Attrs.addAttribute("tune-cpu", TuneCPU); AddedAttr = true; } - if (!Features.empty()) { + if (!Features.empty() && SetTargetFeatures) { llvm::sort(Features); Attrs.addAttribute("target-features", llvm::join(Features, ",")); AddedAttr = true; |
