https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/211059
>From c6251f9eabeae39f21a18851c43db627d704269a Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Wed, 29 Jul 2026 09:34:18 -0400 Subject: [PATCH] [AMDGPU] Add a per-kernel kernarg preload SGPR limit The backend option sets one kernel argument count for all kernels. Kernels can need different preload behavior, and an argument count does not describe the SGPR cost when arguments have different sizes or alignment. Add `amdgpu_kernarg_preload_sgpr_count(N)` for AMDGPU kernels. It limits the SGPRs used by the contiguous preload range, including padding and hidden arguments. A value of zero disables preloading for that kernel. Kernels without the attribute keep using the global option. Allow template-dependent values and validate them after substitution. --- clang/docs/ReleaseNotes.md | 2 + clang/include/clang/Basic/Attr.td | 7 ++++ clang/include/clang/Basic/AttrDocs.td | 41 +++++++++++++++++++ clang/include/clang/Sema/SemaAMDGPU.h | 9 ++++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 6 +++ clang/lib/Sema/SemaAMDGPU.cpp | 36 ++++++++++++++++ clang/lib/Sema/SemaDeclAttr.cpp | 8 ++++ .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 21 ++++++++++ clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu | 20 +++++++++ ...-kernarg-preload-sgpr-count-save-temps.hip | 25 +++++++++++ ...a-attribute-supported-attributes-list.test | 1 + clang/test/SemaCUDA/amdgpu-attrs.cu | 7 ++++ clang/test/SemaOpenCL/amdgpu-attrs.cl | 12 ++++++ .../AMDGPU/AMDGPUPreloadKernelArguments.cpp | 26 ++++++++++-- .../preload-kernargs-sgpr-count-attr.ll | 34 +++++++++++++++ 15 files changed, 252 insertions(+), 3 deletions(-) create mode 100644 clang/test/CodeGenHIP/amdgpu-kernarg-preload-sgpr-count-save-temps.hip create mode 100644 llvm/test/CodeGen/AMDGPU/preload-kernargs-sgpr-count-attr.ll diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index f631d9f858f9f..5b88bafa88323 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -166,6 +166,8 @@ features cannot lower the translation-unit ABI level; ### Attribute Changes in Clang - Clang now properly propagates attributes on class and variable templates to their redeclarations, which will result in redeclarations not interfering with diagnostics. (#GH209812) +- Added the `amdgpu_kernarg_preload_sgpr_count` attribute for AMDGPU kernels to + set a per-kernel SGPR limit for kernel argument preloading. ### Improvements to Clang's diagnostics diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 39c672322d515..a8a8c0bf8b116 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2525,6 +2525,13 @@ def AMDGPUNumVGPR : InheritableAttr { let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; } +def AMDGPUKernargPreloadSGPRCount : InheritableAttr { + let Spellings = [Clang<"amdgpu_kernarg_preload_sgpr_count", 0>]; + let Args = [ExprArgument<"Count">]; + let Documentation = [AMDGPUKernargPreloadSGPRCountDocs]; + let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; +} + def AMDGPUMaxNumWorkGroups : InheritableAttr { let Spellings = [Clang<"amdgpu_max_num_work_groups", 0>]; let Args = [ExprArgument<"MaxNumWorkGroupsX">, ExprArgument<"MaxNumWorkGroupsY", 1>, ExprArgument<"MaxNumWorkGroupsZ", 1>]; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 05e4cb0870652..aa20c32c6f61a 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3574,6 +3574,47 @@ An error will be given if: attributes; - The AMDGPU target backend is unable to create machine code that can meet the request. +}]; +} + +def AMDGPUKernargPreloadSGPRCountDocs : Documentation { + let Category = DocCatAMDGPUAttributes; + let Content = [{ +Clang supports the +``__attribute__((amdgpu_kernarg_preload_sgpr_count(<count>)))`` attribute on +AMDGPU kernel functions. + +Kernel argument preloading can reduce kernel launch latency. Without +preloading, the kernel may need to load explicit kernel arguments from kernel +argument segment memory. On targets that support kernel argument preloading, +the hardware can copy part of the kernel argument segment into user SGPRs +before the kernel starts, avoiding those initial kernel argument memory loads. + +The attribute gives a kernel its own SGPR limit for kernel argument preloading. +This is useful when a translation unit contains kernels that need different +preload behavior. For example, a small hot kernel may benefit from using more +SGPRs for preloading, while another kernel may need preloading disabled to save +user SGPRs. + +The ``<count>`` value is the maximum number of user SGPRs that the backend +should use for the kernel argument preload range. The backend considers +explicit kernel arguments from the start of the kernel argument list and +preloads the longest contiguous prefix that fits within the limit. Padding and +hidden arguments are also included in the limit. The backend may use fewer +SGPRs if an argument is not supported, fewer SGPRs are available, or another +target limit is reached. + +Kernel argument preloading is only supported on AMDGPU targets with the +``kernarg-preload`` target feature. In the current backend, the known +processors with this feature are ``gfx90a``, ``gfx942``, ``gfx950``, +``gfx1250``, and ``gfx1251``. The generic aliases ``gfx9-4-generic`` and +``gfx12-5-generic`` also include the feature. Future processors may also +support it if their feature set includes ``FeatureKernargPreload`` in +``llvm/lib/Target/AMDGPU/AMDGPU.td``. On targets that do not support kernel +argument preloading, this attribute has no effect. + +Passing ``0`` disables kernel argument preloading for the annotated kernel. +Omitting the attribute keeps the default command-line behavior. }]; } diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h index a6205534e0de3..7034ff2817b94 100644 --- a/clang/include/clang/Sema/SemaAMDGPU.h +++ b/clang/include/clang/Sema/SemaAMDGPU.h @@ -64,6 +64,14 @@ class SemaAMDGPU : public SemaBase { void addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max); + AMDGPUKernargPreloadSGPRCountAttr * + CreateAMDGPUKernargPreloadSGPRCountAttr(const AttributeCommonInfo &CI, + Expr *CountExpr); + + void addAMDGPUKernargPreloadSGPRCountAttr(Decl *D, + const AttributeCommonInfo &CI, + Expr *CountExpr); + /// Create an AMDGPUMaxNumWorkGroupsAttr attribute. AMDGPUMaxNumWorkGroupsAttr * CreateAMDGPUMaxNumWorkGroupsAttr(const AttributeCommonInfo &CI, Expr *XExpr, @@ -77,6 +85,7 @@ class SemaAMDGPU : public SemaBase { void handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL); void handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL); void handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL); + void handleAMDGPUKernargPreloadSGPRCountAttr(Decl *D, const ParsedAttr &AL); void handleAMDGPUMaxNumWorkGroupsAttr(Decl *D, const ParsedAttr &AL); void handleAMDGPUFlatWorkGroupSizeAttr(Decl *D, const ParsedAttr &AL); diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 7b37f3f7f9b6e..56c5f8e5b6110 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -377,6 +377,12 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR)); } + if (const auto *Attr = FD->getAttr<AMDGPUKernargPreloadSGPRCountAttr>()) { + unsigned Count = + Attr->getCount()->EvaluateKnownConstInt(M.getContext()).getExtValue(); + F->addFnAttr("amdgpu-kernarg-preload-sgpr-count", llvm::utostr(Count)); + } + if (const auto *Attr = FD->getAttr<AMDGPUMaxNumWorkGroupsAttr>()) { uint32_t X = Attr->getMaxNumWorkGroupsX() ->EvaluateKnownConstInt(M.getContext()) diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 48230fa262d5c..50df5d180c953 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -729,6 +729,42 @@ void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) { AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR)); } +static bool checkAMDGPUKernargPreloadSGPRCountArgument( + Sema &S, Expr *CountExpr, const AMDGPUKernargPreloadSGPRCountAttr &Attr) { + if (S.DiagnoseUnexpandedParameterPack(CountExpr)) + return true; + + if (CountExpr->isValueDependent()) + return false; + + uint32_t Count = 0; + return !S.checkUInt32Argument(Attr, CountExpr, Count); +} + +AMDGPUKernargPreloadSGPRCountAttr * +SemaAMDGPU::CreateAMDGPUKernargPreloadSGPRCountAttr( + const AttributeCommonInfo &CI, Expr *CountExpr) { + ASTContext &Context = getASTContext(); + AMDGPUKernargPreloadSGPRCountAttr TmpAttr(Context, CI, CountExpr); + + if (checkAMDGPUKernargPreloadSGPRCountArgument(SemaRef, CountExpr, TmpAttr)) + return nullptr; + + return ::new (Context) + AMDGPUKernargPreloadSGPRCountAttr(Context, CI, CountExpr); +} + +void SemaAMDGPU::addAMDGPUKernargPreloadSGPRCountAttr( + Decl *D, const AttributeCommonInfo &CI, Expr *CountExpr) { + if (auto *Attr = CreateAMDGPUKernargPreloadSGPRCountAttr(CI, CountExpr)) + D->addAttr(Attr); +} + +void SemaAMDGPU::handleAMDGPUKernargPreloadSGPRCountAttr(Decl *D, + const ParsedAttr &AL) { + addAMDGPUKernargPreloadSGPRCountAttr(D, AL, AL.getArgAsExpr(0)); +} + static bool checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr, Expr *ZExpr, diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 492b125587344..36681a380c460 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7719,6 +7719,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL, case ParsedAttr::AT_AMDGPUNumVGPR: S.AMDGPU().handleAMDGPUNumVGPRAttr(D, AL); break; + case ParsedAttr::AT_AMDGPUKernargPreloadSGPRCount: + S.AMDGPU().handleAMDGPUKernargPreloadSGPRCountAttr(D, AL); + break; case ParsedAttr::AT_AMDGPUMaxNumWorkGroups: S.AMDGPU().handleAMDGPUMaxNumWorkGroupsAttr(D, AL); break; @@ -8609,6 +8612,11 @@ void Sema::ProcessDeclAttributeList( Diag(D->getLocation(), diag::err_attribute_wrong_decl_type) << A << A->isRegularKeywordAttribute() << ExpectedKernelFunction; D->setInvalidDecl(); + } else if (const auto *A = + D->getAttr<AMDGPUKernargPreloadSGPRCountAttr>()) { + Diag(D->getLocation(), diag::err_attribute_wrong_decl_type) + << A << A->isRegularKeywordAttribute() << ExpectedKernelFunction; + D->setInvalidDecl(); } } checkAMDGPUReqdWorkGroupSize(*this, D); diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index f1f97ca125f46..d008d7c570517 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -677,6 +677,20 @@ static void instantiateDependentAMDGPUWavesPerEUAttr( S.AMDGPU().addAMDGPUWavesPerEUAttr(New, Attr, MinExpr, MaxExpr); } +static void instantiateDependentAMDGPUKernargPreloadSGPRCountAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const AMDGPUKernargPreloadSGPRCountAttr &Attr, Decl *New) { + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + + ExprResult Result = S.SubstExpr(Attr.getCount(), TemplateArgs); + if (Result.isInvalid()) + return; + + S.AMDGPU().addAMDGPUKernargPreloadSGPRCountAttr(New, Attr, + Result.getAs<Expr>()); +} + static void instantiateDependentAMDGPUMaxNumWorkGroupsAttr( Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, const AMDGPUMaxNumWorkGroupsAttr &Attr, Decl *New) { @@ -963,6 +977,13 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, *AMDGPUFlatWorkGroupSize, New); } + if (const auto *AMDGPUKernargPreloadSGPRCount = + dyn_cast<AMDGPUKernargPreloadSGPRCountAttr>(TmplAttr)) { + instantiateDependentAMDGPUKernargPreloadSGPRCountAttr( + *this, TemplateArgs, *AMDGPUKernargPreloadSGPRCount, New); + continue; + } + if (const auto *AMDGPUMaxNumWorkGroups = dyn_cast<AMDGPUMaxNumWorkGroupsAttr>(TmplAttr)) { instantiateDependentAMDGPUMaxNumWorkGroupsAttr( diff --git a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu index c1ba7a178cea5..c8af49e1f5a23 100644 --- a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu +++ b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu @@ -12,6 +12,9 @@ // RUN: -check-prefix=NAMD // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \ // RUN: -verify -Wno-deprecated-declarations -o - -x hip %s | FileCheck -check-prefix=NAMD %s +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip \ +// RUN: -ast-dump -ast-dump-filter=template_kernarg_preload_sgpr_count %s \ +// RUN: | FileCheck -check-prefix=AST %s // RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -foffload-uniform-block \ // RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \ @@ -45,6 +48,20 @@ __attribute__((amdgpu_num_vgpr(64))) // expected-no-diagnostics __global__ void num_vgpr_64() { // CHECK: define{{.*}} amdgpu_kernel void @_Z11num_vgpr_64v() [[NUM_VGPR_64:#[0-9]+]] } +__attribute__((amdgpu_kernarg_preload_sgpr_count(2))) // expected-no-diagnostics +__global__ void kernarg_preload_sgpr_count_2() { +// CHECK: define{{.*}} amdgpu_kernel void @_Z28kernarg_preload_sgpr_count_2v() [[KERNARG_PRELOAD_SGPR_COUNT_2:#[0-9]+]] +} + +template<unsigned Count> +__attribute__((amdgpu_kernarg_preload_sgpr_count(Count))) +__global__ void template_kernarg_preload_sgpr_count() {} +template __global__ void template_kernarg_preload_sgpr_count<3>(); +// CHECK: define{{.*}} amdgpu_kernel void @_Z35template_kernarg_preload_sgpr_countILj3EEvv() [[KERNARG_PRELOAD_SGPR_COUNT_3:#[0-9]+]] +// AST-LABEL: FunctionDecl {{.*}} template_kernarg_preload_sgpr_count {{.*}} explicit_instantiation_definition +// AST: AMDGPUKernargPreloadSGPRCountAttr +// AST-NOT: AMDGPUKernargPreloadSGPRCountAttr + __attribute__((amdgpu_max_num_work_groups(32, 4, 2))) // expected-no-diagnostics __global__ void max_num_work_groups_32_4_2() { // CHECK: define{{.*}} amdgpu_kernel void @_Z26max_num_work_groups_32_4_2v() [[MAX_NUM_WORK_GROUPS_32_4_2:#[0-9]+]] @@ -101,6 +118,7 @@ template __global__ void template_a_b_c_max_num_work_groups<32, 4, 2>(); // NAMD-NOT: "amdgpu-waves-per-eu" // NAMD-NOT: "amdgpu-num-vgpr" // NAMD-NOT: "amdgpu-num-sgpr" +// NAMD-NOT: "amdgpu-kernarg-preload-sgpr-count" // NAMD-NOT: "amdgpu-max-num-work-groups" // DEFAULT-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"{{.*}}"uniform-work-group-size" @@ -111,6 +129,8 @@ template __global__ void template_a_b_c_max_num_work_groups<32, 4, 2>(); // CHECK-DAG: attributes [[WAVES_PER_EU_2]] = {{.*}}"amdgpu-waves-per-eu"="2" // CHECK-DAG: attributes [[NUM_SGPR_32]] = {{.*}}"amdgpu-num-sgpr"="32" // CHECK-DAG: attributes [[NUM_VGPR_64]] = {{.*}}"amdgpu-num-vgpr"="64" +// CHECK-DAG: attributes [[KERNARG_PRELOAD_SGPR_COUNT_2]] = {{.*}}"amdgpu-kernarg-preload-sgpr-count"="2" +// CHECK-DAG: attributes [[KERNARG_PRELOAD_SGPR_COUNT_3]] = {{.*}}"amdgpu-kernarg-preload-sgpr-count"="3" // CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_4_2]] = {{.*}}"amdgpu-max-num-workgroups"="32,4,2" // CHECK-DAG: attributes [[MAX_NUM_WORK_GROUPS_32_1_1]] = {{.*}}"amdgpu-max-num-workgroups"="32,1,1" diff --git a/clang/test/CodeGenHIP/amdgpu-kernarg-preload-sgpr-count-save-temps.hip b/clang/test/CodeGenHIP/amdgpu-kernarg-preload-sgpr-count-save-temps.hip new file mode 100644 index 0000000000000..610ab8d6c5b8d --- /dev/null +++ b/clang/test/CodeGenHIP/amdgpu-kernarg-preload-sgpr-count-save-temps.hip @@ -0,0 +1,25 @@ +// REQUIRES: amdgpu-registered-target + +// RUN: rm -rf %t && mkdir -p %t +// RUN: cd %t && %clang --target=x86_64-unknown-linux-gnu --offload-arch=gfx942 \ +// RUN: --cuda-device-only -nogpulib -nogpuinc -save-temps -O2 -c -x hip %s \ +// RUN: -o out.o +// RUN: FileCheck %s --input-file=%t/amdgpu-kernarg-preload-sgpr-count-save-temps-hip-amdgcn-amd-amdhsa-gfx942.s + +extern "C" __attribute__((global, amdgpu_kernarg_preload_sgpr_count(3))) +void preload_sgpr_count_3(int *out, int a, int b) { + *out = a + b; +} + +extern "C" __attribute__((global, amdgpu_kernarg_preload_sgpr_count(0))) +void preload_sgpr_count_0(int *out, int a, int b) { + *out = a + b; +} + +// CHECK-LABEL: .amdhsa_kernel preload_sgpr_count_3 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_length 3 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_offset 0 + +// CHECK-LABEL: .amdhsa_kernel preload_sgpr_count_0 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_length 0 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_offset 0 diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 8bca68e2119e7..8d499666ad29e 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: AMDGPUKernargPreloadSGPRCount (SubjectMatchRule_function) // CHECK-NEXT: AMDGPUMaxNumWorkGroups (SubjectMatchRule_function) // CHECK-NEXT: AMDGPUNumSGPR (SubjectMatchRule_function) // CHECK-NEXT: AMDGPUNumVGPR (SubjectMatchRule_function) diff --git a/clang/test/SemaCUDA/amdgpu-attrs.cu b/clang/test/SemaCUDA/amdgpu-attrs.cu index ee0219696bdf4..41f82b22a3ad9 100644 --- a/clang/test/SemaCUDA/amdgpu-attrs.cu +++ b/clang/test/SemaCUDA/amdgpu-attrs.cu @@ -205,6 +205,13 @@ __global__ void non_cexpr_waves_per_eu_2() {} __attribute__((amdgpu_waves_per_eu(2, ipow2(2)))) __global__ void non_cexpr_waves_per_eu_2_4() {} +// expected-error@+3{{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}} +// expected-note@+4{{in instantiation of}} +template<unsigned long long Count> +__attribute__((amdgpu_kernarg_preload_sgpr_count(Count))) +__global__ void template_kernarg_preload_sgpr_count_too_large() {} +template __global__ void template_kernarg_preload_sgpr_count_too_large<4294967296ULL>(); + __attribute__((amdgpu_max_num_work_groups(32))) __global__ void max_num_work_groups_32() {} diff --git a/clang/test/SemaOpenCL/amdgpu-attrs.cl b/clang/test/SemaOpenCL/amdgpu-attrs.cl index 0e57a88241373..a4e148e43e1f4 100644 --- a/clang/test/SemaOpenCL/amdgpu-attrs.cl +++ b/clang/test/SemaOpenCL/amdgpu-attrs.cl @@ -20,12 +20,17 @@ typedef __attribute__((amdgpu_num_vgpr(64))) struct struct_num_vgpr_64 { // expe int x; float y; } struct_num_vgpr_64; +typedef __attribute__((amdgpu_kernarg_preload_sgpr_count(2))) struct struct_kernarg_preload_sgpr_count_2 { // expected-error {{'amdgpu_kernarg_preload_sgpr_count' attribute only applies to kernel functions}} + int x; + float y; +} struct_kernarg_preload_sgpr_count_2; __attribute__((amdgpu_flat_work_group_size(32, 64))) void func_flat_work_group_size_32_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}} __attribute__((amdgpu_waves_per_eu(2))) void func_waves_per_eu_2() {} // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}} __attribute__((amdgpu_waves_per_eu(2, 4))) void func_waves_per_eu_2_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}} __attribute__((amdgpu_num_sgpr(32))) void func_num_sgpr_32() {} // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}} __attribute__((amdgpu_num_vgpr(64))) void func_num_vgpr_64() {} // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}} +__attribute__((amdgpu_kernarg_preload_sgpr_count(2))) void func_kernarg_preload_sgpr_count_2() {} // expected-error {{'amdgpu_kernarg_preload_sgpr_count' attribute only applies to kernel functions}} __attribute__((amdgpu_flat_work_group_size("ABC", "ABC"))) kernel void kernel_flat_work_group_size_ABC_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}} __attribute__((amdgpu_flat_work_group_size(32, "ABC"))) kernel void kernel_flat_work_group_size_32_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}} @@ -35,6 +40,9 @@ __attribute__((amdgpu_waves_per_eu(2, "ABC"))) kernel void kernel_waves_per_eu_2 __attribute__((amdgpu_waves_per_eu("ABC", 4))) kernel void kernel_waves_per_eu_ABC_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}} __attribute__((amdgpu_num_sgpr("ABC"))) kernel void kernel_num_sgpr_ABC() {} // expected-error {{'amdgpu_num_sgpr' attribute requires an integer constant}} __attribute__((amdgpu_num_vgpr("ABC"))) kernel void kernel_num_vgpr_ABC() {} // expected-error {{'amdgpu_num_vgpr' attribute requires an integer constant}} +__attribute__((amdgpu_kernarg_preload_sgpr_count("ABC"))) kernel void kernel_kernarg_preload_sgpr_count_ABC() {} // expected-error {{'amdgpu_kernarg_preload_sgpr_count' attribute requires an integer constant}} +extern constant int kernarg_preload_sgpr_count; +__attribute__((amdgpu_kernarg_preload_sgpr_count(kernarg_preload_sgpr_count))) kernel void kernel_kernarg_preload_sgpr_count_non_constant() {} // expected-error {{'amdgpu_kernarg_preload_sgpr_count' attribute requires an integer constant}} __attribute__((amdgpu_flat_work_group_size(4294967296, 4294967296))) kernel void kernel_flat_work_group_size_L_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}} __attribute__((amdgpu_flat_work_group_size(32, 4294967296))) kernel void kernel_flat_work_group_size_32_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}} @@ -44,6 +52,7 @@ __attribute__((amdgpu_waves_per_eu(2, 4294967296))) kernel void kernel_waves_per __attribute__((amdgpu_waves_per_eu(4294967296, 4))) kernel void kernel_waves_per_eu_L_4() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}} __attribute__((amdgpu_num_sgpr(4294967296))) kernel void kernel_num_sgpr_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}} __attribute__((amdgpu_num_vgpr(4294967296))) kernel void kernel_num_vgpr_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}} +__attribute__((amdgpu_kernarg_preload_sgpr_count(4294967296))) kernel void kernel_kernarg_preload_sgpr_count_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}} __attribute__((amdgpu_flat_work_group_size(0, 64))) kernel void kernel_flat_work_group_size_0_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute argument is invalid: max must be 0 since min is 0}} __attribute__((amdgpu_waves_per_eu(0, 4))) kernel void kernel_waves_per_eu_0_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute argument is invalid: max must be 0 since min is 0}} @@ -58,6 +67,8 @@ __attribute__((amdgpu_waves_per_eu(2, 4, 8))) kernel void kernel_waves_per_eu_2_ __attribute__((amdgpu_flat_work_group_size(0, 0))) kernel void kernel_flat_work_group_size_0_0() {} __attribute__((amdgpu_waves_per_eu(0))) kernel void kernel_waves_per_eu_0() {} +__attribute__((amdgpu_kernarg_preload_sgpr_count(0))) kernel void kernel_kernarg_preload_sgpr_count_0() {} +__attribute__((amdgpu_kernarg_preload_sgpr_count(2))) kernel void kernel_kernarg_preload_sgpr_count_2() {} __attribute__((amdgpu_waves_per_eu(0, 0))) kernel void kernel_waves_per_eu_0_0() {} __attribute__((amdgpu_num_sgpr(0))) kernel void kernel_num_sgpr_0() {} __attribute__((amdgpu_num_vgpr(0))) kernel void kernel_num_vgpr_0() {} @@ -67,3 +78,4 @@ kernel __attribute__((amdgpu_waves_per_eu(2))) void kernel_waves_per_eu_2() {} kernel __attribute__((amdgpu_waves_per_eu(2, 4))) void kernel_waves_per_eu_2_4() {} kernel __attribute__((amdgpu_num_sgpr(32))) void kernel_num_sgpr_32() {} kernel __attribute__((amdgpu_num_vgpr(64))) void kernel_num_vgpr_64() {} +kernel __attribute__((amdgpu_kernarg_preload_sgpr_count(2))) void kernel_kernarg_preload_sgpr_count_2_alt() {} diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernelArguments.cpp index 7d6e3edc75e1f..59776c6e7bc10 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernelArguments.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernelArguments.cpp @@ -28,6 +28,8 @@ #include "llvm/IR/PassManager.h" #include "llvm/IR/Verifier.h" #include "llvm/Pass.h" +#include <algorithm> +#include <limits> #define DEBUG_TYPE "amdgpu-preload-kernel-arguments" @@ -42,6 +44,9 @@ static cl::opt<bool> cl::desc("Enable preload kernel arguments to SGPRs"), cl::init(true)); +static constexpr StringRef KernargPreloadSGPRCountAttr = + "amdgpu-kernarg-preload-sgpr-count"; + namespace { class AMDGPUPreloadKernelArgumentsLegacy : public ModulePass { @@ -167,8 +172,11 @@ class PreloadKernelArgInfo { } public: - PreloadKernelArgInfo(Function &F, const GCNSubtarget &ST) : F(F), ST(ST) { + PreloadKernelArgInfo(Function &F, const GCNSubtarget &ST, + unsigned MaxNumPreloadSGPRs) + : F(F), ST(ST) { setInitialFreeUserSGPRsCount(); + NumFreeUserSGPRs = std::min(NumFreeUserSGPRs, MaxNumPreloadSGPRs); } // Returns the maximum number of user SGPRs that we have available to preload @@ -291,11 +299,23 @@ static bool markKernelArgsAsInreg(Module &M, const TargetMachine &TM) { F.getCallingConv() != CallingConv::AMDGPU_KERNEL) continue; - PreloadKernelArgInfo PreloadInfo(F, ST); + bool HasKernargPreloadSGPRCountAttr = + F.hasFnAttribute(KernargPreloadSGPRCountAttr); + unsigned MaxNumPreloadSGPRs = + HasKernargPreloadSGPRCountAttr + ? F.getFnAttributeAsParsedInteger(KernargPreloadSGPRCountAttr) + : std::numeric_limits<unsigned>::max(); + if (MaxNumPreloadSGPRs == 0) + continue; + + PreloadKernelArgInfo PreloadInfo(F, ST, MaxNumPreloadSGPRs); uint64_t ExplicitArgOffset = 0; const DataLayout &DL = F.getDataLayout(); const uint64_t BaseOffset = ST.getExplicitKernelArgOffset(); - unsigned NumPreloadsRequested = KernargPreloadCount; + unsigned NumPreloadsRequested = HasKernargPreloadSGPRCountAttr + ? std::numeric_limits<unsigned>::max() + : KernargPreloadCount; + unsigned NumPreloadedExplicitArgs = 0; for (Argument &Arg : F.args()) { // Avoid incompatible attributes and guard against running this pass diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs-sgpr-count-attr.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs-sgpr-count-attr.ll new file mode 100644 index 0000000000000..6e1c257da4894 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs-sgpr-count-attr.ll @@ -0,0 +1,34 @@ +; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -passes=amdgpu-preload-kernel-arguments -amdgpu-kernarg-preload-count=3 < %s | FileCheck %s + +define amdgpu_kernel void @global_default(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c) { +; CHECK-LABEL: define amdgpu_kernel void @global_default( +; CHECK-SAME: ptr addrspace(1) inreg %a, ptr addrspace(1) inreg %b, ptr addrspace(1) inreg %c) + ret void +} + +define amdgpu_kernel void @sgpr_count_five(ptr addrspace(1) %a, i32 %b, ptr addrspace(1) %c) #0 { +; CHECK-LABEL: define amdgpu_kernel void @sgpr_count_five( +; CHECK-SAME: ptr addrspace(1) inreg %a, i32 inreg %b, ptr addrspace(1) %c) + ret void +} + +define amdgpu_kernel void @sgpr_count_zero(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c) #1 { +; CHECK-LABEL: define amdgpu_kernel void @sgpr_count_zero( +; CHECK-SAME: ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c) + ret void +} + +define amdgpu_kernel void @hidden_arg_exceeds_sgpr_count(ptr addrspace(1) %out) #2 { +; CHECK-LABEL: define amdgpu_kernel void @hidden_arg_exceeds_sgpr_count( +; CHECK-SAME: ptr addrspace(1) inreg %out) +; CHECK: call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() + %implicit_arg_ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() + %block_count_x = load i32, ptr addrspace(4) %implicit_arg_ptr + ret void +} + +declare ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr() + +attributes #0 = { "amdgpu-kernarg-preload-sgpr-count"="5" } +attributes #1 = { "amdgpu-kernarg-preload-sgpr-count"="0" } +attributes #2 = { "amdgpu-kernarg-preload-sgpr-count"="2" } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
