From 0c5e40257d7ba46cde95c60afcdc8108ad54131b Mon Sep 17 00:00:00 2001 From: Brendon Cahoon Date: Wed, 18 Dec 2024 14:08:52 -0600 Subject: [PATCH 1/4] [AMDGPU] Do not fold into v_accvpr_mov/write/read In SIFoldOperands, leave copies for moving between agpr and vgpr registers. The register coalescer is able to handle the copies more efficiently than v_accvgpr_mov, v_accvgpr_write, and v_accvgpr_read. Otherwise, the compiler generates unneccesary instructions such as v_accvgpr_mov a0, a0. --- llvm/lib/Target/AMDGPU/SIFoldOperands.cpp | 13 - .../CodeGen/AMDGPU/no-fold-accvgpr-mov.mir | 234 ++++++++++++++++++ .../CodeGen/AMDGPU/no-fold-accvgpr-read.mir | 181 ++++++++++++++ 3 files changed, 415 insertions(+), 13 deletions(-) create mode 100644 llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir create mode 100644 llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir diff --git a/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp b/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp index 4fb5cb066be7c..47e4afea6fa5f 100644 --- a/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp +++ b/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp @@ -1096,21 +1096,8 @@ void SIFoldOperandsImpl::foldOperand( B.addImm(Defs[I].second); } LLVM_DEBUG(dbgs() << "Folded " << *UseMI); - return; } - if (Size != 4) - return; - - Register Reg0 = UseMI->getOperand(0).getReg(); - Register Reg1 = UseMI->getOperand(1).getReg(); - if (TRI->isAGPR(*MRI, Reg0) && TRI->isVGPR(*MRI, Reg1)) - UseMI->setDesc(TII->get(AMDGPU::V_ACCVGPR_WRITE_B32_e64)); - else if (TRI->isVGPR(*MRI, Reg0) && TRI->isAGPR(*MRI, Reg1)) - UseMI->setDesc(TII->get(AMDGPU::V_ACCVGPR_READ_B32_e64)); - else if (ST->hasGFX90AInsts() && TRI->isAGPR(*MRI, Reg0) && - TRI->isAGPR(*MRI, Reg1)) - UseMI->setDesc(TII->get(AMDGPU::V_ACCVGPR_MOV_B32)); return; } diff --git a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir new file mode 100644 index 0000000000000..15d6140b0206d --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir @@ -0,0 +1,234 @@ +# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5 +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs --run-pass si-fold-operands %s -o - | FileCheck %s +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=COALESCE +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=GFX908-COALESCE + +--- +name: test +tracksRegLiveness: true +body: | + ; CHECK-LABEL: name: test + ; CHECK: bb.0: + ; CHECK-NEXT: successors: %bb.1(0x80000000) + ; CHECK-NEXT: liveins: $sgpr4_sgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr4_sgpr5 + ; CHECK-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[COPY]](p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + ; CHECK-NEXT: S_BITCMP1_B32 killed [[S_LOAD_DWORD_IMM]], 0, implicit-def $scc + ; CHECK-NEXT: [[S_MOV_B32_:%[0-9]+]]:sgpr_32 = S_MOV_B32 0 + ; CHECK-NEXT: [[S_CSELECT_B64_:%[0-9]+]]:sreg_64_xexec = S_CSELECT_B64 -1, 0, implicit $scc + ; CHECK-NEXT: [[V_MOV_B32_e32_:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 0, implicit $exec + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:agpr_32 = COPY [[V_MOV_B32_e32_]] + ; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_32 = IMPLICIT_DEF + ; CHECK-NEXT: [[V_CNDMASK_B32_e64_:%[0-9]+]]:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, [[S_CSELECT_B64_]], implicit $exec + ; CHECK-NEXT: [[V_CMP_NE_U32_e64_:%[0-9]+]]:sreg_64_xexec = V_CMP_NE_U32_e64 [[V_CNDMASK_B32_e64_]], 1, implicit $exec + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.1: + ; CHECK-NEXT: successors: %bb.2(0x40000000), %bb.3(0x40000000) + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[PHI:%[0-9]+]]:agpr_32 = PHI [[COPY1]], %bb.0, %24, %bb.3 + ; CHECK-NEXT: [[PHI1:%[0-9]+]]:sreg_32 = PHI [[S_MOV_B32_]], %bb.0, %11, %bb.3 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY [[PHI]] + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:agpr_32 = COPY [[PHI]] + ; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64 = S_MOV_B64 -1 + ; CHECK-NEXT: $vcc = S_AND_B64 $exec, [[V_CMP_NE_U32_e64_]], implicit-def $scc + ; CHECK-NEXT: S_CBRANCH_VCCNZ %bb.3, implicit $vcc + ; CHECK-NEXT: S_BRANCH %bb.2 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.2: + ; CHECK-NEXT: successors: %bb.3(0x80000000) + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[PHI1]], 1, implicit-def dead $scc + ; CHECK-NEXT: [[S_ASHR_I32_:%[0-9]+]]:sreg_32 = S_ASHR_I32 [[PHI1]], 31, implicit-def dead $scc + ; CHECK-NEXT: [[S_AND_B32_:%[0-9]+]]:sreg_32 = S_AND_B32 killed [[S_ASHR_I32_]], killed [[S_OR_B32_]], implicit-def dead $scc + ; CHECK-NEXT: [[REG_SEQUENCE:%[0-9]+]]:vreg_128_align2 = REG_SEQUENCE [[COPY2]], %subreg.sub0, [[V_MOV_B32_e32_]], %subreg.sub1, [[V_MOV_B32_e32_]], %subreg.sub2, [[V_MOV_B32_e32_]], %subreg.sub3 + ; CHECK-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64 = REG_SEQUENCE [[S_MOV_B32_]], %subreg.sub0, [[S_MOV_B32_]], %subreg.sub1 + ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vreg_64_align2 = COPY [[REG_SEQUENCE1]] + ; CHECK-NEXT: [[COPY5:%[0-9]+]]:areg_128_align2 = COPY [[REG_SEQUENCE]] + ; CHECK-NEXT: [[V_MFMA_F32_16X16X16F16_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY4]], [[COPY4]], killed [[COPY5]], 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[S_MOV_B64_1:%[0-9]+]]:sreg_64 = S_MOV_B64 0 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.3: + ; CHECK-NEXT: successors: %bb.4(0x40000000), %bb.1(0x40000000) + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[PHI2:%[0-9]+]]:sreg_32 = PHI [[DEF]], %bb.1, [[S_AND_B32_]], %bb.2 + ; CHECK-NEXT: [[PHI3:%[0-9]+]]:agpr_32 = PHI [[COPY3]], %bb.1, [[V_MFMA_F32_16X16X16F16_e64_]].sub0, %bb.2 + ; CHECK-NEXT: [[PHI4:%[0-9]+]]:sreg_64_xexec = PHI [[S_MOV_B64_]], %bb.1, [[S_MOV_B64_1]], %bb.2 + ; CHECK-NEXT: [[V_CNDMASK_B32_e64_1:%[0-9]+]]:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, [[PHI4]], implicit $exec + ; CHECK-NEXT: [[V_CMP_NE_U32_e64_1:%[0-9]+]]:sreg_64_xexec = V_CMP_NE_U32_e64 [[V_CNDMASK_B32_e64_1]], 1, implicit $exec + ; CHECK-NEXT: $vcc = S_AND_B64 $exec, [[V_CMP_NE_U32_e64_1]], implicit-def $scc + ; CHECK-NEXT: S_CBRANCH_VCCNZ %bb.1, implicit $vcc + ; CHECK-NEXT: S_BRANCH %bb.4 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.4: + ; CHECK-NEXT: successors: %bb.5(0x80000000) + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.5: + ; CHECK-NEXT: S_ENDPGM 0 + ; + ; COALESCE-LABEL: name: test + ; COALESCE: bb.0: + ; COALESCE-NEXT: successors: %bb.1(0x80000000) + ; COALESCE-NEXT: liveins: $sgpr4_sgpr5 + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: [[COPY:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr4_sgpr5 + ; COALESCE-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[COPY]](p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + ; COALESCE-NEXT: S_BITCMP1_B32 [[S_LOAD_DWORD_IMM]], 0, implicit-def $scc + ; COALESCE-NEXT: undef [[S_MOV_B32_:%[0-9]+]].sub0:sgpr_64 = S_MOV_B32 0 + ; COALESCE-NEXT: [[S_CSELECT_B64_:%[0-9]+]]:sreg_64_xexec = S_CSELECT_B64 -1, 0, implicit killed $scc + ; COALESCE-NEXT: undef [[V_MOV_B32_e32_:%[0-9]+]].sub1:vreg_128_align2 = V_MOV_B32_e32 0, implicit $exec + ; COALESCE-NEXT: undef [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]].sub0:areg_128_align2 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; COALESCE-NEXT: [[V_CNDMASK_B32_e64_:%[0-9]+]]:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, [[S_CSELECT_B64_]], implicit $exec + ; COALESCE-NEXT: [[V_CMP_NE_U32_e64_:%[0-9]+]]:sreg_64_xexec = V_CMP_NE_U32_e64 1, [[V_CNDMASK_B32_e64_]], implicit $exec + ; COALESCE-NEXT: [[S_MOV_B32_1:%[0-9]+]]:sreg_32 = S_MOV_B32 0 + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.1: + ; COALESCE-NEXT: successors: %bb.2(0x40000000), %bb.3(0x40000000) + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: [[V_MOV_B32_e32_:%[0-9]+]].sub0:vreg_128_align2 = COPY [[V_ACCVGPR_WRITE_B32_e64_]].sub0 + ; COALESCE-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 -1 + ; COALESCE-NEXT: $vcc = S_AND_B64 $exec, [[V_CMP_NE_U32_e64_]], implicit-def dead $scc + ; COALESCE-NEXT: S_CBRANCH_VCCNZ %bb.3, implicit killed $vcc + ; COALESCE-NEXT: S_BRANCH %bb.2 + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.2: + ; COALESCE-NEXT: successors: %bb.3(0x80000000) + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[S_MOV_B32_1]], 1, implicit-def dead $scc + ; COALESCE-NEXT: [[S_ASHR_I32_:%[0-9]+]]:sreg_32 = S_ASHR_I32 [[S_MOV_B32_1]], 31, implicit-def dead $scc + ; COALESCE-NEXT: [[S_MOV_B32_1:%[0-9]+]]:sreg_32 = S_AND_B32 [[S_ASHR_I32_]], [[S_OR_B32_]], implicit-def dead $scc + ; COALESCE-NEXT: [[V_MOV_B32_e32_:%[0-9]+]].sub2:vreg_128_align2 = COPY [[V_MOV_B32_e32_]].sub1 + ; COALESCE-NEXT: [[V_MOV_B32_e32_:%[0-9]+]].sub3:vreg_128_align2 = COPY [[V_MOV_B32_e32_]].sub1 + ; COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub1:sgpr_64 = COPY [[S_MOV_B32_]].sub0 + ; COALESCE-NEXT: [[COPY1:%[0-9]+]]:vreg_64_align2 = COPY [[S_MOV_B32_]] + ; COALESCE-NEXT: [[COPY2:%[0-9]+]]:areg_128_align2 = COPY [[V_MOV_B32_e32_]] + ; COALESCE-NEXT: [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[COPY2]], 0, 0, 0, implicit $mode, implicit $exec + ; COALESCE-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 0 + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.3: + ; COALESCE-NEXT: successors: %bb.4(0x40000000), %bb.1(0x40000000) + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: [[V_CNDMASK_B32_e64_1:%[0-9]+]]:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, [[S_MOV_B64_]], implicit $exec + ; COALESCE-NEXT: [[V_CMP_NE_U32_e64_1:%[0-9]+]]:sreg_64_xexec = V_CMP_NE_U32_e64 1, [[V_CNDMASK_B32_e64_1]], implicit $exec + ; COALESCE-NEXT: $vcc = S_AND_B64 $exec, [[V_CMP_NE_U32_e64_1]], implicit-def dead $scc + ; COALESCE-NEXT: S_CBRANCH_VCCNZ %bb.1, implicit killed $vcc + ; COALESCE-NEXT: S_BRANCH %bb.4 + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.4: + ; COALESCE-NEXT: successors: %bb.5(0x80000000) + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.5: + ; COALESCE-NEXT: S_ENDPGM 0 + ; + ; GFX908-COALESCE-LABEL: name: test + ; GFX908-COALESCE: bb.0: + ; GFX908-COALESCE-NEXT: successors: %bb.1(0x80000000) + ; GFX908-COALESCE-NEXT: liveins: $sgpr4_sgpr5 + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: [[COPY:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr4_sgpr5 + ; GFX908-COALESCE-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[COPY]](p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + ; GFX908-COALESCE-NEXT: S_BITCMP1_B32 [[S_LOAD_DWORD_IMM]], 0, implicit-def $scc + ; GFX908-COALESCE-NEXT: undef [[S_MOV_B32_:%[0-9]+]].sub0:sgpr_64 = S_MOV_B32 0 + ; GFX908-COALESCE-NEXT: [[S_CSELECT_B64_:%[0-9]+]]:sreg_64_xexec = S_CSELECT_B64 -1, 0, implicit killed $scc + ; GFX908-COALESCE-NEXT: undef [[V_MOV_B32_e32_:%[0-9]+]].sub1:vreg_128_align2 = V_MOV_B32_e32 0, implicit $exec + ; GFX908-COALESCE-NEXT: undef [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]].sub0:areg_128_align2 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; GFX908-COALESCE-NEXT: [[V_CNDMASK_B32_e64_:%[0-9]+]]:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, [[S_CSELECT_B64_]], implicit $exec + ; GFX908-COALESCE-NEXT: [[V_CMP_NE_U32_e64_:%[0-9]+]]:sreg_64_xexec = V_CMP_NE_U32_e64 1, [[V_CNDMASK_B32_e64_]], implicit $exec + ; GFX908-COALESCE-NEXT: [[S_MOV_B32_1:%[0-9]+]]:sreg_32 = S_MOV_B32 0 + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.1: + ; GFX908-COALESCE-NEXT: successors: %bb.2(0x40000000), %bb.3(0x40000000) + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: [[V_MOV_B32_e32_:%[0-9]+]].sub0:vreg_128_align2 = COPY [[V_ACCVGPR_WRITE_B32_e64_]].sub0 + ; GFX908-COALESCE-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 -1 + ; GFX908-COALESCE-NEXT: $vcc = S_AND_B64 $exec, [[V_CMP_NE_U32_e64_]], implicit-def dead $scc + ; GFX908-COALESCE-NEXT: S_CBRANCH_VCCNZ %bb.3, implicit killed $vcc + ; GFX908-COALESCE-NEXT: S_BRANCH %bb.2 + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.2: + ; GFX908-COALESCE-NEXT: successors: %bb.3(0x80000000) + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: [[S_OR_B32_:%[0-9]+]]:sreg_32 = S_OR_B32 [[S_MOV_B32_1]], 1, implicit-def dead $scc + ; GFX908-COALESCE-NEXT: [[S_ASHR_I32_:%[0-9]+]]:sreg_32 = S_ASHR_I32 [[S_MOV_B32_1]], 31, implicit-def dead $scc + ; GFX908-COALESCE-NEXT: [[S_MOV_B32_1:%[0-9]+]]:sreg_32 = S_AND_B32 [[S_ASHR_I32_]], [[S_OR_B32_]], implicit-def dead $scc + ; GFX908-COALESCE-NEXT: [[V_MOV_B32_e32_:%[0-9]+]].sub2:vreg_128_align2 = COPY [[V_MOV_B32_e32_]].sub1 + ; GFX908-COALESCE-NEXT: [[V_MOV_B32_e32_:%[0-9]+]].sub3:vreg_128_align2 = COPY [[V_MOV_B32_e32_]].sub1 + ; GFX908-COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub1:sgpr_64 = COPY [[S_MOV_B32_]].sub0 + ; GFX908-COALESCE-NEXT: [[COPY1:%[0-9]+]]:vreg_64_align2 = COPY [[S_MOV_B32_]] + ; GFX908-COALESCE-NEXT: [[COPY2:%[0-9]+]]:areg_128_align2 = COPY [[V_MOV_B32_e32_]] + ; GFX908-COALESCE-NEXT: [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[COPY2]], 0, 0, 0, implicit $mode, implicit $exec + ; GFX908-COALESCE-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 0 + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.3: + ; GFX908-COALESCE-NEXT: successors: %bb.4(0x40000000), %bb.1(0x40000000) + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: [[V_CNDMASK_B32_e64_1:%[0-9]+]]:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, [[S_MOV_B64_]], implicit $exec + ; GFX908-COALESCE-NEXT: [[V_CMP_NE_U32_e64_1:%[0-9]+]]:sreg_64_xexec = V_CMP_NE_U32_e64 1, [[V_CNDMASK_B32_e64_1]], implicit $exec + ; GFX908-COALESCE-NEXT: $vcc = S_AND_B64 $exec, [[V_CMP_NE_U32_e64_1]], implicit-def dead $scc + ; GFX908-COALESCE-NEXT: S_CBRANCH_VCCNZ %bb.1, implicit killed $vcc + ; GFX908-COALESCE-NEXT: S_BRANCH %bb.4 + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.4: + ; GFX908-COALESCE-NEXT: successors: %bb.5(0x80000000) + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.5: + ; GFX908-COALESCE-NEXT: S_ENDPGM 0 + bb.0: + successors: %bb.1 + liveins: $sgpr4_sgpr5 + + %521:sgpr_64(p4) = COPY $sgpr4_sgpr5 + %655:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM %521(p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + S_BITCMP1_B32 killed %655, 0, implicit-def $scc + %526:sgpr_32 = S_MOV_B32 0 + %690:sreg_64_xexec = S_CSELECT_B64 -1, 0, implicit $scc + %815:vgpr_32 = V_MOV_B32_e32 0, implicit $exec + %658:sreg_32 = IMPLICIT_DEF + %660:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, %690, implicit $exec + %689:sreg_64_xexec = V_CMP_NE_U32_e64 %660, 1, implicit $exec + + bb.1: + successors: %bb.2, %bb.3 + + %125:vgpr_32 = PHI %815, %bb.0, %384, %bb.3 + %130:sreg_32 = PHI %526, %bb.0, %260, %bb.3 + %820:agpr_32 = COPY %125 + %659:sreg_64 = S_MOV_B64 -1 + $vcc = S_AND_B64 $exec, %689, implicit-def $scc + S_CBRANCH_VCCNZ %bb.3, implicit $vcc + S_BRANCH %bb.2 + + bb.2: + successors: %bb.3 + + %665:sreg_32 = S_OR_B32 %130, 1, implicit-def dead $scc + %667:sreg_32 = S_ASHR_I32 %130, 31, implicit-def dead $scc + %131:sreg_32 = S_AND_B32 killed %667, killed %665, implicit-def dead $scc + %685:vreg_128_align2 = REG_SEQUENCE %125, %subreg.sub0, %815, %subreg.sub1, %815, %subreg.sub2, %815, %subreg.sub3 + %671:sreg_64 = REG_SEQUENCE %526, %subreg.sub0, %526, %subreg.sub1 + %673:vreg_64_align2 = COPY %671 + %675:areg_128_align2 = COPY %685 + %672:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %673, %673, killed %675, 0, 0, 0, implicit $mode, implicit $exec + %255:vgpr_32 = COPY %672.sub0 + %663:sreg_64 = S_MOV_B64 0 + + bb.3: + successors: %bb.4, %bb.1 + + %260:sreg_32 = PHI %658, %bb.1, %131, %bb.2 + %821:agpr_32 = PHI %820, %bb.1, %672.sub0, %bb.2 + %389:sreg_64_xexec = PHI %659, %bb.1, %663, %bb.2 + %384:vgpr_32 = COPY %821 + %676:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, %389, implicit $exec + %684:sreg_64_xexec = V_CMP_NE_U32_e64 %676, 1, implicit $exec + $vcc = S_AND_B64 $exec, %684, implicit-def $scc + S_CBRANCH_VCCNZ %bb.1, implicit $vcc + S_BRANCH %bb.4 + + bb.4: + successors: %bb.5 + + bb.5: + S_ENDPGM 0 + +... diff --git a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir new file mode 100644 index 0000000000000..7ea4cf92c0d64 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir @@ -0,0 +1,181 @@ +# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5 +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs --run-pass si-fold-operands %s -o - | FileCheck %s +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=COALESCE +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=GFX908-COALESCE + +--- +name: test +tracksRegLiveness: true +body: | + ; CHECK-LABEL: name: test + ; CHECK: bb.0: + ; CHECK-NEXT: successors: %bb.2(0x40000000), %bb.1(0x40000000) + ; CHECK-NEXT: liveins: $sgpr4_sgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr4_sgpr5 + ; CHECK-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[COPY]](p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + ; CHECK-NEXT: [[S_MOV_B32_:%[0-9]+]]:sgpr_32 = S_MOV_B32 0 + ; CHECK-NEXT: S_BITCMP0_B32 killed [[S_LOAD_DWORD_IMM]], 0, implicit-def $scc + ; CHECK-NEXT: S_CBRANCH_SCC0 %bb.2, implicit $scc + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.1: + ; CHECK-NEXT: successors: %bb.3(0x80000000) + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]]:agpr_32 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; CHECK-NEXT: S_BRANCH %bb.3 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.2: + ; CHECK-NEXT: successors: %bb.3(0x80000000) + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[S_MOV_B32_1:%[0-9]+]]:sgpr_32 = S_MOV_B32 0 + ; CHECK-NEXT: [[V_ACCVGPR_WRITE_B32_e64_1:%[0-9]+]]:agpr_32 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; CHECK-NEXT: [[V_ACCVGPR_WRITE_B32_e64_2:%[0-9]+]]:agpr_32 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; CHECK-NEXT: [[V_ACCVGPR_WRITE_B32_e64_3:%[0-9]+]]:agpr_32 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; CHECK-NEXT: [[V_ACCVGPR_WRITE_B32_e64_4:%[0-9]+]]:agpr_32 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; CHECK-NEXT: [[REG_SEQUENCE:%[0-9]+]]:areg_128_align2 = REG_SEQUENCE [[V_ACCVGPR_WRITE_B32_e64_1]], %subreg.sub0, [[V_ACCVGPR_WRITE_B32_e64_2]], %subreg.sub1, [[V_ACCVGPR_WRITE_B32_e64_3]], %subreg.sub2, [[V_ACCVGPR_WRITE_B32_e64_4]], %subreg.sub3 + ; CHECK-NEXT: [[REG_SEQUENCE1:%[0-9]+]]:sreg_64 = REG_SEQUENCE [[S_MOV_B32_1]], %subreg.sub0, [[S_MOV_B32_1]], %subreg.sub1 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_64_align2 = COPY [[REG_SEQUENCE1]] + ; CHECK-NEXT: [[V_MFMA_F32_16X16X16F16_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[V_MFMA_F32_16X16X16F16_e64_1:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], killed [[V_MFMA_F32_16X16X16F16_e64_]], 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[V_MFMA_F32_16X16X16F16_e64_2:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], killed [[V_MFMA_F32_16X16X16F16_e64_1]], 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[V_MFMA_F32_16X16X16F16_e64_3:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], killed [[V_MFMA_F32_16X16X16F16_e64_2]], 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY [[V_MFMA_F32_16X16X16F16_e64_3]].sub0 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: bb.3: + ; CHECK-NEXT: [[PHI:%[0-9]+]]:agpr_32 = PHI [[V_ACCVGPR_WRITE_B32_e64_]], %bb.1, [[V_MFMA_F32_16X16X16F16_e64_3]].sub0, %bb.2 + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vgpr_32 = COPY [[PHI]] + ; CHECK-NEXT: [[V_CVT_F16_F32_e64_:%[0-9]+]]:vgpr_32 = nofpexcept V_CVT_F16_F32_e64 0, [[COPY3]], 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[V_PACK_B32_F16_e64_:%[0-9]+]]:vgpr_32 = nofpexcept V_PACK_B32_F16_e64 0, killed [[V_CVT_F16_F32_e64_]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[V_MOV_B32_e32_:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 0, implicit $exec + ; CHECK-NEXT: [[REG_SEQUENCE2:%[0-9]+]]:vreg_64_align2 = REG_SEQUENCE [[V_PACK_B32_F16_e64_]], %subreg.sub0, killed [[V_MOV_B32_e32_]], %subreg.sub1 + ; CHECK-NEXT: [[REG_SEQUENCE3:%[0-9]+]]:sgpr_128 = REG_SEQUENCE [[S_MOV_B32_]], %subreg.sub0, [[S_MOV_B32_]], %subreg.sub1, [[S_MOV_B32_]], %subreg.sub2, [[S_MOV_B32_]], %subreg.sub3 + ; CHECK-NEXT: BUFFER_STORE_DWORDX2_OFFSET_exact [[REG_SEQUENCE2]], killed [[REG_SEQUENCE3]], 0, 0, 0, 0, implicit $exec :: (dereferenceable store (s64) into `ptr addrspace(8) null`, align 1, addrspace 8) + ; CHECK-NEXT: S_ENDPGM 0 + ; + ; COALESCE-LABEL: name: test + ; COALESCE: bb.0: + ; COALESCE-NEXT: successors: %bb.2(0x40000000), %bb.1(0x40000000) + ; COALESCE-NEXT: liveins: $sgpr4_sgpr5 + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: [[COPY:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr4_sgpr5 + ; COALESCE-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[COPY]](p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + ; COALESCE-NEXT: undef [[S_MOV_B32_:%[0-9]+]].sub0:sgpr_128 = S_MOV_B32 0 + ; COALESCE-NEXT: S_BITCMP0_B32 [[S_LOAD_DWORD_IMM]], 0, implicit-def $scc + ; COALESCE-NEXT: S_CBRANCH_SCC0 %bb.2, implicit killed $scc + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.1: + ; COALESCE-NEXT: successors: %bb.3(0x80000000) + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: undef [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]].sub0:areg_128_align2 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; COALESCE-NEXT: S_BRANCH %bb.3 + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.2: + ; COALESCE-NEXT: successors: %bb.3(0x80000000) + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub1:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; COALESCE-NEXT: [[COPY1:%[0-9]+]]:vreg_64_align2 = COPY [[S_MOV_B32_]].sub0_sub1 + ; COALESCE-NEXT: [[V_MFMA_F32_16X16X16F16_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; COALESCE-NEXT: [[V_MFMA_F32_16X16X16F16_e64_1:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[V_MFMA_F32_16X16X16F16_e64_]], 0, 0, 0, implicit $mode, implicit $exec + ; COALESCE-NEXT: [[V_MFMA_F32_16X16X16F16_e64_2:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[V_MFMA_F32_16X16X16F16_e64_1]], 0, 0, 0, implicit $mode, implicit $exec + ; COALESCE-NEXT: [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[V_MFMA_F32_16X16X16F16_e64_2]], 0, 0, 0, implicit $mode, implicit $exec + ; COALESCE-NEXT: {{ $}} + ; COALESCE-NEXT: bb.3: + ; COALESCE-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY [[V_ACCVGPR_WRITE_B32_e64_]].sub0 + ; COALESCE-NEXT: [[V_CVT_F16_F32_e32_:%[0-9]+]]:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 [[COPY2]], implicit $mode, implicit $exec + ; COALESCE-NEXT: undef [[V_PACK_B32_F16_e64_:%[0-9]+]].sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, [[V_CVT_F16_F32_e32_]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; COALESCE-NEXT: [[V_PACK_B32_F16_e64_:%[0-9]+]].sub1:vreg_64_align2 = V_MOV_B32_e32 0, implicit $exec + ; COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub1:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub2:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub3:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; COALESCE-NEXT: BUFFER_STORE_DWORDX2_OFFSET_exact [[V_PACK_B32_F16_e64_]], [[S_MOV_B32_]], 0, 0, 0, 0, implicit $exec :: (dereferenceable store (s64) into `ptr addrspace(8) null`, align 1, addrspace 8) + ; COALESCE-NEXT: S_ENDPGM 0 + ; + ; GFX908-COALESCE-LABEL: name: test + ; GFX908-COALESCE: bb.0: + ; GFX908-COALESCE-NEXT: successors: %bb.2(0x40000000), %bb.1(0x40000000) + ; GFX908-COALESCE-NEXT: liveins: $sgpr4_sgpr5 + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: [[COPY:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr4_sgpr5 + ; GFX908-COALESCE-NEXT: [[S_LOAD_DWORD_IMM:%[0-9]+]]:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM [[COPY]](p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + ; GFX908-COALESCE-NEXT: undef [[S_MOV_B32_:%[0-9]+]].sub0:sgpr_128 = S_MOV_B32 0 + ; GFX908-COALESCE-NEXT: S_BITCMP0_B32 [[S_LOAD_DWORD_IMM]], 0, implicit-def $scc + ; GFX908-COALESCE-NEXT: S_CBRANCH_SCC0 %bb.2, implicit killed $scc + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.1: + ; GFX908-COALESCE-NEXT: successors: %bb.3(0x80000000) + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: undef [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]].sub0:areg_128_align2 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; GFX908-COALESCE-NEXT: S_BRANCH %bb.3 + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.2: + ; GFX908-COALESCE-NEXT: successors: %bb.3(0x80000000) + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: undef [[V_ACCVGPR_WRITE_B32_e64_1:%[0-9]+]].sub0:areg_128_align2 = V_ACCVGPR_WRITE_B32_e64 0, implicit $exec + ; GFX908-COALESCE-NEXT: [[V_ACCVGPR_WRITE_B32_e64_1:%[0-9]+]].sub1:areg_128_align2 = COPY [[V_ACCVGPR_WRITE_B32_e64_1]].sub0 + ; GFX908-COALESCE-NEXT: [[V_ACCVGPR_WRITE_B32_e64_1:%[0-9]+]].sub2:areg_128_align2 = COPY [[V_ACCVGPR_WRITE_B32_e64_1]].sub0 + ; GFX908-COALESCE-NEXT: [[V_ACCVGPR_WRITE_B32_e64_1:%[0-9]+]].sub3:areg_128_align2 = COPY [[V_ACCVGPR_WRITE_B32_e64_1]].sub0 + ; GFX908-COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub1:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; GFX908-COALESCE-NEXT: [[COPY1:%[0-9]+]]:vreg_64_align2 = COPY [[S_MOV_B32_]].sub0_sub1 + ; GFX908-COALESCE-NEXT: [[V_MFMA_F32_16X16X16F16_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[V_ACCVGPR_WRITE_B32_e64_1]], 0, 0, 0, implicit $mode, implicit $exec + ; GFX908-COALESCE-NEXT: [[V_MFMA_F32_16X16X16F16_e64_1:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[V_MFMA_F32_16X16X16F16_e64_]], 0, 0, 0, implicit $mode, implicit $exec + ; GFX908-COALESCE-NEXT: [[V_MFMA_F32_16X16X16F16_e64_2:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[V_MFMA_F32_16X16X16F16_e64_1]], 0, 0, 0, implicit $mode, implicit $exec + ; GFX908-COALESCE-NEXT: [[V_ACCVGPR_WRITE_B32_e64_:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 [[COPY1]], [[COPY1]], [[V_MFMA_F32_16X16X16F16_e64_2]], 0, 0, 0, implicit $mode, implicit $exec + ; GFX908-COALESCE-NEXT: {{ $}} + ; GFX908-COALESCE-NEXT: bb.3: + ; GFX908-COALESCE-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY [[V_ACCVGPR_WRITE_B32_e64_]].sub0 + ; GFX908-COALESCE-NEXT: [[V_CVT_F16_F32_e32_:%[0-9]+]]:vgpr_32 = nofpexcept V_CVT_F16_F32_e32 [[COPY2]], implicit $mode, implicit $exec + ; GFX908-COALESCE-NEXT: undef [[V_PACK_B32_F16_e64_:%[0-9]+]].sub0:vreg_64_align2 = nofpexcept V_PACK_B32_F16_e64 0, [[V_CVT_F16_F32_e32_]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; GFX908-COALESCE-NEXT: [[V_PACK_B32_F16_e64_:%[0-9]+]].sub1:vreg_64_align2 = V_MOV_B32_e32 0, implicit $exec + ; GFX908-COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub1:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; GFX908-COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub2:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; GFX908-COALESCE-NEXT: [[S_MOV_B32_:%[0-9]+]].sub3:sgpr_128 = COPY [[S_MOV_B32_]].sub0 + ; GFX908-COALESCE-NEXT: BUFFER_STORE_DWORDX2_OFFSET_exact [[V_PACK_B32_F16_e64_]], [[S_MOV_B32_]], 0, 0, 0, 0, implicit $exec :: (dereferenceable store (s64) into `ptr addrspace(8) null`, align 1, addrspace 8) + ; GFX908-COALESCE-NEXT: S_ENDPGM 0 + bb.0: + successors: %bb.1, %bb.3 + liveins: $sgpr4_sgpr5 + + %259:sgpr_64(p4) = COPY $sgpr4_sgpr5 + %392:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM %259(p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + %264:sgpr_32 = S_MOV_B32 0 + S_BITCMP0_B32 killed %392, 0, implicit-def $scc + S_CBRANCH_SCC0 %bb.1, implicit $scc + + bb.3: + successors: %bb.2 + + %316:sgpr_32 = COPY %264 + %484:vgpr_32 = COPY %316, implicit $exec + S_BRANCH %bb.2 + + bb.1: + successors: %bb.2 + + %396:sgpr_32 = S_MOV_B32 0 + %424:vgpr_32 = COPY %396 + %425:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec + %427:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec + %429:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec + %431:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec + %403:areg_128_align2 = REG_SEQUENCE %425, %subreg.sub0, %427, %subreg.sub1, %429, %subreg.sub2, %431, %subreg.sub3 + %399:sreg_64 = REG_SEQUENCE %396, %subreg.sub0, %396, %subreg.sub1 + %401:vreg_64_align2 = COPY %399 + %400:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %403, 0, 0, 0, implicit $mode, implicit $exec + %404:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %400, 0, 0, 0, implicit $mode, implicit $exec + %407:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %404, 0, 0, 0, implicit $mode, implicit $exec + %410:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %407, 0, 0, 0, implicit $mode, implicit $exec + %48:vgpr_32 = COPY %410.sub0 + %52:vgpr_32 = COPY %48 + + bb.2: + %180:vgpr_32 = PHI %484, %bb.3, %52, %bb.1 + %413:vgpr_32 = nofpexcept V_CVT_F16_F32_e64 0, %180, 0, 0, implicit $mode, implicit $exec + %415:vgpr_32 = nofpexcept V_PACK_B32_F16_e64 0, killed %413, 0, %264, 0, 0, implicit $mode, implicit $exec + %423:vgpr_32 = V_MOV_B32_e32 0, implicit $exec + %422:vreg_64_align2 = REG_SEQUENCE %415, %subreg.sub0, killed %423, %subreg.sub1 + %419:sgpr_128 = REG_SEQUENCE %264, %subreg.sub0, %264, %subreg.sub1, %264, %subreg.sub2, %264, %subreg.sub3 + %420:vreg_64_align2 = COPY %422 + BUFFER_STORE_DWORDX2_OFFSET_exact killed %420, killed %419, %264, 0, 0, 0, implicit $exec :: (dereferenceable store (s64) into `ptr addrspace(8) null`, align 1, addrspace 8) + S_ENDPGM 0 + +... From 32a1a9a1c09d95af1447b6cdee68ca9bc3af93d9 Mon Sep 17 00:00:00 2001 From: Brendon Cahoon Date: Thu, 19 Dec 2024 16:18:24 -0600 Subject: [PATCH 2/4] Update lit tests to address review commits Remove -verify-machineinstrs Renumber registers --- .../CodeGen/AMDGPU/no-fold-accvgpr-mov.mir | 69 ++++++++-------- .../CodeGen/AMDGPU/no-fold-accvgpr-read.mir | 81 ++++++++++--------- 2 files changed, 76 insertions(+), 74 deletions(-) diff --git a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir index 15d6140b0206d..5c83170563e59 100644 --- a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir +++ b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.mir @@ -1,8 +1,9 @@ # NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5 -# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs --run-pass si-fold-operands %s -o - | FileCheck %s -# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=COALESCE -# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=GFX908-COALESCE +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -run-pass si-fold-operands %s -o - | FileCheck %s +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=COALESCE +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=GFX908-COALESCE +... --- name: test tracksRegLiveness: true @@ -177,51 +178,51 @@ body: | successors: %bb.1 liveins: $sgpr4_sgpr5 - %521:sgpr_64(p4) = COPY $sgpr4_sgpr5 - %655:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM %521(p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) - S_BITCMP1_B32 killed %655, 0, implicit-def $scc - %526:sgpr_32 = S_MOV_B32 0 - %690:sreg_64_xexec = S_CSELECT_B64 -1, 0, implicit $scc - %815:vgpr_32 = V_MOV_B32_e32 0, implicit $exec - %658:sreg_32 = IMPLICIT_DEF - %660:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, %690, implicit $exec - %689:sreg_64_xexec = V_CMP_NE_U32_e64 %660, 1, implicit $exec + %0:sgpr_64(p4) = COPY $sgpr4_sgpr5 + %1:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM %0(p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + S_BITCMP1_B32 killed %1, 0, implicit-def $scc + %2:sgpr_32 = S_MOV_B32 0 + %3:sreg_64_xexec = S_CSELECT_B64 -1, 0, implicit $scc + %4:vgpr_32 = V_MOV_B32_e32 0, implicit $exec + %5:sreg_32 = IMPLICIT_DEF + %6:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, %3, implicit $exec + %7:sreg_64_xexec = V_CMP_NE_U32_e64 %6, 1, implicit $exec bb.1: successors: %bb.2, %bb.3 - %125:vgpr_32 = PHI %815, %bb.0, %384, %bb.3 - %130:sreg_32 = PHI %526, %bb.0, %260, %bb.3 - %820:agpr_32 = COPY %125 - %659:sreg_64 = S_MOV_B64 -1 - $vcc = S_AND_B64 $exec, %689, implicit-def $scc + %8:vgpr_32 = PHI %4, %bb.0, %9, %bb.3 + %10:sreg_32 = PHI %2, %bb.0, %11, %bb.3 + %12:agpr_32 = COPY %8 + %13:sreg_64 = S_MOV_B64 -1 + $vcc = S_AND_B64 $exec, %7, implicit-def $scc S_CBRANCH_VCCNZ %bb.3, implicit $vcc S_BRANCH %bb.2 bb.2: successors: %bb.3 - %665:sreg_32 = S_OR_B32 %130, 1, implicit-def dead $scc - %667:sreg_32 = S_ASHR_I32 %130, 31, implicit-def dead $scc - %131:sreg_32 = S_AND_B32 killed %667, killed %665, implicit-def dead $scc - %685:vreg_128_align2 = REG_SEQUENCE %125, %subreg.sub0, %815, %subreg.sub1, %815, %subreg.sub2, %815, %subreg.sub3 - %671:sreg_64 = REG_SEQUENCE %526, %subreg.sub0, %526, %subreg.sub1 - %673:vreg_64_align2 = COPY %671 - %675:areg_128_align2 = COPY %685 - %672:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %673, %673, killed %675, 0, 0, 0, implicit $mode, implicit $exec - %255:vgpr_32 = COPY %672.sub0 - %663:sreg_64 = S_MOV_B64 0 + %14:sreg_32 = S_OR_B32 %10, 1, implicit-def dead $scc + %15:sreg_32 = S_ASHR_I32 %10, 31, implicit-def dead $scc + %16:sreg_32 = S_AND_B32 killed %15, killed %14, implicit-def dead $scc + %17:vreg_128_align2 = REG_SEQUENCE %8, %subreg.sub0, %4, %subreg.sub1, %4, %subreg.sub2, %4, %subreg.sub3 + %18:sreg_64 = REG_SEQUENCE %2, %subreg.sub0, %2, %subreg.sub1 + %19:vreg_64_align2 = COPY %18 + %20:areg_128_align2 = COPY %17 + %21:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %19, %19, killed %20, 0, 0, 0, implicit $mode, implicit $exec + %22:vgpr_32 = COPY %21.sub0 + %23:sreg_64 = S_MOV_B64 0 bb.3: successors: %bb.4, %bb.1 - %260:sreg_32 = PHI %658, %bb.1, %131, %bb.2 - %821:agpr_32 = PHI %820, %bb.1, %672.sub0, %bb.2 - %389:sreg_64_xexec = PHI %659, %bb.1, %663, %bb.2 - %384:vgpr_32 = COPY %821 - %676:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, %389, implicit $exec - %684:sreg_64_xexec = V_CMP_NE_U32_e64 %676, 1, implicit $exec - $vcc = S_AND_B64 $exec, %684, implicit-def $scc + %11:sreg_32 = PHI %5, %bb.1, %16, %bb.2 + %24:agpr_32 = PHI %12, %bb.1, %21.sub0, %bb.2 + %25:sreg_64_xexec = PHI %13, %bb.1, %23, %bb.2 + %9:vgpr_32 = COPY %24 + %26:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, %25, implicit $exec + %27:sreg_64_xexec = V_CMP_NE_U32_e64 %26, 1, implicit $exec + $vcc = S_AND_B64 $exec, %27, implicit-def $scc S_CBRANCH_VCCNZ %bb.1, implicit $vcc S_BRANCH %bb.4 diff --git a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir index 7ea4cf92c0d64..49c0aaf9fb390 100644 --- a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir +++ b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-read.mir @@ -1,8 +1,9 @@ # NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5 -# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs --run-pass si-fold-operands %s -o - | FileCheck %s -# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=COALESCE -# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 -verify-machineinstrs -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=GFX908-COALESCE +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -run-pass si-fold-operands %s -o - | FileCheck %s +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=COALESCE +# RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 -start-before=si-fold-operands -stop-after=register-coalescer %s -o - | FileCheck %s --check-prefixes=GFX908-COALESCE +... --- name: test tracksRegLiveness: true @@ -132,50 +133,50 @@ body: | ; GFX908-COALESCE-NEXT: BUFFER_STORE_DWORDX2_OFFSET_exact [[V_PACK_B32_F16_e64_]], [[S_MOV_B32_]], 0, 0, 0, 0, implicit $exec :: (dereferenceable store (s64) into `ptr addrspace(8) null`, align 1, addrspace 8) ; GFX908-COALESCE-NEXT: S_ENDPGM 0 bb.0: - successors: %bb.1, %bb.3 + successors: %bb.2, %bb.1 liveins: $sgpr4_sgpr5 - %259:sgpr_64(p4) = COPY $sgpr4_sgpr5 - %392:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM %259(p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) - %264:sgpr_32 = S_MOV_B32 0 - S_BITCMP0_B32 killed %392, 0, implicit-def $scc - S_CBRANCH_SCC0 %bb.1, implicit $scc - - bb.3: - successors: %bb.2 - - %316:sgpr_32 = COPY %264 - %484:vgpr_32 = COPY %316, implicit $exec - S_BRANCH %bb.2 + %0:sgpr_64(p4) = COPY $sgpr4_sgpr5 + %1:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM %0(p4), 0, 0 :: (dereferenceable invariant load (s32), align 16, addrspace 4) + %2:sgpr_32 = S_MOV_B32 0 + S_BITCMP0_B32 killed %1, 0, implicit-def $scc + S_CBRANCH_SCC0 %bb.2, implicit $scc bb.1: - successors: %bb.2 + successors: %bb.3 - %396:sgpr_32 = S_MOV_B32 0 - %424:vgpr_32 = COPY %396 - %425:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec - %427:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec - %429:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec - %431:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %424, implicit $exec - %403:areg_128_align2 = REG_SEQUENCE %425, %subreg.sub0, %427, %subreg.sub1, %429, %subreg.sub2, %431, %subreg.sub3 - %399:sreg_64 = REG_SEQUENCE %396, %subreg.sub0, %396, %subreg.sub1 - %401:vreg_64_align2 = COPY %399 - %400:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %403, 0, 0, 0, implicit $mode, implicit $exec - %404:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %400, 0, 0, 0, implicit $mode, implicit $exec - %407:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %404, 0, 0, 0, implicit $mode, implicit $exec - %410:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %401, %401, killed %407, 0, 0, 0, implicit $mode, implicit $exec - %48:vgpr_32 = COPY %410.sub0 - %52:vgpr_32 = COPY %48 + %3:sgpr_32 = COPY %2 + %4:vgpr_32 = COPY %3, implicit $exec + S_BRANCH %bb.3 bb.2: - %180:vgpr_32 = PHI %484, %bb.3, %52, %bb.1 - %413:vgpr_32 = nofpexcept V_CVT_F16_F32_e64 0, %180, 0, 0, implicit $mode, implicit $exec - %415:vgpr_32 = nofpexcept V_PACK_B32_F16_e64 0, killed %413, 0, %264, 0, 0, implicit $mode, implicit $exec - %423:vgpr_32 = V_MOV_B32_e32 0, implicit $exec - %422:vreg_64_align2 = REG_SEQUENCE %415, %subreg.sub0, killed %423, %subreg.sub1 - %419:sgpr_128 = REG_SEQUENCE %264, %subreg.sub0, %264, %subreg.sub1, %264, %subreg.sub2, %264, %subreg.sub3 - %420:vreg_64_align2 = COPY %422 - BUFFER_STORE_DWORDX2_OFFSET_exact killed %420, killed %419, %264, 0, 0, 0, implicit $exec :: (dereferenceable store (s64) into `ptr addrspace(8) null`, align 1, addrspace 8) + successors: %bb.3 + + %5:sgpr_32 = S_MOV_B32 0 + %6:vgpr_32 = COPY %5 + %7:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %6, implicit $exec + %8:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %6, implicit $exec + %9:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %6, implicit $exec + %10:agpr_32 = V_ACCVGPR_WRITE_B32_e64 %6, implicit $exec + %11:areg_128_align2 = REG_SEQUENCE %7, %subreg.sub0, %8, %subreg.sub1, %9, %subreg.sub2, %10, %subreg.sub3 + %12:sreg_64 = REG_SEQUENCE %5, %subreg.sub0, %5, %subreg.sub1 + %13:vreg_64_align2 = COPY %12 + %14:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %13, %13, killed %11, 0, 0, 0, implicit $mode, implicit $exec + %15:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %13, %13, killed %14, 0, 0, 0, implicit $mode, implicit $exec + %16:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %13, %13, killed %15, 0, 0, 0, implicit $mode, implicit $exec + %17:areg_128_align2 = V_MFMA_F32_16X16X16F16_e64 %13, %13, killed %16, 0, 0, 0, implicit $mode, implicit $exec + %18:vgpr_32 = COPY %17.sub0 + %19:vgpr_32 = COPY %18 + + bb.3: + %20:vgpr_32 = PHI %4, %bb.1, %19, %bb.2 + %21:vgpr_32 = nofpexcept V_CVT_F16_F32_e64 0, %20, 0, 0, implicit $mode, implicit $exec + %22:vgpr_32 = nofpexcept V_PACK_B32_F16_e64 0, killed %21, 0, %2, 0, 0, implicit $mode, implicit $exec + %23:vgpr_32 = V_MOV_B32_e32 0, implicit $exec + %24:vreg_64_align2 = REG_SEQUENCE %22, %subreg.sub0, killed %23, %subreg.sub1 + %25:sgpr_128 = REG_SEQUENCE %2, %subreg.sub0, %2, %subreg.sub1, %2, %subreg.sub2, %2, %subreg.sub3 + %26:vreg_64_align2 = COPY %24 + BUFFER_STORE_DWORDX2_OFFSET_exact killed %26, killed %25, %2, 0, 0, 0, implicit $exec :: (dereferenceable store (s64) into `ptr addrspace(8) null`, align 1, addrspace 8) S_ENDPGM 0 ... From 4d72ab22e303d7389521bbc3bfbff37f9f232ec4 Mon Sep 17 00:00:00 2001 From: Brendon Cahoon Date: Thu, 19 Dec 2024 17:44:43 -0600 Subject: [PATCH 3/4] Add end-to-end test --- .../CodeGen/AMDGPU/no-fold-accvgpr-mov.ll | 124 ++++++++++++++++++ 1 file changed, 124 insertions(+) create mode 100644 llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll diff --git a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll new file mode 100644 index 0000000000000..6ccc0aa4aaa3d --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll @@ -0,0 +1,124 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: llc -O3 -mtriple=amdgcn--amdhsa -mcpu=gfx942 %s -o - | FileCheck %s --check-prefixes=GFX942 +; RUN: llc -O3 -mtriple=amdgcn--amdhsa -mcpu=gfx908 %s -o - | FileCheck %s --check-prefixes=GFX908 + +define amdgpu_kernel void @test(i1 %arg, i1 %arg1) { +; GFX942-LABEL: test: +; GFX942: ; %bb.0: ; %bb +; GFX942-NEXT: s_load_dword s0, s[4:5], 0x0 +; GFX942-NEXT: v_mov_b32_e32 v1, 0 +; GFX942-NEXT: s_mov_b32 s2, 0 +; GFX942-NEXT: v_accvgpr_write_b32 a0, v1 +; GFX942-NEXT: s_mov_b32 s3, 0 +; GFX942-NEXT: s_waitcnt lgkmcnt(0) +; GFX942-NEXT: s_bitcmp1_b32 s0, 0 +; GFX942-NEXT: s_cselect_b64 s[0:1], -1, 0 +; GFX942-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1] +; GFX942-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0 +; GFX942-NEXT: s_branch .LBB0_2 +; GFX942-NEXT: .LBB0_1: ; %Flow +; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1 +; GFX942-NEXT: s_andn2_b64 vcc, exec, s[4:5] +; GFX942-NEXT: s_cbranch_vccz .LBB0_4 +; GFX942-NEXT: .LBB0_2: ; %bb2 +; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1 +; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1] +; GFX942-NEXT: s_mov_b64 s[4:5], -1 +; GFX942-NEXT: s_cbranch_vccnz .LBB0_1 +; GFX942-NEXT: ; %bb.3: ; %bb4 +; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1 +; GFX942-NEXT: s_or_b32 s4, s3, 1 +; GFX942-NEXT: s_ashr_i32 s5, s3, 31 +; GFX942-NEXT: s_mov_b32 s3, s2 +; GFX942-NEXT: v_mov_b64_e32 v[4:5], s[2:3] +; GFX942-NEXT: v_accvgpr_read_b32 v0, a0 +; GFX942-NEXT: v_mov_b32_e32 v2, v1 +; GFX942-NEXT: v_mov_b32_e32 v3, v1 +; GFX942-NEXT: v_accvgpr_write_b32 a0, v0 +; GFX942-NEXT: v_accvgpr_write_b32 a1, v1 +; GFX942-NEXT: v_accvgpr_write_b32 a2, v2 +; GFX942-NEXT: v_accvgpr_write_b32 a3, v3 +; GFX942-NEXT: s_and_b32 s3, s5, s4 +; GFX942-NEXT: s_mov_b64 s[4:5], 0 +; GFX942-NEXT: v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[4:5], a[0:3] +; GFX942-NEXT: s_branch .LBB0_1 +; GFX942-NEXT: .LBB0_4: ; %common.ret +; GFX942-NEXT: s_endpgm +; +; GFX908-LABEL: test: +; GFX908: ; %bb.0: ; %bb +; GFX908-NEXT: s_load_dword s0, s[8:9], 0x0 +; GFX908-NEXT: v_mov_b32_e32 v1, 0 +; GFX908-NEXT: s_mov_b32 s2, 0 +; GFX908-NEXT: s_mov_b32 s3, 0 +; GFX908-NEXT: v_accvgpr_write_b32 a0, v1 +; GFX908-NEXT: s_waitcnt lgkmcnt(0) +; GFX908-NEXT: s_bitcmp1_b32 s0, 0 +; GFX908-NEXT: s_cselect_b64 s[0:1], -1, 0 +; GFX908-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1] +; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0 +; GFX908-NEXT: s_branch .LBB0_2 +; GFX908-NEXT: .LBB0_1: ; %Flow +; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1 +; GFX908-NEXT: s_andn2_b64 vcc, exec, s[4:5] +; GFX908-NEXT: s_cbranch_vccz .LBB0_4 +; GFX908-NEXT: .LBB0_2: ; %bb2 +; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1 +; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1] +; GFX908-NEXT: s_mov_b64 s[4:5], -1 +; GFX908-NEXT: s_cbranch_vccnz .LBB0_1 +; GFX908-NEXT: ; %bb.3: ; %bb4 +; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1 +; GFX908-NEXT: s_or_b32 s4, s3, 1 +; GFX908-NEXT: s_ashr_i32 s5, s3, 31 +; GFX908-NEXT: s_mov_b32 s3, s2 +; GFX908-NEXT: v_accvgpr_read_b32 v0, a0 +; GFX908-NEXT: v_mov_b32_e32 v5, s3 +; GFX908-NEXT: v_mov_b32_e32 v4, s2 +; GFX908-NEXT: v_mov_b32_e32 v2, v1 +; GFX908-NEXT: v_mov_b32_e32 v3, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a0, v0 +; GFX908-NEXT: v_accvgpr_write_b32 a1, v1 +; GFX908-NEXT: v_accvgpr_write_b32 a2, v2 +; GFX908-NEXT: v_accvgpr_write_b32 a3, v3 +; GFX908-NEXT: s_and_b32 s3, s5, s4 +; GFX908-NEXT: v_mfma_f32_16x16x16f16 a[0:3], v[4:5], v[4:5], a[0:3] +; GFX908-NEXT: s_mov_b64 s[4:5], 0 +; GFX908-NEXT: s_branch .LBB0_1 +; GFX908-NEXT: .LBB0_4: ; %common.ret +; GFX908-NEXT: s_endpgm +bb: + br label %bb2 + +bb2: + %i = phi { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } [ %i12, %bb4 ], [ zeroinitializer, %bb ] + %i3 = phi i32 [ %i7, %bb4 ], [ 0, %bb ] + br i1 %arg, label %bb4, label %bb13 + +bb4: + %i5 = or i32 %i3, 1 + %i6 = icmp slt i32 %i3, 0 + %i7 = select i1 %i6, i32 %i5, i32 0 + %i8 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 123 + %i9 = insertelement <4 x float> zeroinitializer, float %i8, i32 0 + %i10 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i9, i32 0, i32 0, i32 0) + %i11 = extractelement <4 x float> %i10, i32 0 + %i12 = insertvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } zeroinitializer, float %i11, 123 + br label %bb2 + +bb13: + br i1 %arg1, label %bb14, label %common.ret + +common.ret: + ret void + +bb14: + %i15 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 0 + %i16 = insertelement <4 x float> zeroinitializer, float %i15, i32 0 + br label %common.ret +} + +; Function Attrs: convergent nocallback nofree nosync nounwind willreturn memory(none) +declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg) #0 + +attributes #0 = { convergent nocallback nofree nosync nounwind willreturn memory(none) } From 4a326e8a0fed14077024488d3d3f1d98ceaca34b Mon Sep 17 00:00:00 2001 From: Brendon Cahoon Date: Mon, 6 Jan 2025 14:48:23 -0600 Subject: [PATCH 4/4] Update no-fold-accvgpr-mov.ll test Changed i1 parameter to i32 Don't specify optimization level --- .../CodeGen/AMDGPU/no-fold-accvgpr-mov.ll | 112 +++++++++--------- 1 file changed, 54 insertions(+), 58 deletions(-) diff --git a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll index 6ccc0aa4aaa3d..a9b8663a48dea 100644 --- a/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll +++ b/llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll @@ -1,31 +1,22 @@ ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 -; RUN: llc -O3 -mtriple=amdgcn--amdhsa -mcpu=gfx942 %s -o - | FileCheck %s --check-prefixes=GFX942 -; RUN: llc -O3 -mtriple=amdgcn--amdhsa -mcpu=gfx908 %s -o - | FileCheck %s --check-prefixes=GFX908 +; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 %s -o - | FileCheck %s --check-prefixes=GFX942 +; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 %s -o - | FileCheck %s --check-prefixes=GFX908 -define amdgpu_kernel void @test(i1 %arg, i1 %arg1) { -; GFX942-LABEL: test: -; GFX942: ; %bb.0: ; %bb -; GFX942-NEXT: s_load_dword s0, s[4:5], 0x0 +define amdgpu_kernel void @matmul_kernel(i32 %a0, i32 %a1) { +; GFX942-LABEL: matmul_kernel: +; GFX942: ; %bb.0: ; %entry +; GFX942-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0 ; GFX942-NEXT: v_mov_b32_e32 v1, 0 ; GFX942-NEXT: s_mov_b32 s2, 0 ; GFX942-NEXT: v_accvgpr_write_b32 a0, v1 ; GFX942-NEXT: s_mov_b32 s3, 0 ; GFX942-NEXT: s_waitcnt lgkmcnt(0) -; GFX942-NEXT: s_bitcmp1_b32 s0, 0 +; GFX942-NEXT: s_cmp_lg_u32 s0, 0 ; GFX942-NEXT: s_cselect_b64 s[0:1], -1, 0 ; GFX942-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1] ; GFX942-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0 ; GFX942-NEXT: s_branch .LBB0_2 -; GFX942-NEXT: .LBB0_1: ; %Flow -; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1 -; GFX942-NEXT: s_andn2_b64 vcc, exec, s[4:5] -; GFX942-NEXT: s_cbranch_vccz .LBB0_4 -; GFX942-NEXT: .LBB0_2: ; %bb2 -; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1 -; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1] -; GFX942-NEXT: s_mov_b64 s[4:5], -1 -; GFX942-NEXT: s_cbranch_vccnz .LBB0_1 -; GFX942-NEXT: ; %bb.3: ; %bb4 +; GFX942-NEXT: .LBB0_1: ; %bb2 ; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1 ; GFX942-NEXT: s_or_b32 s4, s3, 1 ; GFX942-NEXT: s_ashr_i32 s5, s3, 31 @@ -39,39 +30,37 @@ define amdgpu_kernel void @test(i1 %arg, i1 %arg1) { ; GFX942-NEXT: v_accvgpr_write_b32 a2, v2 ; GFX942-NEXT: v_accvgpr_write_b32 a3, v3 ; GFX942-NEXT: s_and_b32 s3, s5, s4 -; GFX942-NEXT: s_mov_b64 s[4:5], 0 +; GFX942-NEXT: s_nop 0 ; GFX942-NEXT: v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[4:5], a[0:3] -; GFX942-NEXT: s_branch .LBB0_1 +; GFX942-NEXT: s_cbranch_execz .LBB0_4 +; GFX942-NEXT: .LBB0_2: ; %bb +; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1 +; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1] +; GFX942-NEXT: s_cbranch_vccz .LBB0_1 +; GFX942-NEXT: ; %bb.3: +; GFX942-NEXT: ; implicit-def: $sgpr3 ; GFX942-NEXT: .LBB0_4: ; %common.ret ; GFX942-NEXT: s_endpgm ; -; GFX908-LABEL: test: -; GFX908: ; %bb.0: ; %bb -; GFX908-NEXT: s_load_dword s0, s[8:9], 0x0 +; GFX908-LABEL: matmul_kernel: +; GFX908: ; %bb.0: ; %entry +; GFX908-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0 ; GFX908-NEXT: v_mov_b32_e32 v1, 0 ; GFX908-NEXT: s_mov_b32 s2, 0 ; GFX908-NEXT: s_mov_b32 s3, 0 ; GFX908-NEXT: v_accvgpr_write_b32 a0, v1 ; GFX908-NEXT: s_waitcnt lgkmcnt(0) -; GFX908-NEXT: s_bitcmp1_b32 s0, 0 +; GFX908-NEXT: s_cmp_lg_u32 s0, 0 ; GFX908-NEXT: s_cselect_b64 s[0:1], -1, 0 ; GFX908-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1] ; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0 ; GFX908-NEXT: s_branch .LBB0_2 -; GFX908-NEXT: .LBB0_1: ; %Flow -; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1 -; GFX908-NEXT: s_andn2_b64 vcc, exec, s[4:5] -; GFX908-NEXT: s_cbranch_vccz .LBB0_4 -; GFX908-NEXT: .LBB0_2: ; %bb2 -; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1 -; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1] -; GFX908-NEXT: s_mov_b64 s[4:5], -1 -; GFX908-NEXT: s_cbranch_vccnz .LBB0_1 -; GFX908-NEXT: ; %bb.3: ; %bb4 +; GFX908-NEXT: .LBB0_1: ; %bb2 ; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1 ; GFX908-NEXT: s_or_b32 s4, s3, 1 ; GFX908-NEXT: s_ashr_i32 s5, s3, 31 ; GFX908-NEXT: s_mov_b32 s3, s2 +; GFX908-NEXT: s_nop 3 ; GFX908-NEXT: v_accvgpr_read_b32 v0, a0 ; GFX908-NEXT: v_mov_b32_e32 v5, s3 ; GFX908-NEXT: v_mov_b32_e32 v4, s2 @@ -83,42 +72,49 @@ define amdgpu_kernel void @test(i1 %arg, i1 %arg1) { ; GFX908-NEXT: v_accvgpr_write_b32 a3, v3 ; GFX908-NEXT: s_and_b32 s3, s5, s4 ; GFX908-NEXT: v_mfma_f32_16x16x16f16 a[0:3], v[4:5], v[4:5], a[0:3] -; GFX908-NEXT: s_mov_b64 s[4:5], 0 -; GFX908-NEXT: s_branch .LBB0_1 +; GFX908-NEXT: s_cbranch_execz .LBB0_4 +; GFX908-NEXT: .LBB0_2: ; %bb +; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1 +; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1] +; GFX908-NEXT: s_cbranch_vccz .LBB0_1 +; GFX908-NEXT: ; %bb.3: +; GFX908-NEXT: ; implicit-def: $sgpr3 ; GFX908-NEXT: .LBB0_4: ; %common.ret ; GFX908-NEXT: s_endpgm +entry: + br label %bb + bb: - br label %bb2 + %i = phi { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } [ %i10, %bb2 ], [ zeroinitializer, %entry ] + %i1 = phi i32 [ %i5, %bb2 ], [ 0, %entry ] + %c0 = icmp ne i32 %a0, 0 + br i1 %c0, label %bb2, label %bb11 bb2: - %i = phi { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } [ %i12, %bb4 ], [ zeroinitializer, %bb ] - %i3 = phi i32 [ %i7, %bb4 ], [ 0, %bb ] - br i1 %arg, label %bb4, label %bb13 - -bb4: - %i5 = or i32 %i3, 1 - %i6 = icmp slt i32 %i3, 0 - %i7 = select i1 %i6, i32 %i5, i32 0 - %i8 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 123 - %i9 = insertelement <4 x float> zeroinitializer, float %i8, i32 0 - %i10 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i9, i32 0, i32 0, i32 0) - %i11 = extractelement <4 x float> %i10, i32 0 - %i12 = insertvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } zeroinitializer, float %i11, 123 - br label %bb2 + %i3 = or i32 %i1, 1 + %i4 = icmp slt i32 %i1, 0 + %i5 = select i1 %i4, i32 %i3, i32 0 + %i6 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 123 + %i7 = insertelement <4 x float> zeroinitializer, float %i6, i32 0 + %i8 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i7, i32 0, i32 0, i32 0) + %i9 = extractelement <4 x float> %i8, i32 0 + %i10 = insertvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } zeroinitializer, float %i9, 123 + br label %bb -bb13: - br i1 %arg1, label %bb14, label %common.ret +bb11: + %c1 = icmp ne i32 %a1, 0 + br i1 %c1, label %bb12, label %common.ret common.ret: ret void -bb14: - %i15 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 0 - %i16 = insertelement <4 x float> zeroinitializer, float %i15, i32 0 +bb12: + %i13 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 0 + %i14 = insertelement <4 x float> zeroinitializer, float %i13, i32 0 + %i15 = insertelement <4 x float> %i14, float 0.000000e+00, i32 0 + %i16 = insertelement <4 x float> %i15, float 0.000000e+00, i32 0 br label %common.ret } ; Function Attrs: convergent nocallback nofree nosync nounwind willreturn memory(none) -declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg) #0 - -attributes #0 = { convergent nocallback nofree nosync nounwind willreturn memory(none) } +declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)