tree a4282f10d23b679dbe43558f0406db5399935d5b
parent a80a6b31a721c8d73445c458f1ea2312d9262750
author Valery Pykhtin <valery.pykhtin@amd.com> 1760365031 +0200
committer GitHub <noreply@github.com> 1760365031 +0200
gpgsig -----BEGIN PGP SIGNATURE-----
 
 wsFcBAABCAAQBQJo7QnoCRC1aQ7uu5UhlAAAygEQAH9lFdQVcRO06TWi+lNXK+7d
 6Tx9kymupfMxSwv4EUYZek7LtIdk1g0gHryXoOc+rIGESohiQGO2Xzc8qqE4lBWm
 LHOO0qO48YPqP5kzul2lk/ac3wEKlKHxTVJwuiUHINYrsaDz4H7bMdtU/DLix0ne
 eVxYqD2VnCWxpokBWgyVNIcQO4X1MLZ9balXdMbDZrPotoETa10XBmH2wyDjJ1JA
 GhLHfn+UQbPAsKxmY9U29s7qJKWAf+VSOIKxc4KJz0iI1xp1s1PoZuVL6ezJjcn0
 IO57UBxfiz1h+IGOFmlhrMVDZR3VkrnFD+ZIaUFcgAkMyd44n9Dw82f55a24KYGd
 a5jLQt0g9f+yAweajq9BG7pCBXgHSiiK/pQJ/CVNm7cLuuLdQqS5xWCO0tC8hGXg
 Hh50Eg8hLAE1ePahdIXMJlMmYRTt3Fj8HipPxb6coObwAfw643IKIAAWp7g8mu8X
 TnW1mR6R6LKaWwdf4uX98NLAp3OuqEpUaYw2rm6OqsD6/+At23hNnzwfMlZdI5RH
 AWOjXfk4BcEewErtobrUxqYe91JHObhHIqtTXgK4JQ4jhak4zlMvkjRfbsZgUFkA
 aPqUTAMBQ/PIxngWskhyvOTGFcrB48KgwtMMywI88T0JFA42eENjurWr7wemTJ8C
 OZk/cSOFVD+vRqDq2N/7
 =Ri8O
 -----END PGP SIGNATURE-----
 

[AMDGPU] Add register usage debug printing the point of maximum register pressure. (#161850)

Basically this allows to analyze "why so many VGPRs used?".

It prints all live registers at the point of maximum register pressure
and for each register its defs/uses are dumped.

Currently can be run before and after the scheduler but would be nice if
it can be ran inbetween any passes (not sure this is possible with
legacy pass-manager). Requires debug or built with asserts compiler.

Highly recommended to run with debug info to have debug locations for
instructions.

Example output:

```
*** Register pressure info (VGPRs) for _ZN7ck_tile6ken.... ***
Max pressure is 256 VGPRs at 41780e@BB.18 (LoopHdr BB.16, Depth 1): %9858:vreg_512_align2 = contract V_MFMA_F32_32X32X16_BF16_mac_vgprcd_e64 %10137:vreg_128_align2, %10141:vreg_128_align2, %9858:vreg_512_align2(tied-def 0), 0, 0, 0, implicit $mode, implicit $exec

Live registers with single definition (123 VGPRs):
  %10126:VReg_128_Align2, LiveMask 00000000000000FF (4 VGPRs)
    def 41600r@BB.18 (LoopHdr BB.16, Depth 1): undef %10126.sub0_sub1:vreg_128_align2 = DS_READ_B64_TR_B16 %478:vgpr_32, 15232, 0, implicit $exec :: (load (s64) from %ir.sunkaddr1314, !noalias !60, addrspace 3)
    def 41608r@BB.18 (LoopHdr BB.16, Depth 1): %10126.sub2_sub3:vreg_128_align2 = DS_READ_B64_TR_B16 %478:vgpr_32, 16320, 0, implicit $exec :: (load (s64) from %ir.sunkaddr1315, !noalias !60, addrspace 3)
    use 41848r@BB.18 (LoopHdr BB.16, Depth 1): %9856:vreg_512_align2 = contract V_MFMA_F32_32X32X16_BF16_mac_vgprcd_e64 %10126:vreg_128_align2, %10138:vreg_128_align2, %9856:vreg_512_align2(tied-def 0), 0, 0, 0, implicit $mode, implicit $exec
  %10136:VReg_128_Align2, LiveMask 00000000000000FF (4 VGPRs)
    def 41264r@BB.18 (LoopHdr BB.16, Depth 1): undef %10136.sub0_sub1:vreg_128_align2 = DS_READ_B64_TR_B16 %478:vgpr_32, 2176, 0, implicit $exec :: (load (s64) from %ir.sunkaddr1294, !noalias !60, addrspace 3)
    def 41272r@BB.18 (LoopHdr BB.16, Depth 1): %10136.sub2_sub3:vreg_128_align2 = DS_READ_B64_TR_B16 %478:vgpr_32, 3264, 0, implicit $exec :: (load (s64) from %ir.sunkaddr1295, !noalias !60, addrspace 3)
    use 41788r@BB.18 (LoopHdr BB.16, Depth 1): %9858:vreg_512_align2 = contract V_MFMA_F32_32X32X16_BF16_mac_vgprcd_e64 %10136:vreg_128_align2, %10140:vreg_128_align2, %9858:vreg_512_align2(tied-def 0), 0, 0, 0, implicit $mode, implicit $exec
  %10129:VReg_128_Align2, LiveMask 00000000000000FF (4 VGPRs)
...
Live registers with multiple definitions (133 VGPRs):
  %9856:VReg_512_Align2, LiveMask 00000000FFFFFFFF (16 VGPRs)
    def 16544r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def undef %9856.sub0_sub1:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def 16592r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def %9856.sub2_sub3:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def 16608r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def %9856.sub4_sub5:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def 16656r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def %9856.sub6_sub7:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def 16672r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def %9856.sub8_sub9:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def 16720r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def %9856.sub10_sub11:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def 16736r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def %9856.sub12_sub13:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def 16784r@BB.8: INLINEASM &"v_pk_mul_f32 $0, $1, $2" [sideeffect] [isconvergent] [attdialect], $0:[regdef:VReg_64_Align2], def %9856.sub14_sub15:vreg_512_align2, $1:[reguse:VReg_64_Align2], %4069:vreg_64_align2, $2:[reguse:VReg_64_Align2], %10159:vreg_64_align2, !52
    def use 41828r@BB.18 (LoopHdr BB.16, Depth 1): %9856:vreg_512_align2 = contract V_MFMA_F32_...
...
********** INTERVALS **********
...
********** MACHINEINSTRS **********
# Machine code for function _ZN7ck_tile6kentr...
```