diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index a03b0e44e15f7..b2b7ac88bf594 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{ + 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 2e8d7752c9751..063a15d578b17 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 03ac6b78598fc..08b763c1e7576 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()) + 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..201d867b55047 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 e476c15b35ded..b5bf183c68417 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)