[llvm] [AMDGPU] Infer amdgpu-accum-offset from reachable kernels (PR #218763)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 16:59:01 PDT 2026
https://github.com/JoshuaGrindstaff updated https://github.com/llvm/llvm-project/pull/218763
>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/4] [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/4] 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.
>From a95259c7f51574b6e6c46cb6f0b4ed8193d4f109 Mon Sep 17 00:00:00 2001
From: JoshuaGrindstaff <Joshua.Grindstaff at amd.com>
Date: Tue, 18 Aug 2026 14:00:14 -0500
Subject: [PATCH 3/4] [AMDGPU] Add AMDGPURegistreBudget
Change-Id: Iccbc712f7e2ff659481bfb22f4a99aaab6a8ba2a
---
llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp | 205 ++++-
.../amdgpu-attributor-min-agpr-alloc.ll | 217 ++---
...pu-attributor-register-budget-propagate.ll | 814 ++++++++++++++++++
...ributor-register-budget-work-group-size.ll | 205 +++++
.../AMDGPU/amdgpu-attributor-trap-leaf.ll | 2 +-
5 files changed, 1333 insertions(+), 110 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-propagate.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-work-group-size.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
index 630ffad96e451..5b9f974e70345 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
@@ -1433,6 +1433,199 @@ struct AAAMDGPUMinAGPRAlloc
};
const char AAAMDGPUMinAGPRAlloc::ID = 0;
+/// The register-file split a function has been committed to by its callers.
+/// \p VGPRs and \p AGPRs are both absolute ceilings, so their sum stays within
+/// the total vector register budget of the tightest caller.
+struct RegisterBudgetState {
+ unsigned VGPRs = 0;
+ unsigned AGPRs = 0;
+ bool Unknown = true;
+
+ bool operator==(const RegisterBudgetState &Other) const {
+ return Unknown == Other.Unknown && VGPRs == Other.VGPRs &&
+ AGPRs == Other.AGPRs;
+ }
+ bool operator!=(const RegisterBudgetState &Other) const {
+ return !(*this == Other);
+ }
+
+ /// Combine with the budget of one caller. A function reachable from several
+ /// callers has to fit inside all of their splits, so take the tightest
+ /// ceiling on each axis independently.
+ void merge(const RegisterBudgetState &Other) {
+ assert(!Other.Unknown && "cannot merge an unknown budget");
+ if (Unknown) {
+ *this = Other;
+ return;
+ }
+ VGPRs = std::min(VGPRs, Other.VGPRs);
+ AGPRs = std::min(AGPRs, Other.AGPRs);
+ }
+};
+
+/// An abstract attribute to propagate the register file split a kernel was
+/// compiled with down the call graph to its device functions, emitted as
+/// "amdgpu-register-budget". Entry functions seed the split from their own
+/// vector register budget; every other function inherits the tightest split
+/// over its callers.
+struct AAAMDGPURegisterBudget
+ : public StateWrapper<BooleanState, AbstractAttribute> {
+ using Base = StateWrapper<BooleanState, AbstractAttribute>;
+ AAAMDGPURegisterBudget(const IRPosition &IRP, Attributor &A) : Base(IRP) {}
+
+ static AAAMDGPURegisterBudget &createForPosition(const IRPosition &IRP,
+ Attributor &A) {
+ if (IRP.getPositionKind() == IRPosition::IRP_FUNCTION)
+ return *new (A.Allocator) AAAMDGPURegisterBudget(IRP, A);
+ llvm_unreachable(
+ "AAAMDGPURegisterBudget is only valid for function position");
+ }
+
+ // When we known not all callers are known, we know that any unknown caller
+ // that reaches this function will have a pessimistic amdgpu-agpr-alloc
+ // attribute. This being pessimistic means that the budget for the number of
+ // VGPRs and AGPRs will be split in half for the unknown caller. We also know
+ // that the FlatWorkGroupSize attribute will also be pessimistic for at least
+ // the current function (the one being called in the indirect callsite)
+ unsigned computePessimisticValue(Attributor &A) const {
+ Function *F = getAssociatedFunction();
+ auto &InfoCache = static_cast<AMDGPUInformationCache &>(A.getInfoCache());
+ const GCNSubtarget &ST = InfoCache.TM.getSubtarget<GCNSubtarget>(*F);
+ unsigned MaxWG = ST.getMaxFlatWorkGroupSize();
+ unsigned Occ = std::clamp(ST.getWavesPerEUForWorkGroup(MaxWG), 1u,
+ ST.getMaxWavesPerEU());
+ unsigned Budget =
+ ST.getMaxNumVGPRs(Occ, AMDGPU::getDynamicVGPRBlockSize(*F));
+ return Budget / 2; // 128/2 == 64 on gfx90a at the max work-group size
+ }
+
+ ChangeStatus updateImpl(Attributor &A) override {
+ Function *F = getAssociatedFunction();
+ RegisterBudgetState OldState = Budget;
+
+ // The budget is recomputed from scratch on every update rather than
+ // accumulated, because the seed itself moves: a kernel's split shrinks as
+ // AAAMDGPUMinAGPRAlloc climbs, and callees have to follow it down.
+ if (AMDGPU::isEntryFunctionCC(F->getCallingConv())) {
+ unsigned VGPRBudget = 0;
+ unsigned AGPRBudget = 0;
+ const auto *AGPRAlloc = A.getAAFor<AAAMDGPUMinAGPRAlloc>(
+ *this, IRPosition::function(*F), DepClassTy::OPTIONAL);
+
+ auto &InfoCache = static_cast<AMDGPUInformationCache &>(A.getInfoCache());
+ const GCNSubtarget &ST = InfoCache.TM.getSubtarget<GCNSubtarget>(*F);
+ unsigned MaxRegs = ST.getMaxNumVGPRs(*F);
+ if (!AGPRAlloc || !AGPRAlloc->isValidState()) {
+ VGPRBudget = AGPRBudget = MaxRegs / 2; // pessimistic
+ if (VGPRBudget == computePessimisticValue(A))
+ return indicatePessimisticFixpoint();
+ } else {
+ AGPRBudget = AGPRAlloc->getAssumed();
+ AGPRBudget = alignTo(AGPRBudget, 4);
+ VGPRBudget = MaxRegs - std::min(MaxRegs, AGPRBudget);
+ }
+
+ Budget = {VGPRBudget, AGPRBudget, /*Unknown=*/false};
+ LLVM_DEBUG(dbgs() << "Register budget for " << F->getName() << ": "
+ << VGPRBudget << ", " << AGPRBudget << "\n");
+ } else {
+ RegisterBudgetState Merged;
+
+ auto CheckUse = [&](const Use &U, bool &Follow) {
+ if (auto *CE = dyn_cast<ConstantExpr>(U.getUser())) {
+ if (CE->isCast() && CE->getType()->isPointerTy()) {
+ Follow = true;
+ return true;
+ }
+ }
+ if (isa<SelectInst>(U.getUser()) || isa<PHINode>(U.getUser())) {
+ Follow = true;
+ return true;
+ }
+ AbstractCallSite ACS(&U);
+ const Use *EffectiveUse =
+ ACS && ACS.isCallbackCall() ? &ACS.getCalleeUseForCallback() : &U;
+ if (!ACS || !ACS.isCallee(EffectiveUse))
+ return true;
+
+ Function *Caller = ACS.getInstruction()->getFunction();
+ const auto *CallerAA = A.getAAFor<AAAMDGPURegisterBudget>(
+ *this, IRPosition::function(*Caller), DepClassTy::REQUIRED);
+ if (!CallerAA || !CallerAA->isValidState())
+ return true;
+
+ const RegisterBudgetState &CallerBudget = CallerAA->getBudget();
+ if (!CallerBudget.Unknown)
+ Merged.merge(CallerBudget);
+ return true;
+ };
+
+ A.checkForAllUses(CheckUse, *this, *F);
+
+ // Checks for unknown call sites.
+ bool DummyUAI = false;
+ bool AllCallsitesKnown = A.checkForAllCallSites(
+ [](AbstractCallSite) { return true; }, *this, true, DummyUAI);
+ if (!AllCallsitesKnown && !Merged.Unknown) {
+ if (std::optional<unsigned> PessimisticValue =
+ computePessimisticValue(A)) {
+ Merged.merge(
+ {*PessimisticValue, *PessimisticValue, /*Unknown=*/false});
+ } else
+ return indicatePessimisticFixpoint();
+ }
+ // Stays unknown when no caller contributed a budget, so that functions
+ // outside any kernel's reach are left unconstrained.
+ Budget = Merged;
+ }
+
+ return OldState == Budget ? ChangeStatus::UNCHANGED : ChangeStatus::CHANGED;
+ }
+
+ ChangeStatus manifest(Attributor &A) override {
+ if (Budget.Unknown)
+ return ChangeStatus::UNCHANGED;
+ SmallString<16> Buffer;
+ raw_svector_ostream OS(Buffer);
+ OS << Budget.VGPRs << ',' << Budget.AGPRs;
+ return A.manifestAttrs(
+ getIRPosition(),
+ {Attribute::get(getAssociatedFunction()->getContext(), AttrName,
+ OS.str())},
+ /*ForceReplace=*/true);
+ }
+
+ const RegisterBudgetState &getBudget() const { return Budget; }
+
+ const std::string getAsStr(Attributor *A) const override {
+ if (!getAssumed() || Budget.Unknown)
+ return "unknown";
+ std::string Str;
+ raw_string_ostream OS(Str);
+ OS << AttrName << '=' << Budget.VGPRs << ',' << Budget.AGPRs;
+ return Str;
+ }
+
+ void trackStatistics() const override {}
+
+ StringRef getName() const override { return "AAAMDGPURegisterBudget"; }
+ const char *getIdAddr() const override { return &ID; }
+
+ /// This function should return true if the type of the \p AA is
+ /// AAAMDGPURegisterBudget.
+ static bool classof(const AbstractAttribute *AA) {
+ return AA->getIdAddr() == &ID;
+ }
+
+ static const char ID;
+
+private:
+ RegisterBudgetState Budget;
+
+ static constexpr char AttrName[] = "amdgpu-register-budget";
+};
+
+const char AAAMDGPURegisterBudget::ID = 0;
/// An abstract attribute to propagate the function attribute
/// "amdgpu-cluster-dims" from kernel entry functions to device functions.
@@ -1597,10 +1790,10 @@ static bool runImpl(SetVector<Function *> &Functions, bool IsModulePass,
{&AAAMDAttributes::ID, &AAUniformWorkGroupSize::ID,
&AAPotentialValues::ID, &AAAMDFlatWorkGroupSize::ID,
&AAAMDMaxNumWorkgroups::ID, &AAAMDWavesPerEU::ID,
- &AAAMDGPUMinAGPRAlloc::ID, &AACallEdges::ID, &AAPointerInfo::ID,
- &AAPotentialConstantValues::ID, &AAUnderlyingObjects::ID,
- &AANoAliasAddrSpace::ID, &AAAddressSpace::ID, &AAIndirectCallInfo::ID,
- &AAAMDGPUClusterDims::ID, &AAAlign::ID});
+ &AAAMDGPUMinAGPRAlloc::ID, &AAAMDGPURegisterBudget::ID, &AACallEdges::ID,
+ &AAPointerInfo::ID, &AAPotentialConstantValues::ID,
+ &AAUnderlyingObjects::ID, &AANoAliasAddrSpace::ID, &AAAddressSpace::ID,
+ &AAIndirectCallInfo::ID, &AAAMDGPUClusterDims::ID, &AAAlign::ID});
AttributorConfig AC(CGUpdater);
AC.IsClosedWorldModule = Options.IsClosedWorld;
@@ -1642,8 +1835,10 @@ static bool runImpl(SetVector<Function *> &Functions, bool IsModulePass,
if (!F->isDeclaration() && ST.hasClusters())
A.getOrCreateAAFor<AAAMDGPUClusterDims>(IRPosition::function(*F));
- if (ST.hasGFX90AInsts())
+ if (ST.hasGFX90AInsts()) {
A.getOrCreateAAFor<AAAMDGPUMinAGPRAlloc>(IRPosition::function(*F));
+ A.getOrCreateAAFor<AAAMDGPURegisterBudget>(IRPosition::function(*F));
+ }
for (auto &I : instructions(F)) {
Value *Ptr = nullptr;
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll
index 02eef5ab5e98b..a0d7855b20369 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll
@@ -94,7 +94,7 @@ define amdgpu_kernel void @kernel_uses_asm_virtreg_second_arg() {
define amdgpu_kernel void @kernel_uses_non_agpr_asm() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_non_agpr_asm(
-; CHECK-SAME: ) #[[ATTR0]] {
+; CHECK-SAME: ) #[[ATTR3:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -130,7 +130,7 @@ define amdgpu_kernel void @kernel_uses_asm_physreg_tuple() {
define void @func_uses_asm_virtreg_agpr() {
; CHECK-LABEL: define void @func_uses_asm_virtreg_agpr(
-; CHECK-SAME: ) #[[ATTR1]] {
+; CHECK-SAME: ) #[[ATTR4:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -142,7 +142,7 @@ define void @func_uses_asm_virtreg_agpr() {
define void @func_uses_asm_physreg_agpr() {
; CHECK-LABEL: define void @func_uses_asm_physreg_agpr(
-; CHECK-SAME: ) #[[ATTR1]] {
+; CHECK-SAME: ) #[[ATTR5:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -154,7 +154,7 @@ define void @func_uses_asm_physreg_agpr() {
define void @func_uses_asm_physreg_agpr_tuple() {
; CHECK-LABEL: define void @func_uses_asm_physreg_agpr_tuple(
-; CHECK-SAME: ) #[[ATTR2]] {
+; CHECK-SAME: ) #[[ATTR6:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -179,7 +179,7 @@ define amdgpu_kernel void @kernel_calls_extern() {
define amdgpu_kernel void @kernel_calls_extern_marked_callsite() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_calls_extern_marked_callsite() {
-; CHECK-NEXT: call void @unknown() #[[ATTR35:[0-9]+]]
+; CHECK-NEXT: call void @unknown() #[[ATTR44:[0-9]+]]
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -203,7 +203,7 @@ define amdgpu_kernel void @kernel_calls_indirect(ptr %indirect) {
define amdgpu_kernel void @kernel_calls_indirect_marked_callsite(ptr %indirect) {
; CHECK-LABEL: define amdgpu_kernel void @kernel_calls_indirect_marked_callsite(
; CHECK-SAME: ptr [[INDIRECT:%.*]]) {
-; CHECK-NEXT: call void [[INDIRECT]]() #[[ATTR35]]
+; CHECK-NEXT: call void [[INDIRECT]]() #[[ATTR44]]
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -226,7 +226,7 @@ define amdgpu_kernel void @kernel_transitively_uses_agpr_asm() {
define void @empty() {
; CHECK-LABEL: define void @empty(
-; CHECK-SAME: ) #[[ATTR0]] {
+; CHECK-SAME: ) #[[ATTR7:[0-9]+]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -236,7 +236,7 @@ define void @empty() {
define void @also_empty() {
; CHECK-LABEL: define void @also_empty(
-; CHECK-SAME: ) #[[ATTR0]] {
+; CHECK-SAME: ) #[[ATTR7]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -246,7 +246,7 @@ define void @also_empty() {
define amdgpu_kernel void @kernel_calls_empty() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_calls_empty(
-; CHECK-SAME: ) #[[ATTR0]] {
+; CHECK-SAME: ) #[[ATTR3]] {
; CHECK-NEXT: call void @empty()
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -272,7 +272,7 @@ define amdgpu_kernel void @kernel_calls_non_agpr_and_agpr() {
define amdgpu_kernel void @kernel_calls_generic_intrinsic(ptr %ptr0, ptr %ptr1, i64 %size) {
; CHECK-LABEL: define amdgpu_kernel void @kernel_calls_generic_intrinsic(
-; CHECK-SAME: ptr [[PTR0:%.*]], ptr [[PTR1:%.*]], i64 [[SIZE:%.*]]) #[[ATTR0]] {
+; CHECK-SAME: ptr [[PTR0:%.*]], ptr [[PTR1:%.*]], i64 [[SIZE:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr [[PTR0]], ptr [[PTR1]], i64 [[SIZE]], i1 false)
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -286,7 +286,7 @@ declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>
define amdgpu_kernel void @kernel_calls_mfma.f32.32x32x1f32(ptr addrspace(1) %out, float %a, float %b, <32 x float> %c) {
; CHECK-LABEL: define amdgpu_kernel void @kernel_calls_mfma.f32.32x32x1f32(
-; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], float [[A:%.*]], float [[B:%.*]], <32 x float> [[C:%.*]]) #[[ATTR0]] {
+; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]], float [[A:%.*]], float [[B:%.*]], <32 x float> [[C:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: [[RESULT:%.*]] = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float [[A]], float [[B]], <32 x float> [[C]], i32 0, i32 0, i32 0)
; CHECK-NEXT: store <32 x float> [[RESULT]], ptr addrspace(1) [[OUT]], align 128
; CHECK-NEXT: call void @use_most()
@@ -300,7 +300,7 @@ define amdgpu_kernel void @kernel_calls_mfma.f32.32x32x1f32(ptr addrspace(1) %ou
define amdgpu_kernel void @kernel_calls_workitem_id_x(ptr addrspace(1) %out) {
; CHECK-LABEL: define amdgpu_kernel void @kernel_calls_workitem_id_x(
-; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]]) #[[ATTR0]] {
+; CHECK-SAME: ptr addrspace(1) [[OUT:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: [[RESULT:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
; CHECK-NEXT: store i32 [[RESULT]], ptr addrspace(1) [[OUT]], align 4
; CHECK-NEXT: call void @use_most()
@@ -314,7 +314,7 @@ define amdgpu_kernel void @kernel_calls_workitem_id_x(ptr addrspace(1) %out) {
define amdgpu_kernel void @indirect_calls_none_agpr(i1 %cond) {
; CHECK-LABEL: define amdgpu_kernel void @indirect_calls_none_agpr(
-; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR0]] {
+; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR3]] {
; CHECK-NEXT: [[FPTR:%.*]] = select i1 [[COND]], ptr @empty, ptr @also_empty
; CHECK-NEXT: [[TMP1:%.*]] = icmp eq ptr [[FPTR]], @also_empty
; CHECK-NEXT: br i1 [[TMP1]], label [[TMP2:%.*]], label [[TMP3:%.*]]
@@ -352,7 +352,7 @@ define amdgpu_kernel void @kernel_uses_asm_virtreg_def_struct_0() {
define amdgpu_kernel void @kernel_uses_asm_virtreg_use_struct_1() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_virtreg_use_struct_1(
-; CHECK-SAME: ) #[[ATTR4:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR9:[0-9]+]] {
; CHECK-NEXT: [[DEF:%.*]] = call { i32, <2 x i32> } asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -400,7 +400,7 @@ define amdgpu_kernel void @kernel_uses_asm_virtreg_def_ptr_ty() {
define amdgpu_kernel void @kernel_uses_asm_virtreg_def_vector_ptr_ty() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_virtreg_def_vector_ptr_ty(
-; CHECK-SAME: ) #[[ATTR4]] {
+; CHECK-SAME: ) #[[ATTR9]] {
; CHECK-NEXT: [[DEF:%.*]] = call <2 x ptr> asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -412,7 +412,7 @@ define amdgpu_kernel void @kernel_uses_asm_virtreg_def_vector_ptr_ty() {
define amdgpu_kernel void @kernel_uses_asm_physreg_def_struct_0() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_physreg_def_struct_0(
-; CHECK-SAME: ) #[[ATTR5:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR10:[0-9]+]] {
; CHECK-NEXT: [[DEF:%.*]] = call { i32, i32 } asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -424,7 +424,7 @@ define amdgpu_kernel void @kernel_uses_asm_physreg_def_struct_0() {
define amdgpu_kernel void @kernel_uses_asm_clobber() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_clobber(
-; CHECK-SAME: ) #[[ATTR6:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR11:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -436,7 +436,7 @@ define amdgpu_kernel void @kernel_uses_asm_clobber() {
define amdgpu_kernel void @kernel_uses_asm_clobber_tuple() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_clobber_tuple(
-; CHECK-SAME: ) #[[ATTR7:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR12:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -448,7 +448,7 @@ define amdgpu_kernel void @kernel_uses_asm_clobber_tuple() {
define amdgpu_kernel void @kernel_uses_asm_clobber_oob() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_clobber_oob(
-; CHECK-SAME: ) #[[ATTR8:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR13:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -460,7 +460,7 @@ define amdgpu_kernel void @kernel_uses_asm_clobber_oob() {
define amdgpu_kernel void @kernel_uses_asm_clobber_max() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_clobber_max(
-; CHECK-SAME: ) #[[ATTR8]] {
+; CHECK-SAME: ) #[[ATTR13]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -472,7 +472,7 @@ define amdgpu_kernel void @kernel_uses_asm_clobber_max() {
define amdgpu_kernel void @kernel_uses_asm_physreg_oob() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_physreg_oob(
-; CHECK-SAME: ) #[[ATTR8]] {
+; CHECK-SAME: ) #[[ATTR13]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -484,7 +484,7 @@ define amdgpu_kernel void @kernel_uses_asm_physreg_oob() {
define amdgpu_kernel void @kernel_uses_asm_virtreg_def_max_ty() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_virtreg_def_max_ty(
-; CHECK-SAME: ) #[[ATTR9:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR14:[0-9]+]] {
; CHECK-NEXT: [[DEF:%.*]] = call <32 x i32> asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -496,7 +496,7 @@ define amdgpu_kernel void @kernel_uses_asm_virtreg_def_max_ty() {
define amdgpu_kernel void @kernel_uses_asm_virtreg_use_max_ty() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_virtreg_use_max_ty(
-; CHECK-SAME: ) #[[ATTR9]] {
+; CHECK-SAME: ) #[[ATTR14]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -508,7 +508,7 @@ define amdgpu_kernel void @kernel_uses_asm_virtreg_use_max_ty() {
define amdgpu_kernel void @kernel_uses_asm_virtreg_use_def_max_ty() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_asm_virtreg_use_def_max_ty(
-; CHECK-SAME: ) #[[ATTR9]] {
+; CHECK-SAME: ) #[[ATTR14]] {
; CHECK-NEXT: [[DEF:%.*]] = call <32 x i32> asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -520,7 +520,7 @@ define amdgpu_kernel void @kernel_uses_asm_virtreg_use_def_max_ty() {
define amdgpu_kernel void @vreg_use_exceeds_register_file() {
; CHECK-LABEL: define amdgpu_kernel void @vreg_use_exceeds_register_file(
-; CHECK-SAME: ) #[[ATTR8]] {
+; CHECK-SAME: ) #[[ATTR13]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -532,7 +532,7 @@ define amdgpu_kernel void @vreg_use_exceeds_register_file() {
define amdgpu_kernel void @vreg_def_exceeds_register_file() {
; CHECK-LABEL: define amdgpu_kernel void @vreg_def_exceeds_register_file(
-; CHECK-SAME: ) #[[ATTR8]] {
+; CHECK-SAME: ) #[[ATTR13]] {
; CHECK-NEXT: [[DEF:%.*]] = call <257 x i32> asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -544,7 +544,7 @@ define amdgpu_kernel void @vreg_def_exceeds_register_file() {
define amdgpu_kernel void @multiple() {
; CHECK-LABEL: define amdgpu_kernel void @multiple(
-; CHECK-SAME: ) #[[ATTR9]] {
+; CHECK-SAME: ) #[[ATTR14]] {
; CHECK-NEXT: [[DEF:%.*]] = call { <16 x i32>, <8 x i32>, <8 x i32> } asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -556,7 +556,7 @@ define amdgpu_kernel void @multiple() {
define amdgpu_kernel void @earlyclobber_0() {
; CHECK-LABEL: define amdgpu_kernel void @earlyclobber_0(
-; CHECK-SAME: ) #[[ATTR10:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR15:[0-9]+]] {
; CHECK-NEXT: [[DEF:%.*]] = call <8 x i32> asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -568,7 +568,7 @@ define amdgpu_kernel void @earlyclobber_0() {
define amdgpu_kernel void @earlyclobber_1() {
; CHECK-LABEL: define amdgpu_kernel void @earlyclobber_1(
-; CHECK-SAME: ) #[[ATTR11:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR16:[0-9]+]] {
; CHECK-NEXT: [[DEF:%.*]] = call { <8 x i32>, <16 x i32> } asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -580,7 +580,7 @@ define amdgpu_kernel void @earlyclobber_1() {
define amdgpu_kernel void @physreg_a32__vreg_a256__vreg_a512() {
; CHECK-LABEL: define amdgpu_kernel void @physreg_a32__vreg_a256__vreg_a512(
-; CHECK-SAME: ) #[[ATTR12:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR17:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -592,7 +592,7 @@ define amdgpu_kernel void @physreg_a32__vreg_a256__vreg_a512() {
define amdgpu_kernel void @physreg_def_a32__def_vreg_a256__def_vreg_a512() {
; CHECK-LABEL: define amdgpu_kernel void @physreg_def_a32__def_vreg_a256__def_vreg_a512(
-; CHECK-SAME: ) #[[ATTR12]] {
+; CHECK-SAME: ) #[[ATTR17]] {
; CHECK-NEXT: [[TMP1:%.*]] = call { i32, <8 x i32>, <16 x i32> } asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -604,7 +604,7 @@ define amdgpu_kernel void @physreg_def_a32__def_vreg_a256__def_vreg_a512() {
define amdgpu_kernel void @physreg_def_a32___def_vreg_a512_use_vreg_a256() {
; CHECK-LABEL: define amdgpu_kernel void @physreg_def_a32___def_vreg_a512_use_vreg_a256(
-; CHECK-SAME: ) #[[ATTR13:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR18:[0-9]+]] {
; CHECK-NEXT: [[TMP1:%.*]] = call { i32, <16 x i32> } asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -616,7 +616,7 @@ define amdgpu_kernel void @physreg_def_a32___def_vreg_a512_use_vreg_a256() {
define amdgpu_kernel void @mixed_physreg_vreg_tuples_0() {
; CHECK-LABEL: define amdgpu_kernel void @mixed_physreg_vreg_tuples_0(
-; CHECK-SAME: ) #[[ATTR10]] {
+; CHECK-SAME: ) #[[ATTR15]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -628,7 +628,7 @@ define amdgpu_kernel void @mixed_physreg_vreg_tuples_0() {
define amdgpu_kernel void @mixed_physreg_vreg_tuples_1() {
; CHECK-LABEL: define amdgpu_kernel void @mixed_physreg_vreg_tuples_1(
-; CHECK-SAME: ) #[[ATTR14:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR19:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -640,7 +640,7 @@ define amdgpu_kernel void @mixed_physreg_vreg_tuples_1() {
define amdgpu_kernel void @physreg_raises_limit() {
; CHECK-LABEL: define amdgpu_kernel void @physreg_raises_limit(
-; CHECK-SAME: ) #[[ATTR15:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR20:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -652,7 +652,7 @@ define amdgpu_kernel void @physreg_raises_limit() {
define amdgpu_kernel void @physreg_tuple_alignment_raises_limit() {
; CHECK-LABEL: define amdgpu_kernel void @physreg_tuple_alignment_raises_limit(
-; CHECK-SAME: ) #[[ATTR10]] {
+; CHECK-SAME: ) #[[ATTR15]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -664,7 +664,7 @@ define amdgpu_kernel void @physreg_tuple_alignment_raises_limit() {
define amdgpu_kernel void @align3_virtreg() {
; CHECK-LABEL: define amdgpu_kernel void @align3_virtreg(
-; CHECK-SAME: ) #[[ATTR5]] {
+; CHECK-SAME: ) #[[ATTR10]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -676,7 +676,7 @@ define amdgpu_kernel void @align3_virtreg() {
define amdgpu_kernel void @align3_align4_virtreg() {
; CHECK-LABEL: define amdgpu_kernel void @align3_align4_virtreg(
-; CHECK-SAME: ) #[[ATTR14]] {
+; CHECK-SAME: ) #[[ATTR19]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -688,7 +688,7 @@ define amdgpu_kernel void @align3_align4_virtreg() {
define amdgpu_kernel void @align2_align4_virtreg() {
; CHECK-LABEL: define amdgpu_kernel void @align2_align4_virtreg(
-; CHECK-SAME: ) #[[ATTR14]] {
+; CHECK-SAME: ) #[[ATTR19]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -700,7 +700,7 @@ define amdgpu_kernel void @align2_align4_virtreg() {
define amdgpu_kernel void @kernel_uses_write_register_a55() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_write_register_a55(
-; CHECK-SAME: ) #[[ATTR16:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR21:[0-9]+]] {
; CHECK-NEXT: call void @llvm.write_register.i32(metadata [[META0:![0-9]+]], i32 0)
; CHECK-NEXT: ret void
;
@@ -710,7 +710,7 @@ define amdgpu_kernel void @kernel_uses_write_register_a55() {
define amdgpu_kernel void @kernel_uses_write_register_v55() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_write_register_v55(
-; CHECK-SAME: ) #[[ATTR0]] {
+; CHECK-SAME: ) #[[ATTR3]] {
; CHECK-NEXT: call void @llvm.write_register.i32(metadata [[META1:![0-9]+]], i32 0)
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -722,7 +722,7 @@ define amdgpu_kernel void @kernel_uses_write_register_v55() {
define amdgpu_kernel void @kernel_uses_write_register_a55_57() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_write_register_a55_57(
-; CHECK-SAME: ) #[[ATTR17:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR22:[0-9]+]] {
; CHECK-NEXT: call void @llvm.write_register.i96(metadata [[META2:![0-9]+]], i96 0)
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -734,7 +734,7 @@ define amdgpu_kernel void @kernel_uses_write_register_a55_57() {
define amdgpu_kernel void @kernel_uses_read_register_a55(ptr addrspace(1) %ptr) {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_read_register_a55(
-; CHECK-SAME: ptr addrspace(1) [[PTR:%.*]]) #[[ATTR18:[0-9]+]] {
+; CHECK-SAME: ptr addrspace(1) [[PTR:%.*]]) #[[ATTR23:[0-9]+]] {
; CHECK-NEXT: [[REG:%.*]] = call i32 @llvm.read_register.i32(metadata [[META0]])
; CHECK-NEXT: store i32 [[REG]], ptr addrspace(1) [[PTR]], align 4
; CHECK-NEXT: call void @use_most()
@@ -748,7 +748,7 @@ define amdgpu_kernel void @kernel_uses_read_register_a55(ptr addrspace(1) %ptr)
define amdgpu_kernel void @kernel_uses_read_volatile_register_a55(ptr addrspace(1) %ptr) {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_read_volatile_register_a55(
-; CHECK-SAME: ptr addrspace(1) [[PTR:%.*]]) #[[ATTR19:[0-9]+]] {
+; CHECK-SAME: ptr addrspace(1) [[PTR:%.*]]) #[[ATTR24:[0-9]+]] {
; CHECK-NEXT: [[REG:%.*]] = call i32 @llvm.read_volatile_register.i32(metadata [[META0]])
; CHECK-NEXT: store i32 [[REG]], ptr addrspace(1) [[PTR]], align 4
; CHECK-NEXT: call void @use_most()
@@ -762,7 +762,7 @@ define amdgpu_kernel void @kernel_uses_read_volatile_register_a55(ptr addrspace(
define amdgpu_kernel void @kernel_uses_read_register_a56_59(ptr addrspace(1) %ptr) {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_read_register_a56_59(
-; CHECK-SAME: ptr addrspace(1) [[PTR:%.*]]) #[[ATTR20:[0-9]+]] {
+; CHECK-SAME: ptr addrspace(1) [[PTR:%.*]]) #[[ATTR25:[0-9]+]] {
; CHECK-NEXT: [[REG:%.*]] = call i128 @llvm.read_register.i128(metadata [[META3:![0-9]+]])
; CHECK-NEXT: store i128 [[REG]], ptr addrspace(1) [[PTR]], align 8
; CHECK-NEXT: call void @use_most()
@@ -776,7 +776,7 @@ define amdgpu_kernel void @kernel_uses_read_register_a56_59(ptr addrspace(1) %pt
define amdgpu_kernel void @kernel_uses_write_register_out_of_bounds_a256() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_write_register_out_of_bounds_a256(
-; CHECK-SAME: ) #[[ATTR8]] {
+; CHECK-SAME: ) #[[ATTR13]] {
; CHECK-NEXT: call void @llvm.write_register.i32(metadata [[META4:![0-9]+]], i32 0)
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -788,7 +788,7 @@ define amdgpu_kernel void @kernel_uses_write_register_out_of_bounds_a256() {
define amdgpu_kernel void @kernel_multiple_uses() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_multiple_uses(
-; CHECK-SAME: ) #[[ATTR4]] {
+; CHECK-SAME: ) #[[ATTR9]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void asm sideeffect "
@@ -804,7 +804,7 @@ define amdgpu_kernel void @kernel_multiple_uses() {
define amdgpu_kernel void @kernel_multiple_defs() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_multiple_defs(
-; CHECK-SAME: ) #[[ATTR4]] {
+; CHECK-SAME: ) #[[ATTR9]] {
; CHECK-NEXT: [[TMP1:%.*]] = call i64 asm sideeffect "
; CHECK-NEXT: [[TMP2:%.*]] = call i32 asm sideeffect "
; CHECK-NEXT: [[TMP3:%.*]] = call i128 asm sideeffect "
@@ -820,7 +820,7 @@ define amdgpu_kernel void @kernel_multiple_defs() {
define amdgpu_kernel void @kernel_multiple_use_defs() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_multiple_use_defs(
-; CHECK-SAME: ) #[[ATTR4]] {
+; CHECK-SAME: ) #[[ATTR9]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: [[TMP1:%.*]] = call i128 asm sideeffect "
; CHECK-NEXT: call void @use_most()
@@ -834,7 +834,7 @@ define amdgpu_kernel void @kernel_multiple_use_defs() {
define void @callgraph_b() {
; CHECK-LABEL: define void @callgraph_b(
-; CHECK-SAME: ) #[[ATTR14]] {
+; CHECK-SAME: ) #[[ATTR26:[0-9]+]] {
; CHECK-NEXT: [[TMP1:%.*]] = call <4 x i32> asm sideeffect "
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
@@ -848,7 +848,7 @@ define void @callgraph_b() {
define void @callgraph_c() {
; CHECK-LABEL: define void @callgraph_c(
-; CHECK-SAME: ) #[[ATTR2]] {
+; CHECK-SAME: ) #[[ATTR27:[0-9]+]] {
; CHECK-NEXT: [[TMP1:%.*]] = call i32 asm sideeffect "
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
@@ -862,7 +862,7 @@ define void @callgraph_c() {
define void @callgraph_a(i1 %cond) {
; CHECK-LABEL: define void @callgraph_a(
-; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR14]] {
+; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR26]] {
; CHECK-NEXT: br i1 [[COND]], label [[A:%.*]], label [[B:%.*]]
; CHECK: a:
; CHECK-NEXT: call void @callgraph_b()
@@ -883,9 +883,9 @@ b:
}
-define void @kernel_max_callgraph(i1 %cond) {
-; CHECK-LABEL: define void @kernel_max_callgraph(
-; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR14]] {
+define amdgpu_kernel void @kernel_max_callgraph(i1 %cond) {
+; CHECK-LABEL: define amdgpu_kernel void @kernel_max_callgraph(
+; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR19]] {
; CHECK-NEXT: call void @callgraph_a(i1 [[COND]])
; CHECK-NEXT: ret void
;
@@ -895,7 +895,7 @@ define void @kernel_max_callgraph(i1 %cond) {
define amdgpu_kernel void @kernel_uses_all_virtregs() #1 {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_all_virtregs(
-; CHECK-SAME: ) #[[ATTR21:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR28:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -907,7 +907,7 @@ define amdgpu_kernel void @kernel_uses_all_virtregs() #1 {
define amdgpu_kernel void @kernel_uses_all_virtregs_plus_1() #1 {
; CHECK-LABEL: define amdgpu_kernel void @kernel_uses_all_virtregs_plus_1(
-; CHECK-SAME: ) #[[ATTR21]] {
+; CHECK-SAME: ) #[[ATTR28]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -919,7 +919,7 @@ define amdgpu_kernel void @kernel_uses_all_virtregs_plus_1() #1 {
define void @recursive() {
; CHECK-LABEL: define void @recursive(
-; CHECK-SAME: ) #[[ATTR22:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR29:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: call void @recursive()
@@ -933,7 +933,7 @@ define void @recursive() {
define void @indirect_0() {
; CHECK-LABEL: define void @indirect_0(
-; CHECK-SAME: ) #[[ATTR22]] {
+; CHECK-SAME: ) #[[ATTR30:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -945,7 +945,7 @@ define void @indirect_0() {
define void @indirect_1() {
; CHECK-LABEL: define void @indirect_1(
-; CHECK-SAME: ) #[[ATTR23:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR31:[0-9]+]] {
; CHECK-NEXT: [[TMP1:%.*]] = call <3 x i32> asm sideeffect "
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -957,7 +957,7 @@ define void @indirect_1() {
define amdgpu_kernel void @knowable_indirect_call(i1 %cond) {
; CHECK-LABEL: define amdgpu_kernel void @knowable_indirect_call(
-; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR22]] {
+; CHECK-SAME: i1 [[COND:%.*]]) #[[ATTR32:[0-9]+]] {
; CHECK-NEXT: [[FPTR:%.*]] = select i1 [[COND]], ptr @indirect_0, ptr @indirect_1
; CHECK-NEXT: [[TMP1:%.*]] = icmp eq ptr [[FPTR]], @indirect_1
; CHECK-NEXT: br i1 [[TMP1]], label [[TMP2:%.*]], label [[TMP3:%.*]]
@@ -1021,7 +1021,7 @@ define amdgpu_kernel void @indirect_unknown(ptr %fptr) {
define amdgpu_kernel void @kernel_sanitize_address() sanitize_address {
; CHECK: Function Attrs: sanitize_address
; CHECK-LABEL: define amdgpu_kernel void @kernel_sanitize_address(
-; CHECK-SAME: ) #[[ATTR24:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR33:[0-9]+]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -1032,7 +1032,7 @@ define amdgpu_kernel void @kernel_sanitize_address() sanitize_address {
define amdgpu_kernel void @kernel_sanitize_memory() sanitize_memory {
; CHECK: Function Attrs: sanitize_memory
; CHECK-LABEL: define amdgpu_kernel void @kernel_sanitize_memory(
-; CHECK-SAME: ) #[[ATTR25:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR34:[0-9]+]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -1043,7 +1043,7 @@ define amdgpu_kernel void @kernel_sanitize_memory() sanitize_memory {
define amdgpu_kernel void @kernel_sanitize_thread() sanitize_thread {
; CHECK: Function Attrs: sanitize_thread
; CHECK-LABEL: define amdgpu_kernel void @kernel_sanitize_thread(
-; CHECK-SAME: ) #[[ATTR26:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR35:[0-9]+]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -1054,7 +1054,7 @@ define amdgpu_kernel void @kernel_sanitize_thread() sanitize_thread {
define amdgpu_kernel void @kernel_sanitize_hwaddress() sanitize_hwaddress {
; CHECK: Function Attrs: sanitize_hwaddress
; CHECK-LABEL: define amdgpu_kernel void @kernel_sanitize_hwaddress(
-; CHECK-SAME: ) #[[ATTR27:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR36:[0-9]+]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -1069,7 +1069,7 @@ define amdgpu_kernel void @kernel_sanitize_hwaddress() sanitize_hwaddress {
define amdgpu_kernel void @kernel_sanitize_address_preannotated() #2 {
; CHECK: Function Attrs: sanitize_address
; CHECK-LABEL: define amdgpu_kernel void @kernel_sanitize_address_preannotated(
-; CHECK-SAME: ) #[[ATTR28:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR37:[0-9]+]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -1082,7 +1082,7 @@ define amdgpu_kernel void @kernel_sanitize_address_preannotated() #2 {
define void @sanitized_callee() sanitize_address {
; CHECK: Function Attrs: sanitize_address
; CHECK-LABEL: define void @sanitized_callee(
-; CHECK-SAME: ) #[[ATTR24]] {
+; CHECK-SAME: ) #[[ATTR33]] {
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
;
@@ -1092,7 +1092,7 @@ define void @sanitized_callee() sanitize_address {
define amdgpu_kernel void @kernel_calls_sanitized_callee() {
; CHECK-LABEL: define amdgpu_kernel void @kernel_calls_sanitized_callee(
-; CHECK-SAME: ) #[[ATTR29:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR38:[0-9]+]] {
; CHECK-NEXT: call void @sanitized_callee()
; CHECK-NEXT: call void @use_most()
; CHECK-NEXT: ret void
@@ -1113,42 +1113,51 @@ attributes #2 = { sanitize_address "amdgpu-agpr-alloc"="0" }
!4 = !{!"a256"}
;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR3:[0-9]+]] = { convergent nocallback nocreateundeforpoison nofree nosync nounwind willreturn memory(none) }
-; CHECK: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="4" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR5]] = { "amdgpu-agpr-alloc"="6" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR6]] = { "amdgpu-agpr-alloc"="5" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR7]] = { "amdgpu-agpr-alloc"="14" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR8]] = { "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR9]] = { "amdgpu-agpr-alloc"="32" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR10]] = { "amdgpu-agpr-alloc"="9" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR11]] = { "amdgpu-agpr-alloc"="64" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR12]] = { "amdgpu-agpr-alloc"="49" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR13]] = { "amdgpu-agpr-alloc"="33" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR14]] = { "amdgpu-agpr-alloc"="8" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR15]] = { "amdgpu-agpr-alloc"="13" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR16]] = { "amdgpu-agpr-alloc"="56" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR17]] = { "amdgpu-agpr-alloc"="58" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR18]] = { "amdgpu-agpr-alloc"="56" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR19]] = { "amdgpu-agpr-alloc"="56" }
-; CHECK: attributes #[[ATTR20]] = { "amdgpu-agpr-alloc"="60" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR21]] = { "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" "amdgpu-waves-per-eu"="1,1" }
-; CHECK: attributes #[[ATTR22]] = { "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR23]] = { "amdgpu-agpr-alloc"="3" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR24]] = { sanitize_address "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR25]] = { sanitize_memory "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR26]] = { sanitize_thread "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR27]] = { sanitize_hwaddress "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR28]] = { sanitize_address "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR29]] = { "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR30:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR31:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
-; CHECK: attributes #[[ATTR32:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(read) }
-; CHECK: attributes #[[ATTR33:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR34:[0-9]+]] = { nocallback nounwind }
-; CHECK: attributes #[[ATTR35]] = { "amdgpu-agpr-alloc"="0" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="0,0" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" "amdgpu-register-budget"="124,4" }
+; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" "amdgpu-register-budget"="124,4" }
+; CHECK: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
+; CHECK: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR5]] = { "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" "amdgpu-register-budget"="64,4" }
+; CHECK: attributes #[[ATTR6]] = { "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR7]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="64,0" }
+; CHECK: attributes #[[ATTR8:[0-9]+]] = { convergent nocallback nocreateundeforpoison nofree nosync nounwind willreturn memory(none) }
+; CHECK: attributes #[[ATTR9]] = { "amdgpu-agpr-alloc"="4" "amdgpu-no-wwm" "amdgpu-register-budget"="124,4" }
+; CHECK: attributes #[[ATTR10]] = { "amdgpu-agpr-alloc"="6" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
+; CHECK: attributes #[[ATTR11]] = { "amdgpu-agpr-alloc"="5" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
+; CHECK: attributes #[[ATTR12]] = { "amdgpu-agpr-alloc"="14" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
+; CHECK: attributes #[[ATTR13]] = { "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" "amdgpu-register-budget"="0,256" }
+; CHECK: attributes #[[ATTR14]] = { "amdgpu-agpr-alloc"="32" "amdgpu-no-wwm" "amdgpu-register-budget"="96,32" }
+; CHECK: attributes #[[ATTR15]] = { "amdgpu-agpr-alloc"="9" "amdgpu-no-wwm" "amdgpu-register-budget"="116,12" }
+; CHECK: attributes #[[ATTR16]] = { "amdgpu-agpr-alloc"="64" "amdgpu-no-wwm" "amdgpu-register-budget"="64,64" }
+; CHECK: attributes #[[ATTR17]] = { "amdgpu-agpr-alloc"="49" "amdgpu-no-wwm" "amdgpu-register-budget"="76,52" }
+; CHECK: attributes #[[ATTR18]] = { "amdgpu-agpr-alloc"="33" "amdgpu-no-wwm" "amdgpu-register-budget"="92,36" }
+; CHECK: attributes #[[ATTR19]] = { "amdgpu-agpr-alloc"="8" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
+; CHECK: attributes #[[ATTR20]] = { "amdgpu-agpr-alloc"="13" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
+; CHECK: attributes #[[ATTR21]] = { "amdgpu-agpr-alloc"="56" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" "amdgpu-register-budget"="72,56" }
+; CHECK: attributes #[[ATTR22]] = { "amdgpu-agpr-alloc"="58" "amdgpu-no-wwm" "amdgpu-register-budget"="68,60" }
+; CHECK: attributes #[[ATTR23]] = { "amdgpu-agpr-alloc"="56" "amdgpu-no-wwm" "amdgpu-register-budget"="72,56" }
+; CHECK: attributes #[[ATTR24]] = { "amdgpu-agpr-alloc"="56" "amdgpu-register-budget"="72,56" }
+; CHECK: attributes #[[ATTR25]] = { "amdgpu-agpr-alloc"="60" "amdgpu-no-wwm" "amdgpu-register-budget"="68,60" }
+; CHECK: attributes #[[ATTR26]] = { "amdgpu-agpr-alloc"="8" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
+; CHECK: attributes #[[ATTR27]] = { "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
+; CHECK: attributes #[[ATTR28]] = { "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" "amdgpu-register-budget"="0,256" "amdgpu-waves-per-eu"="1,1" }
+; CHECK: attributes #[[ATTR29]] = { "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR30]] = { "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
+; CHECK: attributes #[[ATTR31]] = { "amdgpu-agpr-alloc"="3" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
+; CHECK: attributes #[[ATTR32]] = { "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
+; CHECK: attributes #[[ATTR33]] = { sanitize_address "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR34]] = { sanitize_memory "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR35]] = { sanitize_thread "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR36]] = { sanitize_hwaddress "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR37]] = { sanitize_address "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
+; CHECK: attributes #[[ATTR38]] = { "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR39:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR40:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
+; CHECK: attributes #[[ATTR41:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(read) }
+; CHECK: attributes #[[ATTR42:[0-9]+]] = { nounwind }
+; CHECK: attributes #[[ATTR43:[0-9]+]] = { nocallback nounwind }
+; CHECK: attributes #[[ATTR44]] = { "amdgpu-agpr-alloc"="0" }
;.
; CHECK: [[META0]] = !{!"a55"}
; CHECK: [[META1]] = !{!"v55"}
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-propagate.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-propagate.ll
new file mode 100644
index 0000000000000..dc16d37b89857
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-propagate.ll
@@ -0,0 +1,814 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --check-attributes --check-globals all --version 6
+; RUN: opt -S -mtriple=amdgpu9.0a-amd-amdhsa -passes=amdgpu-attributor %s | FileCheck %s
+
+; Consolidated tests for AGPR/VGPR split propagation by the AMDGPU attributor.
+; The AGPR demand is propagated *up* the call graph (callee -> caller) via
+; AAAMDGPUMinAGPRAlloc and emitted as "amdgpu-agpr-alloc". The register budget
+; is propagated *down* (caller -> callee) via AAAMDGPURegisterBudget and
+; emitted as "amdgpu-register-budget", which caps a callee's registers so the
+; split the caller was compiled with is runnable with all callees. The
+; attribute is always a "VGPRs,AGPRs" pair of absolute ceilings, and a function
+; reachable from several callers takes the smallest ceiling on each axis
+; independently. The total budget is 128 here (the default flatblocksize/wavesperEU give 128 for 512 register file), so an unknown/indirect caller,
+; having itself been pessimised to a half split, contributes 64,64.
+;
+; Each scenario below is an independent call-graph component (its own copy of
+; the @*_use_most sink) so the scenarios converge independently.
+
+;; ===========================================================================
+;; Scenario 1 (direct): a kernel directly calls an AGPR-hungry leaf and an
+;; all-VGPR leaf. The all-VGPR leaf must be pulled down to the kernel's split
+;; (78,50) so the register allocator cannot grab the whole VGPR file.
+;;
+;; A
+;; / \
+;; B C
+;;
+;; A = k_direct : actual agpr - 0, agpr attribute - 50 register-budget 78,50
+;; B = direct_leaf_hivgpr : actual agpr - 0, agpr attribute - 0 register-budget 78,50
+;; C = direct_leaf_agpr : actual agpr - 50, agpr attribute - 50 register -budget 78,50
+;; ===========================================================================
+
+; Shrink result attribute list by preventing use of most attributes.
+;.
+; CHECK: @fptr = global ptr null
+;.
+define internal void @direct_use_most() {
+; CHECK-LABEL: define internal void @direct_use_most(
+; CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+; A leaf that forces 50 AGPRs by clobbering a49.
+define internal void @direct_leaf_agpr() {
+; CHECK-LABEL: define internal void @direct_leaf_agpr(
+; CHECK-SAME: ) #[[ATTR1:[0-9]+]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @direct_use_most()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a49", "~{a49}"()
+ call void @direct_use_most()
+ ret void
+}
+
+; An all-VGPR leaf with no AGPR references of its own. It must be pulled to the
+; caller's split (78,50) via downward propagation rather than staying unbounded.
+define internal void @direct_leaf_hivgpr(ptr %p) {
+; CHECK-LABEL: define internal void @direct_leaf_hivgpr(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR0]] {
+; CHECK-NEXT: [[V:%.*]] = load volatile <128 x i32>, ptr [[P]], align 512
+; CHECK-NEXT: store volatile <128 x i32> [[V]], ptr [[P]], align 512
+; CHECK-NEXT: call void @direct_use_most()
+; CHECK-NEXT: ret void
+;
+ %v = load volatile <128 x i32>, ptr %p
+ store volatile <128 x i32> %v, ptr %p
+ call void @direct_use_most()
+ ret void
+}
+
+define amdgpu_kernel void @k_direct(ptr %p) {
+; CHECK-LABEL: define amdgpu_kernel void @k_direct(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: call void @direct_leaf_agpr()
+; CHECK-NEXT: call void @direct_leaf_hivgpr(ptr [[P]])
+; CHECK-NEXT: call void @direct_use_most()
+; CHECK-NEXT: ret void
+;
+ call void @direct_leaf_agpr()
+ call void @direct_leaf_hivgpr(ptr %p)
+ call void @direct_use_most()
+ ret void
+}
+
+;; ===========================================================================
+;; Scenario 2 (indirect): the only kernel makes an unknown indirect call, so its
+;; AGPR demand is unbounded and it ends up with no attributes at all.
+;; With "amdgpu-agpr-alloc" absent the consumer falls back to the same half split.
+;; The indirectly-reachable leaves are externally callable, so not all of their
+;; callers are known and they inherit the half-split budget an unknown caller
+;; would impose Upward real demand is still annotated
+;;
+;; A
+;; / \
+;; B C
+;;
+;; C = indirect_leaf_agpr : real agpr 50, no register-budget attribute -> 64,64
+;; B = indirect_leaf_hivgpr : real agpr 0, no register-budget attribute -> 64,64
+;; A = k_indirect : no attributes -> 64,64
+;; ===========================================================================
+
+define internal void @indirect_use_most() {
+; CHECK-LABEL: define internal void @indirect_use_most(
+; CHECK-SAME: ) #[[ATTR2:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+; Leaf that would otherwise force 50 AGPRs, but is only reachable indirectly.
+define void @indirect_leaf_agpr() {
+; CHECK-LABEL: define void @indirect_leaf_agpr(
+; CHECK-SAME: ) #[[ATTR3:[0-9]+]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @indirect_use_most()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a49", "~{a49}"()
+ call void @indirect_use_most()
+ ret void
+}
+
+; All-VGPR leaf, also only reachable indirectly.
+define void @indirect_leaf_hivgpr(ptr %p) {
+; CHECK-LABEL: define void @indirect_leaf_hivgpr(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR2]] {
+; CHECK-NEXT: [[V:%.*]] = load volatile <128 x i32>, ptr [[P]], align 512
+; CHECK-NEXT: store volatile <128 x i32> [[V]], ptr [[P]], align 512
+; CHECK-NEXT: call void @indirect_use_most()
+; CHECK-NEXT: ret void
+;
+ %v = load volatile <128 x i32>, ptr %p
+ store volatile <128 x i32> %v, ptr %p
+ call void @indirect_use_most()
+ ret void
+}
+
+ at fptr = global ptr null
+
+define amdgpu_kernel void @k_indirect() {
+; CHECK-LABEL: define amdgpu_kernel void @k_indirect() {
+; CHECK-NEXT: [[FN:%.*]] = load ptr, ptr @fptr, align 8
+; CHECK-NEXT: call void [[FN]]()
+; CHECK-NEXT: call void @indirect_use_most()
+; CHECK-NEXT: ret void
+;
+ %fn = load ptr, ptr @fptr
+ call void %fn()
+ call void @indirect_use_most()
+ ret void
+}
+
+;; ===========================================================================
+;; Scenario 3 (cycle): This scenario show how a cycle affects the up then
+;; down propogation
+;; use_most_hi : real 0, register-budget 78,50
+;; ping / pong / k_mutual : real 50, register-budget 78,50
+;; use_most_lo : real 0, register-budget 96,32
+;; self_rec / k_self : real 32, register-budget 96,32
+;; ===========================================================================
+
+define internal void @use_most_hi() {
+; CHECK-LABEL: define internal void @use_most_hi(
+; CHECK-SAME: ) #[[ATTR0]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+; Mutual recursion: @ping <-> @pong. Only @ping has a direct AGPR requirement
+; (clobbers a49 -> 50). Propagation around the cycle converges @pong and the
+; kernel to 50 as well.
+define internal void @ping(i1 %c) {
+; CHECK-LABEL: define internal void @ping(
+; CHECK-SAME: i1 [[C:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @use_most_hi()
+; CHECK-NEXT: br i1 [[C]], label %[[REC:.*]], label %[[EXIT:.*]]
+; CHECK: [[REC]]:
+; CHECK-NEXT: call void @pong(i1 [[C]])
+; CHECK-NEXT: br label %[[EXIT]]
+; CHECK: [[EXIT]]:
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a49", "~{a49}"()
+ call void @use_most_hi()
+ br i1 %c, label %rec, label %exit
+rec:
+ call void @pong(i1 %c)
+ br label %exit
+exit:
+ ret void
+}
+
+define internal void @pong(i1 %c) {
+; CHECK-LABEL: define internal void @pong(
+; CHECK-SAME: i1 [[C:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: call void @ping(i1 [[C]])
+; CHECK-NEXT: call void @use_most_hi()
+; CHECK-NEXT: ret void
+;
+ call void @ping(i1 %c)
+ call void @use_most_hi()
+ ret void
+}
+
+define amdgpu_kernel void @k_mutual(i1 %c) {
+; CHECK-LABEL: define amdgpu_kernel void @k_mutual(
+; CHECK-SAME: i1 [[C:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: call void @ping(i1 [[C]])
+; CHECK-NEXT: ret void
+;
+ call void @ping(i1 %c)
+ ret void
+}
+
+; Second, independent component with its own @use_most copy.
+define internal void @use_most_lo() {
+; CHECK-LABEL: define internal void @use_most_lo(
+; CHECK-SAME: ) #[[ATTR4:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+; Self recursion is the simplest cycle: @self_rec calls itself and clobbers
+; a31 -> 32. It must converge to 32 without being dragged to the other
+; component's 64.
+define internal void @self_rec(i1 %c) {
+; CHECK-LABEL: define internal void @self_rec(
+; CHECK-SAME: i1 [[C:%.*]]) #[[ATTR5:[0-9]+]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @use_most_lo()
+; CHECK-NEXT: br i1 [[C]], label %[[REC:.*]], label %[[EXIT:.*]]
+; CHECK: [[REC]]:
+; CHECK-NEXT: call void @self_rec(i1 [[C]])
+; CHECK-NEXT: br label %[[EXIT]]
+; CHECK: [[EXIT]]:
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a31", "~{a31}"()
+ call void @use_most_lo()
+ br i1 %c, label %rec, label %exit
+rec:
+ call void @self_rec(i1 %c)
+ br label %exit
+exit:
+ ret void
+}
+
+define amdgpu_kernel void @k_self(i1 %c) {
+; CHECK-LABEL: define amdgpu_kernel void @k_self(
+; CHECK-SAME: i1 [[C:%.*]]) #[[ATTR5]] {
+; CHECK-NEXT: call void @self_rec(i1 [[C]])
+; CHECK-NEXT: ret void
+;
+ call void @self_rec(i1 %c)
+ ret void
+}
+
+;; ===========================================================================
+;; Scenario 4 (shared callee): regression test for the shared-callee leak.
+;; shared_k1 -> shared_f1 (requires AGPRs)
+;; shared_k1 -> shared_f2
+;; shared_k2 -> shared_f2 (requires no AGPRs)
+;; Only real AGPR demand propagates *up*, so shared_k2 (which reaches only
+;; shared_f2) must NOT inherit an AGPR reservation from shared_k1 just because
+;; they share shared_f2. shared_f2 still gets a downward budget
+;; (amdgpu-register-budget) without claiming a real demand that leaks to k2.
+;;
+;; A B
+;; / \ /
+;; C D
+;; A = shared_k1 : real 0, register-budget 78,50
+;; B = shared_k2 : real 0, register-budget 128,0
+;; C = shared_f1 : real 50, register-budget 78,50
+;; D = shared_f2 : real 0, register-budget 78,0
+;;
+;; ===========================================================================
+
+define internal void @shared_use_most() {
+; CHECK-LABEL: define internal void @shared_use_most(
+; CHECK-SAME: ) #[[ATTR6:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+; shared_f1: forces 50 AGPRs.
+define internal void @shared_f1() {
+; CHECK-LABEL: define internal void @shared_f1(
+; CHECK-SAME: ) #[[ATTR1]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @shared_use_most()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a49", "~{a49}"()
+ call void @shared_use_most()
+ ret void
+}
+
+; shared_f2: shared between shared_k1 and shared_k2, needs no AGPRs of its own.
+define internal void @shared_f2(ptr %p) {
+; CHECK-LABEL: define internal void @shared_f2(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR6]] {
+; CHECK-NEXT: [[V:%.*]] = load volatile <128 x i32>, ptr [[P]], align 512
+; CHECK-NEXT: store volatile <128 x i32> [[V]], ptr [[P]], align 512
+; CHECK-NEXT: call void @shared_use_most()
+; CHECK-NEXT: ret void
+;
+ %v = load volatile <128 x i32>, ptr %p
+ store volatile <128 x i32> %v, ptr %p
+ call void @shared_use_most()
+ ret void
+}
+
+; shared_k1 reaches the AGPR-hungry shared_f1 and the shared shared_f2.
+define amdgpu_kernel void @shared_k1(ptr %p) {
+; CHECK-LABEL: define amdgpu_kernel void @shared_k1(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR1]] {
+; CHECK-NEXT: call void @shared_f1()
+; CHECK-NEXT: call void @shared_f2(ptr [[P]])
+; CHECK-NEXT: ret void
+;
+ call void @shared_f1()
+ call void @shared_f2(ptr %p)
+ ret void
+}
+
+; shared_k2 reaches only shared_f2 and must not inherit shared_k1's split.
+define amdgpu_kernel void @shared_k2(ptr %p) {
+; CHECK-LABEL: define amdgpu_kernel void @shared_k2(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR7:[0-9]+]] {
+; CHECK-NEXT: call void @shared_f2(ptr [[P]])
+; CHECK-NEXT: ret void
+;
+ call void @shared_f2(ptr %p)
+ ret void
+}
+
+;; ;; ===========================================================================
+;; Scenario 5: one kernel reserves no AGPRs at all, so the shared callee gets
+;; no AGPR headroom even though the other kernel reserved 16.
+;;
+;; A B
+;; \ /
+;; C
+;;
+;; A = @noroom_kernel_agpr16 : agpr 16 vgpr 128 - 16 -> 112,16
+;; B = @noroom_kernel_noagpr : agpr 0 vgpr 128 - 0 -> 128,0
+;; C = @noroom_shared 112,16 | 128,0 -> 112,0
+;;
+;; C may use 112 VGPRs and 128 - 128 = 0 AGPRs: any AGPR use would push it past
+;; the 128-register limit when it runs under B's accum_offset.
+;; ===========================================================================
+
+define internal void @noroom_use_most() {
+; CHECK-LABEL: define internal void @noroom_use_most(
+; CHECK-SAME: ) #[[ATTR8:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+; C: shared by both kernels, so it takes the VGPR ceiling from A and the AGPR
+; ceiling from B.
+define internal void @noroom_shared() {
+; CHECK-LABEL: define internal void @noroom_shared(
+; CHECK-SAME: ) #[[ATTR8]] {
+; CHECK-NEXT: call void @noroom_use_most()
+; CHECK-NEXT: ret void
+;
+ call void @noroom_use_most()
+ ret void
+}
+
+define amdgpu_kernel void @noroom_kernel_agpr16() {
+; CHECK-LABEL: define amdgpu_kernel void @noroom_kernel_agpr16(
+; CHECK-SAME: ) #[[ATTR9:[0-9]+]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @noroom_shared()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a15", "~{a15}"()
+ call void @noroom_shared()
+ ret void
+}
+
+define amdgpu_kernel void @noroom_kernel_noagpr() {
+; CHECK-LABEL: define amdgpu_kernel void @noroom_kernel_noagpr(
+; CHECK-SAME: ) #[[ATTR7]] {
+; CHECK-NEXT: call void @noroom_shared()
+; CHECK-NEXT: ret void
+;
+ call void @noroom_shared()
+ ret void
+}
+
+;; ;; ===========================================================================
+;; Scenario 6: both kernels reserve some AGPRs, so the shared callee keeps the
+;; headroom common to both.
+;;
+;; A B
+;; \ /
+;; C
+;;
+;; A = @room_kernel_agpr16 : 16 AGPRs 128 - 16 -> 112,16
+;; B = @room_kernel_agpr12 : 12 AGPRs 128 - 12 -> 116,12
+;; C = @room_shared 112,16 | 116,12 -> 112,12
+;;
+;; C may use 112 VGPRs and 128 - 116 = 12 AGPRs, which it could spill into.
+;; Note this is the smaller of the two kernels' reservations, not the larger.
+;; ===========================================================================
+
+define internal void @room_use_most() {
+; CHECK-LABEL: define internal void @room_use_most(
+; CHECK-SAME: ) #[[ATTR10:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+define internal void @room_shared() {
+; CHECK-LABEL: define internal void @room_shared(
+; CHECK-SAME: ) #[[ATTR10]] {
+; CHECK-NEXT: call void @room_use_most()
+; CHECK-NEXT: ret void
+;
+ call void @room_use_most()
+ ret void
+}
+
+define amdgpu_kernel void @room_kernel_agpr16() {
+; CHECK-LABEL: define amdgpu_kernel void @room_kernel_agpr16(
+; CHECK-SAME: ) #[[ATTR9]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @room_shared()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a15", "~{a15}"()
+ call void @room_shared()
+ ret void
+}
+
+define amdgpu_kernel void @room_kernel_agpr12() {
+; CHECK-LABEL: define amdgpu_kernel void @room_kernel_agpr12(
+; CHECK-SAME: ) #[[ATTR11:[0-9]+]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @room_shared()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a11", "~{a11}"()
+ call void @room_shared()
+ ret void
+}
+
+;; ;; ===========================================================================
+;; Scenario 7: the same merge seen from a single kernel fanning out.
+;;
+;; A
+;; / \
+;; B C
+;; A = @fanout_kernel : agpr 16 | vgpr 128 - 16 = 112,16
+;; B = @fanout_agpr : clobbers a15, so 16 flows up into A
+;; C = @fanout_noagpr : uses no AGPRs of its own, but inheriets 112,16
+;;
+;; ===========================================================================
+
+define internal void @fanout_use_most() {
+; CHECK-LABEL: define internal void @fanout_use_most(
+; CHECK-SAME: ) #[[ATTR12:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP8:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP9:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP10:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP12:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+define internal void @fanout_sink() {
+; CHECK-LABEL: define internal void @fanout_sink(
+; CHECK-SAME: ) #[[ATTR12]] {
+; CHECK-NEXT: call void @fanout_use_most()
+; CHECK-NEXT: ret void
+;
+ call void @fanout_use_most()
+ ret void
+}
+
+define internal void @fanout_agpr() {
+; CHECK-LABEL: define internal void @fanout_agpr(
+; CHECK-SAME: ) #[[ATTR9]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @fanout_sink()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a15", "~{a15}"()
+ call void @fanout_sink()
+ ret void
+}
+
+define internal void @fanout_noagpr(ptr %p) {
+; CHECK-LABEL: define internal void @fanout_noagpr(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR12]] {
+; CHECK-NEXT: [[V:%.*]] = load volatile <128 x i32>, ptr [[P]], align 512
+; CHECK-NEXT: store volatile <128 x i32> [[V]], ptr [[P]], align 512
+; CHECK-NEXT: call void @fanout_sink()
+; CHECK-NEXT: ret void
+;
+ %v = load volatile <128 x i32>, ptr %p
+ store volatile <128 x i32> %v, ptr %p
+ call void @fanout_sink()
+ ret void
+}
+
+define amdgpu_kernel void @fanout_kernel(ptr %p) {
+; CHECK-LABEL: define amdgpu_kernel void @fanout_kernel(
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR9]] {
+; CHECK-NEXT: call void @fanout_agpr()
+; CHECK-NEXT: call void @fanout_noagpr(ptr [[P]])
+; CHECK-NEXT: ret void
+;
+ call void @fanout_agpr()
+ call void @fanout_noagpr(ptr %p)
+ ret void
+}
+;.
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="76,52" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="50" "amdgpu-no-wwm" "amdgpu-register-budget"="76,52" }
+; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="50" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="96,32" }
+; CHECK: attributes #[[ATTR5]] = { "amdgpu-agpr-alloc"="32" "amdgpu-no-wwm" "amdgpu-register-budget"="96,32" }
+; CHECK: attributes #[[ATTR6]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="76,0" }
+; CHECK: attributes #[[ATTR7]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
+; CHECK: attributes #[[ATTR8]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="112,0" }
+; CHECK: attributes #[[ATTR9]] = { "amdgpu-agpr-alloc"="16" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
+; CHECK: attributes #[[ATTR10]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="112,12" }
+; CHECK: attributes #[[ATTR11]] = { "amdgpu-agpr-alloc"="12" "amdgpu-no-wwm" "amdgpu-register-budget"="116,12" }
+; CHECK: attributes #[[ATTR12]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
+; CHECK: attributes #[[ATTR13:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR14:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
+;.
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-work-group-size.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-work-group-size.ll
new file mode 100644
index 0000000000000..2e4ba2d30e445
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-work-group-size.ll
@@ -0,0 +1,205 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --check-attributes --check-globals all --version 6
+; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -passes=amdgpu-attributor %s | FileCheck -check-prefixes=CHECK,GFX90A %s
+; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -passes=amdgpu-attributor %s | FileCheck -check-prefixes=CHECK,GFX942 %s
+
+; Why the downward propagation carries a VGPR budget rather than an AGPR count.
+;
+; The size of the vector register file a wave may use is not a constant: it
+; falls out of the occupancy needed to satisfy the kernel's launch bounds. On
+; gfx90a and gfx942 a max flat work group size of 1024 needs 4 waves, cutting
+; the 512-register file down to 128, while a max of 512 leaves 256. Two kernels
+; sharing a callee can therefore be compiled against completely different
+; totals, and the attributor deliberately widens the callee's flat work group
+; size to the union of its callers ([1,512] U [1,1024] = [1,1024], the default
+; range, so no attribute is emitted on the callee).
+;
+; An AGPR count alone cannot be propagated across that difference, because the
+; same count means different things under different totals. Folding the work
+; group size and the AGPR requirement together into a single VGPR budget per
+; kernel makes the number self-contained, so it can be safely merged in the
+; callee.
+;
+; A B
+; \ /
+; C
+
+;; ===========================================================================
+;; Scenario 1: the AGPR-hungry kernel is the one with the *larger* register
+;; file, so its budget is still tighter than the other kernel's.
+;;
+;; A = @wgs_kernel_512_agpr130 : [1,512], agpr 130 -> 256 - 130 = 126
+;; B = @wgs_kernel_1024_noagpr : [1,1024], agpr 0 -> 128 - 0 = 128
+;; C = @wgs_shared -> 126,128
+;;
+;; C decodes to a ceiling of 126 VGPRs, and 0 AGPRs because its own total is
+;; 128 and the largest budget it must fit under is already 128.
+;; ===========================================================================
+
+; C. Also serves as the sink that keeps the inferred attribute list short by
+; preventing the attributor from proving the implicit inputs are unused.
+define internal void @wgs_shared() {
+; CHECK-LABEL: define internal void @wgs_shared(
+; CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP11:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP12:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP8:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP9:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP10:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+; A: a 512-wide launch gets a 256-register file, and clobbering a129 forces an
+; AGPR requirement of 130, leaving 126 VGPRs.
+define amdgpu_kernel void @wgs_kernel_512_agpr130() #0 {
+; CHECK-LABEL: define amdgpu_kernel void @wgs_kernel_512_agpr130(
+; CHECK-SAME: ) #[[ATTR1:[0-9]+]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @wgs_shared()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a129", "~{a129}"()
+ call void @wgs_shared()
+ ret void
+}
+
+; B: a 1024-wide launch gets a 128-register file and needs no AGPRs.
+define amdgpu_kernel void @wgs_kernel_1024_noagpr() #1 {
+; CHECK-LABEL: define amdgpu_kernel void @wgs_kernel_1024_noagpr(
+; CHECK-SAME: ) #[[ATTR2:[0-9]+]] {
+; CHECK-NEXT: call void @wgs_shared()
+; CHECK-NEXT: ret void
+;
+ call void @wgs_shared()
+ ret void
+}
+
+;; ===========================================================================
+;; Scenario 2: the wide-launch kernel's budget exceeds the narrow-launch
+;; kernel's *total*, so decoding the AGPR ceiling has to saturate.
+;;
+;; A = @sat_kernel_512_agpr56 : [1,512], agpr 56 -> 256 - 56 = 200
+;; B = @sat_kernel_1024_noagpr: [1,1024], agpr 0 -> 128 - 0 = 128
+;; C = @sat_shared -> 128,200
+;;
+;; C decodes to a ceiling of 128 VGPRs. Its own total is 128 and the largest
+;; budget it must fit under is 200, so the AGPR ceiling saturates at 0 rather
+;; than going negative.
+;; ===========================================================================
+
+define internal void @sat_shared() {
+; CHECK-LABEL: define internal void @sat_shared(
+; CHECK-SAME: ) #[[ATTR3:[0-9]+]] {
+; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
+; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
+; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
+; CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.workitem.id.y()
+; CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workitem.id.z()
+; CHECK-NEXT: [[TMP4:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
+; CHECK-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+; CHECK-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
+; CHECK-NEXT: [[TMP11:%.*]] = call i32 @llvm.amdgcn.cluster.id.x()
+; CHECK-NEXT: [[TMP12:%.*]] = call i32 @llvm.amdgcn.cluster.id.y()
+; CHECK-NEXT: [[TMP13:%.*]] = call i32 @llvm.amdgcn.cluster.id.z()
+; CHECK-NEXT: [[TMP7:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: [[TMP8:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+; CHECK-NEXT: [[TMP9:%.*]] = call i64 @llvm.amdgcn.dispatch.id()
+; CHECK-NEXT: [[TMP10:%.*]] = call i32 @llvm.amdgcn.lds.kernel.id()
+; CHECK-NEXT: [[IMPLICIT_ARG_PTR:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+; CHECK-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr [[ALLOCA_CAST]], ptr addrspace(4) [[IMPLICIT_ARG_PTR]], i64 256, i1 false)
+; CHECK-NEXT: ret void
+;
+ %alloca = alloca [256 x i8], addrspace(5)
+ %alloca.cast = addrspacecast ptr addrspace(5) %alloca to ptr
+ call i32 @llvm.amdgcn.workitem.id.x()
+ call i32 @llvm.amdgcn.workitem.id.y()
+ call i32 @llvm.amdgcn.workitem.id.z()
+ call i32 @llvm.amdgcn.workgroup.id.x()
+ call i32 @llvm.amdgcn.workgroup.id.y()
+ call i32 @llvm.amdgcn.workgroup.id.z()
+ call i32 @llvm.amdgcn.cluster.id.x()
+ call i32 @llvm.amdgcn.cluster.id.y()
+ call i32 @llvm.amdgcn.cluster.id.z()
+ call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+ call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
+ call i64 @llvm.amdgcn.dispatch.id()
+ call i32 @llvm.amdgcn.lds.kernel.id()
+ %implicit.arg.ptr = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+ call void @llvm.memcpy.p0.p4(ptr %alloca.cast, ptr addrspace(4) %implicit.arg.ptr, i64 256, i1 false)
+ ret void
+}
+
+define amdgpu_kernel void @sat_kernel_512_agpr56() #0 {
+; CHECK-LABEL: define amdgpu_kernel void @sat_kernel_512_agpr56(
+; CHECK-SAME: ) #[[ATTR4:[0-9]+]] {
+; CHECK-NEXT: call void asm sideeffect "
+; CHECK-NEXT: call void @sat_shared()
+; CHECK-NEXT: ret void
+;
+ call void asm sideeffect "; touch a55", "~{a55}"()
+ call void @sat_shared()
+ ret void
+}
+
+define amdgpu_kernel void @sat_kernel_1024_noagpr() #1 {
+; CHECK-LABEL: define amdgpu_kernel void @sat_kernel_1024_noagpr(
+; CHECK-SAME: ) #[[ATTR2]] {
+; CHECK-NEXT: call void @sat_shared()
+; CHECK-NEXT: ret void
+;
+ call void @sat_shared()
+ ret void
+}
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1,512" }
+attributes #1 = { "amdgpu-flat-work-group-size"="1,1024" }
+;.
+; GFX90A: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="124,0" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="130" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="124,132" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="56" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="200,56" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR5:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR6:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) "target-cpu"="gfx90a" }
+;.
+; GFX942: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="124,0" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="130" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="124,132" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="56" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="200,56" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR5:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR6:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) "target-cpu"="gfx942" }
+;.
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; GFX90A: {{.*}}
+; GFX942: {{.*}}
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll
index 31d196655bc14..1ef6a3596a81b 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll
@@ -58,7 +58,7 @@ attributes #1 = { nocallback "trap-func-name"="handler" }
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
; CHECK: attributes #[[ATTR1:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
; CHECK: attributes #[[ATTR3]] = { "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR4]] = { "trap-func-name"="handler" }
; CHECK: attributes #[[ATTR5]] = { nocallback "trap-func-name"="handler" }
>From 728badb06cd4a162ce4a6972d499acff149caf5a Mon Sep 17 00:00:00 2001
From: JoshuaGrindstaff <Joshua.Grindstaff at amd.com>
Date: Tue, 25 Aug 2026 18:52:19 -0500
Subject: [PATCH 4/4] Change name to accum-offset and removed AGPRbudget
---
llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp | 146 ++++++++--------
...dgpu-attributor-accum-offset-propagate.ll} | 163 +++++++++---------
...ttributor-accum-offset-work-group-size.ll} | 50 +++---
.../amdgpu-attributor-min-agpr-alloc.ll | 60 +++----
.../AMDGPU/amdgpu-attributor-trap-leaf.ll | 2 +-
5 files changed, 209 insertions(+), 212 deletions(-)
rename llvm/test/CodeGen/AMDGPU/{amdgpu-attributor-register-budget-propagate.ll => amdgpu-attributor-accum-offset-propagate.ll} (86%)
rename llvm/test/CodeGen/AMDGPU/{amdgpu-attributor-register-budget-work-group-size.ll => amdgpu-attributor-accum-offset-work-group-size.ll} (77%)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
index 5b9f974e70345..b5aaca72c1154 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
@@ -1433,60 +1433,59 @@ struct AAAMDGPUMinAGPRAlloc
};
const char AAAMDGPUMinAGPRAlloc::ID = 0;
-/// The register-file split a function has been committed to by its callers.
-/// \p VGPRs and \p AGPRs are both absolute ceilings, so their sum stays within
-/// the total vector register budget of the tightest caller.
-struct RegisterBudgetState {
- unsigned VGPRs = 0;
- unsigned AGPRs = 0;
+/// The accum_offset a function has been committed to by its callers, that is,
+/// the ceiling on the number of architectural VGPRs it may allocate.
+struct AccumOffsetState {
+ unsigned Offset = 0;
bool Unknown = true;
- bool operator==(const RegisterBudgetState &Other) const {
- return Unknown == Other.Unknown && VGPRs == Other.VGPRs &&
- AGPRs == Other.AGPRs;
+ bool operator==(const AccumOffsetState &Other) const {
+ return Unknown == Other.Unknown && Offset == Other.Offset;
}
- bool operator!=(const RegisterBudgetState &Other) const {
+ bool operator!=(const AccumOffsetState &Other) const {
return !(*this == Other);
}
- /// Combine with the budget of one caller. A function reachable from several
- /// callers has to fit inside all of their splits, so take the tightest
- /// ceiling on each axis independently.
- void merge(const RegisterBudgetState &Other) {
- assert(!Other.Unknown && "cannot merge an unknown budget");
+ /// Combine with the boundary of one caller. AGPRs are addressed relative to
+ /// accum_offset, so a function reachable from several callers has to fit
+ /// under the lowest boundary any of them committed to.
+ void merge(const AccumOffsetState &Other) {
+ assert(!Other.Unknown && "cannot merge an unknown accum offset");
if (Unknown) {
*this = Other;
return;
}
- VGPRs = std::min(VGPRs, Other.VGPRs);
- AGPRs = std::min(AGPRs, Other.AGPRs);
+ Offset = std::min(Offset, Other.Offset);
}
};
-/// An abstract attribute to propagate the register file split a kernel was
-/// compiled with down the call graph to its device functions, emitted as
-/// "amdgpu-register-budget". Entry functions seed the split from their own
-/// vector register budget; every other function inherits the tightest split
-/// over its callers.
-struct AAAMDGPURegisterBudget
+/// An abstract attribute to propagate the accum_offset a kernel was compiled
+/// with down the call graph to its device functions, emitted as
+/// "amdgpu-accum-offset". Entry functions seed the boundary from their own
+/// vector register budget and AGPR requirement; every other function inherits
+/// the lowest boundary over its callers.
+///
+/// The AGPR side of the split needs no propagation: for a non-entry function
+/// getMaxNumVectorRegs pins the AGPR ceiling to the value of
+/// "amdgpu-agpr-alloc", which AAAMDGPUMinAGPRAlloc has already computed as that
+/// function's own requirement.
+struct AAAMDGPUAccumOffset
: public StateWrapper<BooleanState, AbstractAttribute> {
using Base = StateWrapper<BooleanState, AbstractAttribute>;
- AAAMDGPURegisterBudget(const IRPosition &IRP, Attributor &A) : Base(IRP) {}
+ AAAMDGPUAccumOffset(const IRPosition &IRP, Attributor &A) : Base(IRP) {}
- static AAAMDGPURegisterBudget &createForPosition(const IRPosition &IRP,
- Attributor &A) {
+ static AAAMDGPUAccumOffset &createForPosition(const IRPosition &IRP,
+ Attributor &A) {
if (IRP.getPositionKind() == IRPosition::IRP_FUNCTION)
- return *new (A.Allocator) AAAMDGPURegisterBudget(IRP, A);
- llvm_unreachable(
- "AAAMDGPURegisterBudget is only valid for function position");
+ return *new (A.Allocator) AAAMDGPUAccumOffset(IRP, A);
+ llvm_unreachable("AAAMDGPUAccumOffset is only valid for function position");
}
- // When we known not all callers are known, we know that any unknown caller
- // that reaches this function will have a pessimistic amdgpu-agpr-alloc
- // attribute. This being pessimistic means that the budget for the number of
- // VGPRs and AGPRs will be split in half for the unknown caller. We also know
- // that the FlatWorkGroupSize attribute will also be pessimistic for at least
- // the current function (the one being called in the indirect callsite)
+ // When not all callers are known, any unknown caller that reaches this
+ // function will itself have a pessimistic amdgpu-agpr-alloc attribute, which
+ // splits its register file in half. We also know that the FlatWorkGroupSize
+ // attribute will be pessimistic for at least the current function (the one
+ // being called in the indirect callsite).
unsigned computePessimisticValue(Attributor &A) const {
Function *F = getAssociatedFunction();
auto &InfoCache = static_cast<AMDGPUInformationCache &>(A.getInfoCache());
@@ -1501,14 +1500,13 @@ struct AAAMDGPURegisterBudget
ChangeStatus updateImpl(Attributor &A) override {
Function *F = getAssociatedFunction();
- RegisterBudgetState OldState = Budget;
+ AccumOffsetState OldState = AccumOffset;
- // The budget is recomputed from scratch on every update rather than
- // accumulated, because the seed itself moves: a kernel's split shrinks as
+ // The boundary is recomputed from scratch on every update rather than
+ // accumulated, because the seed itself moves: a kernel's boundary drops as
// AAAMDGPUMinAGPRAlloc climbs, and callees have to follow it down.
if (AMDGPU::isEntryFunctionCC(F->getCallingConv())) {
- unsigned VGPRBudget = 0;
- unsigned AGPRBudget = 0;
+ unsigned Offset = 0;
const auto *AGPRAlloc = A.getAAFor<AAAMDGPUMinAGPRAlloc>(
*this, IRPosition::function(*F), DepClassTy::OPTIONAL);
@@ -1516,20 +1514,18 @@ struct AAAMDGPURegisterBudget
const GCNSubtarget &ST = InfoCache.TM.getSubtarget<GCNSubtarget>(*F);
unsigned MaxRegs = ST.getMaxNumVGPRs(*F);
if (!AGPRAlloc || !AGPRAlloc->isValidState()) {
- VGPRBudget = AGPRBudget = MaxRegs / 2; // pessimistic
- if (VGPRBudget == computePessimisticValue(A))
+ Offset = MaxRegs / 2; // pessimistic
+ if (Offset == computePessimisticValue(A))
return indicatePessimisticFixpoint();
} else {
- AGPRBudget = AGPRAlloc->getAssumed();
- AGPRBudget = alignTo(AGPRBudget, 4);
- VGPRBudget = MaxRegs - std::min(MaxRegs, AGPRBudget);
+ Offset = MaxRegs - std::min(MaxRegs, AGPRAlloc->getAssumed());
}
- Budget = {VGPRBudget, AGPRBudget, /*Unknown=*/false};
- LLVM_DEBUG(dbgs() << "Register budget for " << F->getName() << ": "
- << VGPRBudget << ", " << AGPRBudget << "\n");
+ AccumOffset = {Offset, /*Unknown=*/false};
+ LLVM_DEBUG(dbgs() << "Accum offset for " << F->getName() << ": " << Offset
+ << "\n");
} else {
- RegisterBudgetState Merged;
+ AccumOffsetState Merged;
auto CheckUse = [&](const Use &U, bool &Follow) {
if (auto *CE = dyn_cast<ConstantExpr>(U.getUser())) {
@@ -1549,14 +1545,14 @@ struct AAAMDGPURegisterBudget
return true;
Function *Caller = ACS.getInstruction()->getFunction();
- const auto *CallerAA = A.getAAFor<AAAMDGPURegisterBudget>(
+ const auto *CallerAA = A.getAAFor<AAAMDGPUAccumOffset>(
*this, IRPosition::function(*Caller), DepClassTy::REQUIRED);
if (!CallerAA || !CallerAA->isValidState())
return true;
- const RegisterBudgetState &CallerBudget = CallerAA->getBudget();
- if (!CallerBudget.Unknown)
- Merged.merge(CallerBudget);
+ const AccumOffsetState &CallerOffset = CallerAA->getAccumOffset();
+ if (!CallerOffset.Unknown)
+ Merged.merge(CallerOffset);
return true;
};
@@ -1566,28 +1562,24 @@ struct AAAMDGPURegisterBudget
bool DummyUAI = false;
bool AllCallsitesKnown = A.checkForAllCallSites(
[](AbstractCallSite) { return true; }, *this, true, DummyUAI);
- if (!AllCallsitesKnown && !Merged.Unknown) {
- if (std::optional<unsigned> PessimisticValue =
- computePessimisticValue(A)) {
- Merged.merge(
- {*PessimisticValue, *PessimisticValue, /*Unknown=*/false});
- } else
- return indicatePessimisticFixpoint();
- }
- // Stays unknown when no caller contributed a budget, so that functions
+ if (!AllCallsitesKnown && !Merged.Unknown)
+ Merged.merge({computePessimisticValue(A), /*Unknown=*/false});
+
+ // Stays unknown when no caller contributed a boundary, so that functions
// outside any kernel's reach are left unconstrained.
- Budget = Merged;
+ AccumOffset = Merged;
}
- return OldState == Budget ? ChangeStatus::UNCHANGED : ChangeStatus::CHANGED;
+ return OldState == AccumOffset ? ChangeStatus::UNCHANGED
+ : ChangeStatus::CHANGED;
}
ChangeStatus manifest(Attributor &A) override {
- if (Budget.Unknown)
+ if (AccumOffset.Unknown)
return ChangeStatus::UNCHANGED;
- SmallString<16> Buffer;
+ SmallString<8> Buffer;
raw_svector_ostream OS(Buffer);
- OS << Budget.VGPRs << ',' << Budget.AGPRs;
+ OS << AccumOffset.Offset;
return A.manifestAttrs(
getIRPosition(),
{Attribute::get(getAssociatedFunction()->getContext(), AttrName,
@@ -1595,24 +1587,24 @@ struct AAAMDGPURegisterBudget
/*ForceReplace=*/true);
}
- const RegisterBudgetState &getBudget() const { return Budget; }
+ const AccumOffsetState &getAccumOffset() const { return AccumOffset; }
const std::string getAsStr(Attributor *A) const override {
- if (!getAssumed() || Budget.Unknown)
+ if (!getAssumed() || AccumOffset.Unknown)
return "unknown";
std::string Str;
raw_string_ostream OS(Str);
- OS << AttrName << '=' << Budget.VGPRs << ',' << Budget.AGPRs;
+ OS << AttrName << '=' << AccumOffset.Offset;
return Str;
}
void trackStatistics() const override {}
- StringRef getName() const override { return "AAAMDGPURegisterBudget"; }
+ StringRef getName() const override { return "AAAMDGPUAccumOffset"; }
const char *getIdAddr() const override { return &ID; }
/// This function should return true if the type of the \p AA is
- /// AAAMDGPURegisterBudget.
+ /// AAAMDGPUAccumOffset.
static bool classof(const AbstractAttribute *AA) {
return AA->getIdAddr() == &ID;
}
@@ -1620,12 +1612,12 @@ struct AAAMDGPURegisterBudget
static const char ID;
private:
- RegisterBudgetState Budget;
+ AccumOffsetState AccumOffset;
- static constexpr char AttrName[] = "amdgpu-register-budget";
+ static constexpr char AttrName[] = "amdgpu-accum-offset";
};
-const char AAAMDGPURegisterBudget::ID = 0;
+const char AAAMDGPUAccumOffset::ID = 0;
/// An abstract attribute to propagate the function attribute
/// "amdgpu-cluster-dims" from kernel entry functions to device functions.
@@ -1790,7 +1782,7 @@ static bool runImpl(SetVector<Function *> &Functions, bool IsModulePass,
{&AAAMDAttributes::ID, &AAUniformWorkGroupSize::ID,
&AAPotentialValues::ID, &AAAMDFlatWorkGroupSize::ID,
&AAAMDMaxNumWorkgroups::ID, &AAAMDWavesPerEU::ID,
- &AAAMDGPUMinAGPRAlloc::ID, &AAAMDGPURegisterBudget::ID, &AACallEdges::ID,
+ &AAAMDGPUMinAGPRAlloc::ID, &AAAMDGPUAccumOffset::ID, &AACallEdges::ID,
&AAPointerInfo::ID, &AAPotentialConstantValues::ID,
&AAUnderlyingObjects::ID, &AANoAliasAddrSpace::ID, &AAAddressSpace::ID,
&AAIndirectCallInfo::ID, &AAAMDGPUClusterDims::ID, &AAAlign::ID});
@@ -1837,7 +1829,7 @@ static bool runImpl(SetVector<Function *> &Functions, bool IsModulePass,
if (ST.hasGFX90AInsts()) {
A.getOrCreateAAFor<AAAMDGPUMinAGPRAlloc>(IRPosition::function(*F));
- A.getOrCreateAAFor<AAAMDGPURegisterBudget>(IRPosition::function(*F));
+ A.getOrCreateAAFor<AAAMDGPUAccumOffset>(IRPosition::function(*F));
}
for (auto &I : instructions(F)) {
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-propagate.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-accum-offset-propagate.ll
similarity index 86%
rename from llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-propagate.ll
rename to llvm/test/CodeGen/AMDGPU/amdgpu-attributor-accum-offset-propagate.ll
index dc16d37b89857..45c5d742f71be 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-propagate.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-accum-offset-propagate.ll
@@ -3,30 +3,32 @@
; Consolidated tests for AGPR/VGPR split propagation by the AMDGPU attributor.
; The AGPR demand is propagated *up* the call graph (callee -> caller) via
-; AAAMDGPUMinAGPRAlloc and emitted as "amdgpu-agpr-alloc". The register budget
-; is propagated *down* (caller -> callee) via AAAMDGPURegisterBudget and
-; emitted as "amdgpu-register-budget", which caps a callee's registers so the
-; split the caller was compiled with is runnable with all callees. The
-; attribute is always a "VGPRs,AGPRs" pair of absolute ceilings, and a function
-; reachable from several callers takes the smallest ceiling on each axis
-; independently. The total budget is 128 here (the default flatblocksize/wavesperEU give 128 for 512 register file), so an unknown/indirect caller,
-; having itself been pessimised to a half split, contributes 64,64.
+; AAAMDGPUMinAGPRAlloc and emitted as "amdgpu-agpr-alloc". The accum_offset a
+; kernel was compiled with is propagated *down* (caller -> callee) via
+; AAAMDGPUAccumOffset and emitted as "amdgpu-accum-offset", which caps a
+; callee's architectural VGPRs so the boundary the caller was compiled with is
+; runnable with all callees.
+;
+; The total budget is 128 here (the default flat-work-group-size and
+; waves-per-eu give 128 out of a 512 register file), so an unknown or indirect
+; caller, having itself been pessimised to a half split, contributes 64.
;
; Each scenario below is an independent call-graph component (its own copy of
; the @*_use_most sink) so the scenarios converge independently.
;; ===========================================================================
;; Scenario 1 (direct): a kernel directly calls an AGPR-hungry leaf and an
-;; all-VGPR leaf. The all-VGPR leaf must be pulled down to the kernel's split
-;; (78,50) so the register allocator cannot grab the whole VGPR file.
+;; all-VGPR leaf. The all-VGPR leaf must be pulled down to the kernel's boundary
+;; (78) so the register allocator cannot grab the whole VGPR file.
;;
;; A
;; / \
;; B C
;;
-;; A = k_direct : actual agpr - 0, agpr attribute - 50 register-budget 78,50
-;; B = direct_leaf_hivgpr : actual agpr - 0, agpr attribute - 0 register-budget 78,50
-;; C = direct_leaf_agpr : actual agpr - 50, agpr attribute - 50 register -budget 78,50
+;; A = k_direct : actual agpr 0, agpr attribute 50, accum-offset 78
+;; B = direct_leaf_hivgpr : actual agpr 0, agpr attribute 0, accum-offset 78
+;; C = direct_leaf_agpr : actual agpr 50, agpr attribute 50, accum-offset 78
+;;
;; ===========================================================================
; Shrink result attribute list by preventing use of most attributes.
@@ -89,7 +91,7 @@ define internal void @direct_leaf_agpr() {
}
; An all-VGPR leaf with no AGPR references of its own. It must be pulled to the
-; caller's split (78,50) via downward propagation rather than staying unbounded.
+; caller's boundary (78) by downward propagation rather than staying unbounded.
define internal void @direct_leaf_hivgpr(ptr %p) {
; CHECK-LABEL: define internal void @direct_leaf_hivgpr(
; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR0]] {
@@ -122,17 +124,18 @@ define amdgpu_kernel void @k_direct(ptr %p) {
;; Scenario 2 (indirect): the only kernel makes an unknown indirect call, so its
;; AGPR demand is unbounded and it ends up with no attributes at all.
;; With "amdgpu-agpr-alloc" absent the consumer falls back to the same half split.
-;; The indirectly-reachable leaves are externally callable, so not all of their
-;; callers are known and they inherit the half-split budget an unknown caller
-;; would impose Upward real demand is still annotated
+;; The indirectly-reachable leaves are externally callable, so no known caller
+;; contributes a boundary and they are left without the attribute. Its absence
+;; means the same half split an unknown caller would impose. Upward real demand
+;; is still annotated.
;;
;; A
;; / \
;; B C
;;
-;; C = indirect_leaf_agpr : real agpr 50, no register-budget attribute -> 64,64
-;; B = indirect_leaf_hivgpr : real agpr 0, no register-budget attribute -> 64,64
-;; A = k_indirect : no attributes -> 64,64
+;; C = indirect_leaf_agpr : real agpr 50, no accum-offset attribute -> 64
+;; B = indirect_leaf_hivgpr : real agpr 0, no accum-offset attribute -> 64
+;; A = k_indirect : no attributes -> 64
;; ===========================================================================
define internal void @indirect_use_most() {
@@ -221,12 +224,11 @@ define amdgpu_kernel void @k_indirect() {
}
;; ===========================================================================
-;; Scenario 3 (cycle): This scenario show how a cycle affects the up then
-;; down propogation
-;; use_most_hi : real 0, register-budget 78,50
-;; ping / pong / k_mutual : real 50, register-budget 78,50
-;; use_most_lo : real 0, register-budget 96,32
-;; self_rec / k_self : real 32, register-budget 96,32
+;; Scenario 3 (cycle): shows how a cycle affects the up-then-down propagation.
+;; use_most_hi : real 0, accum-offset 78
+;; ping / pong / k_mutual : real 50, accum-offset 78
+;; use_most_lo : real 0, accum-offset 96
+;; self_rec / k_self : real 32, accum-offset 96
;; ===========================================================================
define internal void @use_most_hi() {
@@ -403,22 +405,25 @@ define amdgpu_kernel void @k_self(i1 %c) {
;; shared_k2 -> shared_f2 (requires no AGPRs)
;; Only real AGPR demand propagates *up*, so shared_k2 (which reaches only
;; shared_f2) must NOT inherit an AGPR reservation from shared_k1 just because
-;; they share shared_f2. shared_f2 still gets a downward budget
-;; (amdgpu-register-budget) without claiming a real demand that leaks to k2.
+;; they share shared_f2. shared_f2 still gets a downward boundary
+;; (amdgpu-accum-offset) without claiming a real demand that leaks to k2.
+;;
+;; A B
+;; / \ /
+;; C D
+;; A = shared_k1 : real 0, accum-offset 78
+;; B = shared_k2 : real 0, accum-offset 128
+;; C = shared_f1 : real 50, accum-offset 78
+;; D = shared_f2 : real 0, accum-offset 78
;;
-;; A B
-;; / \ /
-;; C D
-;; A = shared_k1 : real 0, register-budget 78,50
-;; B = shared_k2 : real 0, register-budget 128,0
-;; C = shared_f1 : real 50, register-budget 78,50
-;; D = shared_f2 : real 0, register-budget 78,0
-;;
+;; D takes the lower of its two callers' boundaries. Its AGPR ceiling is pinned
+;; to its own requirement of 0 by getMaxNumVectorRegs, so B stays within budget
+;; (128 arch + 0 AGPR) even though D runs under B's higher boundary.
;; ===========================================================================
define internal void @shared_use_most() {
; CHECK-LABEL: define internal void @shared_use_most(
-; CHECK-SAME: ) #[[ATTR6:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR0]] {
; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
@@ -474,7 +479,7 @@ define internal void @shared_f1() {
; shared_f2: shared between shared_k1 and shared_k2, needs no AGPRs of its own.
define internal void @shared_f2(ptr %p) {
; CHECK-LABEL: define internal void @shared_f2(
-; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR6]] {
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR0]] {
; CHECK-NEXT: [[V:%.*]] = load volatile <128 x i32>, ptr [[P]], align 512
; CHECK-NEXT: store volatile <128 x i32> [[V]], ptr [[P]], align 512
; CHECK-NEXT: call void @shared_use_most()
@@ -502,7 +507,7 @@ define amdgpu_kernel void @shared_k1(ptr %p) {
; shared_k2 reaches only shared_f2 and must not inherit shared_k1's split.
define amdgpu_kernel void @shared_k2(ptr %p) {
; CHECK-LABEL: define amdgpu_kernel void @shared_k2(
-; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR7:[0-9]+]] {
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR6:[0-9]+]] {
; CHECK-NEXT: call void @shared_f2(ptr [[P]])
; CHECK-NEXT: ret void
;
@@ -518,17 +523,17 @@ define amdgpu_kernel void @shared_k2(ptr %p) {
;; \ /
;; C
;;
-;; A = @noroom_kernel_agpr16 : agpr 16 vgpr 128 - 16 -> 112,16
-;; B = @noroom_kernel_noagpr : agpr 0 vgpr 128 - 0 -> 128,0
-;; C = @noroom_shared 112,16 | 128,0 -> 112,0
+;; A = @noroom_kernel_agpr16 : agpr 16, boundary 128 - 16 -> 112
+;; B = @noroom_kernel_noagpr : agpr 0, boundary 128 - 0 -> 128
+;; C = @noroom_shared min(112, 128) -> 112
;;
-;; C may use 112 VGPRs and 128 - 128 = 0 AGPRs: any AGPR use would push it past
-;; the 128-register limit when it runs under B's accum_offset.
+;; C may use 112 arch VGPRs and 0 AGPRs: any AGPR use would push it past the
+;; 128-register limit when it runs under B's accum_offset.
;; ===========================================================================
define internal void @noroom_use_most() {
; CHECK-LABEL: define internal void @noroom_use_most(
-; CHECK-SAME: ) #[[ATTR8:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR7:[0-9]+]] {
; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
@@ -572,7 +577,7 @@ define internal void @noroom_use_most() {
; ceiling from B.
define internal void @noroom_shared() {
; CHECK-LABEL: define internal void @noroom_shared(
-; CHECK-SAME: ) #[[ATTR8]] {
+; CHECK-SAME: ) #[[ATTR7]] {
; CHECK-NEXT: call void @noroom_use_most()
; CHECK-NEXT: ret void
;
@@ -582,7 +587,7 @@ define internal void @noroom_shared() {
define amdgpu_kernel void @noroom_kernel_agpr16() {
; CHECK-LABEL: define amdgpu_kernel void @noroom_kernel_agpr16(
-; CHECK-SAME: ) #[[ATTR9:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR8:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @noroom_shared()
; CHECK-NEXT: ret void
@@ -594,7 +599,7 @@ define amdgpu_kernel void @noroom_kernel_agpr16() {
define amdgpu_kernel void @noroom_kernel_noagpr() {
; CHECK-LABEL: define amdgpu_kernel void @noroom_kernel_noagpr(
-; CHECK-SAME: ) #[[ATTR7]] {
+; CHECK-SAME: ) #[[ATTR6]] {
; CHECK-NEXT: call void @noroom_shared()
; CHECK-NEXT: ret void
;
@@ -603,24 +608,27 @@ define amdgpu_kernel void @noroom_kernel_noagpr() {
}
;; ;; ===========================================================================
-;; Scenario 6: both kernels reserve some AGPRs, so the shared callee keeps the
-;; headroom common to both.
+;; Scenario 6: both kernels reserve some AGPRs, so both boundaries are below the
+;; budget and the shared callee takes the lower one.
;;
;; A B
;; \ /
;; C
;;
-;; A = @room_kernel_agpr16 : 16 AGPRs 128 - 16 -> 112,16
-;; B = @room_kernel_agpr12 : 12 AGPRs 128 - 12 -> 116,12
-;; C = @room_shared 112,16 | 116,12 -> 112,12
+;; A = @room_kernel_agpr16 : 16 AGPRs, boundary 128 - 16 -> 112
+;; B = @room_kernel_agpr12 : 12 AGPRs, boundary 128 - 12 -> 116
+;; C = @room_shared min(112, 116) -> 112
;;
-;; C may use 112 VGPRs and 128 - 116 = 12 AGPRs, which it could spill into.
-;; Note this is the smaller of the two kernels' reservations, not the larger.
+;; C may use 112 arch VGPRs. Its AGPR ceiling is its own requirement of 0 rather
+;; than the 12 AGPRs both kernels happen to have headroom for, so it cannot use
+;; AGPRs as spill space. That headroom is only recoverable by propagating a
+;; second, independent AGPR ceiling, which this attribute deliberately does not
+;; carry.
;; ===========================================================================
define internal void @room_use_most() {
; CHECK-LABEL: define internal void @room_use_most(
-; CHECK-SAME: ) #[[ATTR10:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR7]] {
; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
@@ -662,7 +670,7 @@ define internal void @room_use_most() {
define internal void @room_shared() {
; CHECK-LABEL: define internal void @room_shared(
-; CHECK-SAME: ) #[[ATTR10]] {
+; CHECK-SAME: ) #[[ATTR7]] {
; CHECK-NEXT: call void @room_use_most()
; CHECK-NEXT: ret void
;
@@ -672,7 +680,7 @@ define internal void @room_shared() {
define amdgpu_kernel void @room_kernel_agpr16() {
; CHECK-LABEL: define amdgpu_kernel void @room_kernel_agpr16(
-; CHECK-SAME: ) #[[ATTR9]] {
+; CHECK-SAME: ) #[[ATTR8]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @room_shared()
; CHECK-NEXT: ret void
@@ -684,7 +692,7 @@ define amdgpu_kernel void @room_kernel_agpr16() {
define amdgpu_kernel void @room_kernel_agpr12() {
; CHECK-LABEL: define amdgpu_kernel void @room_kernel_agpr12(
-; CHECK-SAME: ) #[[ATTR11:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR9:[0-9]+]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @room_shared()
; CHECK-NEXT: ret void
@@ -700,15 +708,15 @@ define amdgpu_kernel void @room_kernel_agpr12() {
;; A
;; / \
;; B C
-;; A = @fanout_kernel : agpr 16 | vgpr 128 - 16 = 112,16
+;; A = @fanout_kernel : agpr 16, boundary 128 - 16 -> 112
;; B = @fanout_agpr : clobbers a15, so 16 flows up into A
-;; C = @fanout_noagpr : uses no AGPRs of its own, but inheriets 112,16
+;; C = @fanout_noagpr : uses no AGPRs of its own, but inherits the boundary 112
;;
;; ===========================================================================
define internal void @fanout_use_most() {
; CHECK-LABEL: define internal void @fanout_use_most(
-; CHECK-SAME: ) #[[ATTR12:[0-9]+]] {
+; CHECK-SAME: ) #[[ATTR7]] {
; CHECK-NEXT: [[ALLOCA:%.*]] = alloca [256 x i8], align 1, addrspace(5)
; CHECK-NEXT: [[ALLOCA_CAST:%.*]] = addrspacecast ptr addrspace(5) [[ALLOCA]] to ptr
; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.workitem.id.x()
@@ -750,7 +758,7 @@ define internal void @fanout_use_most() {
define internal void @fanout_sink() {
; CHECK-LABEL: define internal void @fanout_sink(
-; CHECK-SAME: ) #[[ATTR12]] {
+; CHECK-SAME: ) #[[ATTR7]] {
; CHECK-NEXT: call void @fanout_use_most()
; CHECK-NEXT: ret void
;
@@ -760,7 +768,7 @@ define internal void @fanout_sink() {
define internal void @fanout_agpr() {
; CHECK-LABEL: define internal void @fanout_agpr(
-; CHECK-SAME: ) #[[ATTR9]] {
+; CHECK-SAME: ) #[[ATTR8]] {
; CHECK-NEXT: call void asm sideeffect "
; CHECK-NEXT: call void @fanout_sink()
; CHECK-NEXT: ret void
@@ -772,7 +780,7 @@ define internal void @fanout_agpr() {
define internal void @fanout_noagpr(ptr %p) {
; CHECK-LABEL: define internal void @fanout_noagpr(
-; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR12]] {
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR7]] {
; CHECK-NEXT: [[V:%.*]] = load volatile <128 x i32>, ptr [[P]], align 512
; CHECK-NEXT: store volatile <128 x i32> [[V]], ptr [[P]], align 512
; CHECK-NEXT: call void @fanout_sink()
@@ -786,7 +794,7 @@ define internal void @fanout_noagpr(ptr %p) {
define amdgpu_kernel void @fanout_kernel(ptr %p) {
; CHECK-LABEL: define amdgpu_kernel void @fanout_kernel(
-; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR9]] {
+; CHECK-SAME: ptr [[P:%.*]]) #[[ATTR8]] {
; CHECK-NEXT: call void @fanout_agpr()
; CHECK-NEXT: call void @fanout_noagpr(ptr [[P]])
; CHECK-NEXT: ret void
@@ -796,19 +804,16 @@ define amdgpu_kernel void @fanout_kernel(ptr %p) {
ret void
}
;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="76,52" }
-; CHECK: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="50" "amdgpu-no-wwm" "amdgpu-register-budget"="76,52" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-accum-offset"="78" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-accum-offset"="78" "amdgpu-agpr-alloc"="50" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="50" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="96,32" }
-; CHECK: attributes #[[ATTR5]] = { "amdgpu-agpr-alloc"="32" "amdgpu-no-wwm" "amdgpu-register-budget"="96,32" }
-; CHECK: attributes #[[ATTR6]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="76,0" }
-; CHECK: attributes #[[ATTR7]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
-; CHECK: attributes #[[ATTR8]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="112,0" }
-; CHECK: attributes #[[ATTR9]] = { "amdgpu-agpr-alloc"="16" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
-; CHECK: attributes #[[ATTR10]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="112,12" }
-; CHECK: attributes #[[ATTR11]] = { "amdgpu-agpr-alloc"="12" "amdgpu-no-wwm" "amdgpu-register-budget"="116,12" }
-; CHECK: attributes #[[ATTR12]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
-; CHECK: attributes #[[ATTR13:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-; CHECK: attributes #[[ATTR14:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
+; CHECK: attributes #[[ATTR4]] = { "amdgpu-accum-offset"="96" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR5]] = { "amdgpu-accum-offset"="96" "amdgpu-agpr-alloc"="32" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR6]] = { "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR7]] = { "amdgpu-accum-offset"="112" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR8]] = { "amdgpu-accum-offset"="112" "amdgpu-agpr-alloc"="16" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR9]] = { "amdgpu-accum-offset"="116" "amdgpu-agpr-alloc"="12" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR10:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+; CHECK: attributes #[[ATTR11:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
;.
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-work-group-size.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-accum-offset-work-group-size.ll
similarity index 77%
rename from llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-work-group-size.ll
rename to llvm/test/CodeGen/AMDGPU/amdgpu-attributor-accum-offset-work-group-size.ll
index 2e4ba2d30e445..09361f07a04f1 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-register-budget-work-group-size.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-accum-offset-work-group-size.ll
@@ -2,7 +2,7 @@
; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -passes=amdgpu-attributor %s | FileCheck -check-prefixes=CHECK,GFX90A %s
; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -passes=amdgpu-attributor %s | FileCheck -check-prefixes=CHECK,GFX942 %s
-; Why the downward propagation carries a VGPR budget rather than an AGPR count.
+; Why the downward propagation carries an accum_offset rather than an AGPR count.
;
; The size of the vector register file a wave may use is not a constant: it
; falls out of the occupancy needed to satisfy the kernel's launch bounds. On
@@ -14,10 +14,10 @@
; range, so no attribute is emitted on the callee).
;
; An AGPR count alone cannot be propagated across that difference, because the
-; same count means different things under different totals. Folding the work
-; group size and the AGPR requirement together into a single VGPR budget per
-; kernel makes the number self-contained, so it can be safely merged in the
-; callee.
+; same count means different things under different totals. An accum_offset is
+; an absolute position in the register file, so folding the work group size and
+; the AGPR requirement together into one boundary per kernel makes the number
+; self-contained and safe to merge in the callee.
;
; A B
; \ /
@@ -25,14 +25,14 @@
;; ===========================================================================
;; Scenario 1: the AGPR-hungry kernel is the one with the *larger* register
-;; file, so its budget is still tighter than the other kernel's.
+;; file, so its boundary is still lower than the other kernel's.
;;
;; A = @wgs_kernel_512_agpr130 : [1,512], agpr 130 -> 256 - 130 = 126
;; B = @wgs_kernel_1024_noagpr : [1,1024], agpr 0 -> 128 - 0 = 128
-;; C = @wgs_shared -> 126,128
+;; C = @wgs_shared min(126, 128) -> 126
;;
-;; C decodes to a ceiling of 126 VGPRs, and 0 AGPRs because its own total is
-;; 128 and the largest budget it must fit under is already 128.
+;; C is held to 126 arch VGPRs even though its own total is 128, because it has
+;; to leave A's 130 AGPRs untouched in A's 256-register file.
;; ===========================================================================
; C. Also serves as the sink that keeps the inferred attribute list short by
@@ -105,16 +105,16 @@ define amdgpu_kernel void @wgs_kernel_1024_noagpr() #1 {
}
;; ===========================================================================
-;; Scenario 2: the wide-launch kernel's budget exceeds the narrow-launch
-;; kernel's *total*, so decoding the AGPR ceiling has to saturate.
+;; Scenario 2: the boundary of the kernel with the larger register file sits
+;; above the other kernel's whole total, so it constrains nothing.
;;
;; A = @sat_kernel_512_agpr56 : [1,512], agpr 56 -> 256 - 56 = 200
;; B = @sat_kernel_1024_noagpr: [1,1024], agpr 0 -> 128 - 0 = 128
-;; C = @sat_shared -> 128,200
+;; C = @sat_shared min(200, 128) -> 128
;;
-;; C decodes to a ceiling of 128 VGPRs. Its own total is 128 and the largest
-;; budget it must fit under is 200, so the AGPR ceiling saturates at 0 rather
-;; than going negative.
+;; C keeps a ceiling of 128 arch VGPRs, which is its own total anyway, so A's
+;; boundary of 200 is inert. The 56 AGPRs A reserves are still safe: C uses no
+;; AGPRs of its own, so under A the pair peaks at 128 + 56 <= 256.
;; ===========================================================================
define internal void @sat_shared() {
@@ -184,19 +184,19 @@ define amdgpu_kernel void @sat_kernel_1024_noagpr() #1 {
attributes #0 = { "amdgpu-flat-work-group-size"="1,512" }
attributes #1 = { "amdgpu-flat-work-group-size"="1,1024" }
;.
-; GFX90A: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="124,0" "target-cpu"="gfx90a" }
-; GFX90A: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="130" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="124,132" "target-cpu"="gfx90a" }
-; GFX90A: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx90a" }
-; GFX90A: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx90a" }
-; GFX90A: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="56" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="200,56" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR0]] = { "amdgpu-accum-offset"="126" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR1]] = { "amdgpu-accum-offset"="126" "amdgpu-agpr-alloc"="130" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR2]] = { "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-wwm" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR3]] = { "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "target-cpu"="gfx90a" }
+; GFX90A: attributes #[[ATTR4]] = { "amdgpu-accum-offset"="200" "amdgpu-agpr-alloc"="56" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "target-cpu"="gfx90a" }
; GFX90A: attributes #[[ATTR5:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) "target-cpu"="gfx90a" }
; GFX90A: attributes #[[ATTR6:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) "target-cpu"="gfx90a" }
;.
-; GFX942: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="124,0" "target-cpu"="gfx942" }
-; GFX942: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="130" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="124,132" "target-cpu"="gfx942" }
-; GFX942: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx942" }
-; GFX942: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" "target-cpu"="gfx942" }
-; GFX942: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="56" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "amdgpu-register-budget"="200,56" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR0]] = { "amdgpu-accum-offset"="126" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR1]] = { "amdgpu-accum-offset"="126" "amdgpu-agpr-alloc"="130" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR2]] = { "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-flat-work-group-size"="1,1024" "amdgpu-no-wwm" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR3]] = { "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "target-cpu"="gfx942" }
+; GFX942: attributes #[[ATTR4]] = { "amdgpu-accum-offset"="200" "amdgpu-agpr-alloc"="56" "amdgpu-flat-work-group-size"="1,512" "amdgpu-no-wwm" "target-cpu"="gfx942" }
; GFX942: attributes #[[ATTR5:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) "target-cpu"="gfx942" }
; GFX942: attributes #[[ATTR6:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) "target-cpu"="gfx942" }
;.
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll
index a0d7855b20369..8946a33c3ad1d 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-min-agpr-alloc.ll
@@ -1113,44 +1113,44 @@ attributes #2 = { sanitize_address "amdgpu-agpr-alloc"="0" }
!4 = !{!"a256"}
;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="0,0" }
-; CHECK: attributes #[[ATTR1]] = { "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" "amdgpu-register-budget"="124,4" }
-; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" "amdgpu-register-budget"="124,4" }
-; CHECK: attributes #[[ATTR3]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-accum-offset"="0" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-accum-offset"="127" "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR2]] = { "amdgpu-accum-offset"="126" "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR3]] = { "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR4]] = { "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR5]] = { "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" "amdgpu-register-budget"="64,4" }
+; CHECK: attributes #[[ATTR5]] = { "amdgpu-accum-offset"="64" "amdgpu-agpr-alloc"="1" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR6]] = { "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR7]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="64,0" }
+; CHECK: attributes #[[ATTR7]] = { "amdgpu-accum-offset"="64" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR8:[0-9]+]] = { convergent nocallback nocreateundeforpoison nofree nosync nounwind willreturn memory(none) }
-; CHECK: attributes #[[ATTR9]] = { "amdgpu-agpr-alloc"="4" "amdgpu-no-wwm" "amdgpu-register-budget"="124,4" }
-; CHECK: attributes #[[ATTR10]] = { "amdgpu-agpr-alloc"="6" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
-; CHECK: attributes #[[ATTR11]] = { "amdgpu-agpr-alloc"="5" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
-; CHECK: attributes #[[ATTR12]] = { "amdgpu-agpr-alloc"="14" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
-; CHECK: attributes #[[ATTR13]] = { "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" "amdgpu-register-budget"="0,256" }
-; CHECK: attributes #[[ATTR14]] = { "amdgpu-agpr-alloc"="32" "amdgpu-no-wwm" "amdgpu-register-budget"="96,32" }
-; CHECK: attributes #[[ATTR15]] = { "amdgpu-agpr-alloc"="9" "amdgpu-no-wwm" "amdgpu-register-budget"="116,12" }
-; CHECK: attributes #[[ATTR16]] = { "amdgpu-agpr-alloc"="64" "amdgpu-no-wwm" "amdgpu-register-budget"="64,64" }
-; CHECK: attributes #[[ATTR17]] = { "amdgpu-agpr-alloc"="49" "amdgpu-no-wwm" "amdgpu-register-budget"="76,52" }
-; CHECK: attributes #[[ATTR18]] = { "amdgpu-agpr-alloc"="33" "amdgpu-no-wwm" "amdgpu-register-budget"="92,36" }
-; CHECK: attributes #[[ATTR19]] = { "amdgpu-agpr-alloc"="8" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
-; CHECK: attributes #[[ATTR20]] = { "amdgpu-agpr-alloc"="13" "amdgpu-no-wwm" "amdgpu-register-budget"="112,16" }
-; CHECK: attributes #[[ATTR21]] = { "amdgpu-agpr-alloc"="56" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" "amdgpu-register-budget"="72,56" }
-; CHECK: attributes #[[ATTR22]] = { "amdgpu-agpr-alloc"="58" "amdgpu-no-wwm" "amdgpu-register-budget"="68,60" }
-; CHECK: attributes #[[ATTR23]] = { "amdgpu-agpr-alloc"="56" "amdgpu-no-wwm" "amdgpu-register-budget"="72,56" }
-; CHECK: attributes #[[ATTR24]] = { "amdgpu-agpr-alloc"="56" "amdgpu-register-budget"="72,56" }
-; CHECK: attributes #[[ATTR25]] = { "amdgpu-agpr-alloc"="60" "amdgpu-no-wwm" "amdgpu-register-budget"="68,60" }
-; CHECK: attributes #[[ATTR26]] = { "amdgpu-agpr-alloc"="8" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
-; CHECK: attributes #[[ATTR27]] = { "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
-; CHECK: attributes #[[ATTR28]] = { "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" "amdgpu-register-budget"="0,256" "amdgpu-waves-per-eu"="1,1" }
+; CHECK: attributes #[[ATTR9]] = { "amdgpu-accum-offset"="124" "amdgpu-agpr-alloc"="4" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR10]] = { "amdgpu-accum-offset"="122" "amdgpu-agpr-alloc"="6" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR11]] = { "amdgpu-accum-offset"="123" "amdgpu-agpr-alloc"="5" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR12]] = { "amdgpu-accum-offset"="114" "amdgpu-agpr-alloc"="14" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR13]] = { "amdgpu-accum-offset"="0" "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR14]] = { "amdgpu-accum-offset"="96" "amdgpu-agpr-alloc"="32" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR15]] = { "amdgpu-accum-offset"="119" "amdgpu-agpr-alloc"="9" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR16]] = { "amdgpu-accum-offset"="64" "amdgpu-agpr-alloc"="64" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR17]] = { "amdgpu-accum-offset"="79" "amdgpu-agpr-alloc"="49" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR18]] = { "amdgpu-accum-offset"="95" "amdgpu-agpr-alloc"="33" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR19]] = { "amdgpu-accum-offset"="120" "amdgpu-agpr-alloc"="8" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR20]] = { "amdgpu-accum-offset"="115" "amdgpu-agpr-alloc"="13" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR21]] = { "amdgpu-accum-offset"="72" "amdgpu-agpr-alloc"="56" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR22]] = { "amdgpu-accum-offset"="70" "amdgpu-agpr-alloc"="58" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR23]] = { "amdgpu-accum-offset"="72" "amdgpu-agpr-alloc"="56" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR24]] = { "amdgpu-accum-offset"="72" "amdgpu-agpr-alloc"="56" }
+; CHECK: attributes #[[ATTR25]] = { "amdgpu-accum-offset"="68" "amdgpu-agpr-alloc"="60" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR26]] = { "amdgpu-accum-offset"="64" "amdgpu-agpr-alloc"="8" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR27]] = { "amdgpu-accum-offset"="64" "amdgpu-agpr-alloc"="2" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR28]] = { "amdgpu-accum-offset"="0" "amdgpu-agpr-alloc"="256" "amdgpu-no-wwm" "amdgpu-waves-per-eu"="1,1" }
; CHECK: attributes #[[ATTR29]] = { "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR30]] = { "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
-; CHECK: attributes #[[ATTR31]] = { "amdgpu-agpr-alloc"="3" "amdgpu-no-wwm" "amdgpu-register-budget"="64,8" }
-; CHECK: attributes #[[ATTR32]] = { "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" "amdgpu-register-budget"="120,8" }
+; CHECK: attributes #[[ATTR30]] = { "amdgpu-accum-offset"="64" "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR31]] = { "amdgpu-accum-offset"="64" "amdgpu-agpr-alloc"="3" "amdgpu-no-wwm" }
+; CHECK: attributes #[[ATTR32]] = { "amdgpu-accum-offset"="121" "amdgpu-agpr-alloc"="7" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR33]] = { sanitize_address "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR34]] = { sanitize_memory "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR35]] = { sanitize_thread "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR36]] = { sanitize_hwaddress "amdgpu-no-wwm" }
-; CHECK: attributes #[[ATTR37]] = { sanitize_address "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
+; CHECK: attributes #[[ATTR37]] = { sanitize_address "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR38]] = { "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR39:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
; CHECK: attributes #[[ATTR40:[0-9]+]] = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
diff --git a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll
index 1ef6a3596a81b..777662938fbc7 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdgpu-attributor-trap-leaf.ll
@@ -58,7 +58,7 @@ attributes #1 = { nocallback "trap-func-name"="handler" }
;.
; CHECK: attributes #[[ATTR0:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
; CHECK: attributes #[[ATTR1:[0-9]+]] = { nounwind }
-; CHECK: attributes #[[ATTR2]] = { "amdgpu-agpr-alloc"="0" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" "amdgpu-register-budget"="128,0" }
+; CHECK: attributes #[[ATTR2]] = { "amdgpu-accum-offset"="128" "amdgpu-agpr-alloc"="0" "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR3]] = { "amdgpu-no-cluster-id-x" "amdgpu-no-cluster-id-y" "amdgpu-no-cluster-id-z" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-flat-scratch-init" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-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" "amdgpu-no-wwm" }
; CHECK: attributes #[[ATTR4]] = { "trap-func-name"="handler" }
; CHECK: attributes #[[ATTR5]] = { nocallback "trap-func-name"="handler" }
More information about the llvm-commits
mailing list