[Clang] Propagate target-features if compatible when using mlink-builtin-bitcode

[jmmartinez]

Summary

Buitlins from AMD's device-libs are compiled without specifying a
target-cpu, which results in builtins without the target-features
attribute set.

Before this patch, when linking this builtins with -mlink-builtin-bitcode
the target-features were not propagated in the incoming builtins.

With this patch, the default target features are propagated
if they are compatible with the target-features in the incoming builtin.

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2034

Really it would be less bad if the incompatible functions were not imported rather than the backend pass

2035

do you need to guard against adding the empty attribute? I don't want to see "target-features"=""

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

17

This isn't a real target feature (do we not have a warning on this?)

s/uncompatible/incompatible

[jmmartinez]

clang/lib/CodeGen/CGCall.cpp

2034

I thought it was possible to have functions with incompatible features in the same module.
e.g. one function compiled with some instruction set support, one without, and an ifunc that resolves to one or the other.

Maybe it's not the case in the context of -mlink-builtin-bitcode?

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

17

Only when the feature doesn't map to any "+" or "-" feature.

If I use "foo-insts" I get a warning because there are no +foo-insts / -foo-insts (and it still ends up in the generated llvm-ir)

[jmmartinez]

• Drop incoming builtins with incompatible target-features

[jmmartinez]

clang/lib/CodeGen/CodeGenAction.cpp

282 ↗

(On Diff #554753)

Remove { }

[arsenm]

clang/lib/CodeGen/CGCall.h

398–401 ↗

(On Diff #554753)

i think this should be done in a separate patch, just propagate + append for step 1. There are other edge cases I'm worried about handling with this

[jmmartinez]

• Rollback, drop incompatible functions in a separate commit

[jmmartinez]

• Rebased

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2017

Do you need to guard against empty string?

2018

consume_front

2021

Capitalize

2027

Capitalize

2030–2031

Early return breaks the other features

[jmmartinez]

• Review

[jmmartinez]

clang/lib/CodeGen/CGCall.cpp

2017

I checked and I saw some tests without it. I'm adding the condition for it.

2018

I used drop_front(1) instead, consume_front would only work with "+" and not drop "-".

2030–2031

I did not understand this remark.

If the features are not compatible, we do not add a "target-features" entry in the new "FuncAttrs". Then, the old "target-features" entry is kept in the Function coming from the builtin.

If you think it would be better to set the target-features in FuncAttrs to the old value in any case. If that's the case I've added the following code:

if (EnabledForTarget != EnabledForFunc) {
    FuncAttr.addAttribute(FFeatures);
    return;
}

[arsenm]

clang/lib/CodeGen/CGCall.cpp

2030–2031

You find an incompatible feature and then discontinue processing any further features by early exiting. I expect this to act like an append to any features already present. The incompatibility is at an individual feature level, not the group

2034

The truth is this system isn't really well considered. We don't have real ifunc support and we probably shouldn't be using subtargets for cases with incompatible encodings

[jmmartinez]

• Review

[jmmartinez]

• Capitalize comment

clang/lib/CodeGen/CGCall.cpp

2030–2031

I see. I changed it to match that behaviour.

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2006

can you add a comment about how the original target feature attributes of the function that is incompatible or missing in the target feature are handled.

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

46

The irrelevant function attributes should not be checked

[jmmartinez]

@yaxunl I've addressed your remarks on a GitHub PR: https://github.com/llvm/llvm-project/pull/65938

Thanks for the remarks! It made me realize that I could simplify the code a lot.

clang/lib/CodeGen/CGCall.cpp

static void getTrivialDefaultFunctionAttributes(
 
   for (StringRef Attr : CodeGenOpts.DefaultFunctionAttrs) {
     StringRef Var, Value;
     std::tie(Var, Value) = Attr.split('=');
     FuncAttrs.addAttribute(Var, Value);
   }
 }
 
+static void
+overrideFunctionFeaturesWithTargetFeatures(llvm::AttrBuilder &FuncAttr,
+                                           const llvm::Function &F,
+                                           const TargetOptions &TargetOpts) {
+  auto FFeatures = F.getFnAttribute("target-features");
+
+  SmallVector<StringRef> MergedFeatures(TargetOpts.Features.begin(),
+                                        TargetOpts.Features.end());
+
+  if (FFeatures.isValid()) {
+    const auto &TFeatures = TargetOpts.FeatureMap;
+    for (StringRef Feature : llvm::split(FFeatures.getValueAsString(), ',')) {
+      bool EnabledForFunc = Feature[0] == '+';
+      StringRef Name = Feature.substr(1);
+      auto TEntry = TFeatures.find(Name);
+
+      // if the feature is not set for the target-opts, it must be preserved
+      if (TEntry == TFeatures.end()) {
+        MergedFeatures.push_back(Feature);
+        continue;
+      }
+
+      // if the feature is enabled for one and disabled for the other, they are
+      // not compatible
+      bool EnabledForTarget = TEntry->second;
+      if (EnabledForTarget != EnabledForFunc)
+        return;
+    }
+  }
+
+  FuncAttr.addAttribute("target-features", llvm::join(MergedFeatures, ","));
+}
+
 /// Adds attributes to \p F according to our \p CodeGenOpts and \p LangOpts, as
 /// though we had emitted it ourselves. We remove any attributes on F that
 /// conflict with the attributes we add here.
 static void mergeDefaultFunctionDefinitionAttributes(
     llvm::Function &F, const CodeGenOptions &CodeGenOpts,
     const LangOptions &LangOpts, const TargetOptions &TargetOpts,
     bool WillInternalize) {
 
static void mergeDefaultFunctionDefinitionAttributes(
   } else if (MergedF32 != DenormModeToMergeF32) {
     // Overwrite existing attribute
     FuncAttrs.addAttribute("denormal-fp-math-f32",
                            CodeGenOpts.FP32DenormalMode.str());
   }
 
   F.removeFnAttrs(AttrsToRemove);
   addDenormalModeAttrs(Merged, MergedF32, FuncAttrs);
+
+  overrideFunctionFeaturesWithTargetFeatures(FuncAttrs, F, TargetOpts);
+
   F.addFnAttrs(FuncAttrs);
 }
 
 void clang::CodeGen::mergeDefaultFunctionDefinitionAttributes(
     llvm::Function &F, const CodeGenOptions CodeGenOpts,
     const LangOptions &LangOpts, const TargetOptions &TargetOpts,
     bool WillInternalize) {
 

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

-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-attributes --check-globals --include-generated-funcs --version 2
+// Build two version of the bitcode library, one with a target-cpu set and one without
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx803 -DBITCODE -emit-llvm-bc -o %t-lib.bc %s
+// RUN: %clang_cc1 -triple amdgcn-- -DBITCODE -emit-llvm-bc -o %t-lib.no-cpu.bc %s
+
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm-bc -o %t.bc %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm \
+// RUN:   -mlink-builtin-bitcode %t-lib.bc -o - %t.bc | FileCheck %s --check-prefixes=COMMON,CPU
+
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm-bc -o %t.bc %s
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -emit-llvm \
-// RUN:   -mlink-builtin-bitcode %t-lib.bc -o - %t.bc | FileCheck %s
+// RUN:   -mlink-builtin-bitcode %t-lib.no-cpu.bc -o - %t.bc | FileCheck %s --check-prefixes=COMMON,NOCPU
 
 #ifdef BITCODE
-int foo(void) { return 42; }
+int no_attr(void) { return 42; }
+int __attribute__((target("gfx8-insts"))) attr_in_target(void) { return 42; }
+int __attribute__((target("extended-image-insts"))) attr_not_in_target(void) { return 42; }
+int __attribute__((target("no-gfx9-insts"))) attr_uncompatible(void) { return 42; }
 int x = 12;
 #endif
 
-extern int foo(void);
+extern int no_attr(void);
+extern int attr_in_target(void);
+extern int attr_not_in_target(void);
+extern int attr_uncompatible(void);
 extern int x;
 
-int bar() { return foo() + x; }
-//.
-// CHECK: @x = internal addrspace(1) global i32 12, align 4
-//.
-// CHECK: Function Attrs: noinline nounwind optnone
-// CHECK-LABEL: define dso_local i32 @bar
-// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
-// CHECK-NEXT:  entry:
-// CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
-// CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
-// CHECK-NEXT:    [[CALL:%.*]] = call i32 @foo()
-// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @x to ptr), align 4
-// CHECK-NEXT:    [[ADD:%.*]] = add nsw i32 [[CALL]], [[TMP0]]
-// CHECK-NEXT:    ret i32 [[ADD]]
-//
-//
-// CHECK: Function Attrs: convergent noinline nounwind optnone
-// CHECK-LABEL: define internal i32 @foo
-// CHECK-SAME: () #[[ATTR1:[0-9]+]] {
-// CHECK-NEXT:  entry:
-// CHECK-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
-// CHECK-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
-// CHECK-NEXT:    ret i32 42
+int bar() { return no_attr() + attr_in_target() + attr_not_in_target() + attr_uncompatible() + x; }
+
+// COMMON: @x = internal addrspace(1) global i32 12, align 4
+
+// COMMON-LABEL: define dso_local i32 @bar
+// COMMON-SAME: () #[[ATTR_BAR:[0-9]+]] {
 //
-//.
-// 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" }
-// 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" }
-//.
+// COMMON-LABEL: define internal i32 @no_attr
+// COMMON-SAME: () #[[ATTR_COMPATIBLE:[0-9]+]] {
+
+// COMMON-LABEL: define internal i32 @attr_in_target
+// COMMON-SAME: () #[[ATTR_COMPATIBLE:[0-9]+]] {
+
+// COMMON-LABEL: define internal i32 @attr_not_in_target
+// COMMON-SAME: () #[[ATTR_EXTEND:[0-9]+]] {
+
+// COMMON-LABEL: define internal i32 @attr_uncompatible
+// COMMON-SAME: () #[[ATTR_UNCOMPATIBLE:[0-9]+]] {
+
+// COMMON: attributes #[[ATTR_BAR]] = { 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" }
+// COMMON: attributes #[[ATTR_COMPATIBLE]] = { convergent 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,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
+// COMMON: attributes #[[ATTR_EXTEND]] = { convergent 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,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+extended-image-insts" }
+// CPU: attributes #[[ATTR_UNCOMPATIBLE]] = { 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,-gfx9-insts" }
+// NOCPU: attributes #[[ATTR_UNCOMPATIBLE]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx90a" "target-features"="-gfx9-insts" }

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

 // 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"
+// 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();
   }
 }