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

[Yaxun Liu via cfe-commits]

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

>From b4a30a6c89e59f11368fb71040539d24cacccbfc Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Thu, 7 Dec 2023 11:45:14 -0500
Subject: [PATCH] [AMDGPU] add function attrbute amdgpu-lib-fun

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.
---
 clang/include/clang/Basic/Attr.td                     |  7 +++++++
 clang/include/clang/Basic/AttrDocs.td                 | 11 +++++++++++
 clang/lib/CodeGen/Targets/AMDGPU.cpp                  |  2 ++
 clang/test/CodeGenCUDA/amdgpu-func-attrs.cu           |  8 ++++++++
 clang/test/CodeGenOpenCL/amdgpu-attrs.cl              |  4 ++++
 .../pragma-attribute-supported-attributes-list.test   |  1 +
 6 files changed, 33 insertions(+)

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"
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 707fc8875089f..f9fc7ab9c0f6f 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -4,6 +4,7 @@
 
 // CHECK: #pragma clang attribute supports the following attributes:
 // CHECK-NEXT: AMDGPUFlatWorkGroupSize (SubjectMatchRule_function)
+// CHECK-NEXT: AMDGPULibFun (SubjectMatchRule_function)
 // CHECK-NEXT: AMDGPUNumSGPR (SubjectMatchRule_function)
 // CHECK-NEXT: AMDGPUNumVGPR (SubjectMatchRule_function)
 // CHECK-NEXT: AMDGPUWavesPerEU (SubjectMatchRule_function)