Skip to content

Commit 39ea0ec

Browse files
jrbyrnesjrbyrnes
committed
Update lit tests
Change-Id: If4c7179cd642c38a7c78d545347b8c0e011a0cbf
1 parent 7c8829e commit 39ea0ec

File tree

6 files changed

+1229
-1039
lines changed

6 files changed

+1229
-1039
lines changed

‎llvm/test/CodeGen/AMDGPU/bf16.ll‎

Lines changed: 83 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -468,15 +468,28 @@ define <16 x bfloat> @v_load_global_v16bf16(ptr addrspace(1) %ptr) {
468468
; GFX8-NEXT: s_waitcnt vmcnt(0)
469469
; GFX8-NEXT: s_setpc_b64 s[30:31]
470470
;
471-
; GFX9-LABEL: v_load_global_v16bf16:
472-
; GFX9: ; %bb.0:
473-
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
474-
; GFX9-NEXT: v_mov_b32_e32 v9, v1
475-
; GFX9-NEXT: v_mov_b32_e32 v8, v0
476-
; GFX9-NEXT: global_load_dwordx4 v[0:3], v[8:9], off
477-
; GFX9-NEXT: global_load_dwordx4 v[4:7], v[8:9], off offset:16
478-
; GFX9-NEXT: s_waitcnt vmcnt(0)
479-
; GFX9-NEXT: s_setpc_b64 s[30:31]
471+
; GFX900-LABEL: v_load_global_v16bf16:
472+
; GFX900: ; %bb.0:
473+
; GFX900-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
474+
; GFX900-NEXT: v_mov_b32_e32 v9, v1
475+
; GFX900-NEXT: v_mov_b32_e32 v8, v0
476+
; GFX900-NEXT: global_load_dwordx4 v[0:3], v[8:9], off
477+
; GFX900-NEXT: global_load_dwordx4 v[4:7], v[8:9], off offset:16
478+
; GFX900-NEXT: s_waitcnt vmcnt(0)
479+
; GFX900-NEXT: s_setpc_b64 s[30:31]
480+
;
481+
; GFX950-LABEL: v_load_global_v16bf16:
482+
; GFX950: ; %bb.0:
483+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
484+
; GFX950-NEXT: global_load_dwordx4 v[8:11], v[0:1], off
485+
; GFX950-NEXT: global_load_dwordx4 v[4:7], v[0:1], off offset:16
486+
; GFX950-NEXT: s_waitcnt vmcnt(1)
487+
; GFX950-NEXT: v_mov_b32_e32 v0, v8
488+
; GFX950-NEXT: v_mov_b32_e32 v1, v9
489+
; GFX950-NEXT: v_mov_b32_e32 v2, v10
490+
; GFX950-NEXT: v_mov_b32_e32 v3, v11
491+
; GFX950-NEXT: s_waitcnt vmcnt(0)
492+
; GFX950-NEXT: s_setpc_b64 s[30:31]
480493
;
481494
; GFX10-LABEL: v_load_global_v16bf16:
482495
; GFX10: ; %bb.0:
@@ -619,17 +632,32 @@ define <32 x bfloat> @v_load_global_v32bf16(ptr addrspace(1) %ptr) {
619632
; GFX8-NEXT: s_waitcnt vmcnt(0)
620633
; GFX8-NEXT: s_setpc_b64 s[30:31]
621634
;
622-
; GFX9-LABEL: v_load_global_v32bf16:
623-
; GFX9: ; %bb.0:
624-
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
625-
; GFX9-NEXT: v_mov_b32_e32 v17, v1
626-
; GFX9-NEXT: v_mov_b32_e32 v16, v0
627-
; GFX9-NEXT: global_load_dwordx4 v[0:3], v[16:17], off
628-
; GFX9-NEXT: global_load_dwordx4 v[4:7], v[16:17], off offset:16
629-
; GFX9-NEXT: global_load_dwordx4 v[8:11], v[16:17], off offset:32
630-
; GFX9-NEXT: global_load_dwordx4 v[12:15], v[16:17], off offset:48
631-
; GFX9-NEXT: s_waitcnt vmcnt(0)
632-
; GFX9-NEXT: s_setpc_b64 s[30:31]
635+
; GFX900-LABEL: v_load_global_v32bf16:
636+
; GFX900: ; %bb.0:
637+
; GFX900-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
638+
; GFX900-NEXT: v_mov_b32_e32 v17, v1
639+
; GFX900-NEXT: v_mov_b32_e32 v16, v0
640+
; GFX900-NEXT: global_load_dwordx4 v[0:3], v[16:17], off
641+
; GFX900-NEXT: global_load_dwordx4 v[4:7], v[16:17], off offset:16
642+
; GFX900-NEXT: global_load_dwordx4 v[8:11], v[16:17], off offset:32
643+
; GFX900-NEXT: global_load_dwordx4 v[12:15], v[16:17], off offset:48
644+
; GFX900-NEXT: s_waitcnt vmcnt(0)
645+
; GFX900-NEXT: s_setpc_b64 s[30:31]
646+
;
647+
; GFX950-LABEL: v_load_global_v32bf16:
648+
; GFX950: ; %bb.0:
649+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
650+
; GFX950-NEXT: global_load_dwordx4 v[16:19], v[0:1], off
651+
; GFX950-NEXT: global_load_dwordx4 v[4:7], v[0:1], off offset:16
652+
; GFX950-NEXT: global_load_dwordx4 v[8:11], v[0:1], off offset:32
653+
; GFX950-NEXT: global_load_dwordx4 v[12:15], v[0:1], off offset:48
654+
; GFX950-NEXT: s_waitcnt vmcnt(3)
655+
; GFX950-NEXT: v_mov_b32_e32 v0, v16
656+
; GFX950-NEXT: v_mov_b32_e32 v1, v17
657+
; GFX950-NEXT: v_mov_b32_e32 v2, v18
658+
; GFX950-NEXT: v_mov_b32_e32 v3, v19
659+
; GFX950-NEXT: s_waitcnt vmcnt(0)
660+
; GFX950-NEXT: s_setpc_b64 s[30:31]
633661
;
634662
; GFX10-LABEL: v_load_global_v32bf16:
635663
; GFX10: ; %bb.0:
@@ -877,22 +905,41 @@ define <64 x bfloat> @v_load_global_v64bf16(ptr addrspace(1) %ptr) {
877905
; GFX8-NEXT: s_waitcnt vmcnt(0)
878906
; GFX8-NEXT: s_setpc_b64 s[30:31]
879907
;
880-
; GFX9-LABEL: v_load_global_v64bf16:
881-
; GFX9: ; %bb.0:
882-
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
883-
; GFX9-NEXT: v_mov_b32_e32 v29, v1
884-
; GFX9-NEXT: v_mov_b32_e32 v28, v0
885-
; GFX9-NEXT: global_load_dwordx4 v[0:3], v[28:29], off
886-
; GFX9-NEXT: global_load_dwordx4 v[4:7], v[28:29], off offset:16
887-
; GFX9-NEXT: global_load_dwordx4 v[8:11], v[28:29], off offset:32
888-
; GFX9-NEXT: global_load_dwordx4 v[12:15], v[28:29], off offset:48
889-
; GFX9-NEXT: global_load_dwordx4 v[16:19], v[28:29], off offset:64
890-
; GFX9-NEXT: global_load_dwordx4 v[20:23], v[28:29], off offset:80
891-
; GFX9-NEXT: global_load_dwordx4 v[24:27], v[28:29], off offset:96
892-
; GFX9-NEXT: s_nop 0
893-
; GFX9-NEXT: global_load_dwordx4 v[28:31], v[28:29], off offset:112
894-
; GFX9-NEXT: s_waitcnt vmcnt(0)
895-
; GFX9-NEXT: s_setpc_b64 s[30:31]
908+
; GFX900-LABEL: v_load_global_v64bf16:
909+
; GFX900: ; %bb.0:
910+
; GFX900-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
911+
; GFX900-NEXT: v_mov_b32_e32 v29, v1
912+
; GFX900-NEXT: v_mov_b32_e32 v28, v0
913+
; GFX900-NEXT: global_load_dwordx4 v[0:3], v[28:29], off
914+
; GFX900-NEXT: global_load_dwordx4 v[4:7], v[28:29], off offset:16
915+
; GFX900-NEXT: global_load_dwordx4 v[8:11], v[28:29], off offset:32
916+
; GFX900-NEXT: global_load_dwordx4 v[12:15], v[28:29], off offset:48
917+
; GFX900-NEXT: global_load_dwordx4 v[16:19], v[28:29], off offset:64
918+
; GFX900-NEXT: global_load_dwordx4 v[20:23], v[28:29], off offset:80
919+
; GFX900-NEXT: global_load_dwordx4 v[24:27], v[28:29], off offset:96
920+
; GFX900-NEXT: s_nop 0
921+
; GFX900-NEXT: global_load_dwordx4 v[28:31], v[28:29], off offset:112
922+
; GFX900-NEXT: s_waitcnt vmcnt(0)
923+
; GFX900-NEXT: s_setpc_b64 s[30:31]
924+
;
925+
; GFX950-LABEL: v_load_global_v64bf16:
926+
; GFX950: ; %bb.0:
927+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
928+
; GFX950-NEXT: global_load_dwordx4 v[32:35], v[0:1], off
929+
; GFX950-NEXT: global_load_dwordx4 v[4:7], v[0:1], off offset:16
930+
; GFX950-NEXT: global_load_dwordx4 v[8:11], v[0:1], off offset:32
931+
; GFX950-NEXT: global_load_dwordx4 v[12:15], v[0:1], off offset:48
932+
; GFX950-NEXT: global_load_dwordx4 v[16:19], v[0:1], off offset:64
933+
; GFX950-NEXT: global_load_dwordx4 v[20:23], v[0:1], off offset:80
934+
; GFX950-NEXT: global_load_dwordx4 v[24:27], v[0:1], off offset:96
935+
; GFX950-NEXT: global_load_dwordx4 v[28:31], v[0:1], off offset:112
936+
; GFX950-NEXT: s_waitcnt vmcnt(7)
937+
; GFX950-NEXT: v_mov_b32_e32 v0, v32
938+
; GFX950-NEXT: v_mov_b32_e32 v1, v33
939+
; GFX950-NEXT: v_mov_b32_e32 v2, v34
940+
; GFX950-NEXT: v_mov_b32_e32 v3, v35
941+
; GFX950-NEXT: s_waitcnt vmcnt(0)
942+
; GFX950-NEXT: s_setpc_b64 s[30:31]
896943
;
897944
; GFX10-LABEL: v_load_global_v64bf16:
898945
; GFX10: ; %bb.0:

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx90a.ll‎

Lines changed: 32 additions & 32 deletions
Original file line numberDiff line numberDiff line change
@@ -443,10 +443,10 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
443443
; GFX90A-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24
444444
; GFX90A-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34
445445
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
446-
; GFX90A-NEXT: v_mov_b32_e32 v0, s10
446+
; GFX90A-NEXT: v_mov_b32_e32 v2, s10
447447
; GFX90A-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0
448-
; GFX90A-NEXT: v_mov_b32_e32 v1, s11
449-
; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[12:13], s[12:13] op_sel:[0,1]
448+
; GFX90A-NEXT: v_mov_b32_e32 v3, s11
449+
; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[12:13], s[12:13] op_sel:[0,1]
450450
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
451451
; GFX90A-NEXT: v_accvgpr_write_b32 a0, s0
452452
; GFX90A-NEXT: v_accvgpr_write_b32 a1, s1
@@ -457,7 +457,7 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
457457
; GFX90A-NEXT: v_accvgpr_write_b32 a6, s6
458458
; GFX90A-NEXT: v_accvgpr_write_b32 a7, s7
459459
; GFX90A-NEXT: s_nop 1
460-
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 blgp:3
460+
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[2:3], v[0:1], a[0:7] cbsz:1 abid:2 blgp:3
461461
; GFX90A-NEXT: v_mov_b32_e32 v0, 0
462462
; GFX90A-NEXT: s_nop 7
463463
; GFX90A-NEXT: s_nop 7
@@ -471,10 +471,10 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
471471
; GFX942-NEXT: s_load_dwordx4 s[8:11], s[4:5], 0x24
472472
; GFX942-NEXT: s_load_dwordx2 s[12:13], s[4:5], 0x34
473473
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
474-
; GFX942-NEXT: v_mov_b32_e32 v0, s10
474+
; GFX942-NEXT: v_mov_b32_e32 v2, s10
475475
; GFX942-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x0
476-
; GFX942-NEXT: v_mov_b32_e32 v1, s11
477-
; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[12:13]
476+
; GFX942-NEXT: v_mov_b32_e32 v3, s11
477+
; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
478478
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
479479
; GFX942-NEXT: v_accvgpr_write_b32 a0, s0
480480
; GFX942-NEXT: v_accvgpr_write_b32 a1, s1
@@ -485,7 +485,7 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
485485
; GFX942-NEXT: v_accvgpr_write_b32 a6, s6
486486
; GFX942-NEXT: v_accvgpr_write_b32 a7, s7
487487
; GFX942-NEXT: s_nop 1
488-
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] cbsz:1 abid:2 neg:[1,1,0]
488+
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[2:3], v[0:1], a[0:7] cbsz:1 abid:2 neg:[1,1,0]
489489
; GFX942-NEXT: v_mov_b32_e32 v0, 0
490490
; GFX942-NEXT: s_nop 7
491491
; GFX942-NEXT: s_nop 7
@@ -898,20 +898,20 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_imm(ptr addrspace(1) %arg, d
898898
; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
899899
; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
900900
; GFX90A-NEXT: v_accvgpr_write_b32 a0, 0
901-
; GFX90A-NEXT: v_mov_b32_e32 v2, 0x3ff00000
902-
; GFX90A-NEXT: v_accvgpr_write_b32 a7, v2
901+
; GFX90A-NEXT: v_mov_b32_e32 v0, 0x3ff00000
902+
; GFX90A-NEXT: v_accvgpr_write_b32 a7, v0
903903
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
904-
; GFX90A-NEXT: v_mov_b32_e32 v0, s2
905-
; GFX90A-NEXT: v_mov_b32_e32 v1, s3
904+
; GFX90A-NEXT: v_mov_b32_e32 v2, s2
905+
; GFX90A-NEXT: v_mov_b32_e32 v3, s3
906906
; GFX90A-NEXT: v_accvgpr_mov_b32 a1, a0
907907
; GFX90A-NEXT: v_accvgpr_mov_b32 a2, a0
908908
; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a0
909909
; GFX90A-NEXT: v_accvgpr_mov_b32 a4, a0
910910
; GFX90A-NEXT: v_accvgpr_mov_b32 a5, a0
911911
; GFX90A-NEXT: v_accvgpr_mov_b32 a6, a0
912-
; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]
912+
; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[6:7], s[6:7] op_sel:[0,1]
913913
; GFX90A-NEXT: s_nop 1
914-
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]
914+
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[2:3], v[0:1], a[0:7]
915915
; GFX90A-NEXT: v_mov_b32_e32 v0, 0
916916
; GFX90A-NEXT: s_nop 7
917917
; GFX90A-NEXT: s_nop 7
@@ -925,20 +925,20 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_imm(ptr addrspace(1) %arg, d
925925
; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
926926
; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
927927
; GFX942-NEXT: v_accvgpr_write_b32 a0, 0
928-
; GFX942-NEXT: v_mov_b32_e32 v2, 0x3ff00000
929-
; GFX942-NEXT: v_accvgpr_write_b32 a7, v2
928+
; GFX942-NEXT: v_mov_b32_e32 v0, 0x3ff00000
929+
; GFX942-NEXT: v_accvgpr_write_b32 a7, v0
930930
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
931-
; GFX942-NEXT: v_mov_b32_e32 v0, s2
932-
; GFX942-NEXT: v_mov_b32_e32 v1, s3
931+
; GFX942-NEXT: v_mov_b32_e32 v2, s2
932+
; GFX942-NEXT: v_mov_b32_e32 v3, s3
933933
; GFX942-NEXT: v_accvgpr_mov_b32 a1, a0
934934
; GFX942-NEXT: v_accvgpr_mov_b32 a2, a0
935935
; GFX942-NEXT: v_accvgpr_mov_b32 a3, a0
936936
; GFX942-NEXT: v_accvgpr_mov_b32 a4, a0
937937
; GFX942-NEXT: v_accvgpr_mov_b32 a5, a0
938938
; GFX942-NEXT: v_accvgpr_mov_b32 a6, a0
939-
; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]
939+
; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[6:7]
940940
; GFX942-NEXT: s_nop 1
941-
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]
941+
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[2:3], v[0:1], a[0:7]
942942
; GFX942-NEXT: v_mov_b32_e32 v0, 0
943943
; GFX942-NEXT: s_nop 7
944944
; GFX942-NEXT: s_nop 7
@@ -957,21 +957,21 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_lit(ptr addrspace(1) %
957957
; GFX90A: ; %bb.0: ; %bb
958958
; GFX90A-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
959959
; GFX90A-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
960-
; GFX90A-NEXT: v_mov_b32_e32 v2, 0x405ec000
960+
; GFX90A-NEXT: v_mov_b32_e32 v0, 0x405ec000
961961
; GFX90A-NEXT: v_accvgpr_write_b32 a0, 0
962-
; GFX90A-NEXT: v_accvgpr_write_b32 a1, v2
962+
; GFX90A-NEXT: v_accvgpr_write_b32 a1, v0
963963
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
964-
; GFX90A-NEXT: v_mov_b32_e32 v0, s2
965-
; GFX90A-NEXT: v_mov_b32_e32 v1, s3
964+
; GFX90A-NEXT: v_mov_b32_e32 v2, s2
965+
; GFX90A-NEXT: v_mov_b32_e32 v3, s3
966966
; GFX90A-NEXT: v_accvgpr_mov_b32 a2, a0
967967
; GFX90A-NEXT: v_accvgpr_mov_b32 a3, a1
968968
; GFX90A-NEXT: v_accvgpr_mov_b32 a4, a0
969969
; GFX90A-NEXT: v_accvgpr_mov_b32 a5, a1
970970
; GFX90A-NEXT: v_accvgpr_mov_b32 a6, a0
971971
; GFX90A-NEXT: v_accvgpr_mov_b32 a7, a1
972-
; GFX90A-NEXT: v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]
972+
; GFX90A-NEXT: v_pk_mov_b32 v[0:1], s[6:7], s[6:7] op_sel:[0,1]
973973
; GFX90A-NEXT: s_nop 1
974-
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[0:1], v[2:3], a[0:7]
974+
; GFX90A-NEXT: v_mfma_f64_16x16x4f64 a[0:7], v[2:3], v[0:1], a[0:7]
975975
; GFX90A-NEXT: v_mov_b32_e32 v0, 0
976976
; GFX90A-NEXT: s_nop 7
977977
; GFX90A-NEXT: s_nop 7
@@ -984,21 +984,21 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64_splat_lit(ptr addrspace(1) %
984984
; GFX942: ; %bb.0: ; %bb
985985
; GFX942-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
986986
; GFX942-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
987-
; GFX942-NEXT: v_mov_b32_e32 v2, 0x405ec000
987+
; GFX942-NEXT: v_mov_b32_e32 v0, 0x405ec000
988988
; GFX942-NEXT: v_accvgpr_write_b32 a0, 0
989-
; GFX942-NEXT: v_accvgpr_write_b32 a1, v2
989+
; GFX942-NEXT: v_accvgpr_write_b32 a1, v0
990990
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
991-
; GFX942-NEXT: v_mov_b32_e32 v0, s2
992-
; GFX942-NEXT: v_mov_b32_e32 v1, s3
991+
; GFX942-NEXT: v_mov_b32_e32 v2, s2
992+
; GFX942-NEXT: v_mov_b32_e32 v3, s3
993993
; GFX942-NEXT: v_accvgpr_mov_b32 a2, a0
994994
; GFX942-NEXT: v_accvgpr_mov_b32 a3, a1
995995
; GFX942-NEXT: v_accvgpr_mov_b32 a4, a0
996996
; GFX942-NEXT: v_accvgpr_mov_b32 a5, a1
997997
; GFX942-NEXT: v_accvgpr_mov_b32 a6, a0
998998
; GFX942-NEXT: v_accvgpr_mov_b32 a7, a1
999-
; GFX942-NEXT: v_mov_b64_e32 v[2:3], s[6:7]
999+
; GFX942-NEXT: v_mov_b64_e32 v[0:1], s[6:7]
10001000
; GFX942-NEXT: s_nop 1
1001-
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]
1001+
; GFX942-NEXT: v_mfma_f64_16x16x4_f64 a[0:7], v[2:3], v[0:1], a[0:7]
10021002
; GFX942-NEXT: v_mov_b32_e32 v0, 0
10031003
; GFX942-NEXT: s_nop 7
10041004
; GFX942-NEXT: s_nop 7

0 commit comments

Comments
 (0)