[clang][CodeGen] Fix GPU-specific attributes being dropped by bitcode linking

[Pierre-vh]

Summary

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

[arsenm]

clang/test/CodeGenCUDA/Inputs/ocml-sample-target-attrs.cl

3

Sound really be ulong

clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu

35

Also should make sure target-cpu was set

43

ulong

[Pierre-vh]

target-cpu wasn't set so I tried something a bit different, but now I still need to remove the target-cpu check in the old test.
For some reason filecheck doesn't match it but it's there?

/home/pierre/work/trunk/llvm-project/clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu:150:28: error: IEEEF32-PSZF64-DYNFULL: expected string not found in input
// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}} }
                           ^
<stdin>:260:427: note: scanning from here
attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx803" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" "uniform-work-group-size"="true" }
                                                                                                                                                                                                                                                                                                                                                                                                                                          ^
<stdin>:260:427: note: with "$FUNCATTR" equal to "1"
attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx803" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" "uniform-work-group-size"="true" }
                                                                                                                                                                                                                                                                                                                                                                                                                                          ^
<stdin>:261:85: note: possible intended match here
attributes #1 = { convergent mustprogress nofree norecurse nosync nounwind willreturn memory(none) "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx803" }

[Pierre-vh]

Fix check lines, I think it's just FileCheck weirdness

[arsenm]

clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu

35

Did this previously receive the target-features spam implied by the target?

[yaxunl]

LGTM. Thanks

clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu

36

Do we know why internalize keeps the target-cpu attribute but non-internalize does not?

[Pierre-vh]

clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu

35

Did this previously receive the target-features spam implied by the target?

I think it did, the attributes were filled with things like "+gfx9-insts", etc.

Do we know why internalize keeps the target-cpu attribute but non-internalize does not?

PropagateAttrs is only set for -mlink-builtin-bitcode, see CompilerInvocation.cpp@1888 (where OPT_mlink_builtin_bitcode is processed)

clang/lib/CodeGen/CGCall.cpp

 
 /// 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);
+  GetCPUAndFeaturesAttributes(GlobalDecl(), FuncAttrs,
+                              /*AddTargetFeatures=*/false);
 
   if (!WillInternalize && F.isInterposable()) {
     // Do not promote "dynamic" denormal-fp-math to this translation unit's
     // setting for weak functions that won't be internalized. The user has no
     // real control for how builtin bitcode is linked, so we shouldn't assume
     // later copies will use a consistent mode.
     F.addFnAttrs(FuncAttrs);
     return;

clang/lib/CodeGen/CodeGenModule.h

void UpdateMultiVersionNames(GlobalDecl GD, const FunctionDecl *FD,
                                StringRef &CurName);
 
   llvm::Constant *
   GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty, LangAS AddrSpace,
                         const VarDecl *D,
                         ForDefinition_t IsForDefinition = NotForDefinition);
 
   bool GetCPUAndFeaturesAttributes(GlobalDecl GD,
-                                   llvm::AttrBuilder &AttrBuilder);
+                                   llvm::AttrBuilder &AttrBuilder,
+                                   bool SetTargetFeatures = true);
   void setNonAliasAttributes(GlobalDecl GD, llvm::GlobalObject *GO);
 
   /// Set function attributes for a function declaration.
   void SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
                              bool IsIncompleteFunction, bool IsThunk);
 
   void EmitGlobalDefinition(GlobalDecl D, llvm::GlobalValue *GV = nullptr);
 

clang/lib/CodeGen/CodeGenModule.cpp

if (CodeGenOpts.KeepStaticConsts && D && isa<VarDecl>(D)) {
     const auto *VD = cast<VarDecl>(D);
     if (VD->getType().isConstQualified() &&
         VD->getStorageDuration() == SD_Static)
       addUsedOrCompilerUsedGlobal(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.
   StringRef TargetCPU = getTarget().getTargetOpts().CPU;
   StringRef TuneCPU = getTarget().getTargetOpts().TuneCPU;
   std::vector<std::string> Features;
   const auto *FD = dyn_cast_or_null<FunctionDecl>(GD.getDecl());
   FD = FD ? FD->getMostRecentDecl() : FD;
bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
   if (!TargetCPU.empty()) {
     Attrs.addAttribute("target-cpu", TargetCPU);
     AddedAttr = true;
   }
   if (!TuneCPU.empty()) {
     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;
   }
 
   return AddedAttr;
 }
 

clang/test/CodeGenCUDA/Inputs/ocml-sample-target-attrs.cl

+typedef unsigned long ulong;
+
+__attribute__((target("gfx11-insts")))
+ulong do_intrin_stuff(void)
+{
+  return __builtin_amdgcn_s_sendmsg_rtnl(0x0);
+}

clang/test/CodeGenCUDA/link-builtin-bitcode-denormal-fp-mode.cu

 // We should not be littering call sites with the attribute
 // Everything should use the default ieee with no explicit attribute
 
 // FIXME: Should check-not "denormal-fp-math" within the denormal-fp-math-f32
 // lines.
 
 // Default mode relies on the implicit check-not for the denormal-fp-math.
 
-// PSZ: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
-// PSZ: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
-// PSZ: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// PSZ: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign"
+// PSZ-SAME: "target-cpu"="gfx803"
+// PSZ: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// PSZ-SAME: "target-cpu"="gfx803"
+// PSZ: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// PSZ-SAME: "target-cpu"="gfx803"
 
 // FIXME: Should check-not "denormal-fp-math" within the line
-// IEEEF64-PSZF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
-// IEEEF64-PSZF32: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
-// IEEEF64-PSZF32: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign" {{.*}} "target-cpu"="gfx803" {{.*}} }
+// IEEEF64-PSZF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
+// IEEEF64-PSZF32: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
+// IEEEF64-PSZF32: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math-f32"="preserve-sign,preserve-sign"
+// IEEEF64-PSZF32-SAME: "target-cpu"="gfx803"
 
 // IEEEF32-PSZF64-DYNF32: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}}  }
 // implicit check-not
 // implicit check-not
 
 
-// IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}} }
-// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}}  }
-// IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"  {{.*}} "target-cpu"="gfx803" {{.*}} }
-
+// IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"
+// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
+// IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"
+// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
+// IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee"
+// IEEEF32-PSZF64-DYNFULL-SAME: "target-cpu"="gfx803"
 
 // -mlink-bitcode-file doesn't internalize or propagate attributes.
 // NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$KERNELATTR]] = { {{.*}} "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" {{.*}} "target-cpu"="gfx803" {{.*}} }
 // NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$FUNCATTR]] = { {{.*}} "denormal-fp-math"="dynamic,dynamic" {{.*}} }
 // NOINTERNALIZE-IEEEF32-PSZF64-DYNFULL: #[[$WEAK_FUNCATTR]] = { {{.*}} "denormal-fp-math"="dynamic,dynamic" {{.*}} }

clang/test/CodeGenCUDA/link-builtin-bitcode-gpu-attrs-preserved.cu

+// Verify the behavior of the +gfxN-insts in the way that
+// rocm-device-libs should be built with. e.g. If the device libraries has a function
+// with "+gfx11-insts", that attribute should still be present after linking and not
+// overwritten with the current target's settings.
+
+// This is important because at this time, many device-libs functions that are only
+// available on some GPUs put an attribute such as "+gfx11-insts" so that
+// AMDGPURemoveIncompatibleFunctions can detect & remove them if needed.
+
+// Build the fake device library in the way rocm-device-libs should be built.
+//
+// RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa\
+// RUN:   -mcode-object-version=none -emit-llvm-bc \
+// RUN:   %S/Inputs/ocml-sample-target-attrs.cl -o %t.bc
+
+// Check the default behavior
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
+// RUN:   -mlink-builtin-bitcode %t.bc \
+// RUN:   -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,INTERNALIZE
+
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx1101 -fcuda-is-device \
+// RUN:   -mlink-builtin-bitcode %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,INTERNALIZE
+
+// Check the case where no internalization is performed
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
+// RUN:   -fcuda-is-device -mlink-bitcode-file %t.bc -emit-llvm %s -o -  | FileCheck %s --check-prefixes=CHECK,NOINTERNALIZE
+
+// Check the case where no internalization is performed
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx1101 \
+// RUN:   -fcuda-is-device -mlink-bitcode-file %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,NOINTERNALIZE
+
+
+// CHECK: define {{.*}} i64 @do_intrin_stuff() #[[ATTR:[0-9]+]]
+// INTERNALIZE: attributes #[[ATTR]] = {{.*}} "target-cpu"="gfx{{.*}}" "target-features"="+gfx11-insts"
+// NOINTERNALIZE: attributes #[[ATTR]] = {{.*}} "target-features"="+gfx11-insts"
+
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+
+typedef unsigned long ulong;
+
+extern "C" {
+__device__ ulong do_intrin_stuff(void);
+
+__global__ void kernel_f16(ulong* out) {
+    *out = do_intrin_stuff();
+  }
+}