[all-commits] [llvm/llvm-project] 23431b: [clang][CodeGen] Fix GPU-specific attributes being...

[Pierre van Houtryve via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 23431b52460328c554ad244fd7b50ecb751cec31
      https://github.com/llvm/llvm-project/commit/23431b52460328c554ad244fd7b50ecb751cec31
  Author: pvanhout <pierre.vanhoutryve at amd.com>
  Date:   2023-06-07 (Wed, 07 Jun 2023)

  Changed paths:
    M clang/lib/CodeGen/CGCall.cpp
    M clang/lib/CodeGen/CodeGenModule.cpp
    M clang/lib/CodeGen/CodeGenModule.h
    A clang/test/CodeGenCUDA/Inputs/ocml-sample-target-attrs.cl
    M clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu
    A clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu

  Log Message:
  -----------
  [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