Skip to content

[AMDGPU] Reflect amdgpu-waves-per-eu attribute minimum occupancy to RegPressure analysis (used in machine licm, machine scheduler, and register allocation) #167390

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
+32 −43
Open

[AMDGPU] Reflect amdgpu-waves-per-eu attribute minimum occupancy to RegPressure analysis (used in machine licm, machine scheduler, and register allocation) #167390

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Comment on lines +3762 to +3763
Copy link
Contributor

@arsenm arsenm Nov 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This should not be using the low level parsing utility and directly use getWavesPerEU

Copy link
Contributor Author

@hidekisaito hidekisaito Nov 11, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Contributor

@arsenm arsenm Nov 11, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Contributor

@shiltian shiltian Nov 11, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yeah, IIUC the attribute is guaranteed ABI.

Copy link
Contributor

@arsenm arsenm Nov 11, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It's not quite ABI but essentially

Copy link
Contributor

@arsenm arsenm Nov 11, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

but it's important there is no distinction between default value and attribute set to default

MinOcc = ST.getWavesPerEU(F).first;
}
Copy link
Contributor

@arsenm arsenm Nov 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Contributor Author

@hidekisaito hidekisaito Nov 11, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

@arsenm arsenm Nov 11, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

switch (RC->getID()) {
default:
return AMDGPUGenRegisterInfo::getRegPressureLimit(RC, MF);
Expand Down
70 changes: 27 additions & 43 deletions llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
Loading