Skip to content

Commit e35cd55

Browse files
shiltianrampitec
shiltian
and
rampitec
committed
[AMDGPU] Add support for v_prng_b32 on gfx1250
Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
1 parent 602d43c commit e35cd55

22 files changed

+585
-4
lines changed

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

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@
44

55
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
66

7+
typedef unsigned int uint;
78
typedef half __attribute__((ext_vector_type(2))) half2;
89

910
// CHECK-LABEL: @test_setprio_inc_wg(
@@ -42,6 +43,24 @@ void test_s_wait_tensorcnt() {
4243
__builtin_amdgcn_s_wait_tensorcnt(0);
4344
}
4445

46+
// CHECK-LABEL: @test_prng_b32(
47+
// CHECK-NEXT: entry:
48+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
49+
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
50+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
51+
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
52+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
53+
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
54+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
55+
// CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.prng.b32(i32 [[TMP0]])
56+
// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
57+
// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(1) [[TMP2]], align 4
58+
// CHECK-NEXT: ret void
59+
//
60+
void test_prng_b32(global uint* out, uint a) {
61+
*out = __builtin_amdgcn_prng_b32(a);
62+
}
63+
4564
// CHECK-LABEL: @test_tanh_f32(
4665
// CHECK-NEXT: entry:
4766
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)

‎llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4007,7 +4007,8 @@ SDValue AMDGPUTargetLowering::performIntrinsicWOChainCombine(
40074007
case Intrinsic::amdgcn_rcp_legacy:
40084008
case Intrinsic::amdgcn_rsq_legacy:
40094009
case Intrinsic::amdgcn_rsq_clamp:
4010-
case Intrinsic::amdgcn_tanh: {
4010+
case Intrinsic::amdgcn_tanh:
4011+
case Intrinsic::amdgcn_prng_b32: {
40114012
// FIXME: This is probably wrong. If src is an sNaN, it won't be quieted
40124013
SDValue Src = N->getOperand(1);
40134014
return Src.isUndef() ? Src : SDValue();

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1148,6 +1148,7 @@ defm V_MOV_B64 : VOP1_Real_FULL <GFX1250Gen, 0x1d>;
11481148
defm V_TANH_F32 : VOP1_Real_FULL<GFX1250Gen, 0x01e>;
11491149
defm V_TANH_F16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x01f>;
11501150
defm V_TANH_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
1151+
defm V_PRNG_B32 : VOP1_Real_FULL<GFX1250Gen, 0x04b>;
11511152
defm V_CVT_F32_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072, "v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;
11521153
defm V_CVT_PK_F16_FP8 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x075>;
11531154
defm V_CVT_PK_F16_BF8 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x076>;

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.prng.ll‎

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,7 @@
11
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN %s
22
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx950 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
3+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefixes=GCN,SDAG %s
4+
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefix=GCN %s
35

46
declare i32 @llvm.amdgcn.prng.b32(i32) #0
57

@@ -27,6 +29,13 @@ define amdgpu_kernel void @prng_b32_constant_100(ptr addrspace(1) %out) #1 {
2729
ret void
2830
}
2931

32+
; GCN-LABEL: {{^}}prng_undef_i32:
33+
; SDAG-NOT: v_prng_b32
34+
define amdgpu_kernel void @prng_undef_i32(ptr addrspace(1) %out) #1 {
35+
%prng = call i32 @llvm.amdgcn.prng.b32(i32 undef)
36+
store i32 %prng, ptr addrspace(1) %out, align 4
37+
ret void
38+
}
3039

3140
attributes #0 = { nounwind readnone }
32-
attributes #1 = { nounwind }
41+
attributes #1 = { nounwind }

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

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -163,6 +163,51 @@ v_tanh_bf16 v5, src_scc
163163
v_tanh_bf16 v127, 0x8000
164164
// GFX1250: v_tanh_bf16_e32 v127, 0x8000 ; encoding: [0xff,0x94,0xfe,0x7e,0x00,0x80,0x00,0x00]
165165

166+
v_prng_b32 v5, v1
167+
// GFX1250: v_prng_b32_e32 v5, v1 ; encoding: [0x01,0x97,0x0a,0x7e]
168+
169+
v_prng_b32 v5, v255
170+
// GFX1250: v_prng_b32_e32 v5, v255 ; encoding: [0xff,0x97,0x0a,0x7e]
171+
172+
v_prng_b32 v5, s1
173+
// GFX1250: v_prng_b32_e32 v5, s1 ; encoding: [0x01,0x96,0x0a,0x7e]
174+
175+
v_prng_b32 v5, s105
176+
// GFX1250: v_prng_b32_e32 v5, s105 ; encoding: [0x69,0x96,0x0a,0x7e]
177+
178+
v_prng_b32 v5, vcc_lo
179+
// GFX1250: v_prng_b32_e32 v5, vcc_lo ; encoding: [0x6a,0x96,0x0a,0x7e]
180+
181+
v_prng_b32 v5, vcc_hi
182+
// GFX1250: v_prng_b32_e32 v5, vcc_hi ; encoding: [0x6b,0x96,0x0a,0x7e]
183+
184+
v_prng_b32 v5, ttmp15
185+
// GFX1250: v_prng_b32_e32 v5, ttmp15 ; encoding: [0x7b,0x96,0x0a,0x7e]
186+
187+
v_prng_b32 v5, m0
188+
// GFX1250: v_prng_b32_e32 v5, m0 ; encoding: [0x7d,0x96,0x0a,0x7e]
189+
190+
v_prng_b32 v5, exec_lo
191+
// GFX1250: v_prng_b32_e32 v5, exec_lo ; encoding: [0x7e,0x96,0x0a,0x7e]
192+
193+
v_prng_b32 v5, exec_hi
194+
// GFX1250: v_prng_b32_e32 v5, exec_hi ; encoding: [0x7f,0x96,0x0a,0x7e]
195+
196+
v_prng_b32 v5, null
197+
// GFX1250: v_prng_b32_e32 v5, null ; encoding: [0x7c,0x96,0x0a,0x7e]
198+
199+
v_prng_b32 v5, -1
200+
// GFX1250: v_prng_b32_e32 v5, -1 ; encoding: [0xc1,0x96,0x0a,0x7e]
201+
202+
v_prng_b32 v5, 0.5
203+
// GFX1250: v_prng_b32_e32 v5, 0.5 ; encoding: [0xf0,0x96,0x0a,0x7e]
204+
205+
v_prng_b32 v5, src_scc
206+
// GFX1250: v_prng_b32_e32 v5, src_scc ; encoding: [0xfd,0x96,0x0a,0x7e]
207+
208+
v_prng_b32 v255, 0xaf123456
209+
// GFX1250: v_prng_b32_e32 v255, 0xaf123456 ; encoding: [0xff,0x96,0xfe,0x7f,0x56,0x34,0x12,0xaf]
210+
166211
v_rcp_bf16 v5, v1
167212
// GFX1250: v_rcp_bf16_e32 v5, v1 ; encoding: [0x01,0xf3,0x0a,0x7e]
168213

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

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -169,6 +169,51 @@ v_tanh_bf16 v127, 0x8000
169169
v_tanh_bf16 v5.h, v1.h
170170
// GFX1250: v_tanh_bf16_e32 v5.h, v1.h ; encoding: [0x81,0x95,0x0a,0x7f]
171171

172+
v_prng_b32 v5, v1
173+
// GFX1250: v_prng_b32_e32 v5, v1 ; encoding: [0x01,0x97,0x0a,0x7e]
174+
175+
v_prng_b32 v5, v255
176+
// GFX1250: v_prng_b32_e32 v5, v255 ; encoding: [0xff,0x97,0x0a,0x7e]
177+
178+
v_prng_b32 v5, s1
179+
// GFX1250: v_prng_b32_e32 v5, s1 ; encoding: [0x01,0x96,0x0a,0x7e]
180+
181+
v_prng_b32 v5, s105
182+
// GFX1250: v_prng_b32_e32 v5, s105 ; encoding: [0x69,0x96,0x0a,0x7e]
183+
184+
v_prng_b32 v5, vcc_lo
185+
// GFX1250: v_prng_b32_e32 v5, vcc_lo ; encoding: [0x6a,0x96,0x0a,0x7e]
186+
187+
v_prng_b32 v5, vcc_hi
188+
// GFX1250: v_prng_b32_e32 v5, vcc_hi ; encoding: [0x6b,0x96,0x0a,0x7e]
189+
190+
v_prng_b32 v5, ttmp15
191+
// GFX1250: v_prng_b32_e32 v5, ttmp15 ; encoding: [0x7b,0x96,0x0a,0x7e]
192+
193+
v_prng_b32 v5, m0
194+
// GFX1250: v_prng_b32_e32 v5, m0 ; encoding: [0x7d,0x96,0x0a,0x7e]
195+
196+
v_prng_b32 v5, exec_lo
197+
// GFX1250: v_prng_b32_e32 v5, exec_lo ; encoding: [0x7e,0x96,0x0a,0x7e]
198+
199+
v_prng_b32 v5, exec_hi
200+
// GFX1250: v_prng_b32_e32 v5, exec_hi ; encoding: [0x7f,0x96,0x0a,0x7e]
201+
202+
v_prng_b32 v5, null
203+
// GFX1250: v_prng_b32_e32 v5, null ; encoding: [0x7c,0x96,0x0a,0x7e]
204+
205+
v_prng_b32 v5, -1
206+
// GFX1250: v_prng_b32_e32 v5, -1 ; encoding: [0xc1,0x96,0x0a,0x7e]
207+
208+
v_prng_b32 v5, 0.5
209+
// GFX1250: v_prng_b32_e32 v5, 0.5 ; encoding: [0xf0,0x96,0x0a,0x7e]
210+
211+
v_prng_b32 v5, src_scc
212+
// GFX1250: v_prng_b32_e32 v5, src_scc ; encoding: [0xfd,0x96,0x0a,0x7e]
213+
214+
v_prng_b32 v255, 0xaf123456
215+
// GFX1250: v_prng_b32_e32 v255, 0xaf123456 ; encoding: [0xff,0x96,0xfe,0x7f,0x56,0x34,0x12,0xaf]
216+
172217
v_rcp_bf16 v5, v1
173218
// GFX1250: v_rcp_bf16_e32 v5, v1 ; encoding: [0x01,0xf3,0x0a,0x7e]
174219

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

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -170,6 +170,58 @@ v_tanh_bf16 v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 bound_ctrl:0 f
170170
// GFX1250: v_tanh_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 fi:1 ; encoding: [0xfa,0x94,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
171171
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
172172

173+
v_prng_b32 v5, v1 quad_perm:[3,2,1,0]
174+
// GFX1250: v_prng_b32_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x1b,0x00,0xff]
175+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
176+
177+
v_prng_b32 v5, v1 quad_perm:[0,1,2,3]
178+
// GFX1250: v_prng_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0xe4,0x00,0xff]
179+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
180+
181+
v_prng_b32 v5, v1 row_mirror
182+
// GFX1250: v_prng_b32_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x40,0x01,0xff]
183+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
184+
185+
v_prng_b32 v5, v1 row_half_mirror
186+
// GFX1250: v_prng_b32_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x41,0x01,0xff]
187+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
188+
189+
v_prng_b32 v5, v1 row_shl:1
190+
// GFX1250: v_prng_b32_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x01,0x01,0xff]
191+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
192+
193+
v_prng_b32 v5, v1 row_shl:15
194+
// GFX1250: v_prng_b32_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x0f,0x01,0xff]
195+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
196+
197+
v_prng_b32 v5, v1 row_shr:1
198+
// GFX1250: v_prng_b32_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x11,0x01,0xff]
199+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
200+
201+
v_prng_b32 v5, v1 row_shr:15
202+
// GFX1250: v_prng_b32_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x1f,0x01,0xff]
203+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
204+
205+
v_prng_b32 v5, v1 row_ror:1
206+
// GFX1250: v_prng_b32_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x21,0x01,0xff]
207+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
208+
209+
v_prng_b32 v5, v1 row_ror:15
210+
// GFX1250: v_prng_b32_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x2f,0x01,0xff]
211+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
212+
213+
v_prng_b32 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
214+
// GFX1250: v_prng_b32_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x50,0x01,0xff]
215+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
216+
217+
v_prng_b32 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
218+
// GFX1250: v_prng_b32_dpp v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1 ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x5f,0x01,0x01]
219+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
220+
221+
v_prng_b32 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
222+
// GFX1250: v_prng_b32_dpp v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x60,0x09,0x13]
223+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
224+
173225
v_rcp_bf16 v5, v1 quad_perm:[3,2,1,0]
174226
// GFX1250: v_rcp_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xf2,0x0a,0x7e,0x01,0x1b,0x00,0xff]
175227
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU

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

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -178,6 +178,58 @@ v_tanh_bf16 v5.h, v1.h quad_perm:[3,2,1,0]
178178
// GFX1250: v_tanh_bf16_dpp v5.h, v1.h quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7f,0x81,0x1b,0x00,0xff]
179179
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
180180

181+
v_prng_b32 v5, v1 quad_perm:[3,2,1,0]
182+
// GFX1250: v_prng_b32_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x1b,0x00,0xff]
183+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
184+
185+
v_prng_b32 v5, v1 quad_perm:[0,1,2,3]
186+
// GFX1250: v_prng_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0xe4,0x00,0xff]
187+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
188+
189+
v_prng_b32 v5, v1 row_mirror
190+
// GFX1250: v_prng_b32_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x40,0x01,0xff]
191+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
192+
193+
v_prng_b32 v5, v1 row_half_mirror
194+
// GFX1250: v_prng_b32_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x41,0x01,0xff]
195+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
196+
197+
v_prng_b32 v5, v1 row_shl:1
198+
// GFX1250: v_prng_b32_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x01,0x01,0xff]
199+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
200+
201+
v_prng_b32 v5, v1 row_shl:15
202+
// GFX1250: v_prng_b32_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x0f,0x01,0xff]
203+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
204+
205+
v_prng_b32 v5, v1 row_shr:1
206+
// GFX1250: v_prng_b32_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x11,0x01,0xff]
207+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
208+
209+
v_prng_b32 v5, v1 row_shr:15
210+
// GFX1250: v_prng_b32_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x1f,0x01,0xff]
211+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
212+
213+
v_prng_b32 v5, v1 row_ror:1
214+
// GFX1250: v_prng_b32_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x21,0x01,0xff]
215+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
216+
217+
v_prng_b32 v5, v1 row_ror:15
218+
// GFX1250: v_prng_b32_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x2f,0x01,0xff]
219+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
220+
221+
v_prng_b32 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
222+
// GFX1250: v_prng_b32_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x50,0x01,0xff]
223+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
224+
225+
v_prng_b32 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
226+
// GFX1250: v_prng_b32_dpp v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1 ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x5f,0x01,0x01]
227+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
228+
229+
v_prng_b32 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
230+
// GFX1250: v_prng_b32_dpp v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 ; encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x60,0x09,0x13]
231+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
232+
181233
v_rcp_bf16 v5, v1 quad_perm:[3,2,1,0]
182234
// GFX1250: v_rcp_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xf2,0x0a,0x7e,0x01,0x1b,0x00,0xff]
183235
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,18 @@ v_tanh_bf16 v127, v127 dpp8:[0,0,0,0,0,0,0,0] fi:0
3838
// GFX1250: v_tanh_bf16_dpp v127, v127 dpp8:[0,0,0,0,0,0,0,0] ; encoding: [0xe9,0x94,0xfe,0x7e,0x7f,0x00,0x00,0x00]
3939
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
4040

41+
v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
42+
// GFX1250: v_prng_b32_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding: [0xe9,0x96,0x0a,0x7e,0x01,0x77,0x39,0x05]
43+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
44+
45+
v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1
46+
// GFX1250: v_prng_b32_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1 ; encoding: [0xea,0x96,0x0a,0x7e,0x01,0x77,0x39,0x05]
47+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
48+
49+
v_prng_b32 v255, v255 dpp8:[0,0,0,0,0,0,0,0] fi:0
50+
// GFX1250: v_prng_b32_dpp v255, v255 dpp8:[0,0,0,0,0,0,0,0] ; encoding: [0xe9,0x96,0xfe,0x7f,0xff,0x00,0x00,0x00]
51+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
52+
4153
v_rcp_bf16 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
4254
// GFX1250: v_rcp_bf16_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding: [0xe9,0xf2,0x0a,0x7e,0x01,0x77,0x39,0x05]
4355
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,18 @@ v_tanh_bf16 v5.h, v1.h dpp8:[7,6,5,4,3,2,1,0]
4646
// GFX1250: v_tanh_bf16_dpp v5.h, v1.h dpp8:[7,6,5,4,3,2,1,0] ; encoding: [0xe9,0x94,0x0a,0x7f,0x81,0x77,0x39,0x05]
4747
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
4848

49+
v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
50+
// GFX1250: v_prng_b32_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding: [0xe9,0x96,0x0a,0x7e,0x01,0x77,0x39,0x05]
51+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
52+
53+
v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1
54+
// GFX1250: v_prng_b32_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1 ; encoding: [0xea,0x96,0x0a,0x7e,0x01,0x77,0x39,0x05]
55+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
56+
57+
v_prng_b32 v255, v255 dpp8:[0,0,0,0,0,0,0,0] fi:0
58+
// GFX1250: v_prng_b32_dpp v255, v255 dpp8:[0,0,0,0,0,0,0,0] ; encoding: [0xe9,0x96,0xfe,0x7f,0xff,0x00,0x00,0x00]
59+
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
60+
4961
v_rcp_bf16 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
5062
// GFX1250: v_rcp_bf16_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding: [0xe9,0xf2,0x0a,0x7e,0x01,0x77,0x39,0x05]
5163
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU

0 commit comments

Comments
 (0)