llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Matt Arsenault (arsenm) <details> <summary>Changes</summary> We currently handle the case where an MFMA is copied to an AGPR, but not the case where the MFMA is copying from an AGPR. --- Patch is 34.43 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/153020.diff 2 Files Affected: - (added) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir (+261) - (modified) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll (+387) ``````````diff diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir new file mode 100644 index 0000000000000..7fdc8c0d8019b --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir @@ -0,0 +1,261 @@ +# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5 +# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -run-pass=greedy,amdgpu-rewrite-agpr-copy-mfma -verify-machineinstrs -o - %s | FileCheck %s + +--- +name: test_rewrite_mfma_copy_from_agpr_physreg +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $agpr0_agpr1 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_physreg + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $agpr0_agpr1 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_64_align2 = COPY $agpr0_agpr1 + ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:vreg_64_align2 = COPY $agpr0_agpr1 + %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0, implicit $mode, implicit $exec + GLOBAL_STORE_DWORDX2 %0, %4, 0, 0, implicit $exec :: (store (s64), addrspace 1) + SI_RETURN +... + +--- +name: test_rewrite_mfma_copy_from_agpr_unrewritable_use +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_unrewritable_use + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1) + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]] + ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3866633 /* reguse:VReg_64_Align2 */, [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]] + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1) + %4:vreg_128_align2 = COPY %3 + %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec + INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3866633 /* reguse:VReg_64_Align2 */, %5 + SI_RETURN +... + +--- +name: test_rewrite_mfma_copy_from_agpr_src2_subreg_use +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_src2_subreg_use + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1) + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]] + ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1) + %4:vreg_128_align2 = COPY %3 + %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec + GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64), addrspace 1) + SI_RETURN +... + +--- +name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1) + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]] + ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit $exec :: (store (s128), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1) + %4:vreg_128_align2 = COPY %3 + %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec + GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1) + SI_RETURN +... + +# A-to-V copy is performed subregister at a time instead. +--- +name: test_rewrite_mfma_copy_from_agpr_split_copy +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_split_copy + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1) + ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]].sub0 + ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub1:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]].sub1 + ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1) + undef %4.sub0:vreg_64_align2 = COPY %3.sub0 + %4.sub1:vreg_64_align2 = COPY %3.sub1 + %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4, 0, 0, 0, implicit $mode, implicit $exec + GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64), addrspace 1) + SI_RETURN +... + +--- +name: test_rewrite_mfma_copy_from_agpr_copyback +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_copyback + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1) + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]] + ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_e64_:%[0-9]+]]:areg_64_align2 = V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: [[COPY4:%[0-9]+]]:areg_64_align2 = COPY [[V_MFMA_F64_4X4X4F64_e64_]] + ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY4]], 0, 0, implicit $exec :: (store (s64), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1) + %4:vreg_64_align2 = COPY %3 + %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4, 0, 0, 0, implicit $mode, implicit $exec + %6:areg_64_align2 = COPY %5 + GLOBAL_STORE_DWORDX2 %0, %6, 0, 0, implicit $exec :: (store (s64), addrspace 1) + SI_RETURN +... + +# There is a read of the coy-from-agpr in the dst operand of the MFMA. +--- +name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use_imm_src2 +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use_imm_src2 + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1) + ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]] + ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit $exec :: (store (s128), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1) + %4:vreg_128_align2 = COPY %3 + %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec + GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1) + SI_RETURN +... + +# Degenerate case. Copy from AGPR to VGPR is dead undef subreg def +--- +name: test_rewrite_mfma_copy_from_agpr_undef_vdst_subreg_use_imm_src2 +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_undef_vdst_subreg_use_imm_src2 + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1) + ; CHECK-NEXT: dead [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]] + ; CHECK-NEXT: undef [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s128), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1) + %4:vreg_128_align2 = COPY %3 + undef %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec + GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1) + SI_RETURN +... + +# Degenerate case. Copy from AGPR to VGPR is dead, but same register +# is redefined as whole register. +--- +name: test_rewrite_mfma_copy_from_agpr_to_vdst_def_imm_src2 +tracksRegLiveness: true +body: | + bb.0: + liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + + ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_to_vdst_def_imm_src2 + ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5 + ; CHECK-NEXT: {{ $}} + ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5 + ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1 + ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3 + ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1) + ; CHECK-NEXT: dead [[COPY3:%[0-9]+]]:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]] + ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec + ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1) + ; CHECK-NEXT: SI_RETURN + %0:vreg_64_align2 = COPY $vgpr4_vgpr5 + %1:av_64_align2 = COPY $vgpr0_vgpr1 + %2:av_64_align2 = COPY $vgpr2_vgpr3 + %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1) + %4:vreg_64_align2 = COPY %3 + %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec + GLOBAL_STORE_DWORDX2 %0, %4, 0, 0, implicit $exec :: (store (s64), addrspace 1) + SI_RETURN +... diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll index 58ec41e3928bd..81613f69c982b 100644 --- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll +++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll @@ -522,7 +522,394 @@ define void @test_rewrite_mfma_subreg_insert2(double %arg0, double %arg1, ptr ad ret void } +define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class(ptr addrspace(1) %arg0, ptr addrspace(1) %arg1) #0 { +; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class: +; CHECK: ; %bb.0: +; CHECK-NEXT: ;;#ASMSTART +; CHECK-NEXT: ; def a[0:31] +; CHECK-NEXT: ;;#ASMEND +; CHECK-NEXT: v_accvgpr_write_b32 a32, v0 +; CHECK-NEXT: v_accvgpr_read_b32 v63, a31 +; CHECK-NEXT: v_accvgpr_read_b32 v62, a30 +; CHECK-NEXT: v_accvgpr_read_b32 v61, a29 +; CHECK-NEXT: v_accvgpr_read_b32 v60, a28 +; CHECK-NEXT: v_accvgpr_read_b32 v59, a27 +; CHECK-NEXT: v_accvgpr_read_b32 v58, a26 +; CHECK-NEXT: v_accvgpr_read_b32 v57, a25 +; CHECK-NEXT: v_accvgpr_read_b32 v56, a24 +; CHECK-NEXT: v_accvgpr_read_b32 v55, a23 +; CHECK-NEXT: v_accvgpr_read_b32 v54, a22 +; CHECK-NEXT: v_accvgpr_read_b32 v53, a21 +; CHECK-NEXT: v_accvgpr_read_b32 v52, a20 +; CHECK-NEXT: v_accvgpr_read_b32 v51, a19 +; CHECK-NEXT: v_accvgpr_read_b32 v50, a18 +; CHECK-NEXT: v_accvgpr_read_b32 v49, a17 +; CHECK-NEXT: v_accvgpr_read_b32 v48, a16 +; CHECK-NEXT: v_accvgpr_read_b32 v47, a15 +; CHECK-NEXT: v_accvgpr_read_b32 v46, a14 +; CHECK-NEXT: v_accvgpr_read_b32 v45, a13 +; CHECK-NEXT: v_accvgpr_read_b32 v44, a12 +; CHECK-NEXT: v_accvgpr_read_b32 v43, a11 +; CHECK-NEXT: v_accvgpr_read_b32 v42, a10 +; CHECK-NEXT: v_accvgpr_read_b32 v41, a9 +; CHECK-NEXT: v_accvgpr_read_b32 v40, a8 +; CHECK-NEXT: v_accvgpr_read_b32 v39, a7 +; CHECK-NEXT: v_accvgpr_read_b32 v38, a6 +; CHECK-NEXT: v_accvgpr_read_b32 v37, a5 +; CHECK-NEXT: v_accvgpr_read_b32 v36, a4 +; CHECK-NEXT: v_accvgpr_read_b32 v35, a3 +; CHECK-NEXT: v_accvgpr_read_b32 v34, a2 +; CHECK-NEXT: v_accvgpr_read_b32 v33, a1 +; CHECK-NEXT: v_accvgpr_read_b32 v32, a0 +; CHECK-NEXT: v_accvgpr_write_b32 a0, 2.0 +; CHECK-NEXT: v_accvgpr_write_b32 a1, 4.0 +; CHECK-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0 +; CHECK-NEXT: s_nop 0 +; CHECK-NEXT: v_mfma_f32_32x32x1_2b_f32 v[0:31], a0, a1, v[32:63] +; CHECK-NEXT: v_accvgpr_write_b32 a0, v32 +; CHECK-NEXT: v_accvgpr_write_b32 a1, v33 +; CHECK-NEXT: v_accvgpr_write_b32 a2, v34 +; CHECK-NEXT: v_accvgpr_write_b32 a3, v35 +; CHECK-NEXT: v_accvgpr_write_b32 a4, v36 +; CHECK-NEXT: v_accvgpr_write_b32 a5, v37 +; CHECK-NEXT: v_accvgpr_write_b32 a6, v38 +; CHECK-NEXT: v_accvgpr_write_b32 a7, v39 +; CHECK-NEXT: v_accvgpr_write_b32 a8, v40 +; CHECK-NEXT: v_accvgpr_write_b32 a9, v41 +; CHECK-NEXT: v_accvgpr_write_b32 a10, v42 +; CHECK-NEXT: v_accvgpr_write_b32 a11, v43 +; CHECK-NEXT: v_accvgpr_write_b32 a12, v44 +; CHECK-NEXT: v_accvgpr_write_b32 a13, v45 +; CHECK-NEXT: v_accvgpr_write_b32 a14, v46 +; CHECK-NEXT: v_accvgpr_write_b32 a15, v47 +; CHECK-NEXT: v_accvgpr_write_b32 a16, v48 +; CHECK-NEXT: v_accvgpr_write_b32 a17, v49 +; CHECK-NEXT: v_accvgpr_write_b32 a18, v50 +; CHECK-NEXT: v_accvgpr_write_b32 a19, v51 +; CHECK-NEXT: v_accvgpr_write_b32 a20, v52 +; CHECK-NEXT: v_accvgpr_write_b32 a21, v53 +; CHECK-NEXT: v_accvgpr_write_b32 a22, v54 +; CHECK-NEXT: v_accvgpr_write_b32 a23, v55 +; CHECK-NEXT: v_accvgpr_write_b32 a24, v56 +; CHECK-NEXT: v_accvgpr_write_b32 a25, v57 +; CHECK-NEXT: v_accvgpr_write_b32 a26, v58 +; CHECK-NEXT: v_accvgpr_write_b32 a27, v59 +; CHECK-NEXT: v_accvgpr_write_b32 a28, v60 +; CHECK-NEXT: v_accvgpr_write_b32 a29, v61 +; CHECK-NEXT: v_accvgpr_write_b32 a30, v62 +; CHECK-NEXT: v_accvgpr_write_b32 a31, v63 +; CHECK-NEXT: v_accvgpr_read_b32 v32, a32 +; CHECK-NEXT: v_mov_b32_e32 v33, 0x41000000 +; CHECK-NEXT: v_and_b32_e32 v32, 0x3ff, v32 +; CHECK-NEXT: v_lshlrev_b32_e32 v32, 7, v32 +; CHECK-NEXT: s_waitcnt lgkmcnt(0) +; CHECK-NEXT: global_store_dwordx4 v32, v[28:31], s[0:1] offset:112 +; CHECK-NEXT: global_store_dwordx4 v32, v[24:27], s[0:1] offset:96 +; CHECK-NEXT: global_store_dwordx4 v32, v[20:23], s[0:1] offset:80 +; CHECK-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:64 +; CHECK-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 +; CHECK-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 +; CHECK-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 +; CHECK-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] +; CHECK-NEXT: v_mov_b32_e32 v34, 0x41800000 +; CHECK-NEXT: s_nop 0 +; CHECK-NEXT: v_accvgpr_read_b32 v0, a0 +; CHECK-NEXT: v_accvgpr_read_b32 v1, a1 +; CHECK-NEXT: v_accvgpr_read_b3... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/153020 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
