[clang] [AMDGPU] add function attrbute amdgpu-lib-fun (PR #74737)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jan 12 14:21:37 PST 2024
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/74737
>From 4264e7e9c7f655f134623d113ba9dccc5564f4c3 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 | 8 ++
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 | 89 ++++++++++---------
...a-attribute-supported-attributes-list.test | 1 +
6 files changed, 78 insertions(+), 41 deletions(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index a03b0e44e15f7d..b2b7ac88bf5943 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -446,6 +446,7 @@ def TargetX86 : TargetArch<["x86"]>;
def TargetAnyX86 : TargetArch<["x86", "x86_64"]>;
def TargetWebAssembly : TargetArch<["wasm32", "wasm64"]>;
def TargetNVPTX : TargetArch<["nvptx", "nvptx64"]>;
+def TargetAMDGPU : TargetArch<["r600", "amdgcn"]>;
def TargetWindows : TargetSpec {
let OSes = ["Win32"];
}
@@ -2028,6 +2029,13 @@ def AMDGPUNumVGPR : InheritableAttr {
let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
}
+def AMDGPULibFun : InheritableAttr, TargetSpecificAttr<TargetAMDGPU>{
+ 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 2e8d7752c9751e..063a15d578b178 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -2702,6 +2702,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 while targeting
+AMDGPU to indicate that it will be handled specially by the backend.
+A library function will not be optimized out by standard LLVM passes and can be
+used to resolve function calls generated by the backend. These functions will
+not be emitted by the backend if they are not used.
+ }];
+}
+
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 03ac6b78598fc8..08b763c1e7576c 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 89add87919c12d..e319cd4809e0dd 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 b0dfc97b53b2c5..201d867b55047b 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-attrs.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-attrs.cl
@@ -1,105 +1,107 @@
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify -o - %s | FileCheck -check-prefix=X86 %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu tahiti -O0 -emit-llvm -verify -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -O0 -emit-llvm -verify -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -verify=x86 -o - %s | FileCheck -check-prefix=X86 %s
-__attribute__((amdgpu_flat_work_group_size(0, 0))) // expected-no-diagnostics
+// expected-no-diagnostics
+
+__attribute__((amdgpu_flat_work_group_size(0, 0)))
kernel void flat_work_group_size_0_0() {}
-__attribute__((amdgpu_waves_per_eu(0))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(0)))
kernel void waves_per_eu_0() {}
-__attribute__((amdgpu_waves_per_eu(0, 0))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(0, 0)))
kernel void waves_per_eu_0_0() {}
-__attribute__((amdgpu_num_sgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_num_sgpr(0)))
kernel void num_sgpr0() {}
-__attribute__((amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_num_vgpr(0)))
kernel void num_vgpr0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_sgpr(0)))
kernel void flat_work_group_size_0_0_num_sgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_num_vgpr(0)))
kernel void flat_work_group_size_0_0_num_vgpr_0() {}
-__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0)))
kernel void waves_per_eu_0_num_sgpr_0() {}
-__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0)))
kernel void waves_per_eu_0_num_vgpr_0() {}
-__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0)))
kernel void waves_per_eu_0_0_num_sgpr_0() {}
-__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0)))
kernel void waves_per_eu_0_0_num_vgpr_0() {}
-__attribute__((amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_num_sgpr(0), amdgpu_num_vgpr(0)))
kernel void num_sgpr_0_num_vgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0_num_sgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_vgpr(0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0_num_vgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_sgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_vgpr(0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_vgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0_num_sgpr_0_num_vgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(0, 0), amdgpu_waves_per_eu(0, 0), amdgpu_num_sgpr(0), amdgpu_num_vgpr(0)))
kernel void flat_work_group_size_0_0_waves_per_eu_0_0_num_sgpr_0_num_vgpr_0() {}
-__attribute__((amdgpu_flat_work_group_size(32, 64))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(32, 64)))
kernel void flat_work_group_size_32_64() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64() [[FLAT_WORK_GROUP_SIZE_32_64:#[0-9]+]]
}
-__attribute__((amdgpu_waves_per_eu(2))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(2)))
kernel void waves_per_eu_2() {
// CHECK: define{{.*}} amdgpu_kernel void @waves_per_eu_2() [[WAVES_PER_EU_2:#[0-9]+]]
}
-__attribute__((amdgpu_waves_per_eu(2, 4))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(2, 4)))
kernel void waves_per_eu_2_4() {
// CHECK: define{{.*}} amdgpu_kernel void @waves_per_eu_2_4() [[WAVES_PER_EU_2_4:#[0-9]+]]
}
-__attribute__((amdgpu_num_sgpr(32))) // expected-no-diagnostics
+__attribute__((amdgpu_num_sgpr(32)))
kernel void num_sgpr_32() {
// CHECK: define{{.*}} amdgpu_kernel void @num_sgpr_32() [[NUM_SGPR_32:#[0-9]+]]
}
-__attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__attribute__((amdgpu_num_vgpr(64)))
kernel void num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @num_vgpr_64() [[NUM_VGPR_64:#[0-9]+]]
}
-__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2)))
kernel void flat_work_group_size_32_64_waves_per_eu_2() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2:#[0-9]+]]
}
-__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4)))
kernel void flat_work_group_size_32_64_waves_per_eu_2_4() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4:#[0-9]+]]
}
-__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_sgpr(32))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_sgpr(32)))
kernel void flat_work_group_size_32_64_num_sgpr_32() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64_num_sgpr_32() [[FLAT_WORK_GROUP_SIZE_32_64_NUM_SGPR_32:#[0-9]+]]
}
-__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_vgpr(64)))
kernel void flat_work_group_size_32_64_num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_NUM_VGPR_64:#[0-9]+]]
}
-__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32)))
kernel void waves_per_eu_2_num_sgpr_32() {
// CHECK: define{{.*}} amdgpu_kernel void @waves_per_eu_2_num_sgpr_32() [[WAVES_PER_EU_2_NUM_SGPR_32:#[0-9]+]]
}
-__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64)))
kernel void waves_per_eu_2_num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @waves_per_eu_2_num_vgpr_64() [[WAVES_PER_EU_2_NUM_VGPR_64:#[0-9]+]]
}
-__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32)))
kernel void waves_per_eu_2_4_num_sgpr_32() {
// CHECK: define{{.*}} amdgpu_kernel void @waves_per_eu_2_4_num_sgpr_32() [[WAVES_PER_EU_2_4_NUM_SGPR_32:#[0-9]+]]
}
-__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64)))
kernel void waves_per_eu_2_4_num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @waves_per_eu_2_4_num_vgpr_64() [[WAVES_PER_EU_2_4_NUM_VGPR_64:#[0-9]+]]
}
-__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
kernel void num_sgpr_32_num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @num_sgpr_32_num_vgpr_64() [[NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]]
}
@@ -121,20 +123,20 @@ kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_VGPR_64:#[0-9]+]]
}
-__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
kernel void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]]
}
-__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64))) // expected-no-diagnostics
+__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))
kernel void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() {
// CHECK: define{{.*}} amdgpu_kernel void @flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() [[FLAT_WORK_GROUP_SIZE_32_64_WAVES_PER_EU_2_4_NUM_SGPR_32_NUM_VGPR_64:#[0-9]+]]
}
-__attribute__((reqd_work_group_size(32, 2, 1))) // expected-no-diagnostics
+__attribute__((reqd_work_group_size(32, 2, 1)))
kernel void reqd_work_group_size_32_2_1() {
// CHECK: define{{.*}} amdgpu_kernel void @reqd_work_group_size_32_2_1() [[FLAT_WORK_GROUP_SIZE_64_64:#[0-9]+]]
}
-__attribute__((reqd_work_group_size(32, 2, 1), amdgpu_flat_work_group_size(16, 128))) // expected-no-diagnostics
+__attribute__((reqd_work_group_size(32, 2, 1), amdgpu_flat_work_group_size(16, 128)))
kernel void reqd_work_group_size_32_2_1_flat_work_group_size_16_128() {
// CHECK: define{{.*}} amdgpu_kernel void @reqd_work_group_size_32_2_1_flat_work_group_size_16_128() [[FLAT_WORK_GROUP_SIZE_16_128:#[0-9]+]]
}
@@ -147,12 +149,16 @@ kernel void default_kernel() {
// CHECK: define{{.*}} amdgpu_kernel void @default_kernel() [[DEFAULT_KERNEL_ATTRS:#[0-9]+]]
}
+__attribute__((amdgpu_lib_fun)) void lib_fun() { // x86-warning {{unknown attribute 'amdgpu_lib_fun' ignored}}
+// 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"
// X86-NOT: "amdgpu-waves-per-eu"
// X86-NOT: "amdgpu-num-vgpr"
// X86-NOT: "amdgpu-num-sgpr"
+// X86-NOT: "amdgpu-lib-fun"
// CHECK-NOT: "amdgpu-implicitarg-num-bytes"
// CHECK-NOT: "amdgpu-flat-work-group-size"="0,0"
@@ -191,3 +197,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 e476c15b35ded9..b5bf183c684175 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)
More information about the cfe-commits
mailing list