[clang][CodeGen] Fix GPU-specific attributes being dropped by bitcode linking
authorpvanhout <pierre.vanhoutryve@amd.com>
Tue, 6 Jun 2023 11:32:00 +0000 (13:32 +0200)
committerpvanhout <pierre.vanhoutryve@amd.com>
Wed, 7 Jun 2023 13:51:52 +0000 (15:51 +0200)
commit23431b52460328c554ad244fd7b50ecb751cec31
tree1ba3db4336580769601403717bace5c9f08fbc7a
parentdcc8f9490f790e7ecbe08a655abec272422df163
[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
clang/lib/CodeGen/CGCall.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/CodeGen/CodeGenModule.h
clang/test/CodeGenCUDA/Inputs/ocml-sample-target-attrs.cl [new file with mode: 0644]
clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu [new file with mode: 0644]