https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/211059
>From d9efde198c88c61a38dedce199e21f8dafeaed58 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" <[email protected]> Date: Thu, 30 Jul 2026 09:45:45 -0400 Subject: [PATCH] [AMDGPU] Select a kernel argument preload range Users usually know which kernel arguments are used on hot paths, but raw SGPR counts do not directly describe those arguments. Add an AMDGPU kernel attribute that selects an inclusive source argument range. Derive the hardware offset and length from the ABI layout, and use the same offset for the descriptor and compatibility prologue. --- clang/docs/ReleaseNotes.md | 2 + clang/include/clang/Basic/Attr.td | 7 + clang/include/clang/Basic/AttrDocs.td | 56 ++++++++ .../clang/Basic/DiagnosticSemaKinds.td | 5 + clang/include/clang/Sema/SemaAMDGPU.h | 8 ++ clang/lib/CodeGen/Targets/AMDGPU.cpp | 10 ++ clang/lib/Sema/SemaAMDGPU.cpp | 64 +++++++++ clang/lib/Sema/SemaDeclAttr.cpp | 7 + .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 25 ++++ clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu | 32 +++++ ...mdgpu-kernarg-preload-range-save-temps.hip | 52 +++++++ ...a-attribute-supported-attributes-list.test | 1 + clang/test/SemaCUDA/amdgpu-attrs.cu | 9 +- clang/test/SemaOpenCL/amdgpu-attrs.cl | 16 +++ llvm/lib/Target/AMDGPU/AMDGPU.h | 19 +++ llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp | 11 +- llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp | 13 +- llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp | 14 +- .../AMDGPU/AMDGPULowerKernelArguments.cpp | 10 +- .../AMDGPU/AMDGPUPreloadKernArgProlog.cpp | 2 +- .../AMDGPU/AMDGPUPreloadKernelArguments.cpp | 131 +++++++++++++++--- llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp | 23 ++- llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 63 ++++++--- .../Target/AMDGPU/SIMachineFunctionInfo.cpp | 2 + .../lib/Target/AMDGPU/SIMachineFunctionInfo.h | 10 ++ .../mir-kernarg-preload-offset-roundtrip.mir | 20 +++ .../AMDGPU/preload-kernargs-range-attr.ll | 22 +++ .../AMDGPU/preload-kernargs-range-errors.ll | 38 +++++ .../CodeGen/AMDGPU/preload-kernargs-range.ll | 32 +++++ 29 files changed, 632 insertions(+), 72 deletions(-) create mode 100644 clang/test/CodeGenHIP/amdgpu-kernarg-preload-range-save-temps.hip create mode 100644 llvm/test/CodeGen/AMDGPU/mir-kernarg-preload-offset-roundtrip.mir create mode 100644 llvm/test/CodeGen/AMDGPU/preload-kernargs-range-attr.ll create mode 100644 llvm/test/CodeGen/AMDGPU/preload-kernargs-range-errors.ll create mode 100644 llvm/test/CodeGen/AMDGPU/preload-kernargs-range.ll diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index f631d9f858f9f..c01663f67b015 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` attribute for AMDGPU kernels to select an + inclusive range of explicit kernel arguments to preload. ### Improvements to Clang's diagnostics diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 39c672322d515..815471bd94229 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 AMDGPUKernargPreload : InheritableAttr { + let Spellings = [Clang<"amdgpu_kernarg_preload", 0>]; + let Args = [ExprArgument<"FirstArg">, ExprArgument<"LastArg">]; + let Documentation = [AMDGPUKernargPreloadDocs]; + 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..2520bf235170a 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -3574,6 +3574,62 @@ An error will be given if: attributes; - The AMDGPU target backend is unable to create machine code that can meet the request. +}]; +} + +def AMDGPUKernargPreloadDocs : Documentation { + let Category = DocCatAMDGPUAttributes; + let Content = [{ +Clang supports the +``[[clang::amdgpu_kernarg_preload(<first>, <last>)]]`` and +``__attribute__((amdgpu_kernarg_preload(<first>, <last>)))`` attributes on +AMDGPU kernel functions. + +In HIP, the standard C++ spelling follows the function name:: + + __global__ void kernel(int *out, int value) + [[clang::amdgpu_kernarg_preload(1, 1)]]; + +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. +This attribute is a performance hint and is only applied when optimization is +enabled. For example, compile with ``-O2`` to apply the hint. At ``-O0``, the +hint is ignored. + +The ``<first>`` and ``<last>`` values select an inclusive range of explicit +kernel arguments to preload. Argument indices are zero based. For example, +``amdgpu_kernarg_preload(2, 4)`` preloads the third, fourth, and fifth +arguments. The backend derives the hardware offset and number of SGPRs from +the selected arguments. ABI padding within the selected range is also +preloaded. + +The SGPR budget for preloading is limited and is shared with other uniform +values. Selecting a range lets the user spend that budget on +performance-sensitive kernel arguments. The derived preload offset and SGPR +cap avoid preloading earlier or later arguments that are unused or less +useful, leaving more SGPRs available for other uniform values. + +Both indices must name explicit kernel arguments, and ``<first>`` must not be +greater than ``<last>``. The complete selected range must fit in the user SGPRs +available for kernel argument preloading. Support for a source-language +aggregate argument depends on how the target ABI lowers it. An aggregate +lowered to a supported scalar or vector type can be preloaded. The backend +reports an error if a selected argument is lowered as an LLVM aggregate or +with the LLVM ``byref`` parameter attribute. + +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. + +Omitting the attribute keeps the default command-line behavior. }]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index d9d0d485f16ac..d85be0b6b76e8 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3734,6 +3734,11 @@ def err_swift_abi_parameter_wrong_type : Error< def err_attribute_argument_invalid : Error< "%0 attribute argument is invalid: %select{max must be 0 since min is 0|" "min must not be greater than max}1">; +def err_attribute_amdgpu_kernarg_preload_invalid_range : Error< + "%0 attribute requires the first argument index to be no greater than " + "the last argument index">; +def err_attribute_amdgpu_kernarg_preload_index_out_of_range : Error< + "%0 attribute argument index %1 is out of range for this function">; def err_attribute_amdgpu_flat_work_group_size_mismatch : Error< "'amdgpu_flat_work_group_size' attribute must match " "'reqd_work_group_size' product">; diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h index a6205534e0de3..efc4a5da17767 100644 --- a/clang/include/clang/Sema/SemaAMDGPU.h +++ b/clang/include/clang/Sema/SemaAMDGPU.h @@ -64,6 +64,13 @@ class SemaAMDGPU : public SemaBase { void addAMDGPUWavesPerEUAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Min, Expr *Max); + AMDGPUKernargPreloadAttr * + CreateAMDGPUKernargPreloadAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *FirstArgExpr, Expr *LastArgExpr); + + void addAMDGPUKernargPreloadAttr(Decl *D, const AttributeCommonInfo &CI, + Expr *FirstArgExpr, Expr *LastArgExpr); + /// Create an AMDGPUMaxNumWorkGroupsAttr attribute. AMDGPUMaxNumWorkGroupsAttr * CreateAMDGPUMaxNumWorkGroupsAttr(const AttributeCommonInfo &CI, Expr *XExpr, @@ -77,6 +84,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 handleAMDGPUKernargPreloadAttr(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..8b9afec525174 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -377,6 +377,16 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( F->addFnAttr("amdgpu-num-vgpr", llvm::utostr(NumVGPR)); } + if (const auto *Attr = FD->getAttr<AMDGPUKernargPreloadAttr>()) { + unsigned FirstArg = Attr->getFirstArg() + ->EvaluateKnownConstInt(M.getContext()) + .getExtValue(); + unsigned LastArg = + Attr->getLastArg()->EvaluateKnownConstInt(M.getContext()).getExtValue(); + F->addFnAttr("amdgpu-kernarg-preload-first-arg", llvm::utostr(FirstArg)); + F->addFnAttr("amdgpu-kernarg-preload-last-arg", llvm::utostr(LastArg)); + } + 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..6cc8f3c3d6420 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -729,6 +729,70 @@ void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) { AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR)); } +static bool +checkAMDGPUKernargPreloadArguments(Sema &S, Decl *D, Expr *FirstArgExpr, + Expr *LastArgExpr, + const AMDGPUKernargPreloadAttr &Attr) { + if (S.DiagnoseUnexpandedParameterPack(FirstArgExpr) || + S.DiagnoseUnexpandedParameterPack(LastArgExpr)) + return true; + + if (FirstArgExpr->isValueDependent() || LastArgExpr->isValueDependent()) + return false; + + uint32_t FirstArg = 0; + if (!S.checkUInt32Argument(Attr, FirstArgExpr, FirstArg, 0)) + return true; + + uint32_t LastArg = 0; + if (!S.checkUInt32Argument(Attr, LastArgExpr, LastArg, 1)) + return true; + + if (FirstArg > LastArg) { + S.Diag(Attr.getLocation(), + diag::err_attribute_amdgpu_kernarg_preload_invalid_range) + << &Attr; + return true; + } + + const auto *FD = dyn_cast<FunctionDecl>(D); + if (FD && LastArg >= FD->getNumParams()) { + S.Diag(LastArgExpr->getBeginLoc(), + diag::err_attribute_amdgpu_kernarg_preload_index_out_of_range) + << &Attr << LastArg << LastArgExpr->getSourceRange(); + return true; + } + + return false; +} + +AMDGPUKernargPreloadAttr *SemaAMDGPU::CreateAMDGPUKernargPreloadAttr( + Decl *D, const AttributeCommonInfo &CI, Expr *FirstArgExpr, + Expr *LastArgExpr) { + ASTContext &Context = getASTContext(); + AMDGPUKernargPreloadAttr TmpAttr(Context, CI, FirstArgExpr, LastArgExpr); + + if (checkAMDGPUKernargPreloadArguments(SemaRef, D, FirstArgExpr, LastArgExpr, + TmpAttr)) + return nullptr; + + return ::new (Context) + AMDGPUKernargPreloadAttr(Context, CI, FirstArgExpr, LastArgExpr); +} + +void SemaAMDGPU::addAMDGPUKernargPreloadAttr(Decl *D, + const AttributeCommonInfo &CI, + Expr *FirstArgExpr, + Expr *LastArgExpr) { + if (auto *Attr = + CreateAMDGPUKernargPreloadAttr(D, CI, FirstArgExpr, LastArgExpr)) + D->addAttr(Attr); +} + +void SemaAMDGPU::handleAMDGPUKernargPreloadAttr(Decl *D, const ParsedAttr &AL) { + addAMDGPUKernargPreloadAttr(D, AL, AL.getArgAsExpr(0), AL.getArgAsExpr(1)); +} + 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..997d0fb26b5dd 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_AMDGPUKernargPreload: + S.AMDGPU().handleAMDGPUKernargPreloadAttr(D, AL); + break; case ParsedAttr::AT_AMDGPUMaxNumWorkGroups: S.AMDGPU().handleAMDGPUMaxNumWorkGroupsAttr(D, AL); break; @@ -8609,6 +8612,10 @@ 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<AMDGPUKernargPreloadAttr>()) { + 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..3b93bf4559016 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -677,6 +677,24 @@ static void instantiateDependentAMDGPUWavesPerEUAttr( S.AMDGPU().addAMDGPUWavesPerEUAttr(New, Attr, MinExpr, MaxExpr); } +static void instantiateDependentAMDGPUKernargPreloadAttr( + Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, + const AMDGPUKernargPreloadAttr &Attr, Decl *New) { + EnterExpressionEvaluationContext Unevaluated( + S, Sema::ExpressionEvaluationContext::ConstantEvaluated); + + ExprResult FirstArgResult = S.SubstExpr(Attr.getFirstArg(), TemplateArgs); + if (FirstArgResult.isInvalid()) + return; + + ExprResult LastArgResult = S.SubstExpr(Attr.getLastArg(), TemplateArgs); + if (LastArgResult.isInvalid()) + return; + + S.AMDGPU().addAMDGPUKernargPreloadAttr( + New, Attr, FirstArgResult.getAs<Expr>(), LastArgResult.getAs<Expr>()); +} + static void instantiateDependentAMDGPUMaxNumWorkGroupsAttr( Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs, const AMDGPUMaxNumWorkGroupsAttr &Attr, Decl *New) { @@ -963,6 +981,13 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, *AMDGPUFlatWorkGroupSize, New); } + if (const auto *AMDGPUKernargPreload = + dyn_cast<AMDGPUKernargPreloadAttr>(TmplAttr)) { + instantiateDependentAMDGPUKernargPreloadAttr(*this, TemplateArgs, + *AMDGPUKernargPreload, 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..e95527fededbb 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 %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,30 @@ __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]+]] } +extern "C" __global__ void kernarg_preload_1_2 + [[clang::amdgpu_kernarg_preload(1, 2)]] (int a, int b, int c) { // expected-no-diagnostics +// CHECK: define{{.*}} amdgpu_kernel void @kernarg_preload_1_2({{.*}}) [[KERNARG_PRELOAD_1_2:#[0-9]+]] +} + +struct KernargScalarizedAggregate { + int value; +}; + +extern "C" __attribute__((amdgpu_kernarg_preload(0, 0))) +__global__ void kernarg_preload_scalarized_aggregate( + KernargScalarizedAggregate arg) { +// CHECK: define{{.*}} amdgpu_kernel void @kernarg_preload_scalarized_aggregate(i32{{.*}}) [[KERNARG_PRELOAD_SCALARIZED_AGGREGATE:#[0-9]+]] +} + +template<unsigned First, unsigned Last> +__attribute__((amdgpu_kernarg_preload(First, Last))) +__global__ void template_kernarg_preload(int a, int b, int c, int d) {} +template __global__ void template_kernarg_preload<2, 3>(int, int, int, int); +// CHECK: define{{.*}} amdgpu_kernel void @{{.*template_kernarg_preload.*}}({{.*}}) [[KERNARG_PRELOAD_2_3:#[0-9]+]] +// AST-LABEL: FunctionDecl {{.*}} template_kernarg_preload {{.*}} explicit_instantiation_definition +// AST: AMDGPUKernargPreloadAttr +// AST-NOT: AMDGPUKernargPreloadAttr + __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 +128,8 @@ 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-first-arg" +// NAMD-NOT: "amdgpu-kernarg-preload-last-arg" // 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 +140,9 @@ 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_1_2]] = {{.*}}"amdgpu-kernarg-preload-first-arg"="1"{{.*}}"amdgpu-kernarg-preload-last-arg"="2" +// CHECK-DAG: attributes [[KERNARG_PRELOAD_SCALARIZED_AGGREGATE]] = {{.*}}"amdgpu-kernarg-preload-first-arg"="0"{{.*}}"amdgpu-kernarg-preload-last-arg"="0" +// CHECK-DAG: attributes [[KERNARG_PRELOAD_2_3]] = {{.*}}"amdgpu-kernarg-preload-first-arg"="2"{{.*}}"amdgpu-kernarg-preload-last-arg"="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-range-save-temps.hip b/clang/test/CodeGenHIP/amdgpu-kernarg-preload-range-save-temps.hip new file mode 100644 index 0000000000000..82295b331a072 --- /dev/null +++ b/clang/test/CodeGenHIP/amdgpu-kernarg-preload-range-save-temps.hip @@ -0,0 +1,52 @@ +// 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-range-save-temps-hip-amdgcn-amd-amdhsa-gfx942.s +// RUN: %clang --target=x86_64-unknown-linux-gnu --offload-arch=gfx942 \ +// RUN: --cuda-device-only -nogpulib -nogpuinc -O0 -S -x hip %s \ +// RUN: -o %t/unoptimized.s +// RUN: FileCheck %s --check-prefix=O0 --input-file=%t/unoptimized.s +// RUN: not %clang --target=x86_64-unknown-linux-gnu --offload-arch=gfx942 \ +// RUN: --cuda-device-only -nogpulib -nogpuinc -O2 -S -x hip \ +// RUN: -DUNSUPPORTED_AGGREGATE %s -o %t/unsupported.s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=ERROR + +extern "C" __attribute__((global, amdgpu_kernarg_preload(1, 2))) +void preload_range_1_2(int *out, int a, int b) { + *out = a + b; +} + +extern "C" __attribute__((global, amdgpu_kernarg_preload(0, 0))) +void preload_range_0_0(int *out, int a, int b) { + *out = a + b; +} + +#ifdef UNSUPPORTED_AGGREGATE +struct UnsupportedAggregate { + int value[2]; +}; + +extern "C" __attribute__((global, amdgpu_kernarg_preload(0, 0))) +void preload_unsupported_aggregate(UnsupportedAggregate arg) {} +#endif + +// CHECK-LABEL: .amdhsa_kernel preload_range_1_2 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_length 2 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_offset 2 + +// CHECK-LABEL: .amdhsa_kernel preload_range_0_0 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_length 2 +// CHECK: .amdhsa_user_sgpr_kernarg_preload_offset 0 + +// O0-LABEL: .amdhsa_kernel preload_range_1_2 +// O0: .amdhsa_user_sgpr_kernarg_preload_length 0 +// O0: .amdhsa_user_sgpr_kernarg_preload_offset 0 + +// O0-LABEL: .amdhsa_kernel preload_range_0_0 +// O0: .amdhsa_user_sgpr_kernarg_preload_length 0 +// O0: .amdhsa_user_sgpr_kernarg_preload_offset 0 + +// ERROR: error: unsupported argument in amdgpu kernarg preload range diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 8bca68e2119e7..862adb3e6a227 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: AMDGPUKernargPreload (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..0618b5cc8a875 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 Last> +__attribute__((amdgpu_kernarg_preload(0, Last))) +__global__ void template_kernarg_preload_too_large(int x) {} +template __global__ void template_kernarg_preload_too_large<4294967296ULL>(int); + __attribute__((amdgpu_max_num_work_groups(32))) __global__ void max_num_work_groups_32() {} @@ -324,5 +331,3 @@ template<unsigned b> __attribute__((amdgpu_max_num_work_groups(32, 1, b))) __global__ void template_32_1_b_max_num_work_groups() {} template __global__ void template_32_1_b_max_num_work_groups<0>(); - - diff --git a/clang/test/SemaOpenCL/amdgpu-attrs.cl b/clang/test/SemaOpenCL/amdgpu-attrs.cl index 0e57a88241373..ccd6ab0c3e978 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(0, 0))) struct struct_kernarg_preload_0_0 { // expected-error {{'amdgpu_kernarg_preload' attribute only applies to kernel functions}} + int x; + float y; +} struct_kernarg_preload_0_0; __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(0, 0))) void func_kernarg_preload_0_0(int x) {} // expected-error {{'amdgpu_kernarg_preload' 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("ABC", 0))) kernel void kernel_kernarg_preload_ABC(int x) {} // expected-error {{'amdgpu_kernarg_preload' attribute requires parameter 0 to be an integer constant}} +extern constant int kernarg_preload_index; +__attribute__((amdgpu_kernarg_preload(0, kernarg_preload_index))) kernel void kernel_kernarg_preload_non_constant(int x) {} // expected-error {{'amdgpu_kernarg_preload' attribute requires parameter 1 to be 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(4294967296, 4294967296))) kernel void kernel_kernarg_preload_L(int x) {} // 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,12 @@ __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(0, 0))) kernel void kernel_kernarg_preload_0_0(int x) {} +__attribute__((amdgpu_kernarg_preload(1, 2))) kernel void kernel_kernarg_preload_1_2(int x, int y, int z) {} +__attribute__((amdgpu_kernarg_preload(2, 1))) kernel void kernel_kernarg_preload_bad_range(int x, int y, int z) {} // expected-error {{'amdgpu_kernarg_preload' attribute requires the first argument index to be no greater than the last argument index}} +__attribute__((amdgpu_kernarg_preload(0, 1))) kernel void kernel_kernarg_preload_bad_index(int x) {} // expected-error {{'amdgpu_kernarg_preload' attribute argument index 1 is out of range for this function}} +__attribute__((amdgpu_kernarg_preload(0))) kernel void kernel_kernarg_preload_too_few(int x) {} // expected-error {{'amdgpu_kernarg_preload' attribute requires exactly 2 arguments}} +__attribute__((amdgpu_kernarg_preload(0, 0, 0))) kernel void kernel_kernarg_preload_too_many(int x) {} // expected-error {{'amdgpu_kernarg_preload' attribute requires exactly 2 arguments}} __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 +82,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(0, 0))) void kernel_kernarg_preload_0_0_alt(int x) {} diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h index c72fa69aa1419..cadf32160b5e3 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPU.h +++ b/llvm/lib/Target/AMDGPU/AMDGPU.h @@ -10,16 +10,21 @@ #ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPU_H #define LLVM_LIB_TARGET_AMDGPU_AMDGPU_H +#include "llvm/ADT/StringRef.h" #include "llvm/Analysis/CGSCCPassManager.h" #include "llvm/CodeGen/MachinePassManager.h" #include "llvm/IR/PassManager.h" #include "llvm/Pass.h" #include "llvm/Support/AMDGPUAddrSpace.h" +#include "llvm/Support/Alignment.h" #include "llvm/Support/CodeGen.h" +#include <cstdint> namespace llvm { class AMDGPUTargetMachine; +class Argument; +class DataLayout; class LazyCallGraph; class GCNTargetMachine; class TargetMachine; @@ -606,6 +611,20 @@ struct AMDGPUUniformIntrinsicCombinePass }; namespace AMDGPU { +struct KernArgLayout { + uint64_t Begin; + uint64_t End; + Align Alignment; +}; + +KernArgLayout getKernArgLayout(const Argument &Arg, const DataLayout &DL, + uint64_t CurrentOffset); + +inline constexpr StringLiteral KernargPreloadFirstArgAttr = + "amdgpu-kernarg-preload-first-arg"; +inline constexpr StringLiteral KernargPreloadLastArgAttr = + "amdgpu-kernarg-preload-last-arg"; + enum TargetIndex { TI_CONSTDATA_START, TI_SCRATCH_RSRC_DWORD0, diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp index e1800c4ebded2..1227ca36de29d 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp @@ -866,9 +866,14 @@ AMDGPUAsmPrinter::getAmdhsaKernelDescriptor(const MachineFunction &MF, static_cast<uint64_t>(PGM_Rsrc3) == 0); KernelDescriptor.compute_pgm_rsrc3 = CurrentProgramInfo.ComputePGMRSrc3; - KernelDescriptor.kernarg_preload = MCConstantExpr::create( - AMDGPU::hasKernargPreload(STM) ? Info->getNumKernargPreloadedSGPRs() : 0, - Ctx); + unsigned KernargPreload = 0; + if (AMDGPU::hasKernargPreload(STM)) { + KernargPreload = Info->getNumKernargPreloadedSGPRs(); + KernargPreload |= Info->getKernargPreloadOffset() + << amdhsa::KERNARG_PRELOAD_SPEC_OFFSET_SHIFT; + } + KernelDescriptor.kernarg_preload = + MCConstantExpr::create(KernargPreload, Ctx); return KernelDescriptor; } diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp index 3c78bb1dd7be3..9aa986bc17b78 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp @@ -597,16 +597,13 @@ bool AMDGPUCallLowering::lowerFormalArgumentsKernel( } const bool IsByRef = Arg.hasByRefAttr(); - Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType(); - unsigned AllocSize = DL.getTypeAllocSize(ArgTy); - if (AllocSize == 0) + AMDGPU::KernArgLayout Layout = + AMDGPU::getKernArgLayout(Arg, DL, ExplicitArgOffset); + if (Layout.Begin == Layout.End) continue; - MaybeAlign ParamAlign = IsByRef ? Arg.getParamAlign() : std::nullopt; - Align ABIAlign = DL.getValueOrABITypeAlignment(ParamAlign, ArgTy); - - uint64_t ArgOffset = alignTo(ExplicitArgOffset, ABIAlign) + BaseOffset; - ExplicitArgOffset = alignTo(ExplicitArgOffset, ABIAlign) + AllocSize; + uint64_t ArgOffset = Layout.Begin + BaseOffset; + ExplicitArgOffset = Layout.End; if (Arg.use_empty()) { ++i; diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp index 962988ff97e39..fa0aa0d5a63fa 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp @@ -1231,16 +1231,12 @@ void AMDGPUTargetLowering::analyzeFormalArgumentsCompute( unsigned InIndex = 0; for (const Argument &Arg : Fn.args()) { - const bool IsByRef = Arg.hasByRefAttr(); Type *BaseArgTy = Arg.getType(); - Type *MemArgTy = IsByRef ? Arg.getParamByRefType() : BaseArgTy; - Align Alignment = DL.getValueOrABITypeAlignment( - IsByRef ? Arg.getParamAlign() : std::nullopt, MemArgTy); - MaxAlign = std::max(Alignment, MaxAlign); - uint64_t AllocSize = DL.getTypeAllocSize(MemArgTy); - - uint64_t ArgOffset = alignTo(ExplicitArgOffset, Alignment) + ExplicitOffset; - ExplicitArgOffset = alignTo(ExplicitArgOffset, Alignment) + AllocSize; + AMDGPU::KernArgLayout Layout = + AMDGPU::getKernArgLayout(Arg, DL, ExplicitArgOffset); + MaxAlign = std::max(Layout.Alignment, MaxAlign); + uint64_t ArgOffset = Layout.Begin + ExplicitOffset; + ExplicitArgOffset = Layout.End; // We're basically throwing away everything passed into us and starting over // to get accurate in-memory offsets. The "PartOffset" is completely useless diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp index d4f6532bcec67..3024f9bf36fff 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp @@ -235,14 +235,12 @@ static bool lowerKernelArguments(Function &F, const TargetMachine &TM, for (Argument &Arg : F.args()) { const bool IsByRef = Arg.hasByRefAttr(); Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType(); - MaybeAlign ParamAlign = IsByRef ? Arg.getParamAlign() : std::nullopt; - Align ABITypeAlign = DL.getValueOrABITypeAlignment(ParamAlign, ArgTy); uint64_t Size = DL.getTypeSizeInBits(ArgTy); - uint64_t AllocSize = DL.getTypeAllocSize(ArgTy); - - uint64_t EltOffset = alignTo(ExplicitArgOffset, ABITypeAlign) + BaseOffset; - ExplicitArgOffset = alignTo(ExplicitArgOffset, ABITypeAlign) + AllocSize; + AMDGPU::KernArgLayout Layout = + AMDGPU::getKernArgLayout(Arg, DL, ExplicitArgOffset); + uint64_t EltOffset = Layout.Begin + BaseOffset; + ExplicitArgOffset = Layout.End; // Skip inreg arguments which should be preloaded. if (Arg.use_empty() || Arg.hasInRegAttr()) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernArgProlog.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernArgProlog.cpp index cf8f12c4704fc..36c2cd3b985ca 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernArgProlog.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernArgProlog.cpp @@ -183,7 +183,7 @@ void AMDGPUPreloadKernArgProlog::addBackCompatLoads( MachineBasicBlock *BackCompatMBB, Register KernArgSegmentPtr, unsigned NumKernArgPreloadSGPRs) { Register KernArgPreloadSGPR = MFI.getArgInfo().FirstKernArgPreloadReg; - unsigned Offset = ST.getExplicitKernelArgOffset(); + unsigned Offset = MFI.getKernargPreloadOffset() * 4; // Fill all user SGPRs used for kernarg preloading with sequential data from // the kernarg segment while (NumKernArgPreloadSGPRs > 0) { diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPreloadKernelArguments.cpp index 7d6e3edc75e1f..99c934966315a 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 "llvm/Support/MathExtras.h" +#include <algorithm> #define DEBUG_TYPE "amdgpu-preload-kernel-arguments" @@ -182,6 +184,10 @@ class PreloadKernelArgInfo { return ExplicitArgOffset <= NumFreeUserSGPRs * 4; } + bool canPreloadKernArgRange(uint64_t Begin, uint64_t End) { + return End - Begin <= NumFreeUserSGPRs * 4; + } + // Try to allocate SGPRs to preload hidden kernel arguments. void tryAllocHiddenArgPreloadSGPRs(uint64_t ImplicitArgsBaseOffset, @@ -291,54 +297,145 @@ static bool markKernelArgsAsInreg(Module &M, const TargetMachine &TM) { F.getCallingConv() != CallingConv::AMDGPU_KERNEL) continue; + bool HasFirstPreloadArg = + F.hasFnAttribute(AMDGPU::KernargPreloadFirstArgAttr); + bool HasLastPreloadArg = + F.hasFnAttribute(AMDGPU::KernargPreloadLastArgAttr); + if (HasFirstPreloadArg != HasLastPreloadArg) { + F.getContext().emitError("incomplete amdgpu kernarg preload range"); + continue; + } + bool HasKernargPreloadRange = HasFirstPreloadArg; + uint64_t FirstPreloadArg = HasKernargPreloadRange + ? F.getFnAttributeAsParsedInteger( + AMDGPU::KernargPreloadFirstArgAttr) + : 0; + uint64_t LastPreloadArg = + HasKernargPreloadRange + ? F.getFnAttributeAsParsedInteger(AMDGPU::KernargPreloadLastArgAttr) + : 0; + if (HasKernargPreloadRange && + (!isUInt<32>(FirstPreloadArg) || !isUInt<32>(LastPreloadArg))) { + F.getContext().emitError( + "amdgpu kernarg preload argument index exceeds the 32-bit limit"); + continue; + } + if (HasKernargPreloadRange && + (FirstPreloadArg > LastPreloadArg || LastPreloadArg >= F.arg_size())) { + F.getContext().emitError("invalid amdgpu kernarg preload argument range"); + continue; + } + PreloadKernelArgInfo PreloadInfo(F, ST); uint64_t ExplicitArgOffset = 0; const DataLayout &DL = F.getDataLayout(); const uint64_t BaseOffset = ST.getExplicitKernelArgOffset(); unsigned NumPreloadsRequested = KernargPreloadCount; + uint64_t PreloadRangeBegin = 0; + uint64_t PreloadRangeEnd = 0; + bool FailedPreloadRange = false; + SmallVector<Argument *, 8> RequestedPreloadArgs; + unsigned NumPreloadedExplicitArgs = 0; for (Argument &Arg : F.args()) { + unsigned ArgNo = Arg.getArgNo(); // Avoid incompatible attributes and guard against running this pass // twice. // // TODO: Preload byref kernel arguments - if (Arg.hasByRefAttr() || Arg.hasNestAttr() || - Arg.hasAttribute("amdgpu-hidden-argument")) + if (Arg.hasAttribute("amdgpu-hidden-argument")) break; + bool IsByRef = Arg.hasByRefAttr(); + Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType(); + AMDGPU::KernArgLayout Layout = + AMDGPU::getKernArgLayout(Arg, DL, ExplicitArgOffset); + ExplicitArgOffset = Layout.End; + + bool InRequestedRange = HasKernargPreloadRange && + ArgNo >= FirstPreloadArg && + ArgNo <= LastPreloadArg; + if (HasKernargPreloadRange && !InRequestedRange) + continue; + // Inreg may be pre-existing on some arguments, try to preload these. - if (NumPreloadsRequested == 0 && !Arg.hasInRegAttr()) + if (!HasKernargPreloadRange && NumPreloadsRequested == 0 && + !Arg.hasInRegAttr()) break; - // FIXME: Preload aggregates. - if (Arg.getType()->isAggregateType()) + // FIXME: Preload aggregates and byref arguments. + if (IsByRef || Arg.hasNestAttr() || ArgTy->isAggregateType()) { + if (HasKernargPreloadRange) { + F.getContext().emitError( + "unsupported argument in amdgpu kernarg preload range"); + } + FailedPreloadRange = HasKernargPreloadRange; break; + } - Type *ArgTy = Arg.getType(); - Align ABITypeAlign = DL.getABITypeAlign(ArgTy); - uint64_t AllocSize = DL.getTypeAllocSize(ArgTy); - ExplicitArgOffset = alignTo(ExplicitArgOffset, ABITypeAlign) + AllocSize; + if (HasKernargPreloadRange) { + if (ArgNo == FirstPreloadArg) + PreloadRangeBegin = alignDown(Layout.Begin + BaseOffset, 4); + if (ArgNo == LastPreloadArg) + PreloadRangeEnd = alignTo(ExplicitArgOffset + BaseOffset, 4); + } - if (!PreloadInfo.canPreloadKernArgAtOffset(ExplicitArgOffset)) + if (!HasKernargPreloadRange && + !PreloadInfo.canPreloadKernArgAtOffset(ExplicitArgOffset)) break; - Arg.addAttr(Attribute::InReg); - NumPreloadedExplicitArgs++; - if (NumPreloadsRequested > 0) - NumPreloadsRequested--; + if (HasKernargPreloadRange) { + RequestedPreloadArgs.push_back(&Arg); + } else { + bool HadInRegAttr = Arg.hasInRegAttr(); + Arg.addAttr(Attribute::InReg); + Changed |= !HadInRegAttr; + NumPreloadedExplicitArgs++; + if (NumPreloadsRequested > 0) + NumPreloadsRequested--; + } + } + + if (FailedPreloadRange) + continue; + + if (HasKernargPreloadRange && + RequestedPreloadArgs.size() != LastPreloadArg - FirstPreloadArg + 1) { + F.getContext().emitError( + "amdgpu kernarg preload range must contain explicit arguments"); + continue; + } + + if (HasKernargPreloadRange && !isUInt<9>(PreloadRangeBegin / 4)) { + F.getContext().emitError( + "amdgpu kernarg preload offset exceeds the hardware limit"); + continue; + } + + if (HasKernargPreloadRange && !PreloadInfo.canPreloadKernArgRange( + PreloadRangeBegin, PreloadRangeEnd)) { + F.getContext().emitError( + "amdgpu kernarg preload range exceeds available user SGPRs"); + continue; + } + + if (HasKernargPreloadRange) { + for (Argument *Arg : RequestedPreloadArgs) { + bool HadInRegAttr = Arg->hasInRegAttr(); + Arg->addAttr(Attribute::InReg); + Changed |= !HadInRegAttr; + } } // Only try preloading hidden arguments if we can successfully preload the // last explicit argument. - if (NumPreloadedExplicitArgs == F.arg_size()) { + if (!HasKernargPreloadRange && NumPreloadedExplicitArgs == F.arg_size()) { uint64_t ImplicitArgsBaseOffset = alignTo(ExplicitArgOffset, ST.getAlignmentForImplicitArgPtr()) + BaseOffset; PreloadInfo.tryAllocHiddenArgPreloadSGPRs(ImplicitArgsBaseOffset, FunctionsToErase); } - - Changed |= NumPreloadedExplicitArgs > 0; } Changed |= !FunctionsToErase.empty(); diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp index 87515d22ed422..600cdeccfbdfe 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.cpp @@ -12,6 +12,7 @@ //===----------------------------------------------------------------------===// #include "AMDGPUSubtarget.h" +#include "AMDGPU.h" #include "AMDGPUCallLowering.h" #include "AMDGPUInstructionSelector.h" #include "AMDGPULegalizerInfo.h" @@ -32,6 +33,17 @@ using namespace llvm; #define DEBUG_TYPE "amdgpu-subtarget" +AMDGPU::KernArgLayout AMDGPU::getKernArgLayout(const Argument &Arg, + const DataLayout &DL, + uint64_t CurrentOffset) { + const bool IsByRef = Arg.hasByRefAttr(); + Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType(); + Align Alignment = DL.getValueOrABITypeAlignment( + IsByRef ? Arg.getParamAlign() : std::nullopt, ArgTy); + uint64_t Begin = alignTo(CurrentOffset, Alignment); + return {Begin, Begin + DL.getTypeAllocSize(ArgTy), Alignment}; +} + // Returns the maximum per-workgroup LDS allocation size (in bytes) that still // allows the given function to achieve an occupancy of NWaves waves per // SIMD / EU, taking into account only the function's *maximum* workgroup size. @@ -378,13 +390,10 @@ uint64_t AMDGPUSubtarget::getExplicitKernArgSize(const Function &F, if (Arg.hasAttribute("amdgpu-hidden-argument")) continue; - const bool IsByRef = Arg.hasByRefAttr(); - Type *ArgTy = IsByRef ? Arg.getParamByRefType() : Arg.getType(); - Align Alignment = DL.getValueOrABITypeAlignment( - IsByRef ? Arg.getParamAlign() : std::nullopt, ArgTy); - uint64_t AllocSize = DL.getTypeAllocSize(ArgTy); - ExplicitArgBytes = alignTo(ExplicitArgBytes, Alignment) + AllocSize; - MaxAlign = std::max(MaxAlign, Alignment); + AMDGPU::KernArgLayout Layout = + AMDGPU::getKernArgLayout(Arg, DL, ExplicitArgBytes); + ExplicitArgBytes = Layout.End; + MaxAlign = std::max(MaxAlign, Layout.Alignment); } return ExplicitArgBytes; diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index 78bbb7f2d6146..e8cce277252c1 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -3118,8 +3118,7 @@ void SITargetLowering::allocateHSAUserSGPRs(CCState &CCInfo, // these from the dispatch pointer. } -// Allocate pre-loaded kernel arguemtns. Arguments to be preloading must be -// sequential starting from the first argument. +// Allocate a sequential range of preloaded kernel arguments. void SITargetLowering::allocatePreloadKernArgSGPRs( CCState &CCInfo, SmallVectorImpl<CCValAssign> &ArgLocs, const SmallVectorImpl<ISD::InputArg> &Ins, MachineFunction &MF, @@ -3127,26 +3126,41 @@ void SITargetLowering::allocatePreloadKernArgSGPRs( Function &F = MF.getFunction(); unsigned LastExplicitArgOffset = Subtarget->getExplicitKernelArgOffset(); GCNUserSGPRUsageInfo &SGPRInfo = Info.getUserSGPRInfo(); + bool HasPreloadRange = F.hasFnAttribute(AMDGPU::KernargPreloadFirstArgAttr) && + F.hasFnAttribute(AMDGPU::KernargPreloadLastArgAttr); + uint64_t FirstPreloadArg = + HasPreloadRange + ? F.getFnAttributeAsParsedInteger(AMDGPU::KernargPreloadFirstArgAttr) + : 0; + uint64_t LastPreloadArg = + HasPreloadRange + ? F.getFnAttributeAsParsedInteger(AMDGPU::KernargPreloadLastArgAttr) + : 0; bool InPreloadSequence = true; unsigned InIdx = 0; bool AlignedForImplictArgs = false; unsigned ImplicitArgOffset = 0; for (auto &Arg : F.args()) { - if (!InPreloadSequence || !Arg.hasInRegAttr()) + if (!InPreloadSequence) break; unsigned ArgIdx = Arg.getArgNo(); - // Don't preload non-original args or parts not in the current preload - // sequence. + unsigned FirstInIdx = InIdx; if (InIdx < Ins.size() && (!Ins[InIdx].isOrigArg() || Ins[InIdx].getOrigArgIndex() != ArgIdx)) break; + while (InIdx < Ins.size() && Ins[InIdx].isOrigArg() && + Ins[InIdx].getOrigArgIndex() == ArgIdx) + ++InIdx; - for (; InIdx < Ins.size() && Ins[InIdx].isOrigArg() && - Ins[InIdx].getOrigArgIndex() == ArgIdx; - InIdx++) { - assert(ArgLocs[ArgIdx].isMemLoc()); - auto &ArgLoc = ArgLocs[InIdx]; + if (HasPreloadRange && ArgIdx < FirstPreloadArg) + continue; + if ((HasPreloadRange && ArgIdx > LastPreloadArg) || !Arg.hasInRegAttr()) + break; + + for (unsigned PartIdx = FirstInIdx; PartIdx < InIdx; ++PartIdx) { + assert(ArgLocs[PartIdx].isMemLoc()); + CCValAssign &ArgLoc = ArgLocs[PartIdx]; const Align KernelArgBaseAlign = Align(16); unsigned ArgOffset = ArgLoc.getLocMemOffset(); Align Alignment = commonAlignment(KernelArgBaseAlign, ArgOffset); @@ -3165,18 +3179,29 @@ void SITargetLowering::allocatePreloadKernArgSGPRs( ArgOffset += ImplicitArgOffset; } + unsigned ArgDwordOffset = alignDown(ArgOffset, 4); + bool IsFirstPreload = Info.getNumKernargPreloadedSGPRs() == 0; + if (IsFirstPreload) { + Info.setKernargPreloadOffset(ArgDwordOffset / 4); + LastExplicitArgOffset = ArgDwordOffset; + } + // Arg is preloaded into the previous SGPR. - if (ArgLoc.getLocVT().getStoreSize() < 4 && Alignment < 4) { - assert(InIdx >= 1 && "No previous SGPR"); - Info.getArgInfo().PreloadKernArgs[InIdx].Regs.push_back( - Info.getArgInfo().PreloadKernArgs[InIdx - 1].Regs[0]); + if (!IsFirstPreload && ArgLoc.getLocVT().getStoreSize() < 4 && + Alignment < 4 && ArgDwordOffset < LastExplicitArgOffset) { + assert(PartIdx >= 1 && "No previous SGPR"); + Info.getArgInfo().PreloadKernArgs[PartIdx].Regs.push_back( + Info.getArgInfo().PreloadKernArgs[PartIdx - 1].Regs[0]); continue; } - unsigned Padding = ArgOffset - LastExplicitArgOffset; - unsigned PaddingSGPRs = alignTo(Padding, 4) / 4; + unsigned PaddingSGPRs = (ArgDwordOffset - LastExplicitArgOffset) / 4; // Check for free user SGPRs for preloading. if (PaddingSGPRs + NumAllocSGPRs > SGPRInfo.getNumFreeUserSGPRs()) { + if (HasPreloadRange) { + F.getContext().emitError( + "amdgpu kernarg preload range exceeds available user SGPRs"); + } InPreloadSequence = false; break; } @@ -3184,8 +3209,8 @@ void SITargetLowering::allocatePreloadKernArgSGPRs( // Preload this argument. const TargetRegisterClass *RC = TRI.getSGPRClassForBitWidth(NumAllocSGPRs * 32); - SmallVectorImpl<MCRegister> *PreloadRegs = - Info.addPreloadedKernArg(TRI, RC, NumAllocSGPRs, InIdx, PaddingSGPRs); + SmallVectorImpl<MCRegister> *PreloadRegs = Info.addPreloadedKernArg( + TRI, RC, NumAllocSGPRs, PartIdx, PaddingSGPRs); if (PreloadRegs->size() > 1) RC = &AMDGPU::SGPR_32RegClass; @@ -3195,7 +3220,7 @@ void SITargetLowering::allocatePreloadKernArgSGPRs( CCInfo.AllocateReg(Reg); } - LastExplicitArgOffset = NumAllocSGPRs * 4 + ArgOffset; + LastExplicitArgOffset = ArgDwordOffset + NumAllocSGPRs * 4; } } } diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp index 59971923e4a5d..298fb68c11118 100644 --- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp @@ -756,6 +756,7 @@ yaml::SIMachineFunctionInfo::SIMachineFunctionInfo( DynamicVGPRBlockSize(MFI.getDynamicVGPRBlockSize()), ScratchReservedForDynamicVGPRs(MFI.getScratchReservedForDynamicVGPRs()), NumKernargPreloadSGPRs(MFI.getNumKernargPreloadedSGPRs()), + KernargPreloadOffset(MFI.getKernargPreloadOffset()), MinNumAGPRs(MFI.getMinNumAGPRs()) { for (Register Reg : MFI.getSGPRSpillPhysVGPRs()) SpillPhysVGPRS.push_back(regToString(Reg, TRI)); @@ -810,6 +811,7 @@ bool SIMachineFunctionInfo::initializeBaseYamlFields( DynamicVGPRBlockSize = *YamlMFI.DynamicVGPRBlockSize; UserSGPRInfo.allocKernargPreloadSGPRs(YamlMFI.NumKernargPreloadSGPRs); + KernargPreloadOffset = YamlMFI.KernargPreloadOffset; if (YamlMFI.ScavengeFI) { auto FIOrErr = YamlMFI.ScavengeFI->getFI(MF.getFrameInfo()); diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h index 7374e996837a5..326eefd0e5d0e 100644 --- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h +++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h @@ -308,6 +308,7 @@ struct SIMachineFunctionInfo final : public yaml::MachineFunctionInfo { unsigned ScratchReservedForDynamicVGPRs = 0; unsigned NumKernargPreloadSGPRs = 0; + unsigned KernargPreloadOffset = 0; unsigned MinNumAGPRs = ~0u; @@ -367,6 +368,7 @@ template <> struct MappingTraits<SIMachineFunctionInfo> { YamlIO.mapOptional("scratchReservedForDynamicVGPRs", MFI.ScratchReservedForDynamicVGPRs, 0); YamlIO.mapOptional("numKernargPreloadSGPRs", MFI.NumKernargPreloadSGPRs, 0); + YamlIO.mapOptional("kernargPreloadOffset", MFI.KernargPreloadOffset, 0); YamlIO.mapOptional("isWholeWaveFunction", MFI.IsWholeWaveFunction, false); YamlIO.mapOptional("minNumAGPRs", MFI.MinNumAGPRs, ~0u); } @@ -500,6 +502,9 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunctionInfo, // will apply to all wavefronts of the grid. GCNUserSGPRUsageInfo UserSGPRInfo; + // Dword offset in the kernel argument segment where preloading starts. + unsigned KernargPreloadOffset = 0; + // Feature bits required for inputs passed in system SGPRs. bool WorkGroupIDX : 1; // Always initialized. bool WorkGroupIDY : 1; @@ -1022,6 +1027,11 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunctionInfo, return UserSGPRInfo.getNumKernargPreloadSGPRs(); } + unsigned getKernargPreloadOffset() const { return KernargPreloadOffset; } + void setKernargPreloadOffset(unsigned Offset) { + KernargPreloadOffset = Offset; + } + unsigned getNumWaveDispatchSGPRs() const { return NumWaveDispatchSGPRs; } void setNumWaveDispatchSGPRs(unsigned Count) { NumWaveDispatchSGPRs = Count; } diff --git a/llvm/test/CodeGen/AMDGPU/mir-kernarg-preload-offset-roundtrip.mir b/llvm/test/CodeGen/AMDGPU/mir-kernarg-preload-offset-roundtrip.mir new file mode 100644 index 0000000000000..a08c1a1418094 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/mir-kernarg-preload-offset-roundtrip.mir @@ -0,0 +1,20 @@ +# RUN: llc -mtriple=amdgpu9.42-amd-amdhsa -run-pass=none -o - %s | FileCheck %s + +--- | + define amdgpu_kernel void @kernarg_preload_offset() { + ret void + } +... +--- +name: kernarg_preload_offset +machineFunctionInfo: + numKernargPreloadSGPRs: 1 + kernargPreloadOffset: 2 +body: | + bb.0: + S_ENDPGM 0 +... + +# CHECK-LABEL: name: kernarg_preload_offset +# CHECK: numKernargPreloadSGPRs: 1 +# CHECK: kernargPreloadOffset: 2 diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs-range-attr.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs-range-attr.ll new file mode 100644 index 0000000000000..b1849b9e66020 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs-range-attr.ll @@ -0,0 +1,22 @@ +; 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 @middle_range(ptr addrspace(1) %a, i32 %b, ptr addrspace(1) %c, i32 %d) #0 { +; CHECK-LABEL: define amdgpu_kernel void @middle_range( +; CHECK-SAME: ptr addrspace(1) %a, i32 inreg %b, ptr addrspace(1) inreg %c, i32 %d) + ret void +} + +define amdgpu_kernel void @single_argument(i32 %a, i32 %b, i32 %c) #1 { +; CHECK-LABEL: define amdgpu_kernel void @single_argument( +; CHECK-SAME: i32 %a, i32 %b, i32 inreg %c) + ret void +} + +attributes #0 = { "amdgpu-kernarg-preload-first-arg"="1" "amdgpu-kernarg-preload-last-arg"="2" } +attributes #1 = { "amdgpu-kernarg-preload-first-arg"="2" "amdgpu-kernarg-preload-last-arg"="2" } diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs-range-errors.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs-range-errors.ll new file mode 100644 index 0000000000000..12fdc002223fa --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs-range-errors.ll @@ -0,0 +1,38 @@ +; RUN: not opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -passes=amdgpu-preload-kernel-arguments -disable-output < %s 2>&1 | FileCheck %s + +; CHECK: error: invalid amdgpu kernarg preload argument range +define amdgpu_kernel void @invalid_range(i32 %arg) #0 { + ret void +} + +; CHECK: error: incomplete amdgpu kernarg preload range +define amdgpu_kernel void @incomplete_range(i32 %arg) #1 { + ret void +} + +; CHECK: error: amdgpu kernarg preload range exceeds available user SGPRs +define amdgpu_kernel void @range_too_large(i512 %arg) #2 { + ret void +} + +; CHECK: error: amdgpu kernarg preload offset exceeds the hardware limit +define amdgpu_kernel void @offset_too_large([2048 x i8] %unused, i32 %arg) #3 { + ret void +} + +; CHECK: error: unsupported argument in amdgpu kernarg preload range +define amdgpu_kernel void @unsupported_aggregate([2 x i32] %arg) #4 { + ret void +} + +; CHECK: error: amdgpu kernarg preload argument index exceeds the 32-bit limit +define amdgpu_kernel void @index_too_large(i32 %arg) #5 { + ret void +} + +attributes #0 = { "amdgpu-kernarg-preload-first-arg"="1" "amdgpu-kernarg-preload-last-arg"="0" } +attributes #1 = { "amdgpu-kernarg-preload-first-arg"="0" } +attributes #2 = { "amdgpu-kernarg-preload-first-arg"="0" "amdgpu-kernarg-preload-last-arg"="0" } +attributes #3 = { "amdgpu-kernarg-preload-first-arg"="1" "amdgpu-kernarg-preload-last-arg"="1" } +attributes #4 = { "amdgpu-kernarg-preload-first-arg"="0" "amdgpu-kernarg-preload-last-arg"="0" } +attributes #5 = { "amdgpu-kernarg-preload-first-arg"="4294967296" "amdgpu-kernarg-preload-last-arg"="4294967296" } diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs-range.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs-range.ll new file mode 100644 index 0000000000000..1a6b7d03e16f3 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs-range.ll @@ -0,0 +1,32 @@ +; RUN: llc -mtriple=amdgpu9.42-amd-amdhsa -asm-verbose=0 < %s | FileCheck %s + +define amdgpu_kernel void @range_with_padding(i64 %unused, i32 inreg %hot0, i64 inreg %hot1, i32 %unused2) #0 { +; CHECK-LABEL: range_with_padding: +; CHECK: s_load_dwordx4 {{.*}}, 0x8 +; CHECK: .amdhsa_kernel range_with_padding +; CHECK: .amdhsa_user_sgpr_kernarg_preload_length 4 +; CHECK: .amdhsa_user_sgpr_kernarg_preload_offset 2 + ret void +} + +define amdgpu_kernel void @subdword_range_start(i8 %unused, i8 inreg %hot, ptr addrspace(1) %out) #1 { +; CHECK-LABEL: subdword_range_start: +; CHECK: s_load_dword {{.*}}, 0x0 +; CHECK: .amdhsa_kernel subdword_range_start +; CHECK: .amdhsa_user_sgpr_kernarg_preload_length 1 +; CHECK: .amdhsa_user_sgpr_kernarg_preload_offset 0 + store i8 %hot, ptr addrspace(1) %out + ret void +} + +define amdgpu_kernel void @range_after_byref(ptr addrspace(4) byref([8 x i32]) align 32 %unused, i32 inreg %hot) #1 { +; CHECK-LABEL: range_after_byref: +; CHECK: s_load_dword {{.*}}, 0x20 +; CHECK: .amdhsa_kernel range_after_byref +; CHECK: .amdhsa_user_sgpr_kernarg_preload_length 1 +; CHECK: .amdhsa_user_sgpr_kernarg_preload_offset 8 + ret void +} + +attributes #0 = { "amdgpu-kernarg-preload-first-arg"="1" "amdgpu-kernarg-preload-last-arg"="2" } +attributes #1 = { "amdgpu-kernarg-preload-first-arg"="1" "amdgpu-kernarg-preload-last-arg"="1" } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
