[clang] [AMDGPU] add function attrbute amdgpu-lib-fun (PR #74737)

[via cfe-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang-codegen

Author: Yaxun (Sam) Liu (yxsamliu)

<details>
<summary>Changes</summary>

Add a function attribute "amdgpu-lib-fun" to indicate that the function needs special handling in backend. Basically it will not be internalized so that it will not be removed by DCE after internalization. This is to keep the library functions that are not called by users' code but will be called by instructions generated by LLVM passes or instruction selection, e.g. sanitizers or lowering of 128 bit integer divisioin.

---
Full diff: https://github.com/llvm/llvm-project/pull/74737.diff


5 Files Affected:

- (modified) clang/include/clang/Basic/Attr.td (+7) 
- (modified) clang/include/clang/Basic/AttrDocs.td (+11) 
- (modified) clang/lib/CodeGen/Targets/AMDGPU.cpp (+2) 
- (modified) clang/test/CodeGenCUDA/amdgpu-func-attrs.cu (+8) 
- (modified) clang/test/CodeGenOpenCL/amdgpu-attrs.cl (+4) 


``````````diff
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 121ed203829ce..676faddd2d1ac 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -2011,6 +2011,13 @@ def AMDGPUNumVGPR : InheritableAttr {
   let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
 }
 
+def AMDGPULibFun : InheritableAttr {
+  let Spellings = [Clang<"amdgpu_lib_fun">];
+  let Documentation = [AMDGPULibFunDocs];
+  let Subjects = SubjectList<[Function]>;
+  let SimpleHandler = 1;
+}
+
 def AMDGPUKernelCall : DeclOrTypeAttr {
   let Spellings = [Clang<"amdgpu_kernel">];
   let Documentation = [Undocumented];
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 88f7c65e6e847..c2c77d4d1d817 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -2693,6 +2693,17 @@ An error will be given if:
   }];
 }
 
+def AMDGPULibFunDocs : Documentation {
+  let Category = DocCatAMDGPUAttributes;
+  let Content = [{
+The ``amdgpu_lib_fun`` attribute can be applied to a function for AMDGPU target
+to indicate it is a library function which are handled specially in backend.
+An AMDGPU library function is not internalized and can be used to fullfill
+calls generated by LLVM passes or instruction selection. Unused AMDGPU library
+functions will be eliminated by the backend.
+  }];
+}
+
 def DocCatCallingConvs : DocumentationCategory<"Calling Conventions"> {
   let Content = [{
 Clang supports several different calling conventions, depending on the target
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index b654e3f12af8d..1a56ee3692d07 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -356,6 +356,8 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
     if (NumVGPR != 0)
       F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR));
   }
+  if (FD->getAttr<AMDGPULibFunAttr>())
+    F->addFnAttr("amdgpu-lib-fun");
 }
 
 /// Emits control constants used to change per-architecture behaviour in the
diff --git a/clang/test/CodeGenCUDA/amdgpu-func-attrs.cu b/clang/test/CodeGenCUDA/amdgpu-func-attrs.cu
index 89add87919c12..e319cd4809e0d 100644
--- a/clang/test/CodeGenCUDA/amdgpu-func-attrs.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-func-attrs.cu
@@ -8,6 +8,9 @@
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
 // RUN:     -o - -x hip %s -munsafe-fp-atomics \
 // RUN:     | FileCheck -check-prefix=NO-UNSAFE-FP-ATOMICS %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
+// RUN:     | FileCheck %s
 
 #include "Inputs/cuda.h"
 
@@ -15,8 +18,13 @@ __device__ void test() {
 // UNSAFE-FP-ATOMICS: define{{.*}} void @_Z4testv() [[ATTR:#[0-9]+]]
 }
 
+__attribute__((amdgpu_lib_fun)) __device__ void lib_fun() {
+// CHECK: define{{.*}} void @_Z7lib_funv() [[LIB_FUN:#[0-9]+]]
+}
+
 
 // Make sure this is silently accepted on other targets.
 // NO-UNSAFE-FP-ATOMICS-NOT: "amdgpu-unsafe-fp-atomics"
 
 // UNSAFE-FP-ATOMICS-DAG: attributes [[ATTR]] = {{.*}}"amdgpu-unsafe-fp-atomics"="true"
+// CHECK-DAG: attributes [[LIB_FUN]] = {{.*}}"amdgpu-lib-fun"
diff --git a/clang/test/CodeGenOpenCL/amdgpu-attrs.cl b/clang/test/CodeGenOpenCL/amdgpu-attrs.cl
index b0dfc97b53b2c..bce7739c7a429 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-attrs.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-attrs.cl
@@ -147,6 +147,9 @@ kernel void default_kernel() {
 // CHECK: define{{.*}} amdgpu_kernel void @default_kernel() [[DEFAULT_KERNEL_ATTRS:#[0-9]+]]
 }
 
+__attribute__((amdgpu_lib_fun)) void lib_fun() {
+// CHECK: define{{.*}} void @lib_fun() [[LIB_FUN:#[0-9]+]]
+}
 
 // Make sure this is silently accepted on other targets.
 // X86-NOT: "amdgpu-flat-work-group-size"
@@ -191,3 +194,4 @@ kernel void default_kernel() {
 
 // CHECK-DAG: attributes [[A_FUNCTION]] = {{.*}}
 // CHECK-DAG: attributes [[DEFAULT_KERNEL_ATTRS]] = {{.*}} "amdgpu-flat-work-group-size"="1,256"
+// CHECK-DAG: attributes [[LIB_FUN]] = {{.*}} "amdgpu-lib-fun"

``````````

</details>


https://github.com/llvm/llvm-project/pull/74737