Skip to content

Commit e11d28f

Browse files
shiltianrampitec
shiltian
and
rampitec
authored
[AMDGPU] Add support for v_permlane16_swap_b32 on gfx1250 (llvm#149518)
Co-authored-by: Mekhanoshin, Stanislav <[email protected]>
1 parent 7e0ae01 commit e11d28f

File tree

8 files changed

+270
-3
lines changed

8 files changed

+270
-3
lines changed

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl

Lines changed: 47 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
66

77
typedef unsigned int uint;
8+
typedef unsigned int __attribute__((ext_vector_type(2))) uint2;
89
typedef half __attribute__((ext_vector_type(2))) half2;
910

1011
// CHECK-LABEL: @test_setprio_inc_wg(
@@ -368,6 +369,52 @@ void test_cvt_pk_f16_bf8(global half2* out, short a)
368369
out[0] = __builtin_amdgcn_cvt_pk_f16_bf8(a);
369370
}
370371

372+
// CHECK-LABEL: @test_permlane16_swap(
373+
// CHECK-NEXT: entry:
374+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
375+
// CHECK-NEXT: [[OLD_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
376+
// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
377+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
378+
// CHECK-NEXT: [[OLD_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OLD_ADDR]] to ptr
379+
// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
380+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
381+
// CHECK-NEXT: store i32 [[OLD:%.*]], ptr [[OLD_ADDR_ASCAST]], align 4
382+
// CHECK-NEXT: store i32 [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
383+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[OLD_ADDR_ASCAST]], align 4
384+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
385+
// CHECK-NEXT: [[TMP2:%.*]] = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 [[TMP0]], i32 [[TMP1]], i1 false, i1 false)
386+
// CHECK-NEXT: [[TMP3:%.*]] = extractvalue { i32, i32 } [[TMP2]], 0
387+
// CHECK-NEXT: [[TMP4:%.*]] = extractvalue { i32, i32 } [[TMP2]], 1
388+
// CHECK-NEXT: [[TMP5:%.*]] = insertelement <2 x i32> poison, i32 [[TMP3]], i64 0
389+
// CHECK-NEXT: [[TMP6:%.*]] = insertelement <2 x i32> [[TMP5]], i32 [[TMP4]], i64 1
390+
// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
391+
// CHECK-NEXT: store <2 x i32> [[TMP6]], ptr addrspace(1) [[TMP7]], align 8
392+
// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr [[OLD_ADDR_ASCAST]], align 4
393+
// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
394+
// CHECK-NEXT: [[TMP10:%.*]] = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 [[TMP8]], i32 [[TMP9]], i1 true, i1 false)
395+
// CHECK-NEXT: [[TMP11:%.*]] = extractvalue { i32, i32 } [[TMP10]], 0
396+
// CHECK-NEXT: [[TMP12:%.*]] = extractvalue { i32, i32 } [[TMP10]], 1
397+
// CHECK-NEXT: [[TMP13:%.*]] = insertelement <2 x i32> poison, i32 [[TMP11]], i64 0
398+
// CHECK-NEXT: [[TMP14:%.*]] = insertelement <2 x i32> [[TMP13]], i32 [[TMP12]], i64 1
399+
// CHECK-NEXT: [[TMP15:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
400+
// CHECK-NEXT: store <2 x i32> [[TMP14]], ptr addrspace(1) [[TMP15]], align 8
401+
// CHECK-NEXT: [[TMP16:%.*]] = load i32, ptr [[OLD_ADDR_ASCAST]], align 4
402+
// CHECK-NEXT: [[TMP17:%.*]] = load i32, ptr [[SRC_ADDR_ASCAST]], align 4
403+
// CHECK-NEXT: [[TMP18:%.*]] = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 [[TMP16]], i32 [[TMP17]], i1 false, i1 true)
404+
// CHECK-NEXT: [[TMP19:%.*]] = extractvalue { i32, i32 } [[TMP18]], 0
405+
// CHECK-NEXT: [[TMP20:%.*]] = extractvalue { i32, i32 } [[TMP18]], 1
406+
// CHECK-NEXT: [[TMP21:%.*]] = insertelement <2 x i32> poison, i32 [[TMP19]], i64 0
407+
// CHECK-NEXT: [[TMP22:%.*]] = insertelement <2 x i32> [[TMP21]], i32 [[TMP20]], i64 1
408+
// CHECK-NEXT: [[TMP23:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
409+
// CHECK-NEXT: store <2 x i32> [[TMP22]], ptr addrspace(1) [[TMP23]], align 8
410+
// CHECK-NEXT: ret void
411+
//
412+
void test_permlane16_swap(global uint2* out, uint old, uint src) {
413+
*out = __builtin_amdgcn_permlane16_swap(old, src, false, false);
414+
*out = __builtin_amdgcn_permlane16_swap(old, src, true, false);
415+
*out = __builtin_amdgcn_permlane16_swap(old, src, false, true);
416+
}
417+
371418
// CHECK-LABEL: @test_cvt_f32_fp8_e5m3(
372419
// CHECK-NEXT: entry:
373420
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)

‎llvm/lib/Target/AMDGPU/VOP1Instructions.td

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1080,6 +1080,13 @@ multiclass VOP1_Real_FULL_t16_and_fake16_gfx1250<
10801080
VOP1_Real_FULL_with_name<GFX1250Gen, op, opName#"_fake16", asmName>;
10811081
}
10821082

1083+
multiclass VOP1_Real_OpSelIsDPP_gfx1250<bits<9> op> : VOP1_Real_e32<GFX1250Gen, op> {
1084+
defvar ps = !cast<VOP_Pseudo>(NAME#"_e64");
1085+
def _e64_gfx1250 :
1086+
VOP3_Real_Gen<ps, GFX1250Gen>,
1087+
VOP3OpSelIsDPP_gfx12<{0, 1, 1, op{6-0}}, ps.Pfl>;
1088+
}
1089+
10831090
defm V_CVT_F32_FP8 : VOP1_Real_FULL_with_name<GFX12Not12_50Gen, 0x06c, "V_CVT_F32_FP8_OP_SEL", "v_cvt_f32_fp8">;
10841091
defm V_CVT_F32_FP8 : VOP1_Real_FULL_with_name<GFX1250Gen, 0x06c, "V_CVT_F32_FP8_gfx1250", "v_cvt_f32_fp8">;
10851092

@@ -1147,6 +1154,7 @@ defm V_MOV_B64 : VOP1_Real_FULL <GFX1250Gen, 0x1d>;
11471154

11481155
defm V_TANH_F32 : VOP1_Real_FULL<GFX1250Gen, 0x01e>;
11491156
defm V_TANH_F16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x01f>;
1157+
defm V_PERMLANE16_SWAP_B32 : VOP1_Real_OpSelIsDPP_gfx1250<0x049>;
11501158
defm V_TANH_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
11511159
defm V_PRNG_B32 : VOP1_Real_FULL<GFX1250Gen, 0x04b>;
11521160
defm V_CVT_F32_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072, "v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;

‎llvm/lib/Target/AMDGPU/VOPInstructions.td

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -331,10 +331,19 @@ class VOP3OpSel_gfx9 <bits<10> op, VOPProfile P> : VOP3e_vi <op, P> {
331331

332332
// Special case for v_permlane16_swap_b32/v_permlane32_swap_b32
333333
// op_sel[0]/op_sel[1] are treated as bound_ctrl and fi dpp operands.
334-
class VOP3OpSelIsDPP_gfx9 <bits<10> op, VOPProfile P> : VOP3e_vi <op, P> {
334+
class VOP3OpSelIsDPP_base {
335335
bits<1> fi;
336336
bits<1> bound_ctrl;
337+
}
338+
339+
class VOP3OpSelIsDPP_gfx9 <bits<10> op, VOPProfile P> : VOP3OpSelIsDPP_base, VOP3e_vi <op, P> {
340+
// OPSEL[0] specifies FI
341+
let Inst{11} = fi;
342+
// OPSEL[1] specifies BOUND_CTRL
343+
let Inst{12} = bound_ctrl;
344+
}
337345

346+
class VOP3OpSelIsDPP_gfx12 <bits<10> op, VOPProfile P> : VOP3OpSelIsDPP_base, VOP3e_gfx11_gfx12 <op, P> {
338347
// OPSEL[0] specifies FI
339348
let Inst{11} = fi;
340349
// OPSEL[1] specifies BOUND_CTRL

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.permlane16.swap.ll

Lines changed: 151 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,8 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
2-
; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 < %s | FileCheck -check-prefix=GCN %s
3-
; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 < %s | FileCheck -check-prefix=GCN %s
2+
; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 < %s | FileCheck -check-prefix=GFX950 %s
3+
; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 < %s | FileCheck -check-prefix=GFX950 %s
4+
; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s | FileCheck -check-prefix=GFX1250 %s
5+
; RUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 < %s | FileCheck -check-prefix=GFX1250 %s
46

57
; RUN: not --crash llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR-SDAG %s
68
; RUN: not llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -filetype=null %s 2>&1 | FileCheck -check-prefix=ERR-GISEL %s
@@ -17,6 +19,18 @@ define { i32, i32 } @v_permlane16_swap_b32_vv(i32 %vdst_old, i32 %src0_old) {
1719
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
1820
; GCN-NEXT: v_permlane16_swap_b32_e32 v0, v1
1921
; GCN-NEXT: s_setpc_b64 s[30:31]
22+
; GFX950-LABEL: v_permlane16_swap_b32_vv:
23+
; GFX950: ; %bb.0:
24+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
25+
; GFX950-NEXT: v_permlane16_swap_b32_e32 v0, v1
26+
; GFX950-NEXT: s_setpc_b64 s[30:31]
27+
;
28+
; GFX1250-LABEL: v_permlane16_swap_b32_vv:
29+
; GFX1250: ; %bb.0:
30+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
31+
; GFX1250-NEXT: s_wait_kmcnt 0x0
32+
; GFX1250-NEXT: v_permlane16_swap_b32_e32 v0, v1
33+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
2034
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 %src0_old, i1 false, i1 false)
2135
ret { i32, i32 } %v
2236
}
@@ -29,6 +43,22 @@ define { i32, i32 } @v_permlane16_swap_b32_vi(i32 %vdst_old) {
2943
; GCN-NEXT: s_nop 1
3044
; GCN-NEXT: v_permlane16_swap_b32_e32 v0, v1
3145
; GCN-NEXT: s_setpc_b64 s[30:31]
46+
; GFX950-LABEL: v_permlane16_swap_b32_vi:
47+
; GFX950: ; %bb.0:
48+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
49+
; GFX950-NEXT: v_mov_b32_e32 v1, 1
50+
; GFX950-NEXT: s_nop 1
51+
; GFX950-NEXT: v_permlane16_swap_b32_e32 v0, v1
52+
; GFX950-NEXT: s_setpc_b64 s[30:31]
53+
;
54+
; GFX1250-LABEL: v_permlane16_swap_b32_vi:
55+
; GFX1250: ; %bb.0:
56+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
57+
; GFX1250-NEXT: s_wait_kmcnt 0x0
58+
; GFX1250-NEXT: v_mov_b32_e32 v1, 1
59+
; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
60+
; GFX1250-NEXT: v_permlane16_swap_b32_e32 v0, v1
61+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
3262
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 1, i1 false, i1 false)
3363
ret { i32, i32 } %v
3464
}
@@ -41,6 +71,22 @@ define { i32, i32 } @v_permlane16_swap_b32_vl(i32 %vdst_old) {
4171
; GCN-NEXT: s_nop 1
4272
; GCN-NEXT: v_permlane16_swap_b32_e32 v0, v1
4373
; GCN-NEXT: s_setpc_b64 s[30:31]
74+
; GFX950-LABEL: v_permlane16_swap_b32_vl:
75+
; GFX950: ; %bb.0:
76+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
77+
; GFX950-NEXT: v_mov_b32_e32 v1, 0xc1d1
78+
; GFX950-NEXT: s_nop 1
79+
; GFX950-NEXT: v_permlane16_swap_b32_e32 v0, v1
80+
; GFX950-NEXT: s_setpc_b64 s[30:31]
81+
;
82+
; GFX1250-LABEL: v_permlane16_swap_b32_vl:
83+
; GFX1250: ; %bb.0:
84+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
85+
; GFX1250-NEXT: s_wait_kmcnt 0x0
86+
; GFX1250-NEXT: v_mov_b32_e32 v1, 0xc1d1
87+
; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
88+
; GFX1250-NEXT: v_permlane16_swap_b32_e32 v0, v1
89+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
4490
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 49617, i1 false, i1 false)
4591
ret { i32, i32 } %v
4692
}
@@ -54,6 +100,23 @@ define { i32, i32 } @v_permlane16_swap_b32_iv(i32 %src0_old) {
54100
; GCN-NEXT: s_nop 1
55101
; GCN-NEXT: v_permlane16_swap_b32_e32 v0, v1
56102
; GCN-NEXT: s_setpc_b64 s[30:31]
103+
; GFX950-LABEL: v_permlane16_swap_b32_iv:
104+
; GFX950: ; %bb.0:
105+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
106+
; GFX950-NEXT: v_mov_b32_e32 v1, v0
107+
; GFX950-NEXT: v_mov_b32_e32 v0, 1
108+
; GFX950-NEXT: s_nop 1
109+
; GFX950-NEXT: v_permlane16_swap_b32_e32 v0, v1
110+
; GFX950-NEXT: s_setpc_b64 s[30:31]
111+
;
112+
; GFX1250-LABEL: v_permlane16_swap_b32_iv:
113+
; GFX1250: ; %bb.0:
114+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
115+
; GFX1250-NEXT: s_wait_kmcnt 0x0
116+
; GFX1250-NEXT: v_dual_mov_b32 v1, v0 :: v_dual_mov_b32 v0, 1
117+
; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
118+
; GFX1250-NEXT: v_permlane16_swap_b32_e32 v0, v1
119+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
57120
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 1, i32 %src0_old, i1 false, i1 false)
58121
ret { i32, i32 } %v
59122
}
@@ -67,6 +130,23 @@ define { i32, i32 } @v_permlane16_swap_b32_ss(i32 inreg %vdst_old, i32 inreg %sr
67130
; GCN-NEXT: s_nop 1
68131
; GCN-NEXT: v_permlane16_swap_b32_e32 v0, v1
69132
; GCN-NEXT: s_setpc_b64 s[30:31]
133+
; GFX950-LABEL: v_permlane16_swap_b32_ss:
134+
; GFX950: ; %bb.0:
135+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
136+
; GFX950-NEXT: v_mov_b32_e32 v0, s0
137+
; GFX950-NEXT: v_mov_b32_e32 v1, s1
138+
; GFX950-NEXT: s_nop 1
139+
; GFX950-NEXT: v_permlane16_swap_b32_e32 v0, v1
140+
; GFX950-NEXT: s_setpc_b64 s[30:31]
141+
;
142+
; GFX1250-LABEL: v_permlane16_swap_b32_ss:
143+
; GFX1250: ; %bb.0:
144+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
145+
; GFX1250-NEXT: s_wait_kmcnt 0x0
146+
; GFX1250-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
147+
; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
148+
; GFX1250-NEXT: v_permlane16_swap_b32_e32 v0, v1
149+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
70150
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 %src0_old, i1 false, i1 false)
71151
ret { i32, i32 } %v
72152
}
@@ -80,6 +160,23 @@ define { i32, i32 } @v_permlane16_swap_b32_sv(i32 inreg %vdst_old, i32 %src0_old
80160
; GCN-NEXT: s_nop 1
81161
; GCN-NEXT: v_permlane16_swap_b32_e32 v0, v1
82162
; GCN-NEXT: s_setpc_b64 s[30:31]
163+
; GFX950-LABEL: v_permlane16_swap_b32_sv:
164+
; GFX950: ; %bb.0:
165+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
166+
; GFX950-NEXT: v_mov_b32_e32 v1, v0
167+
; GFX950-NEXT: v_mov_b32_e32 v0, s0
168+
; GFX950-NEXT: s_nop 1
169+
; GFX950-NEXT: v_permlane16_swap_b32_e32 v0, v1
170+
; GFX950-NEXT: s_setpc_b64 s[30:31]
171+
;
172+
; GFX1250-LABEL: v_permlane16_swap_b32_sv:
173+
; GFX1250: ; %bb.0:
174+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
175+
; GFX1250-NEXT: s_wait_kmcnt 0x0
176+
; GFX1250-NEXT: v_dual_mov_b32 v1, v0 :: v_dual_mov_b32 v0, s0
177+
; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
178+
; GFX1250-NEXT: v_permlane16_swap_b32_e32 v0, v1
179+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
83180
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 %src0_old, i1 false, i1 false)
84181
ret { i32, i32 } %v
85182
}
@@ -92,6 +189,22 @@ define { i32, i32 } @v_permlane16_swap_b32_vs(i32 %vdst_old, i32 inreg %src0_old
92189
; GCN-NEXT: s_nop 1
93190
; GCN-NEXT: v_permlane16_swap_b32_e32 v0, v1
94191
; GCN-NEXT: s_setpc_b64 s[30:31]
192+
; GFX950-LABEL: v_permlane16_swap_b32_vs:
193+
; GFX950: ; %bb.0:
194+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
195+
; GFX950-NEXT: v_mov_b32_e32 v1, s0
196+
; GFX950-NEXT: s_nop 1
197+
; GFX950-NEXT: v_permlane16_swap_b32_e32 v0, v1
198+
; GFX950-NEXT: s_setpc_b64 s[30:31]
199+
;
200+
; GFX1250-LABEL: v_permlane16_swap_b32_vs:
201+
; GFX1250: ; %bb.0:
202+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
203+
; GFX1250-NEXT: s_wait_kmcnt 0x0
204+
; GFX1250-NEXT: v_mov_b32_e32 v1, s0
205+
; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
206+
; GFX1250-NEXT: v_permlane16_swap_b32_e32 v0, v1
207+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
95208
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 %src0_old, i1 false, i1 false)
96209
ret { i32, i32 } %v
97210
}
@@ -102,6 +215,18 @@ define { i32, i32 } @v_permlane16_swap_b32_vv_fi(i32 %vdst_old, i32 %src0_old) {
102215
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
103216
; GCN-NEXT: v_permlane16_swap_b32_e64 v0, v1 fi:1
104217
; GCN-NEXT: s_setpc_b64 s[30:31]
218+
; GFX950-LABEL: v_permlane16_swap_b32_vv_fi:
219+
; GFX950: ; %bb.0:
220+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
221+
; GFX950-NEXT: v_permlane16_swap_b32_e64 v0, v1 fi:1
222+
; GFX950-NEXT: s_setpc_b64 s[30:31]
223+
;
224+
; GFX1250-LABEL: v_permlane16_swap_b32_vv_fi:
225+
; GFX1250: ; %bb.0:
226+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
227+
; GFX1250-NEXT: s_wait_kmcnt 0x0
228+
; GFX1250-NEXT: v_permlane16_swap_b32_e64 v0, v1 fi:1
229+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
105230
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 %src0_old, i1 true, i1 false)
106231
ret { i32, i32 } %v
107232
}
@@ -112,6 +237,18 @@ define { i32, i32 } @v_permlane16_swap_b32_vv_bc(i32 %vdst_old, i32 %src0_old) {
112237
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
113238
; GCN-NEXT: v_permlane16_swap_b32_e64 v0, v1 bound_ctrl:1
114239
; GCN-NEXT: s_setpc_b64 s[30:31]
240+
; GFX950-LABEL: v_permlane16_swap_b32_vv_bc:
241+
; GFX950: ; %bb.0:
242+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
243+
; GFX950-NEXT: v_permlane16_swap_b32_e64 v0, v1 bound_ctrl:1
244+
; GFX950-NEXT: s_setpc_b64 s[30:31]
245+
;
246+
; GFX1250-LABEL: v_permlane16_swap_b32_vv_bc:
247+
; GFX1250: ; %bb.0:
248+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
249+
; GFX1250-NEXT: s_wait_kmcnt 0x0
250+
; GFX1250-NEXT: v_permlane16_swap_b32_e64 v0, v1 bound_ctrl:1
251+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
115252
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 %src0_old, i1 false, i1 true)
116253
ret { i32, i32 } %v
117254
}
@@ -122,6 +259,18 @@ define { i32, i32 } @v_permlane16_swap_b32_vv_fi_bc(i32 %vdst_old, i32 %src0_old
122259
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
123260
; GCN-NEXT: v_permlane16_swap_b32_e64 v0, v1 bound_ctrl:1 fi:1
124261
; GCN-NEXT: s_setpc_b64 s[30:31]
262+
; GFX950-LABEL: v_permlane16_swap_b32_vv_fi_bc:
263+
; GFX950: ; %bb.0:
264+
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
265+
; GFX950-NEXT: v_permlane16_swap_b32_e64 v0, v1 bound_ctrl:1 fi:1
266+
; GFX950-NEXT: s_setpc_b64 s[30:31]
267+
;
268+
; GFX1250-LABEL: v_permlane16_swap_b32_vv_fi_bc:
269+
; GFX1250: ; %bb.0:
270+
; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
271+
; GFX1250-NEXT: s_wait_kmcnt 0x0
272+
; GFX1250-NEXT: v_permlane16_swap_b32_e64 v0, v1 bound_ctrl:1 fi:1
273+
; GFX1250-NEXT: s_set_pc_i64 s[30:31]
125274
%v = call { i32, i32 } @llvm.amdgcn.permlane16.swap(i32 %vdst_old, i32 %src0_old, i1 true, i1 true)
126275
ret { i32, i32 } %v
127276
}

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -627,3 +627,9 @@ v_cvt_f32_fp8_e32 v1, 3
627627

628628
v_cvt_f32_fp8_e32 v1, v3
629629
// GFX1250: v_cvt_f32_fp8_e32 v1, v3 ; encoding: [0x03,0xd9,0x02,0x7e]
630+
631+
v_permlane16_swap_b32 v1, v2
632+
// GFX1250: v_permlane16_swap_b32_e32 v1, v2 ; encoding: [0x02,0x93,0x02,0x7e]
633+
634+
v_permlane16_swap_b32_e32 v1, v2
635+
// GFX1250: v_permlane16_swap_b32_e32 v1, v2 ; encoding: [0x02,0x93,0x02,0x7e]

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -663,3 +663,9 @@ v_cvt_f32_fp8_e32 v1, 3
663663

664664
v_cvt_f32_fp8_e32 v1, v3
665665
// GFX1250: v_cvt_f32_fp8_e32 v1, v3 ; encoding: [0x03,0xd9,0x02,0x7e]
666+
667+
v_permlane16_swap_b32 v1, v2
668+
// GFX1250: v_permlane16_swap_b32_e32 v1, v2 ; encoding: [0x02,0x93,0x02,0x7e]
669+
670+
v_permlane16_swap_b32_e32 v1, v2
671+
// GFX1250: v_permlane16_swap_b32_e32 v1, v2 ; encoding: [0x02,0x93,0x02,0x7e]

‎llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -720,3 +720,24 @@ v_cvt_pk_f16_fp8 v1, v150 op_sel:[1]
720720

721721
v_cvt_pk_f16_fp8 v1, s2 op_sel:[1]
722722
// GFX1250: v_cvt_pk_f16_fp8 v1, s2 op_sel:[1,0] ; encoding: [0x01,0x08,0xf5,0xd5,0x02,0x00,0x00,0x00]
723+
724+
v_permlane16_swap_b32_e64 v1, v2
725+
// GFX1250: v_permlane16_swap_b32_e64 v1, v2 ; encoding: [0x01,0x00,0xc9,0xd5,0x02,0x01,0x00,0x00]
726+
727+
v_permlane16_swap_b32 v1, v2 bound_ctrl:0
728+
// GFX1250: v_permlane16_swap_b32_e64 v1, v2 ; encoding: [0x01,0x00,0xc9,0xd5,0x02,0x01,0x00,0x00]
729+
730+
v_permlane16_swap_b32 v1, v2 fi:0
731+
// GFX1250: v_permlane16_swap_b32_e64 v1, v2 ; encoding: [0x01,0x00,0xc9,0xd5,0x02,0x01,0x00,0x00]
732+
733+
v_permlane16_swap_b32 v1, v2 bound_ctrl:1
734+
// GFX1250: v_permlane16_swap_b32_e64 v1, v2 bound_ctrl:1 ; encoding: [0x01,0x10,0xc9,0xd5,0x02,0x01,0x00,0x00]
735+
736+
v_permlane16_swap_b32 v1, v2 fi:1
737+
// GFX1250: v_permlane16_swap_b32_e64 v1, v2 fi:1 ; encoding: [0x01,0x08,0xc9,0xd5,0x02,0x01,0x00,0x00]
738+
739+
v_permlane16_swap_b32 v1, v2 bound_ctrl:1 fi:1
740+
// GFX1250: v_permlane16_swap_b32_e64 v1, v2 bound_ctrl:1 fi:1 ; encoding: [0x01,0x18,0xc9,0xd5,0x02,0x01,0x00,0x00]
741+
742+
v_permlane16_swap_b32_e64 v1, v2 bound_ctrl:1 fi:1
743+
// GFX1250: v_permlane16_swap_b32_e64 v1, v2 bound_ctrl:1 fi:1 ; encoding: [0x01,0x18,0xc9,0xd5,0x02,0x01,0x00,0x00]

0 commit comments

Comments
 (0)