AMDGPU: Set more mov flags on V_ACCVGPR_{READ|WRITE}_B32

[arsenm]

Summary

This fixes extra copies when materializing constants in AGPRs. This
made it a lot harder to trigger the spilling in spill-agpr.ll

[rampitec]

I would not be so sure it is trivially materializable and cheap as move in a general case given the limitation on the sources and 3 cycle hazard after it. But let's pretend it is.

[arsenm]

In D82996#2126680, @rampitec wrote:

I would not be so sure it is trivially materializable and cheap as move in a general case given the limitation on the sources and 3 cycle hazard after it. But let's pretend it is.

It is itself, the hazards are a different question. The source constraints also don't matter

d2e74fad20bf8cf66ff20a43fe2934d71e046528

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

bool SIInstrInfo::isReallyTriviallyReMaterializable(const MachineInstr &MI,
                                                     AliasAnalysis *AA) const {
   // TODO: The generic check fails for VALU instructions that should be
   // rematerializable due to implicit reads of exec. We really want all of the
   // generic logic for this except for this.
   switch (MI.getOpcode()) {
   case AMDGPU::V_MOV_B32_e32:
   case AMDGPU::V_MOV_B32_e64:
   case AMDGPU::V_MOV_B64_PSEUDO:
+  case AMDGPU::V_ACCVGPR_READ_B32:
+  case AMDGPU::V_ACCVGPR_WRITE_B32:
     // No implicit operands.
     return MI.getNumOperands() == MI.getDesc().getNumOperands();
   default:
     return false;
   }
 }
 
 bool SIInstrInfo::areLoadsFromSameBasePtr(SDNode *Load0, SDNode *Load1,

llvm/lib/Target/AMDGPU/VOP3PInstructions.td

 def VOPProfileMAI_F32_V2I16_X4  : VOPProfileMAI<VOP_V4F32_V2I16_V2I16_V4F32,   AISrc_128_b32,  ADst_128>;
 def VOPProfileMAI_F32_V2I16_X16 : VOPProfileMAI<VOP_V16F32_V2I16_V2I16_V16F32, AISrc_512_b32,  ADst_512>;
 def VOPProfileMAI_F32_V2I16_X32 : VOPProfileMAI<VOP_V32F32_V2I16_V2I16_V32F32, AISrc_1024_b32, ADst_1024>;
 def VOPProfileMAI_F32_V4F16_X4  : VOPProfileMAI<VOP_V4F32_V4F16_V4F16_V4F32,   AISrc_128_b32,  ADst_128,  AVSrc_64>;
 def VOPProfileMAI_F32_V4F16_X16 : VOPProfileMAI<VOP_V16F32_V4F16_V4F16_V16F32, AISrc_512_b32,  ADst_512,  AVSrc_64>;
 def VOPProfileMAI_F32_V4F16_X32 : VOPProfileMAI<VOP_V32F32_V4F16_V4F16_V32F32, AISrc_1024_b32, ADst_1024, AVSrc_64>;
 
 let Predicates = [HasMAIInsts] in {
+
+let isAsCheapAsAMove = 1, isReMaterializable = 1 in {
 def V_ACCVGPR_READ_B32  : VOP3Inst<"v_accvgpr_read_b32",  VOPProfileAccRead>;
 def V_ACCVGPR_WRITE_B32 : VOP3Inst<"v_accvgpr_write_b32", VOPProfileAccWrite> {
   let isMoveImm = 1;
 }
+}
 
 // FP32 denorm mode is respected, rounding mode is not. Exceptions are not supported.
 let isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 in {
 def V_MFMA_F32_4X4X1F32    : VOP3Inst<"v_mfma_f32_4x4x1f32",    VOPProfileMAI_F32_F32_X4,    int_amdgcn_mfma_f32_4x4x1f32>;
 def V_MFMA_F32_4X4X4F16    : VOP3Inst<"v_mfma_f32_4x4x4f16",    VOPProfileMAI_F32_V4F16_X4,  int_amdgcn_mfma_f32_4x4x4f16>;
 def V_MFMA_I32_4X4X4I8     : VOP3Inst<"v_mfma_i32_4x4x4i8",     VOPProfileMAI_I32_I32_X4,    int_amdgcn_mfma_i32_4x4x4i8>;
 def V_MFMA_F32_4X4X2BF16   : VOP3Inst<"v_mfma_f32_4x4x2bf16",   VOPProfileMAI_F32_V2I16_X4,  int_amdgcn_mfma_f32_4x4x2bf16>;
 def V_MFMA_F32_16X16X1F32  : VOP3Inst<"v_mfma_f32_16x16x1f32",  VOPProfileMAI_F32_F32_X16,   int_amdgcn_mfma_f32_16x16x1f32>;

llvm/test/CodeGen/AMDGPU/agpr-remat.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GFX908 %s
+
+; Make sure there are no v_accvgpr_read_b32 copying back and forth
+; between AGPR and VGPR.
+define amdgpu_kernel void @remat_constant_voids_spill(i32 addrspace(1)* %p) #1 {
+; GFX908-LABEL: remat_constant_voids_spill:
+; GFX908:       ; %bb.0:
+; GFX908-NEXT:    v_accvgpr_write_b32 a1, 1
+; GFX908-NEXT:    v_accvgpr_write_b32 a5, 6
+; GFX908-NEXT:    v_accvgpr_write_b32 a6, 7
+; GFX908-NEXT:    v_accvgpr_write_b32 a7, 8
+; GFX908-NEXT:    v_accvgpr_write_b32 a0, 9
+; GFX908-NEXT:    v_accvgpr_write_b32 a2, 2
+; GFX908-NEXT:    v_accvgpr_write_b32 a3, 3
+; GFX908-NEXT:    v_accvgpr_write_b32 a4, 4
+; GFX908-NEXT:    ;;#ASMSTART
+; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    v_accvgpr_write_b32 a1, 5
+; GFX908-NEXT:    ;;#ASMSTART
+; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_endpgm
+  call void asm sideeffect "", "a,a,a,a"(i32 1, i32 2, i32 3, i32 4)
+  call void asm sideeffect "", "a,a,a,a,a"(i32 5, i32 6, i32 7, i32 8, i32 9)
+  ret void
+}
+
+define void @remat_regcopy_avoids_spill(i32 %v0, i32 %v1, i32 %v2, i32 %v3, i32 %v4, i32 %v5, i32 %v6, i32 %v7, i32 %v8, i32 %v9, i32 %v10) #1 {
+; GFX908-LABEL: remat_regcopy_avoids_spill:
+; GFX908:       ; %bb.0:
+; GFX908-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX908-NEXT:    v_accvgpr_write_b32 a1, v0
+; GFX908-NEXT:    v_accvgpr_write_b32 a2, v1
+; GFX908-NEXT:    v_accvgpr_write_b32 a3, v2
+; GFX908-NEXT:    v_accvgpr_write_b32 a4, v3
+; GFX908-NEXT:    ;;#ASMSTART
+; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    v_accvgpr_write_b32 a0, v8
+; GFX908-NEXT:    v_accvgpr_write_b32 a1, v4
+; GFX908-NEXT:    v_accvgpr_write_b32 a2, v5
+; GFX908-NEXT:    v_accvgpr_write_b32 a3, v6
+; GFX908-NEXT:    v_accvgpr_write_b32 a4, v7
+; GFX908-NEXT:    ;;#ASMSTART
+; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_setpc_b64 s[30:31]
+  call void asm sideeffect "", "a,a,a,a"(i32 %v0, i32 %v1, i32 %v2, i32 %v3)
+  call void asm sideeffect "", "a,a,a,a,a"(i32 %v4, i32 %v5, i32 %v6, i32 %v7, i32 %v8)
+  ret void
+}
+
+attributes #1 = { nounwind "amdgpu-num-vgpr"="8" }

llvm/test/CodeGen/AMDGPU/spill-agpr.ll

st:
   store <4 x float> %mai.2, <4 x float> addrspace(1)* %gep2
   ret void
 }
 
 ; GCN-LABEL: {{^}}max_10_vgprs_used_9a:
 ; A2M-DAG:    s_mov_b32 s{{[0-9]+}}, SCRATCH_RSRC_DWORD0
 ; A2M-DAG:    s_mov_b32 s{{[0-9]+}}, SCRATCH_RSRC_DWORD1
 ; A2V-NOT:    SCRATCH_RSRC
-; GFX908-DAG: v_accvgpr_read_b32 v[[VSPILL:[0-9]+]], a0
-; A2M:        buffer_store_dword v[[VSPILL]], off, s[{{[0-9:]+}}], 0 offset:[[FI:[0-9]+]] ; 4-byte Folded Spill
-; A2M:        buffer_load_dword v[[VSPILL:[0-9]+]], off, s[{{[0-9:]+}}], 0 offset:[[FI]] ; 4-byte Folded Reload
-; GFX908:     v_accvgpr_write_b32 a{{[0-9]+}}, v[[VSPILL]]
-; A2V:        ScratchSize: 0
-define amdgpu_kernel void @max_10_vgprs_used_9a(i32 addrspace(1)* %p) #1 {
-  %tid = call i32 @llvm.amdgcn.workitem.id.x()
-  call void asm sideeffect "", "a,a,a,a"(i32 1, i32 2, i32 3, i32 4)
-  call void asm sideeffect "", "a,a,a,a,a"(i32 5, i32 6, i32 7, i32 8, i32 9)
+
+; A2V: v_accvgpr_read_b32 v[[VSPILL:[0-9]+]], a{{[0-9]+}}
+; A2V: v_accvgpr_write_b32 a{{[0-9]+}}, v[[VSPILL]]
+; A2V: ScratchSize: 0
+
+; A2M: buffer_store_dword v[[VSPILLSTORE:[0-9]+]], off, s[{{[0-9:]+}}], 0 offset:[[FI:[0-9]+]] ; 4-byte Folded Spill
+; A2M: buffer_load_dword v[[VSPILL_RELOAD:[0-9]+]], off, s[{{[0-9:]+}}], 0 offset:[[FI]] ; 4-byte Folded Reload
+; A2M: v_accvgpr_write_b32 a{{[0-9]+}}, v[[VSPILL_RELOAD]]
+define amdgpu_kernel void @max_10_vgprs_used_9a() #1 {
+  %v0 = load volatile i32, i32 addrspace(3)* undef
+  %v1 = load volatile i32, i32 addrspace(3)* undef
+  %v2 = load volatile i32, i32 addrspace(3)* undef
+  %v3 = load volatile i32, i32 addrspace(3)* undef
+  %v4 = load volatile i32, i32 addrspace(3)* undef
+  %v5 = load volatile i32, i32 addrspace(3)* undef
+  %v6 = load volatile i32, i32 addrspace(3)* undef
+  %v7 = load volatile i32, i32 addrspace(3)* undef
+  call void asm sideeffect "", "a,a,a,a,~{v0},~{v1},~{v2},~{v3},~{v4},~{v5},~{v6}"(i32 %v0, i32 %v1, i32 %v2, i32 %v3)
+  %v8 = load volatile i32, i32 addrspace(3)* undef
+  call void asm sideeffect "", "a,a,a,a,a"(i32 %v4, i32 %v5, i32 %v6, i32 %v7, i32 %v8)
   ret void
 }
 
 ; GCN-LABEL: {{^}}max_32regs_mfma32:
 ; A2M-DAG:    s_mov_b32 s{{[0-9]+}}, SCRATCH_RSRC_DWORD0
 ; A2M-DAG:    s_mov_b32 s{{[0-9]+}}, SCRATCH_RSRC_DWORD1
 ; A2V-NOT:    SCRATCH_RSRC
 ; GFX908-DAG: v_accvgpr_read_b32 v[[VSPILL:[0-9]+]], a0