Skip to content

fix(tensilelite): fix VGPR occupancy calculation on gfx11 #9708

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
mgehre-amd wants to merge 8 commits into from
+3,478 −3,270
Open

fix(tensilelite): fix VGPR occupancy calculation on gfx11 #9708

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
ee20bdd
fix(tensilelite): use per-SIMD VGPR file for occupancy on gfx11
mgehre-amd Jul 22, 2026
7d14ceb
fix(tensilelite): guard gfx1250 onto legacy occupancy model
mgehre-amd Jul 23, 2026
051d119
fix(tensilelite): derive the occupancy multiplier from the wavefront …
mgehre-amd Jul 24, 2026
0b0c80a
test(tensilelite): exercise the real doubleVgpr value on gfx11
mgehre-amd Oct 5, 2026
4ab568f
perf(hipblaslt): hold gfx11 occupancy at its develop values
mgehre-amd Oct 6, 2026
6f1368c
test(tensilelite): give the HIP occupancy stub its wavefront size
mgehre-amd Oct 21, 2026
cd97ac1
fix(tensilelite): cap the gfx11 VGPR pool at the modeled budget
mgehre-amd Oct 22, 2026
9f65876
perf(hipblaslt): re-pin gfx11 MaxOccupancy for the capped VGPR pool
mgehre-amd Oct 23, 2026
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
66 changes: 33 additions & 33 deletions .../asm_full/gfx1103/GridBased/gfx1103_Cijk_Ailk_Bjlk_BBS_BH_Bias_AuxB_HAS_SAV_UserArgs.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -178,7 +178,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -366,7 +366,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -554,7 +554,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -742,7 +742,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -1306,7 +1306,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -1494,7 +1494,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -1682,7 +1682,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -1870,7 +1870,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -2058,7 +2058,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -2622,7 +2622,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -2810,7 +2810,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -2998,7 +2998,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -3186,7 +3186,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -3374,7 +3374,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 6
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -3562,7 +3562,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 6
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -3750,7 +3750,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 6
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -3938,7 +3938,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -4126,7 +4126,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -4314,7 +4314,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -4502,7 +4502,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -4690,7 +4690,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -5442,7 +5442,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -5630,7 +5630,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -5818,7 +5818,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 9
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -6006,7 +6006,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 6
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -6194,7 +6194,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 6
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -6382,7 +6382,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 6
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -6570,7 +6570,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -6758,7 +6758,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -8262,7 +8262,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -8450,7 +8450,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -8638,7 +8638,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down Expand Up @@ -8826,7 +8826,7 @@
MatrixInstM: 16
MatrixInstN: 16
MatrixInstruction: [16, 16, 16, 1]
MaxOccupancy: 40
MaxOccupancy: 5
NoLdsWriteCode: false
NoReject: false
NoTailLoop: false
Expand Down
Loading
Loading