[AMDGPU] Reflect amdgpu-waves-per-eu attribute minimum occupancy to RegPressure analysis (used in machine licm, machine scheduler, and register allocation) #167390
Conversation
…egPressure analysis (used in machine licm, machine scheduler, and register allocation)
|
@llvm/pr-subscribers-backend-amdgpu Author: None (hidekisaito) ChangesFull diff: https://github.com/llvm/llvm-project/pull/167390.diff 2 Files Affected:
diff --git a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
index a6c1af24e13e9..ed04d6bf713c7 100644
--- a/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
@@ -3758,6 +3758,11 @@ bool SIRegisterInfo::shouldCoalesce(MachineInstr *MI,
unsigned SIRegisterInfo::getRegPressureLimit(const TargetRegisterClass *RC,
MachineFunction &MF) const {
unsigned MinOcc = ST.getOccupancyWithWorkGroupSizes(MF).first;
+ Function &F = MF.getFunction();
+ if (AMDGPU::getIntegerPairAttribute(F, "amdgpu-waves-per-eu", true) !=
+ std::nullopt) {
+ MinOcc = ST.getWavesPerEU(F).first;
+ }
switch (RC->getID()) {
default:
return AMDGPUGenRegisterInfo::getRegPressureLimit(RC, MF);
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 ebbeab94066d6..b34f17e28afb2 100644
--- a/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll
+++ b/llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll
@@ -375,64 +375,48 @@ define void @v32_asm_def_use(float %v0, float %v1) #4 {
; GFX908-NEXT: ;;#ASMSTART
; GFX908-NEXT: ; def v[0:31] a[0:15]
; GFX908-NEXT: ;;#ASMEND
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a15
-; GFX908-NEXT: ;;#ASMSTART
-; GFX908-NEXT: ; def v32
-; GFX908-NEXT: ;;#ASMEND
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a31, v35
+; GFX908-NEXT: v_accvgpr_read_b32 v32, a15
; GFX908-NEXT: v_accvgpr_read_b32 v35, a14
-; GFX908-NEXT: s_nop 1
+; GFX908-NEXT: v_accvgpr_read_b32 v36, a13
+; GFX908-NEXT: v_accvgpr_write_b32 a31, v32
+; GFX908-NEXT: v_accvgpr_read_b32 v32, a12
; GFX908-NEXT: v_accvgpr_write_b32 a30, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a13
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a29, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a12
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a28, v35
+; GFX908-NEXT: v_accvgpr_write_b32 a29, v36
+; GFX908-NEXT: v_accvgpr_write_b32 a28, v32
; GFX908-NEXT: v_accvgpr_read_b32 v35, a11
-; GFX908-NEXT: s_nop 1
+; GFX908-NEXT: v_accvgpr_read_b32 v36, a10
+; GFX908-NEXT: v_accvgpr_read_b32 v32, a9
; GFX908-NEXT: v_accvgpr_write_b32 a27, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a10
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a26, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a9
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a25, v35
+; GFX908-NEXT: v_accvgpr_write_b32 a26, v36
+; GFX908-NEXT: v_accvgpr_write_b32 a25, v32
; GFX908-NEXT: v_accvgpr_read_b32 v35, a8
-; GFX908-NEXT: s_nop 1
+; GFX908-NEXT: v_accvgpr_read_b32 v36, a7
+; GFX908-NEXT: v_accvgpr_read_b32 v32, a6
; GFX908-NEXT: v_accvgpr_write_b32 a24, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a7
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a23, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a6
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a22, v35
+; GFX908-NEXT: v_accvgpr_write_b32 a23, v36
+; GFX908-NEXT: v_accvgpr_write_b32 a22, v32
; GFX908-NEXT: v_accvgpr_read_b32 v35, a5
-; GFX908-NEXT: s_nop 1
+; GFX908-NEXT: v_accvgpr_read_b32 v36, a4
+; GFX908-NEXT: v_accvgpr_read_b32 v32, a3
; GFX908-NEXT: v_accvgpr_write_b32 a21, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a4
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a20, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a3
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a19, v35
+; GFX908-NEXT: v_accvgpr_write_b32 a20, v36
+; GFX908-NEXT: v_accvgpr_write_b32 a19, v32
; GFX908-NEXT: v_accvgpr_read_b32 v35, a2
-; GFX908-NEXT: s_nop 1
+; GFX908-NEXT: v_accvgpr_read_b32 v36, a1
+; GFX908-NEXT: v_accvgpr_read_b32 v32, a0
; GFX908-NEXT: v_accvgpr_write_b32 a18, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a1
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a17, v35
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a0
-; GFX908-NEXT: s_nop 1
-; GFX908-NEXT: v_accvgpr_write_b32 a16, v35
+; GFX908-NEXT: v_accvgpr_write_b32 a17, v36
+; GFX908-NEXT: v_accvgpr_write_b32 a16, v32
+; GFX908-NEXT: ;;#ASMSTART
+; GFX908-NEXT: ; def v32
+; GFX908-NEXT: ;;#ASMEND
; GFX908-NEXT: ;;#ASMSTART
; GFX908-NEXT: ; copy
; GFX908-NEXT: ;;#ASMEND
-; GFX908-NEXT: v_accvgpr_read_b32 v35, a1
+; GFX908-NEXT: v_accvgpr_read_b32 v37, a1
; GFX908-NEXT: v_mfma_f32_16x16x1f32 a[0:15], v34, v33, a[16:31]
; GFX908-NEXT: s_nop 0
-; GFX908-NEXT: v_accvgpr_write_b32 a32, v35
+; GFX908-NEXT: v_accvgpr_write_b32 a32, v37
; GFX908-NEXT: ;;#ASMSTART
; GFX908-NEXT: ; copy
; GFX908-NEXT: ;;#ASMEND
|
hidekisaito
requested review from
dstutt,
jayfoad,
jrbyrnes,
lucas-rami,
nhaehnle and
perlfu
| if (AMDGPU::getIntegerPairAttribute(F, "amdgpu-waves-per-eu", true) != | ||
| std::nullopt) { |
There was a problem hiding this comment.
This should not be using the low level parsing utility and directly use getWavesPerEU
There was a problem hiding this comment.
That's exactly the point of
getWavesPerEU() using default minimum occupancy estimate based on single WG per CU and getPressure[Set]Limit() using default minimum occupancy estimate based on multiple WGs per CU
There was a problem hiding this comment.
Not sure I follow, but you shouldn't be trying to glean meaning from there being an explicit attribute or not. getWavesPerEU should be the authoritative range to use
There was a problem hiding this comment.
yeah, IIUC the attribute is guaranteed ABI.
There was a problem hiding this comment.
It's not quite ABI but essentially
There was a problem hiding this comment.
but it's important there is no distinction between default value and attribute set to default
| if (AMDGPU::getIntegerPairAttribute(F, "amdgpu-waves-per-eu", true) != | ||
| std::nullopt) { | ||
| MinOcc = ST.getWavesPerEU(F).first; | ||
| } |
There was a problem hiding this comment.
Why isn't this pressure limit computed with RegisterClassInfo? I would assume all of the pressure information is implied by the register counts, which are already set by the occupancy, and wouldn't require manual intervention in terms of the occupancy
There was a problem hiding this comment.
All I can say is that this is the value used by machine licm, pre-RA machine scheduler, and RA ---- and the minimum occupancy computed here, without this patch, is different from what you can get from MachineFunctionInfo's WavesPerEU (that comes from getWavesPerEU). Such inconsistency is certainly not desirable ---- but I need those who are in the experts in the area address that bigger problem. Trying to be consistent one way or the other way (probably one becoming default and choosing the other under an option), I expect a lot of GPU ASM churn to happen.
This patch is just addressing a small but obvious problem while waiting for the bigger problem to be discussed/resolved. I'm fine if others don't like this patch going in ---- we can keep discussing how to resolve the inconsistency issue.
There was a problem hiding this comment.
I'm confused by the state of the MachineFunctionInfo field. It's a scalar, but it really should be a range. It changes mid-scheduling, so I'm not sure how to reconcile it. I'd hope it stays within the bounds of getWavesPerEU
4 participants
I'd like to have a discussion about getWavesPerEU() using default minimum occupancy estimate based on single WG per CU and getPressure[Set]Limit() using default minimum occupancy estimate based on multiple WGs per CU, but that may take time to resolve. As a starter, this patch addresses getPressure[Set]Limit() ignoring amdgpu-waves-per-eu attribute (e.g., using much higher occupancy estimate when the kernel attribute says single occupancy)