https://github.com/shiltian updated
https://github.com/llvm/llvm-project/pull/149450
>From 487881ca26a91b76e24e10e152aa98fedf19414a Mon Sep 17 00:00:00 2001
From: Shilei Tian <i...@tianshilei.me>
Date: Fri, 18 Jul 2025 00:26:15 -0400
Subject: [PATCH] [AMDGPU] Add support for `v_prng_b32` on gfx1250
Co-authored-by: Mekhanoshin, Stanislav <stanislav.mekhanos...@amd.com>
---
.../CodeGenOpenCL/builtins-amdgcn-gfx1250.cl | 19 +++++++
llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp | 3 +-
llvm/lib/Target/AMDGPU/VOP1Instructions.td | 1 +
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.prng.ll | 11 +++-
llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s | 45 ++++++++++++++++
llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s | 45 ++++++++++++++++
.../MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s | 52 +++++++++++++++++++
llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s | 52 +++++++++++++++++++
.../MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s | 12 +++++
llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s | 12 +++++
.../gfx1250_asm_vop3_from_vop1-fake16.s | 36 +++++++++++++
.../MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s | 36 +++++++++++++
.../gfx1250_asm_vop3_from_vop1_dpp16-fake16.s | 44 ++++++++++++++++
.../AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s | 44 ++++++++++++++++
.../gfx1250_asm_vop3_from_vop1_dpp8-fake16.s | 4 ++
.../AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s | 4 ++
.../Disassembler/AMDGPU/gfx1250_dasm_vop1.txt | 45 ++++++++++++++++
.../AMDGPU/gfx1250_dasm_vop1_dpp16.txt | 39 ++++++++++++++
.../AMDGPU/gfx1250_dasm_vop1_dpp8.txt | 9 ++++
.../AMDGPU/gfx1250_dasm_vop3_from_vop1.txt | 36 +++++++++++++
.../gfx1250_dasm_vop3_from_vop1_dpp16.txt | 33 ++++++++++++
.../gfx1250_dasm_vop3_from_vop1_dpp8.txt | 7 ++-
22 files changed, 585 insertions(+), 4 deletions(-)
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
index a9ea17642d6ad..d42e51d04ab9d 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
@@ -4,6 +4,7 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+typedef unsigned int uint;
typedef half __attribute__((ext_vector_type(2))) half2;
// CHECK-LABEL: @test_setprio_inc_wg(
@@ -42,6 +43,24 @@ void test_s_wait_tensorcnt() {
__builtin_amdgcn_s_wait_tensorcnt(0);
}
+// CHECK-LABEL: @test_prng_b32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[OUT_ADDR]] to ptr
+// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[A_ADDR]] to ptr
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]],
align 8
+// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.prng.b32(i32 [[TMP0]])
+// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(1) [[TMP2]], align 4
+// CHECK-NEXT: ret void
+//
+void test_prng_b32(global uint* out, uint a) {
+ *out = __builtin_amdgcn_prng_b32(a);
+}
+
// CHECK-LABEL: @test_tanh_f32(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
index 3d040fb705a8d..84e4fa1dc84aa 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp
@@ -4007,7 +4007,8 @@ SDValue
AMDGPUTargetLowering::performIntrinsicWOChainCombine(
case Intrinsic::amdgcn_rcp_legacy:
case Intrinsic::amdgcn_rsq_legacy:
case Intrinsic::amdgcn_rsq_clamp:
- case Intrinsic::amdgcn_tanh: {
+ case Intrinsic::amdgcn_tanh:
+ case Intrinsic::amdgcn_prng_b32: {
// FIXME: This is probably wrong. If src is an sNaN, it won't be quieted
SDValue Src = N->getOperand(1);
return Src.isUndef() ? Src : SDValue();
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index 1bbbb610305e9..3ee90857b34b8 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -1148,6 +1148,7 @@ defm V_MOV_B64 : VOP1_Real_FULL <GFX1250Gen, 0x1d>;
defm V_TANH_F32 : VOP1_Real_FULL<GFX1250Gen, 0x01e>;
defm V_TANH_F16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x01f>;
defm V_TANH_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
+defm V_PRNG_B32 : VOP1_Real_FULL<GFX1250Gen, 0x04b>;
defm V_CVT_F32_BF16 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072,
"v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;
defm V_CVT_PK_F16_FP8 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x075>;
defm V_CVT_PK_F16_BF8 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x076>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.prng.ll
b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.prng.ll
index 2faf375a97a86..e3cfbd41b923d 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.prng.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.prng.ll
@@ -1,5 +1,7 @@
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx950 -verify-machineinstrs <
%s | FileCheck -check-prefixes=GCN %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx950 -verify-machineinstrs <
%s | FileCheck -check-prefix=GCN %s
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck
-check-prefixes=GCN,SDAG %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck
-check-prefix=GCN %s
declare i32 @llvm.amdgcn.prng.b32(i32) #0
@@ -27,6 +29,13 @@ define amdgpu_kernel void @prng_b32_constant_100(ptr
addrspace(1) %out) #1 {
ret void
}
+; GCN-LABEL: {{^}}prng_undef_i32:
+; SDAG-NOT: v_prng_b32
+define amdgpu_kernel void @prng_undef_i32(ptr addrspace(1) %out) #1 {
+ %prng = call i32 @llvm.amdgcn.prng.b32(i32 undef)
+ store i32 %prng, ptr addrspace(1) %out, align 4
+ ret void
+}
attributes #0 = { nounwind readnone }
-attributes #1 = { nounwind }
\ No newline at end of file
+attributes #1 = { nounwind }
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
index 279bb262bff04..5f310a9954ad0 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
@@ -163,6 +163,51 @@ v_tanh_bf16 v5, src_scc
v_tanh_bf16 v127, 0x8000
// GFX1250: v_tanh_bf16_e32 v127, 0x8000 ; encoding:
[0xff,0x94,0xfe,0x7e,0x00,0x80,0x00,0x00]
+v_prng_b32 v5, v1
+// GFX1250: v_prng_b32_e32 v5, v1 ; encoding:
[0x01,0x97,0x0a,0x7e]
+
+v_prng_b32 v5, v255
+// GFX1250: v_prng_b32_e32 v5, v255 ; encoding:
[0xff,0x97,0x0a,0x7e]
+
+v_prng_b32 v5, s1
+// GFX1250: v_prng_b32_e32 v5, s1 ; encoding:
[0x01,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, s105
+// GFX1250: v_prng_b32_e32 v5, s105 ; encoding:
[0x69,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, vcc_lo
+// GFX1250: v_prng_b32_e32 v5, vcc_lo ; encoding:
[0x6a,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, vcc_hi
+// GFX1250: v_prng_b32_e32 v5, vcc_hi ; encoding:
[0x6b,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, ttmp15
+// GFX1250: v_prng_b32_e32 v5, ttmp15 ; encoding:
[0x7b,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, m0
+// GFX1250: v_prng_b32_e32 v5, m0 ; encoding:
[0x7d,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, exec_lo
+// GFX1250: v_prng_b32_e32 v5, exec_lo ; encoding:
[0x7e,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, exec_hi
+// GFX1250: v_prng_b32_e32 v5, exec_hi ; encoding:
[0x7f,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, null
+// GFX1250: v_prng_b32_e32 v5, null ; encoding:
[0x7c,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, -1
+// GFX1250: v_prng_b32_e32 v5, -1 ; encoding:
[0xc1,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, 0.5
+// GFX1250: v_prng_b32_e32 v5, 0.5 ; encoding:
[0xf0,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, src_scc
+// GFX1250: v_prng_b32_e32 v5, src_scc ; encoding:
[0xfd,0x96,0x0a,0x7e]
+
+v_prng_b32 v255, 0xaf123456
+// GFX1250: v_prng_b32_e32 v255, 0xaf123456 ; encoding:
[0xff,0x96,0xfe,0x7f,0x56,0x34,0x12,0xaf]
+
v_rcp_bf16 v5, v1
// GFX1250: v_rcp_bf16_e32 v5, v1 ; encoding:
[0x01,0xf3,0x0a,0x7e]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
index 76272d25d92d4..aa2e028f661e1 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
@@ -169,6 +169,51 @@ v_tanh_bf16 v127, 0x8000
v_tanh_bf16 v5.h, v1.h
// GFX1250: v_tanh_bf16_e32 v5.h, v1.h ; encoding:
[0x81,0x95,0x0a,0x7f]
+v_prng_b32 v5, v1
+// GFX1250: v_prng_b32_e32 v5, v1 ; encoding:
[0x01,0x97,0x0a,0x7e]
+
+v_prng_b32 v5, v255
+// GFX1250: v_prng_b32_e32 v5, v255 ; encoding:
[0xff,0x97,0x0a,0x7e]
+
+v_prng_b32 v5, s1
+// GFX1250: v_prng_b32_e32 v5, s1 ; encoding:
[0x01,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, s105
+// GFX1250: v_prng_b32_e32 v5, s105 ; encoding:
[0x69,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, vcc_lo
+// GFX1250: v_prng_b32_e32 v5, vcc_lo ; encoding:
[0x6a,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, vcc_hi
+// GFX1250: v_prng_b32_e32 v5, vcc_hi ; encoding:
[0x6b,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, ttmp15
+// GFX1250: v_prng_b32_e32 v5, ttmp15 ; encoding:
[0x7b,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, m0
+// GFX1250: v_prng_b32_e32 v5, m0 ; encoding:
[0x7d,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, exec_lo
+// GFX1250: v_prng_b32_e32 v5, exec_lo ; encoding:
[0x7e,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, exec_hi
+// GFX1250: v_prng_b32_e32 v5, exec_hi ; encoding:
[0x7f,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, null
+// GFX1250: v_prng_b32_e32 v5, null ; encoding:
[0x7c,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, -1
+// GFX1250: v_prng_b32_e32 v5, -1 ; encoding:
[0xc1,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, 0.5
+// GFX1250: v_prng_b32_e32 v5, 0.5 ; encoding:
[0xf0,0x96,0x0a,0x7e]
+
+v_prng_b32 v5, src_scc
+// GFX1250: v_prng_b32_e32 v5, src_scc ; encoding:
[0xfd,0x96,0x0a,0x7e]
+
+v_prng_b32 v255, 0xaf123456
+// GFX1250: v_prng_b32_e32 v255, 0xaf123456 ; encoding:
[0xff,0x96,0xfe,0x7f,0x56,0x34,0x12,0xaf]
+
v_rcp_bf16 v5, v1
// GFX1250: v_rcp_bf16_e32 v5, v1 ; encoding:
[0x01,0xf3,0x0a,0x7e]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
index 0a8ee84561d33..e1cd2e3043693 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
@@ -170,6 +170,58 @@ v_tanh_bf16 v127, -|v127| row_xmask:15 row_mask:0x3
bank_mask:0x0 bound_ctrl:0 f
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32 v5, v1 quad_perm:[3,2,1,0]
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 quad_perm:[0,1,2,3]
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_mirror
+// GFX1250: v_prng_b32_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ;
encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_half_mirror
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shl:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shl:15
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shr:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shr:15
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_ror:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_ror:15
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16 v5, v1 quad_perm:[3,2,1,0]
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
index d4afb9d9b2d9a..c1d3238b65cbd 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
@@ -178,6 +178,58 @@ v_tanh_bf16 v5.h, v1.h quad_perm:[3,2,1,0]
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32 v5, v1 quad_perm:[3,2,1,0]
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 quad_perm:[0,1,2,3]
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_mirror
+// GFX1250: v_prng_b32_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ;
encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_half_mirror
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shl:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shl:15
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shr:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_shr:15
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_ror:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_ror:15
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16 v5, v1 quad_perm:[3,2,1,0]
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
index a7cb6bf8de69c..100e9f92ff58b 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
@@ -38,6 +38,18 @@ v_tanh_bf16 v127, v127 dpp8:[0,0,0,0,0,0,0,0] fi:0
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v255, v255 dpp8:[0,0,0,0,0,0,0,0] fi:0
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
index 6acab7edc0d49..2ae103545443c 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
@@ -46,6 +46,18 @@ v_tanh_bf16 v5.h, v1.h dpp8:[7,6,5,4,3,2,1,0]
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v5, v1 dpp8:[7,6,5,4,3,2,1,0] fi:1
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32 v255, v255 dpp8:[0,0,0,0,0,0,0,0] fi:0
+// 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]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16 v5, v1 dpp8:[7,6,5,4,3,2,1,0]
// 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]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
index 7486d849253e8..9c6a9127d82e4 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
@@ -127,6 +127,42 @@ v_cvt_f32_fp8 v1, v3 byte_sel:1 clamp
v_cvt_f32_fp8 v1, v3 byte_sel:2 clamp
// GFX1250: v_cvt_f32_fp8_e64 v1, v3 byte_sel:2 clamp ; encoding:
[0x01,0x88,0xec,0xd5,0x03,0x01,0x00,0x00]
+v_prng_b32_e64 v5, v1
+// GFX1250: v_prng_b32_e64 v5, v1 ; encoding:
[0x05,0x00,0xcb,0xd5,0x01,0x01,0x00,0x00]
+
+v_prng_b32_e64 v5, v255
+// GFX1250: v_prng_b32_e64 v5, v255 ; encoding:
[0x05,0x00,0xcb,0xd5,0xff,0x01,0x00,0x00]
+
+v_prng_b32_e64 v5, s1
+// GFX1250: v_prng_b32_e64 v5, s1 ; encoding:
[0x05,0x00,0xcb,0xd5,0x01,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, s105
+// GFX1250: v_prng_b32_e64 v5, s105 ; encoding:
[0x05,0x00,0xcb,0xd5,0x69,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, vcc_lo
+// GFX1250: v_prng_b32_e64 v5, vcc_lo ; encoding:
[0x05,0x00,0xcb,0xd5,0x6a,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, vcc_hi
+// GFX1250: v_prng_b32_e64 v5, vcc_hi ; encoding:
[0x05,0x00,0xcb,0xd5,0x6b,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, ttmp15
+// GFX1250: v_prng_b32_e64 v5, ttmp15 ; encoding:
[0x05,0x00,0xcb,0xd5,0x7b,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, m0
+// GFX1250: v_prng_b32_e64 v5, m0 ; encoding:
[0x05,0x00,0xcb,0xd5,0x7d,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, exec_lo
+// GFX1250: v_prng_b32_e64 v5, exec_lo ; encoding:
[0x05,0x00,0xcb,0xd5,0x7e,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, exec_hi
+// GFX1250: v_prng_b32_e64 v5, exec_hi ; encoding:
[0x05,0x00,0xcb,0xd5,0x7f,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, null
+// GFX1250: v_prng_b32_e64 v5, null ; encoding:
[0x05,0x00,0xcb,0xd5,0x7c,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, -1
+// GFX1250: v_prng_b32_e64 v5, -1 ; encoding:
[0x05,0x00,0xcb,0xd5,0xc1,0x00,0x00,0x00]
+
v_tanh_f32_e64 v5, v1
// GFX1250: v_tanh_f32_e64 v5, v1 ; encoding:
[0x05,0x00,0x9e,0xd5,0x01,0x01,0x00,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
index b59b8b31e2d5f..2f57d1c331c42 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
@@ -223,6 +223,42 @@ v_tanh_f16_e64 v255, -|0x8000| clamp div:2
v_tanh_f16 v5.l, v128.h
// GFX1250: v_tanh_f16_e64 v5.l, v128.h op_sel:[1,0] ; encoding:
[0x05,0x08,0x9f,0xd5,0x80,0x01,0x00,0x00]
+v_prng_b32_e64 v5, v1
+// GFX1250: v_prng_b32_e64 v5, v1 ; encoding:
[0x05,0x00,0xcb,0xd5,0x01,0x01,0x00,0x00]
+
+v_prng_b32_e64 v5, v255
+// GFX1250: v_prng_b32_e64 v5, v255 ; encoding:
[0x05,0x00,0xcb,0xd5,0xff,0x01,0x00,0x00]
+
+v_prng_b32_e64 v5, s1
+// GFX1250: v_prng_b32_e64 v5, s1 ; encoding:
[0x05,0x00,0xcb,0xd5,0x01,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, s105
+// GFX1250: v_prng_b32_e64 v5, s105 ; encoding:
[0x05,0x00,0xcb,0xd5,0x69,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, vcc_lo
+// GFX1250: v_prng_b32_e64 v5, vcc_lo ; encoding:
[0x05,0x00,0xcb,0xd5,0x6a,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, vcc_hi
+// GFX1250: v_prng_b32_e64 v5, vcc_hi ; encoding:
[0x05,0x00,0xcb,0xd5,0x6b,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, ttmp15
+// GFX1250: v_prng_b32_e64 v5, ttmp15 ; encoding:
[0x05,0x00,0xcb,0xd5,0x7b,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, m0
+// GFX1250: v_prng_b32_e64 v5, m0 ; encoding:
[0x05,0x00,0xcb,0xd5,0x7d,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, exec_lo
+// GFX1250: v_prng_b32_e64 v5, exec_lo ; encoding:
[0x05,0x00,0xcb,0xd5,0x7e,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, exec_hi
+// GFX1250: v_prng_b32_e64 v5, exec_hi ; encoding:
[0x05,0x00,0xcb,0xd5,0x7f,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, null
+// GFX1250: v_prng_b32_e64 v5, null ; encoding:
[0x05,0x00,0xcb,0xd5,0x7c,0x00,0x00,0x00]
+
+v_prng_b32_e64 v5, -1
+// GFX1250: v_prng_b32_e64 v5, -1 ; encoding:
[0x05,0x00,0xcb,0xd5,0xc1,0x00,0x00,0x00]
+
v_rcp_bf16_e64 v5, v1
// GFX1250: v_rcp_bf16_e64 v5, v1 ; encoding:
[0x05,0x00,0xf9,0xd5,0x01,0x01,0x00,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
index f7f20f46161ce..29bb842b529b7 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
@@ -170,6 +170,50 @@ v_tanh_bf16_e64_dpp v255, -|v255| clamp div:2 row_xmask:15
row_mask:0x3 bank_mas
// GFX1250: v_tanh_bf16_e64_dpp v255, -|v255| clamp div:2 row_xmask:15
row_mask:0x3 bank_mask:0x0 fi:1 ; encoding:
[0xff,0x81,0xca,0xd5,0xfa,0x00,0x00,0x38,0xff,0x6f,0x05,0x30]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32_e64_dpp v5, v1 quad_perm:[3,2,1,0]
+// GFX1250: v_prng_b32_e64_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1b,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 quad_perm:[0,1,2,3]
+// GFX1250: v_prng_b32_e64_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_mirror
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_half_mirror
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_half_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x41,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shl:1
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x01,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shl:15
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x0f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shr:1
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x11,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shr:15
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_ror:1
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x21,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_ror:15
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x2f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x50,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16_e64_dpp v5, v1 quad_perm:[3,2,1,0]
// GFX1250: v_rcp_bf16_e64_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xf9,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1b,0x00,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
index e1241b01ccae1..7df92751c38d1 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
@@ -178,6 +178,50 @@ v_tanh_bf16_e64_dpp v5.h, v128.h quad_perm:[3,2,1,0]
// GFX1250: v_tanh_bf16_e64_dpp v5.h, v128.h op_sel:[1,1] quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x48,0xca,0xd5,0xfa,0x00,0x00,0x00,0x80,0x1b,0x00,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32_e64_dpp v5, v1 quad_perm:[3,2,1,0]
+// GFX1250: v_prng_b32_e64_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1b,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 quad_perm:[0,1,2,3]
+// GFX1250: v_prng_b32_e64_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_mirror
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_half_mirror
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_half_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x41,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shl:1
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x01,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shl:15
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x0f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shr:1
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x11,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_shr:15
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_ror:1
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x21,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_ror:15
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x2f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_prng_b32_e64_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
+// GFX1250: v_prng_b32_e64_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x50,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16_e64_dpp v5, v1 quad_perm:[3,2,1,0]
// GFX1250: v_rcp_bf16_e64_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xf9,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1b,0x00,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
index 0106175301d20..d235aeb9f3e62 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
@@ -50,6 +50,10 @@ v_tanh_bf16_e64_dpp v255, -|v255| clamp div:2
dpp8:[0,0,0,0,0,0,0,0] fi:0
// GFX1250: v_tanh_bf16_e64_dpp v255, -|v255| clamp div:2
dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xff,0x81,0xca,0xd5,0xe9,0x00,0x00,0x38,0xff,0x00,0x00,0x00]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_prng_b32_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0xcb,0xd5,0xe9,0x00,0x00,0x00,0x01,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0]
// GFX1250: v_rcp_bf16_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0xf9,0xd5,0xe9,0x00,0x00,0x00,0x01,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
index 93b86f3ffb841..f25e2a5882436 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
@@ -58,6 +58,10 @@ v_tanh_bf16_e64_dpp v5.h, v128.h dpp8:[7,6,5,4,3,2,1,0]
// GFX1250: v_tanh_bf16_e64_dpp v5.h, v128.h op_sel:[1,1]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x48,0xca,0xd5,0xe9,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_prng_b32_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_prng_b32_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0xcb,0xd5,0xe9,0x00,0x00,0x00,0x01,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_rcp_bf16_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0]
// GFX1250: v_rcp_bf16_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0xf9,0xd5,0xe9,0x00,0x00,0x00,0x01,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
index 5f37ba91e071b..aa968b2bb2bee 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
@@ -200,6 +200,51 @@
0x81,0x95,0x0a,0x7f
# GFX1250-REAL16: v_tanh_bf16_e32 v5.h, v1.h ; encoding:
[0x81,0x95,0x0a,0x7f]
+0xff,0x96,0xfe,0x7f,0x56,0x34,0x12,0xaf
+# GFX1250: v_prng_b32_e32 v255, 0xaf123456 ; encoding:
[0xff,0x96,0xfe,0x7f,0x56,0x34,0x12,0xaf]
+
+0xc1,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, -1 ; encoding:
[0xc1,0x96,0x0a,0x7e]
+
+0xf0,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, 0.5 ; encoding:
[0xf0,0x96,0x0a,0x7e]
+
+0x7f,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, exec_hi ; encoding:
[0x7f,0x96,0x0a,0x7e]
+
+0x7e,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, exec_lo ; encoding:
[0x7e,0x96,0x0a,0x7e]
+
+0x7d,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, m0 ; encoding:
[0x7d,0x96,0x0a,0x7e]
+
+0x7c,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, null ; encoding:
[0x7c,0x96,0x0a,0x7e]
+
+0x01,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, s1 ; encoding:
[0x01,0x96,0x0a,0x7e]
+
+0x69,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, s105 ; encoding:
[0x69,0x96,0x0a,0x7e]
+
+0xfd,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, src_scc ; encoding:
[0xfd,0x96,0x0a,0x7e]
+
+0x7b,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, ttmp15 ; encoding:
[0x7b,0x96,0x0a,0x7e]
+
+0x01,0x97,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, v1 ; encoding:
[0x01,0x97,0x0a,0x7e]
+
+0xff,0x97,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, v255 ; encoding:
[0xff,0x97,0x0a,0x7e]
+
+0x6b,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, vcc_hi ; encoding:
[0x6b,0x96,0x0a,0x7e]
+
+0x6a,0x96,0x0a,0x7e
+# GFX1250: v_prng_b32_e32 v5, vcc_lo ; encoding:
[0x6a,0x96,0x0a,0x7e]
+
0xff,0xf2,0xfe,0x7e,0x00,0x80,0x00,0x00
# GFX1250-REAL16: v_rcp_bf16_e32 v127.l, 0x8000 ; encoding:
[0xff,0xf2,0xfe,0x7e,0x00,0x80,0x00,0x00]
# GFX1250-FAKE16: v_rcp_bf16_e32 v127, 0x8000 ; encoding:
[0xff,0xf2,0xfe,0x7e,0x00,0x80,0x00,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
index 57bee2766ce44..913a2a916ff62 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
@@ -162,6 +162,45 @@
0xfa,0x94,0x0a,0x7f,0x81,0x1b,0x00,0xff
# GFX1250-REAL16: 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]
+0xfa,0x96,0x0a,0x7e,0x01,0xe4,0x00,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x1b,0x00,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x41,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x40,0x01,0xff
+# GFX1250: v_prng_b32_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ;
encoding: [0xfa,0x96,0x0a,0x7e,0x01,0x40,0x01,0xff]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x21,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x2f,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x50,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x5f,0x01,0x01
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x01,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x0f,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x11,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x1f,0x01,0xff
+# 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]
+
+0xfa,0x96,0x0a,0x7e,0x01,0x60,0x09,0x13
+# 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]
+
0xfa,0xf2,0xfe,0x7e,0x7f,0x6f,0x35,0x30
# GFX1250-REAL16: v_rcp_bf16_dpp v127.l, -|v127.l| row_xmask:15 row_mask:0x3
bank_mask:0x0 fi:1 ; encoding: [0xfa,0xf2,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
# GFX1250-FAKE16: v_rcp_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3
bank_mask:0x0 fi:1 ; encoding: [0xfa,0xf2,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
index 28ec6b11b4de3..4afe44e241bf3 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
@@ -43,6 +43,15 @@
# GFX1250-REAL16: 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]
# GFX1250-FAKE16: v_add_f64_e32 v[156:157], v[129:130], v[187:188] ; encoding:
[0x81,0x77,0x39,0x05]
+0xe9,0x96,0xfe,0x7f,0xff,0x00,0x00,0x00
+# 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]
+
+0xe9,0x96,0x0a,0x7e,0x01,0x77,0x39,0x05
+# 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]
+
+0xea,0x96,0x0a,0x7e,0x01,0x77,0x39,0x05
+# 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]
+
0xe9,0xf2,0xfe,0x7e,0x7f,0x00,0x00,0x00
# GFX1250-REAL16: v_rcp_bf16_dpp v127.l, v127.l dpp8:[0,0,0,0,0,0,0,0] ;
encoding: [0xe9,0xf2,0xfe,0x7e,0x7f,0x00,0x00,0x00]
# GFX1250-FAKE16: v_rcp_bf16_dpp v127, v127 dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xe9,0xf2,0xfe,0x7e,0x7f,0x00,0x00,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
index 5004762729701..1cf3b8807d044 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
@@ -175,6 +175,42 @@
# GFX1250-REAL16: v_tanh_bf16_e64 v5.l, v128.h op_sel:[1,0] ; encoding:
[0x05,0x08,0xca,0xd5,0x80,0x01,0x00,0x00]
# GFX1250-FAKE16: v_tanh_bf16_e64 v5, v128 ; encoding:
[0x05,0x00,0xca,0xd5,0x80,0x01,0x00,0x00]
+0x05,0x00,0xcb,0xd5,0xc1,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, -1 ; encoding:
[0x05,0x00,0xcb,0xd5,0xc1,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x7f,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, exec_hi ; encoding:
[0x05,0x00,0xcb,0xd5,0x7f,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x7e,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, exec_lo ; encoding:
[0x05,0x00,0xcb,0xd5,0x7e,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x7d,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, m0 ; encoding:
[0x05,0x00,0xcb,0xd5,0x7d,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x7c,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, null ; encoding:
[0x05,0x00,0xcb,0xd5,0x7c,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x01,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, s1 ; encoding:
[0x05,0x00,0xcb,0xd5,0x01,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x69,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, s105 ; encoding:
[0x05,0x00,0xcb,0xd5,0x69,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x7b,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, ttmp15 ; encoding:
[0x05,0x00,0xcb,0xd5,0x7b,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x01,0x01,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, v1 ; encoding:
[0x05,0x00,0xcb,0xd5,0x01,0x01,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0xff,0x01,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, v255 ; encoding:
[0x05,0x00,0xcb,0xd5,0xff,0x01,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x6b,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, vcc_hi ; encoding:
[0x05,0x00,0xcb,0xd5,0x6b,0x00,0x00,0x00]
+
+0x05,0x00,0xcb,0xd5,0x6a,0x00,0x00,0x00
+# GFX1250: v_prng_b32_e64 v5, vcc_lo ; encoding:
[0x05,0x00,0xcb,0xd5,0x6a,0x00,0x00,0x00]
+
0xff,0x81,0xf9,0xd5,0xff,0x00,0x00,0x38,0x00,0x80,0x00,0x00
# GFX1250-REAL16: v_rcp_bf16_e64 v255.l, -|0x8000| clamp div:2 ; encoding:
[0xff,0x81,0xf9,0xd5,0xff,0x00,0x00,0x38,0x00,0x80,0x00,0x00]
# GFX1250-FAKE16: v_rcp_bf16_e64 v255, -|0x8000| clamp div:2 ; encoding:
[0xff,0x81,0xf9,0xd5,0xff,0x00,0x00,0x38,0x00,0x80,0x00,0x00]
diff --git
a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
index de908b95d94f9..83a647ad7c658 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
@@ -104,6 +104,39 @@
# GFX1250-REAL16: v_tanh_f16_e64_dpp v5.h, v128.h op_sel:[1,1]
quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x48,0x9f,0xd5,0xfa,0x00,0x00,0x00,0x80,0x1b,0x00,0xff]
# GFX1250-FAKE16: v_tanh_f16_e64_dpp v5, v128 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x9f,0xd5,0xfa,0x00,0x00,0x00,0x80,0x1b,0x00,0xff]
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0xe4,0x00,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0xe4,0x00,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1b,0x00,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1b,0x00,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x41,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_half_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x41,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x40,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x40,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x21,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x21,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x2f,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x2f,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x50,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x50,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x01,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x01,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x0f,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x0f,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x11,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x11,0x01,0xff]
+
+0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1f,0x01,0xff
+# GFX1250: v_prng_b32_e64_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ;
encoding: [0x05,0x00,0xcb,0xd5,0xfa,0x00,0x00,0x00,0x01,0x1f,0x01,0xff]
+
0xff,0x81,0xfb,0xd5,0xfa,0x00,0x00,0x38,0xff,0x6f,0x05,0x30
# GFX1250-REAL16: v_rsq_bf16_e64_dpp v255.l, -|v255.l| clamp div:2
row_xmask:15 row_mask:0x3 bank_mask:0x0 fi:1 ; encoding:
[0xff,0x81,0xfb,0xd5,0xfa,0x00,0x00,0x38,0xff,0x6f,0x05,0x30]
# GFX1250-FAKE16: v_rsq_bf16_e64_dpp v255, -|v255| clamp div:2 row_xmask:15
row_mask:0x3 bank_mask:0x0 fi:1 ; encoding:
[0xff,0x81,0xfb,0xd5,0xfa,0x00,0x00,0x38,0xff,0x6f,0x05,0x30]
diff --git
a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
index cfe7173c383b3..ef5ede4d1d453 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
@@ -1,6 +1,6 @@
# NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py
UTC_ARGS: --version 5
-# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250-REAL16 %s
-# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250-FAKE16 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-REAL16 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-FAKE16 %s
0xff,0x81,0x9e,0xd5,0xe9,0x00,0x00,0x38,0xff,0x00,0x00,0x00
# GFX1250: v_tanh_f32_e64_dpp v255, -|v255| clamp div:2 dpp8:[0,0,0,0,0,0,0,0]
; encoding: [0xff,0x81,0x9e,0xd5,0xe9,0x00,0x00,0x38,0xff,0x00,0x00,0x00]
@@ -34,6 +34,9 @@
# GFX1250-REAL16: v_tanh_f16_e64_dpp v5.h, v128.h op_sel:[1,1]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x48,0x9f,0xd5,0xe9,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
# GFX1250-FAKE16: v_tanh_f16_e64_dpp v5, v128 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x9f,0xd5,0xe9,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
+0x05,0x00,0xcb,0xd5,0xe9,0x00,0x00,0x00,0x01,0x77,0x39,0x05
+# GFX1250: v_prng_b32_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0xcb,0xd5,0xe9,0x00,0x00,0x00,0x01,0x77,0x39,0x05]
+
0xff,0x81,0xfb,0xd5,0xe9,0x00,0x00,0x38,0xff,0x00,0x00,0x00
# GFX1250-REAL16: v_rsq_bf16_e64_dpp v255.l, -|v255.l| clamp div:2
dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xff,0x81,0xfb,0xd5,0xe9,0x00,0x00,0x38,0xff,0x00,0x00,0x00]
# GFX1250-FAKE16: v_rsq_bf16_e64_dpp v255, -|v255| clamp div:2
dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xff,0x81,0xfb,0xd5,0xe9,0x00,0x00,0x38,0xff,0x00,0x00,0x00]
_______________________________________________
llvm-branch-commits mailing list
llvm-branch-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
