[llvm] [AMDGPU] Changed vector register calculation for gfx90a+ architecture by adding amdgpu-accum-offset attribute so that the number of available architectural VGPR can be limited (PR #218745)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 11:49:51 PDT 2026
https://github.com/JoshuaGrindstaff updated https://github.com/llvm/llvm-project/pull/218745
>From 2affcec22f15d70bdd4a1be7f6362389e467211a Mon Sep 17 00:00:00 2001
From: JoshuaGrindstaff <Joshua.Grindstaff at amd.com>
Date: Wed, 19 Aug 2026 18:05:04 -0500
Subject: [PATCH 1/2] [AMDGPU] Added Accum_Offset
Change-Id: Id9d0fdd3c560437d5254cd30f29efd4b91cd1276
---
llvm/docs/AMDGPUUsage.rst | 20 ++-
llvm/lib/Target/AMDGPU/GCNSubtarget.cpp | 26 ++-
...ee-registers-assertion-after-ra-failure.ll | 2 +-
.../AMDGPU/agpr-copy-no-free-registers.ll | 4 +-
.../AMDGPU/amdgpu-accum-offset-attr.ll | 149 ++++++++++++++++++
.../AMDGPU/amdgpu-no-agprs-violations.ll | 2 +-
llvm/test/CodeGen/AMDGPU/amdgpu-num-agpr.ll | 74 ++++-----
.../AMDGPU/amdgpu-prepare-agpr-alloc.mir | 4 +-
.../AMDGPU/copy-vgpr-clobber-spill-vgpr.mir | 2 +-
.../AMDGPU/large-avgpr-assign-last.mir | 2 +-
.../AMDGPU/preload-implicit-kernargs.ll | 2 +-
llvm/test/CodeGen/AMDGPU/preload-kernargs.ll | 2 +-
llvm/test/CodeGen/AMDGPU/smfmac_no_agprs.ll | 2 +-
.../CodeGen/AMDGPU/spill-regpressure-less.mir | 2 +-
.../CodeGen/AMDGPU/vgpr-agpr-limit-gfx90a.ll | 10 +-
15 files changed, 239 insertions(+), 64 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/amdgpu-accum-offset-attr.ll
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index e65baa0f48b78..d319b59b6debe 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -2673,8 +2673,8 @@ The AMDGPU backend supports the following LLVM IR attributes.
of the allocation granularity (4). The minimum value is interpreted as the
minimum required number of AGPRs for the function to allocate (that is, the
function requires no more than min registers). If only one value is specified,
- it is interpreted as the minimum register budget. The maximum will restrict
- allocation to use no more than max AGPRs.
+ it is interpreted as the minimum and maximum register budget. The maximum
+ will restrict allocation to use no more than max AGPRs.
The values may be ignored if satisfying it would violate other allocation
constraints.
@@ -2686,6 +2686,22 @@ The AMDGPU backend supports the following LLVM IR attributes.
This is only relevant on targets with AGPRs which support accum_offset (gfx90a+).
+ "amdgpu-accum-offset"="n" Bounds the number of architectural VGPRs the function may allocate,
+ and therefore the position of the boundary between architectural VGPRs
+ and AGPRs in the register file.
+
+ The value is rounded down (the inverse of amdgpu-agpr-alloc)
+ to a multiple of the allocation granularity (4) and clamped
+ to the addressable number of architectural VGPRs. Values below
+ the granularity are treated as 4, the smallest representable offset.
+ as 4, the smallest representable offset.
+
+ The behavior is undefined if a function which allocates more
+ architectural VGPRs than this bound is reached through any
+ function marked with a lower value of this attribute.
+
+ This is only relevant on targets with accum_offset (gfx90a+).
+
"amdgpu-sgpr-hazard-boundary-cull" Enable insertion of SGPR hazard cull sequences at function call boundaries.
Cull sequence reduces future hazard waits, but has a performance cost.
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
index 1c0e718bd8d97..ce9cece3909e0 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
@@ -659,39 +659,49 @@ GCNSubtarget::getMaxNumVectorRegs(const Function &F) const {
// VGPRs and AGPRs. Hence, in an entry function without calls and without
// AGPRs used within it, it is possible to use the whole vector register
// budget for VGPRs.
- //
- // TODO: it shall be possible to estimate maximum AGPR/VGPR pressure and split
- // register file accordingly.
if (hasGFX90AInsts()) {
unsigned MinNumAGPRs = 0;
+ unsigned VGPRCap = ~0u;
const unsigned TotalNumAGPRs = AMDGPU::AGPR_32RegClass.getNumRegs();
const std::pair<unsigned, unsigned> DefaultNumAGPR = {~0u, ~0u};
+ const unsigned DefaultAccumOffset = ~0u;
// TODO: The lower bound should probably force the number of required
// registers up, overriding amdgpu-waves-per-eu.
std::tie(MinNumAGPRs, MaxNumAGPRs) =
AMDGPU::getIntegerPairAttribute(F, "amdgpu-agpr-alloc", DefaultNumAGPR,
/*OnlyFirstRequired=*/true);
-
+ VGPRCap = F.getFnAttributeAsParsedInteger("amdgpu-accum-offset",DefaultAccumOffset);
+ if (VGPRCap == DefaultAccumOffset) {
+ MaxNumVGPRs = MaxVectorRegs / 2;
+ }
+ else {
+ MaxNumVGPRs = alignDown(VGPRCap, 4);
+ }
if (MinNumAGPRs == DefaultNumAGPR.first) {
- // Default to splitting half the registers if AGPRs are required.
- MinNumAGPRs = MaxNumAGPRs = MaxVectorRegs / 2;
+ MaxNumAGPRs = MaxVectorRegs / 2;
+ MinNumAGPRs = 0;
} else {
// Align to accum_offset's allocation granularity.
MinNumAGPRs = alignTo(MinNumAGPRs, 4);
MinNumAGPRs = std::min(MinNumAGPRs, TotalNumAGPRs);
+ // Since we can't know for sure how many availible AGPRs we have, an unset MaxNumAGPRs does not imply that we can use MaxVectorRegs - MaxNumVGPRs number of AGPRs
+ if (MaxNumAGPRs == DefaultNumAGPR.second && !AMDGPU::isEntryFunctionCC(F.getCallingConv()))
+ MaxNumAGPRs = MinNumAGPRs;
}
-
// Clamp values to be inbounds of our limits, and ensure min <= max.
MaxNumAGPRs = std::min(std::max(MinNumAGPRs, MaxNumAGPRs), MaxVectorRegs);
MinNumAGPRs = std::min({MinNumAGPRs, TotalNumAGPRs, MaxNumAGPRs});
- MaxNumVGPRs = std::min(MaxVectorRegs - MinNumAGPRs, NumArchVGPRs);
+ MaxNumVGPRs =
+ std::min({MaxVectorRegs - MinNumAGPRs, NumArchVGPRs, MaxNumVGPRs});
MaxNumAGPRs = std::min(MaxVectorRegs - MaxNumVGPRs, MaxNumAGPRs);
+ LLVM_DEBUG(dbgs() << "MaxNumVGPRs: " << MaxNumVGPRs << ", MaxNumAGPRs: "
+ << MaxNumAGPRs << " (" << F.getName() << ")\n");
assert(MaxNumVGPRs + MaxNumAGPRs <= MaxVectorRegs &&
MaxNumAGPRs <= TotalNumAGPRs && MaxNumVGPRs <= NumArchVGPRs &&
"invalid register counts");
diff --git a/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers-assertion-after-ra-failure.ll b/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers-assertion-after-ra-failure.ll
index d35f04c64ac89..5402106b7e01e 100644
--- a/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers-assertion-after-ra-failure.ll
+++ b/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers-assertion-after-ra-failure.ll
@@ -17,6 +17,6 @@ define void @no_free_vgprs_at_agpr_to_agpr_copy(float %v0, float %v1) #0 {
declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32 immarg, i32 immarg, i32 immarg) #1
declare noundef i32 @llvm.amdgcn.workitem.id.x() #2
-attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-waves-per-eu"="6,6" }
+attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" "amdgpu-waves-per-eu"="6,6" }
attributes #1 = { convergent nocallback nofree nosync nounwind willreturn memory(none) }
attributes #2 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
diff --git a/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll b/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll
index 41c3e0a07f5f5..4d547b70c8cc2 100644
--- a/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll
+++ b/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll
@@ -1161,6 +1161,6 @@ declare i32 @llvm.amdgcn.workitem.id.x() #2
attributes #0 = { "amdgpu-waves-per-eu"="6,6" nounwind }
attributes #1 = { convergent nounwind readnone willreturn }
attributes #2 = { nounwind readnone willreturn }
-attributes #3 = { "amdgpu-waves-per-eu"="7,7" "amdgpu-agpr-alloc"="0" nounwind }
+attributes #3 = { "amdgpu-waves-per-eu"="7,7" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" nounwind }
attributes #4 = { "amdgpu-waves-per-eu"="6,6" "amdgpu-flat-work-group-size"="1024,1024" nounwind }
-attributes #5 = { "amdgpu-waves-per-eu"="6,6" "amdgpu-agpr-alloc"="0" nounwind }
+attributes #5 = { "amdgpu-waves-per-eu"="6,6" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-accum-offset-attr.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-accum-offset-attr.ll
new file mode 100644
index 0000000000000..4ec856a2158d6
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-accum-offset-attr.ll
@@ -0,0 +1,149 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -filetype=null %s 2>&1 | FileCheck --implicit-check-not=warning -check-prefix=WARN %s
+
+; Consumer test for the "amdgpu-accum-offset" register-file split
+; cap applied in GCNSubtarget::getMaxNumVectorRegs.
+
+;no amdgpu-agrp-alloc or amdgpu- accum-offset
+; Budget 128, VGPR ceiling 64, AGPR ceiling 64
+
+; no amdgpu-agpr-alloc attributes
+; Budget 64, VGPR ceiling 48, AGPR Ceiling 16
+; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'split_96_32': desired occupancy was 8, final occupancy is 7
+; WARN: warning: inline asm clobber list contains reserved registers: v48 at line 12
+; WARN: warning: inline asm clobber list contains reserved registers: a16 at line 14
+define amdgpu_kernel void @split_96_32() #0 {
+ call void asm sideeffect "; c $0","~{v47}"(), !srcloc !{i32 11}
+ call void asm sideeffect "; c $0","~{v48}"(), !srcloc !{i32 12}
+ call void asm sideeffect "; c $0","~{a15}"(), !srcloc !{i32 13}
+ call void asm sideeffect "; c $0","~{a16}"(), !srcloc !{i32 14}
+ ret void
+}
+attributes #0 = {"amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-accum-offset"="48" }
+
+; Budget 64, VGPR ceiling 32, AGPR ceiling 16.
+; WARN: warning: inline asm clobber list contains reserved registers: v32 at line 22
+; WARN: warning: inline asm clobber list contains reserved registers: a16 at line 24
+define amdgpu_kernel void @split_32_16() #1 {
+ call void asm sideeffect "; c $0","~{v31}"(), !srcloc !{i32 21}
+ call void asm sideeffect "; c $0","~{v32}"(), !srcloc !{i32 22}
+ call void asm sideeffect "; c $0","~{a15}"(), !srcloc !{i32 23}
+ call void asm sideeffect "; c $0","~{a16}"(), !srcloc !{i32 24}
+ ret void
+}
+attributes #1 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="16,16" "amdgpu-accum-offset"="32" }
+
+; Budget 64, VGPR ceiling 24, AGPR ceiling 8.
+; WARN: warning: inline asm clobber list contains reserved registers: v24 at line 32
+; WARN: warning: inline asm clobber list contains reserved registers: a8 at line 34
+define amdgpu_kernel void @split_24_8() #2 {
+ call void asm sideeffect "; c $0","~{v23}"(), !srcloc !{i32 31}
+ call void asm sideeffect "; c $0","~{v24}"(), !srcloc !{i32 32}
+ call void asm sideeffect "; c $0","~{a7}"(), !srcloc !{i32 33}
+ call void asm sideeffect "; c $0","~{a8}"(), !srcloc !{i32 34}
+ ret void
+}
+attributes #2 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="8,8" "amdgpu-accum-offset"="24" }
+
+; Budget 64, VGPR ceiling 16, AGPR ceiling 4
+; WARN: warning: inline asm clobber list contains reserved registers: v16 at line 42
+; WARN: warning: inline asm clobber list contains reserved registers: a4 at line 44
+define amdgpu_kernel void @split_16_32() #3 {
+ call void asm sideeffect "; c $0","~{v15}"(), !srcloc !{i32 41}
+ call void asm sideeffect "; c $0","~{v16}"(), !srcloc !{i32 42}
+ call void asm sideeffect "; c $0","~{a3}"(), !srcloc !{i32 43}
+ call void asm sideeffect "; c $0","~{a4}"(), !srcloc !{i32 44}
+ ret void
+}
+attributes #3 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="4,4" "amdgpu-accum-offset"="16" }
+
+; Budget 64, VGPR ceiling 24, AGPR ceiling 32
+; WARN: warning: inline asm clobber list contains reserved registers: v24 at line 52
+; WARN: warning: inline asm clobber list contains reserved registers: a32 at line 54
+define amdgpu_kernel void @split_24_40() #4 {
+ call void asm sideeffect "; c $0","~{v23}"(), !srcloc !{i32 51}
+ call void asm sideeffect "; c $0","~{v24}"(), !srcloc !{i32 52}
+ call void asm sideeffect "; c $0","~{a31}"(), !srcloc !{i32 53}
+ call void asm sideeffect "; c $0","~{a32}"(), !srcloc !{i32 54}
+ ret void
+}
+attributes #4 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-accum-offset"="24" }
+
+; Budget 128, VGPR ceiling 64, AGPR ceiling 64.
+; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'pessimistic_split_64_64': desired occupancy was 4, final occupancy is 3
+; WARN: warning: inline asm clobber list contains reserved registers: v64 at line 62
+; WARN: warning: inline asm clobber list contains reserved registers: a64 at line 64
+define amdgpu_kernel void @pessimistic_split_64_64() #5 {
+ call void asm sideeffect "; c $0","~{v63}"(), !srcloc !{i32 61}
+ call void asm sideeffect "; c $0","~{v64}"(), !srcloc !{i32 62}
+ call void asm sideeffect "; c $0","~{a63}"(), !srcloc !{i32 63}
+ call void asm sideeffect "; c $0","~{a64}"(), !srcloc !{i32 64}
+ ret void
+}
+attributes #5 = {"amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="256,256"}
+
+; absent amdgpu-accum-offset attribute
+; Budget 128, VGPR ceiling 64, AGPR ceiling 16
+; WARN: warning: inline asm clobber list contains reserved registers: v64 at line 72
+; WARN: warning: inline asm clobber list contains reserved registers: a16 at line 74
+define amdgpu_kernel void @absent_accum_offset_64_16() #6 {
+ call void asm sideeffect "; c $0","~{v63}"(), !srcloc !{i32 71}
+ call void asm sideeffect "; c $0","~{v64}"(), !srcloc !{i32 72}
+ call void asm sideeffect "; c $0","~{a15}"(), !srcloc !{i32 73}
+ call void asm sideeffect "; c $0","~{a16}"(), !srcloc !{i32 74}
+ ret void
+}
+
+attributes #6 = {"amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="256,256" "amdgpu-agpr-alloc"="16,16" }
+
+; Budget 128, VGPR ceiling 64, AGPR ceiling 32
+; WARN: warning: inline asm clobber list contains reserved registers: v64 at line 82
+; WARN: warning: inline asm clobber list contains reserved registers: a32 at line 84
+define amdgpu_kernel void @agpr_alloc_16_16_64_32() #7 {
+ call void asm sideeffect "; c $0","~{v63}"(), !srcloc !{i32 81}
+ call void asm sideeffect "; c $0","~{v64}"(), !srcloc !{i32 82}
+ call void asm sideeffect "; c $0","~{a31}"(), !srcloc !{i32 83}
+ call void asm sideeffect "; c $0","~{a32}"(), !srcloc !{i32 84}
+ ret void
+}
+attributes #7 = {"amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="256,256" "amdgpu-agpr-alloc"="16,32" }
+
+; Budget 128, VGPR ceiling 56, AGPR ceiling 72
+; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'agpr_alloc_70_58_70': desired occupancy was 4, final occupancy is 3
+; WARN: warning: inline asm clobber list contains reserved registers: v56 at line 92
+; WARN: warning: inline asm clobber list contains reserved registers: a72 at line 94
+define amdgpu_kernel void @agpr_alloc_70_58_70() #8 {
+ call void asm sideeffect "; c $0","~{v55}"(), !srcloc !{i32 91}
+ call void asm sideeffect "; c $0","~{v56}"(), !srcloc !{i32 92}
+ call void asm sideeffect "; c $0","~{a61}"(), !srcloc !{i32 93}
+ call void asm sideeffect "; c $0","~{a72}"(), !srcloc !{i32 94}
+ ret void
+}
+attributes #8 = {"amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="256,256" "amdgpu-agpr-alloc"="72" }
+
+;more restraining attribute takes priority over amdgpu-accum-offset attribute
+; Budget 128, VGPR ceiling 80, AGPR Ceiling 48
+; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'agpr_alloc_48_80_48': desired occupancy was 4, final occupancy is 3
+; WARN: warning: inline asm clobber list contains reserved registers: v80 at line 102
+; WARN: warning: inline asm clobber list contains reserved registers: a48 at line 104
+define amdgpu_kernel void @agpr_alloc_48_80_48() #9 {
+ call void asm sideeffect "; c $0","~{v79}"(), !srcloc !{i32 101}
+ call void asm sideeffect "; c $0","~{v80}"(), !srcloc !{i32 102}
+ call void asm sideeffect "; c $0","~{a47}"(), !srcloc !{i32 103}
+ call void asm sideeffect "; c $0","~{a48}"(), !srcloc !{i32 104}
+ ret void
+}
+attributes #9 = {"amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="256,256" "amdgpu-accum-offset"="100" "amdgpu-agpr-alloc"="48" }
+
+;lower amdgpu-accum-offset attribute could allow agpr more space
+; Budget 128, VGPR ceiling 100, AGPR ceiling 26
+; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'agpr_alloc_10_26_100_26': desired occupancy was 4, final occupancy is 3
+; WARN: warning: inline asm clobber list contains reserved registers: v100 at line 112
+; WARN: warning: inline asm clobber list contains reserved registers: a26 at line 114
+define amdgpu_kernel void @agpr_alloc_10_26_100_26() #10 {
+ call void asm sideeffect "; c $0","~{v99}"(), !srcloc !{i32 111}
+ call void asm sideeffect "; c $0","~{v100}"(), !srcloc !{i32 112}
+ call void asm sideeffect "; c $0","~{a25}"(), !srcloc !{i32 113}
+ call void asm sideeffect "; c $0","~{a26}"(), !srcloc !{i32 114}
+ ret void
+}
+attributes #10 = {"amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="256,256" "amdgpu-accum-offset"="100" "amdgpu-agpr-alloc"="10,26" }
\ No newline at end of file
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-no-agprs-violations.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-no-agprs-violations.ll
index c8f75373a55dd..4d80978090124 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-no-agprs-violations.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-no-agprs-violations.ll
@@ -47,4 +47,4 @@ define amdgpu_kernel void @kernel_calls_mfma.f32.32x32x1f32(ptr addrspace(1) %ou
ret void
}
-attributes #0 = { "amdgpu-agpr-alloc"="0" }
+attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-num-agpr.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-num-agpr.ll
index bb4f2a68b182a..facec987d05bc 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-num-agpr.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-num-agpr.ll
@@ -15,7 +15,7 @@ define amdgpu_kernel void @min_num_agpr_0_0__amdgpu_no_agpr() #0 {
ret void
}
-attributes #0 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,0" }
+attributes #0 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,0" "amdgpu-accum-offset"="256" }
; Check parse of single entry 0
@@ -26,7 +26,7 @@ define amdgpu_kernel void @min_num_agpr_0__amdgpu_no_agpr() #1 {
ret void
}
-attributes #1 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0" }
+attributes #1 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
; Undefined use
@@ -35,7 +35,7 @@ define amdgpu_kernel void @min_num_agpr_1_1() #2 {
ret void
}
-attributes #2 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="1,1" }
+attributes #2 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="1,1" "amdgpu-accum-offset"="256" }
; Check parse of single entry 4, interpreted as the minimum. Total budget is 64.
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_4__amdgpu_no_agpr': desired occupancy was 8, final occupancy is 7
@@ -48,7 +48,7 @@ define amdgpu_kernel void @min_num_agpr_4__amdgpu_no_agpr() #3 {
ret void
}
-attributes #3 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="4" }
+attributes #3 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="4" "amdgpu-accum-offset"="256" }
; Allocation granularity requires rounding this to use 4 AGPRs, so the
@@ -68,7 +68,7 @@ define amdgpu_kernel void @min_num_agpr_occ8_1_1() #4 {
ret void
}
-attributes #4 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="1,1" }
+attributes #4 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="1,1" "amdgpu-accum-offset"="256" }
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_64_64__amdgpu_no_agpr': desired occupancy was 8, final occupancy is 7
@@ -79,7 +79,7 @@ define amdgpu_kernel void @min_num_agpr_64_64__amdgpu_no_agpr() #5 {
ret void
}
-attributes #5 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="64,64" }
+attributes #5 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="64,64" "amdgpu-accum-offset"="256" }
; No free VGPRs
; WARN: warning: inline asm clobber list contains reserved registers: v0 at line 7
@@ -89,7 +89,7 @@ define amdgpu_kernel void @min_num_agpr_64_64() #6 {
ret void
}
-attributes #6 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="64,64" }
+attributes #6 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="64,64" "amdgpu-accum-offset"="256" }
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_63_64': desired occupancy was 8, final occupancy is 7
; WARN: warning: inline asm clobber list contains reserved registers: v0 at line 8
@@ -103,7 +103,7 @@ define amdgpu_kernel void @min_num_agpr_63_64() #7 {
ret void
}
-attributes #7 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="63,64" }
+attributes #7 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="63,64" "amdgpu-accum-offset"="256" }
; No-op value.
@@ -113,7 +113,7 @@ define amdgpu_kernel void @min_num_agpr_occ8_0_64() #8 {
ret void
}
-attributes #8 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,64" }
+attributes #8 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,64" "amdgpu-accum-offset"="256" }
; Register budget 64
@@ -128,7 +128,7 @@ define amdgpu_kernel void @min_num_agpr_occ8_11_59() #9 {
ret void
}
-attributes #9 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="11,59" }
+attributes #9 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="11,59" "amdgpu-accum-offset"="256" }
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ8_12_59': desired occupancy was 8, final occupancy is 7
@@ -142,7 +142,7 @@ define amdgpu_kernel void @min_num_agpr_occ8_12_59() #10 {
ret void
}
-attributes #10 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="12,59" }
+attributes #10 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="12,59" "amdgpu-accum-offset"="256" }
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ8_12_20': desired occupancy was 8, final occupancy is 7
@@ -156,7 +156,7 @@ define amdgpu_kernel void @min_num_agpr_occ8_12_20() #11 {
ret void
}
-attributes #11 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="12,20" }
+attributes #11 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="12,20" "amdgpu-accum-offset"="256" }
; WARN: warning: inline asm clobber list contains reserved registers: a20 at line 13
@@ -170,7 +170,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_12_20() #12 {
ret void
}
-attributes #12 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="12,20" }
+attributes #12 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="12,20" "amdgpu-accum-offset"="256" }
; WARN: warning: inline asm clobber list contains reserved registers: a20 at line 14
define amdgpu_kernel void @min_num_agpr_occ1_13_20() #13 {
@@ -185,7 +185,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_13_20() #13 {
ret void
}
-attributes #13 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="13,20" }
+attributes #13 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="13,20" "amdgpu-accum-offset"="256" }
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ2_13_20': desired occupancy was 2, final occupancy is 1
@@ -203,7 +203,7 @@ define amdgpu_kernel void @min_num_agpr_occ2_13_20() #14 {
ret void
}
-attributes #14 = { "amdgpu-waves-per-eu"="2,2" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="13,20" }
+attributes #14 = { "amdgpu-waves-per-eu"="2,2" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="13,20" "amdgpu-accum-offset"="256" }
; Test maximum exceeds the hardware limit.
@@ -213,7 +213,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_13_257() #15 {
ret void
}
-attributes #15 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="13,257" }
+attributes #15 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="13,257" "amdgpu-accum-offset"="256" }
; Test min and max exceeds the hardware limit.
@@ -223,7 +223,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_257_257() #16 {
ret void
}
-attributes #16 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="257,257" }
+attributes #16 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="257,257" "amdgpu-accum-offset"="256" }
; Test round up hits the hardware limit
@@ -233,7 +233,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_255_255() #17 {
ret void
}
-attributes #17 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="255,255" }
+attributes #17 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="255,255" "amdgpu-accum-offset"="256" }
; Test round up hits the hardware limit
@@ -243,7 +243,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_253_259() #18 {
ret void
}
-attributes #18 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="253,259" }
+attributes #18 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="253,259" "amdgpu-accum-offset"="256" }
; With a minimum of 0, we are not required to allocate any AGPRs
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ4_0_64': desired occupancy was 4, final occupancy is 2
@@ -262,7 +262,7 @@ define amdgpu_kernel void @min_num_agpr_occ4_0_64() #19 {
ret void
}
-attributes #19 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,64" }
+attributes #19 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,64" "amdgpu-accum-offset"="256" }
; With a non-0 minimum, we must allocate at least 4 AGPRs. The rest of
@@ -285,7 +285,7 @@ define amdgpu_kernel void @min_num_agpr_occ4_1_64() #20 {
ret void
}
-attributes #20 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="1,64" }
+attributes #20 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="1,64" "amdgpu-accum-offset"="256" }
; 128 vector registers
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ4_32_64': desired occupancy was 4, final occupancy is 3
@@ -299,7 +299,7 @@ define amdgpu_kernel void @min_num_agpr_occ4_32_64() #21 {
ret void
}
-attributes #21 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="32,64" }
+attributes #21 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="32,64" "amdgpu-accum-offset"="256" }
; Evenly partition the 128 vector registers
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ4_64_64': desired occupancy was 4, final occupancy is 3
@@ -313,7 +313,7 @@ define amdgpu_kernel void @min_num_agpr_occ4_64_64() #22 {
ret void
}
-attributes #22 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="64,64" }
+attributes #22 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="64,64" "amdgpu-accum-offset"="256" }
; We are not required to allocate any AGPRs, but they are available
; with a budget of 512 vector registers. We are artificially limiting
@@ -327,7 +327,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_0_64() #23 {
ret void
}
-attributes #23 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,64" }
+attributes #23 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,64" "amdgpu-accum-offset"="256" }
; WARN: warning: inline asm clobber list contains reserved registers: a68 at line 25
define amdgpu_kernel void @min_num_agpr_occ1_0_68() #24 {
@@ -337,7 +337,7 @@ define amdgpu_kernel void @min_num_agpr_occ1_0_68() #24 {
ret void
}
-attributes #24 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,68" }
+attributes #24 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" "amdgpu-agpr-alloc"="0,68" "amdgpu-accum-offset"="256" }
; The total vector register budget is 128, claim more than that for
@@ -352,7 +352,7 @@ define amdgpu_kernel void @min_num_agpr_occ10__min_agpr_129() #25 {
ret void
}
-attributes #25 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="129" }
+attributes #25 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="129" "amdgpu-accum-offset"="256" }
; Check for another assertion, request beyond the budget.
@@ -366,7 +366,7 @@ define amdgpu_kernel void @min_num_agpr_occ10__min_agpr_129_129() #26 {
ret void
}
-attributes #26 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="129,129" }
+attributes #26 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="129,129" "amdgpu-accum-offset"="256" }
; The total vector register budget is 128, claim all of it for AGPRs.
@@ -381,7 +381,7 @@ define amdgpu_kernel void @min_num_agpr_occ10__min_agpr_128() #27 {
ret void
}
-attributes #27 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="128" }
+attributes #27 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="128" "amdgpu-accum-offset"="256" }
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ10__min_agpr_257': desired occupancy was 8, final occupancy is 3
; WARN: warning: inline asm clobber list contains reserved registers: a128 at line 29
@@ -393,7 +393,7 @@ define amdgpu_kernel void @min_num_agpr_occ10__min_agpr_257() #28 {
ret void
}
-attributes #28 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="257" }
+attributes #28 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="257" "amdgpu-accum-offset"="256" }
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ10__min_agpr_257_257': desired occupancy was 8, final occupancy is 3
; WARN: warning: inline asm clobber list contains reserved registers: a128 at line 30
@@ -405,7 +405,7 @@ define amdgpu_kernel void @min_num_agpr_occ10__min_agpr_257_257() #29 {
ret void
}
-attributes #29 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="257,257" }
+attributes #29 = { "amdgpu-waves-per-eu"="8,10" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="257,257" "amdgpu-accum-offset"="256" }
; The total vector register budget is 96
@@ -421,7 +421,7 @@ define amdgpu_kernel void @min_num_agpr_occ5__min_agpr_8_256() #30 {
ret void
}
-attributes #30 = { "amdgpu-waves-per-eu"="5,5" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="8,256" }
+attributes #30 = { "amdgpu-waves-per-eu"="5,5" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="8,256" "amdgpu-accum-offset"="256" }
; The total vector register budget is 96
; WARN: warning: <unknown>:0:0: failed to meet occupancy target given by 'amdgpu-waves-per-eu' in 'min_num_agpr_occ5__min_agpr_8': desired occupancy was 5, final occupancy is 4
@@ -443,7 +443,7 @@ define amdgpu_kernel void @min_num_agpr_occ5__min_agpr_8_no_agpr_references() #3
ret void
}
-attributes #31 = { "amdgpu-waves-per-eu"="5,5" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="8" }
+attributes #31 = { "amdgpu-waves-per-eu"="5,5" "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="8" "amdgpu-accum-offset"="256" }
; register budget 256
@@ -458,7 +458,7 @@ define amdgpu_kernel void @min_num_agpr_occ2__min_agpr_93() #33 {
ret void
}
-attributes #33 = { "amdgpu-waves-per-eu"="2,2" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="93" }
+attributes #33 = { "amdgpu-waves-per-eu"="2,2" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="93" "amdgpu-accum-offset"="256" }
; register budget 512, no warnings and fully allocated
define amdgpu_kernel void @min_num_agpr_occ1__min_agpr_93() #34 {
@@ -467,7 +467,7 @@ define amdgpu_kernel void @min_num_agpr_occ1__min_agpr_93() #34 {
ret void
}
-attributes #34 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="93" }
+attributes #34 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="93" "amdgpu-accum-offset"="256" }
; register budget 256
; WARN: warning: inline asm clobber list contains reserved registers: a96 at line 36
@@ -478,7 +478,7 @@ define amdgpu_kernel void @min_num_agpr_occ1__min_agpr_93_93() #35 {
ret void
}
-attributes #35 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="93,93" }
+attributes #35 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="93,93" "amdgpu-accum-offset"="256" }
; register budget 512, fully allocated and no warnings.
define amdgpu_kernel void @min_num_agpr_occ1__min_agpr_256() #36 {
@@ -487,7 +487,7 @@ define amdgpu_kernel void @min_num_agpr_occ1__min_agpr_256() #36 {
ret void
}
-attributes #36 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="256" }
+attributes #36 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="256" "amdgpu-accum-offset"="256" }
; register budget 512, fully allocated and no warnings.
define amdgpu_kernel void @min_num_agpr_occ1__min_agpr_256_256() #37 {
@@ -496,7 +496,7 @@ define amdgpu_kernel void @min_num_agpr_occ1__min_agpr_256_256() #37 {
ret void
}
-attributes #37 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="256,256" }
+attributes #37 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="256,256" "amdgpu-accum-offset"="256" }
; register budget 512, fully allocated and no warnings.
define amdgpu_kernel void @occ1_min_agpr_no_attr() #38 {
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-prepare-agpr-alloc.mir b/llvm/test/CodeGen/AMDGPU/amdgpu-prepare-agpr-alloc.mir
index 878abd30567df..074c8f480ca30 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-prepare-agpr-alloc.mir
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-prepare-agpr-alloc.mir
@@ -9,7 +9,7 @@
}
; Attribute is ignored for gfx90a
- define void @no_agprs() "amdgpu-agpr-alloc"="0,0" {
+ define void @no_agprs() "amdgpu-agpr-alloc"="0,0" "amdgpu-accum-offset"="256" {
ret void
}
@@ -17,7 +17,7 @@
ret void
}
- define void @func64_no_agprs() "amdgpu-agpr-alloc"="0,0" {
+ define void @func64_no_agprs() "amdgpu-agpr-alloc"="0,0" "amdgpu-accum-offset"="256" {
ret void
}
diff --git a/llvm/test/CodeGen/AMDGPU/copy-vgpr-clobber-spill-vgpr.mir b/llvm/test/CodeGen/AMDGPU/copy-vgpr-clobber-spill-vgpr.mir
index 0130b13c56d0e..fd66806fdf520 100644
--- a/llvm/test/CodeGen/AMDGPU/copy-vgpr-clobber-spill-vgpr.mir
+++ b/llvm/test/CodeGen/AMDGPU/copy-vgpr-clobber-spill-vgpr.mir
@@ -333,7 +333,7 @@
ret void
}
- attributes #0 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-agpr-alloc"="0" }
+ attributes #0 = { "amdgpu-waves-per-eu"="4,4" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
...
---
diff --git a/llvm/test/CodeGen/AMDGPU/large-avgpr-assign-last.mir b/llvm/test/CodeGen/AMDGPU/large-avgpr-assign-last.mir
index fcb8461963dbb..a7c484a38a1ac 100644
--- a/llvm/test/CodeGen/AMDGPU/large-avgpr-assign-last.mir
+++ b/llvm/test/CodeGen/AMDGPU/large-avgpr-assign-last.mir
@@ -7,7 +7,7 @@
unreachable
}
- attributes #0 = { "amdgpu-agpr-alloc"="0,0" }
+ attributes #0 = { "amdgpu-agpr-alloc"="0,0" "amdgpu-accum-offset"="256" }
...
diff --git a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
index db7f998270f16..03f6578d8f928 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-implicit-kernargs.ll
@@ -1152,7 +1152,7 @@ define amdgpu_kernel void @preload_block_count_z_workgroup_size_z_remainder_z(pt
ret void
}
-attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
+attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
; GFX1250-FAKE16: {{.*}}
; GFX1250-REAL16: {{.*}}
diff --git a/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll b/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll
index 0d817f6497f1a..f9b9c2e37cfdd 100644
--- a/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll
+++ b/llvm/test/CodeGen/AMDGPU/preload-kernargs.ll
@@ -1755,4 +1755,4 @@ define amdgpu_kernel void @ptr1_i8_trailing_unused(ptr addrspace(1) inreg %out,
ret void
}
-attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
+attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
diff --git a/llvm/test/CodeGen/AMDGPU/smfmac_no_agprs.ll b/llvm/test/CodeGen/AMDGPU/smfmac_no_agprs.ll
index a2c5db0d2646d..b333d852c1276 100644
--- a/llvm/test/CodeGen/AMDGPU/smfmac_no_agprs.ll
+++ b/llvm/test/CodeGen/AMDGPU/smfmac_no_agprs.ll
@@ -52,4 +52,4 @@ entry:
}
declare <4 x i32> @llvm.amdgcn.smfmac.i32.16x16x64.i8(<2 x i32>, <4 x i32>, <4 x i32>, i32, i32 immarg, i32 immarg)
-attributes #0 = { "amdgpu-agpr-alloc"="0" }
+attributes #0 = { "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
diff --git a/llvm/test/CodeGen/AMDGPU/spill-regpressure-less.mir b/llvm/test/CodeGen/AMDGPU/spill-regpressure-less.mir
index 403d24ab2e04f..0c73749d05641 100644
--- a/llvm/test/CodeGen/AMDGPU/spill-regpressure-less.mir
+++ b/llvm/test/CodeGen/AMDGPU/spill-regpressure-less.mir
@@ -6,7 +6,7 @@
ret void
}
- attributes #0 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-agpr-alloc"="0" }
+ attributes #0 = { "amdgpu-waves-per-eu"="8,8" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
...
---
diff --git a/llvm/test/CodeGen/AMDGPU/vgpr-agpr-limit-gfx90a.ll b/llvm/test/CodeGen/AMDGPU/vgpr-agpr-limit-gfx90a.ll
index 12ffc853188e7..3eb68944200e4 100644
--- a/llvm/test/CodeGen/AMDGPU/vgpr-agpr-limit-gfx90a.ll
+++ b/llvm/test/CodeGen/AMDGPU/vgpr-agpr-limit-gfx90a.ll
@@ -1082,7 +1082,7 @@ define amdgpu_kernel void @k256_w8_no_agprs() #2569 {
}
attributes #2568 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="8" }
-attributes #2569 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="8" "amdgpu-agpr-alloc"="0" }
+attributes #2569 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="8" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256"}
; GCN-LABEL: {{^}}k256_w4:
; GFX90A: NumVgprs: 64
@@ -1104,7 +1104,7 @@ define amdgpu_kernel void @k256_w4_no_agprs() #2565 {
}
attributes #2564 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="4" }
-attributes #2565 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="4" "amdgpu-agpr-alloc"="0" }
+attributes #2565 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="4" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
; GCN-LABEL: {{^}}k256_w2:
; GFX90A: NumVgprs: 128
@@ -1126,7 +1126,7 @@ define amdgpu_kernel void @k256_w2_no_agprs() #2563 {
}
attributes #2562 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="2" }
-attributes #2563 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="2" "amdgpu-agpr-alloc"="0" }
+attributes #2563 = { nounwind "amdgpu-flat-work-group-size"="256,256" "amdgpu-waves-per-eu"="2" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
; GCN-LABEL: {{^}}k256_w1:
; GFX90A: NumVgprs: 256
@@ -1213,7 +1213,7 @@ define void @f512_no_agpr_ub() #513 {
}
attributes #512 = { nounwind "amdgpu-flat-work-group-size"="512,512" }
-attributes #513 = { nounwind "amdgpu-flat-work-group-size"="512,512" "amdgpu-agpr-alloc"="0" }
+attributes #513 = { nounwind "amdgpu-flat-work-group-size"="512,512" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
; GCN-LABEL: {{^}}k1024:
; GFX90A: NumVgprs: 64
@@ -1300,4 +1300,4 @@ define void @f1024_call_no_agprs_ub() #1025 {
}
attributes #1024 = { nounwind "amdgpu-flat-work-group-size"="1024,1024" }
-attributes #1025 = { nounwind "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="0" }
+attributes #1025 = { nounwind "amdgpu-flat-work-group-size"="1024,1024" "amdgpu-agpr-alloc"="0" "amdgpu-accum-offset"="256" }
>From c5cc7eb90c7584d5ad647e01b6a11dd3c8a6be9d Mon Sep 17 00:00:00 2001
From: JoshuaGrindstaff <Joshua.Grindstaff at amd.com>
Date: Tue, 25 Aug 2026 13:38:21 -0500
Subject: [PATCH 2/2] fix formatting
---
llvm/lib/Target/AMDGPU/GCNSubtarget.cpp | 19 +++++++++++--------
1 file changed, 11 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
index ce9cece3909e0..b8de20094a48a 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
@@ -665,30 +665,33 @@ GCNSubtarget::getMaxNumVectorRegs(const Function &F) const {
const unsigned TotalNumAGPRs = AMDGPU::AGPR_32RegClass.getNumRegs();
const std::pair<unsigned, unsigned> DefaultNumAGPR = {~0u, ~0u};
- const unsigned DefaultAccumOffset = ~0u;
+ const unsigned DefaultAccumOffset = ~0u;
// TODO: The lower bound should probably force the number of required
// registers up, overriding amdgpu-waves-per-eu.
std::tie(MinNumAGPRs, MaxNumAGPRs) =
AMDGPU::getIntegerPairAttribute(F, "amdgpu-agpr-alloc", DefaultNumAGPR,
/*OnlyFirstRequired=*/true);
- VGPRCap = F.getFnAttributeAsParsedInteger("amdgpu-accum-offset",DefaultAccumOffset);
+ VGPRCap = F.getFnAttributeAsParsedInteger("amdgpu-accum-offset",
+ DefaultAccumOffset);
if (VGPRCap == DefaultAccumOffset) {
MaxNumVGPRs = MaxVectorRegs / 2;
- }
- else {
+ } else {
MaxNumVGPRs = alignDown(VGPRCap, 4);
}
if (MinNumAGPRs == DefaultNumAGPR.first) {
- MaxNumAGPRs = MaxVectorRegs / 2;
- MinNumAGPRs = 0;
+ MaxNumAGPRs = MaxVectorRegs / 2;
+ MinNumAGPRs = 0;
} else {
// Align to accum_offset's allocation granularity.
MinNumAGPRs = alignTo(MinNumAGPRs, 4);
MinNumAGPRs = std::min(MinNumAGPRs, TotalNumAGPRs);
- // Since we can't know for sure how many availible AGPRs we have, an unset MaxNumAGPRs does not imply that we can use MaxVectorRegs - MaxNumVGPRs number of AGPRs
- if (MaxNumAGPRs == DefaultNumAGPR.second && !AMDGPU::isEntryFunctionCC(F.getCallingConv()))
+ // Since we can't know for sure how many availible AGPRs we have, an unset
+ // MaxNumAGPRs does not imply that we can use MaxVectorRegs - MaxNumVGPRs
+ // number of AGPRs
+ if (MaxNumAGPRs == DefaultNumAGPR.second &&
+ !AMDGPU::isEntryFunctionCC(F.getCallingConv()))
MaxNumAGPRs = MinNumAGPRs;
}
// Clamp values to be inbounds of our limits, and ensure min <= max.
More information about the llvm-commits
mailing list