Skip to content

Commit d613fe2

Browse files
pravinjagtaparsenm
pravinjagtap
authored and
arsenm
committed
AMDGPU: Add v_prng_b32 instruction for gfx950
Rand num instruction for stochastic rounding.
1 parent 54f5e02 commit d613fe2

File tree

16 files changed

+252
-3
lines changed

16 files changed

+252
-3
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -522,5 +522,7 @@ TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64, "V4fiV2iV4fs",
522522
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64, "V4fiV2iV4fs", "nc", "gfx12-insts,wavefrontsize64")
523523
TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64, "V4fiV2iV4fs", "nc", "gfx12-insts,wavefrontsize64")
524524

525+
TARGET_BUILTIN(__builtin_amdgcn_prng_b32, "UiUi", "nc", "prng-inst")
526+
525527
#undef BUILTIN
526528
#undef TARGET_BUILTIN

‎clang/test/CodeGenOpenCL/amdgpu-features.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -89,7 +89,7 @@
8989
// GFX941: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
9090
// GFX942: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
9191
// GFX9_4_Generic: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
92-
// GFX950: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
92+
// GFX950: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+mai-insts,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
9393
// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
9494
// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
9595
// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950-err.cl

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
1+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -emit-llvm \
2+
// RUN: -verify -o - %s
3+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -emit-llvm \
4+
// RUN: -verify -o - %s
5+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx940 -emit-llvm \
6+
// RUN: -verify -o - %s
7+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1200 -emit-llvm \
8+
// RUN: -verify -o - %s
9+
10+
11+
// REQUIRES: amdgpu-registered-target
12+
13+
typedef unsigned int uint;
14+
void test_prng_b32(global uint* out, uint a) {
15+
*out = __builtin_amdgcn_prng_b32(a); // expected-error{{'__builtin_amdgcn_prng_b32' needs target feature prng-inst}}
16+
}

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

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
2+
// RUN: %clang_cc1 -cl-std=CL1.2 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx950 -emit-llvm -o - %s | FileCheck %s
3+
// REQUIRES: amdgpu-registered-target
4+
5+
typedef unsigned int uint;
6+
7+
// CHECK-LABEL: @test_prng_b32(
8+
// CHECK-NEXT: entry:
9+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
10+
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
11+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
12+
// CHECK-NEXT: store i32 [[A:%.*]], ptr addrspace(5) [[A_ADDR]], align 4
13+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[A_ADDR]], align 4
14+
// CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.prng.b32(i32 [[TMP0]])
15+
// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
16+
// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(1) [[TMP2]], align 4
17+
// CHECK-NEXT: ret void
18+
//
19+
void test_prng_b32(global uint* out, uint a) {
20+
*out = __builtin_amdgcn_prng_b32(a);
21+
}

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -594,6 +594,10 @@ def int_amdgcn_ds_ordered_swap : AMDGPUDSOrderedIntrinsic;
594594
def int_amdgcn_ds_append : AMDGPUDSAppendConsumedIntrinsic;
595595
def int_amdgcn_ds_consume : AMDGPUDSAppendConsumedIntrinsic;
596596

597+
def int_amdgcn_prng_b32 : DefaultAttrsIntrinsic<
598+
[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]
599+
>, ClangBuiltin<"__builtin_amdgcn_prng_b32">;
600+
597601
} // TargetPrefix = "amdgcn"
598602

599603
// New-style image intrinsics

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

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -978,6 +978,12 @@ def FeatureVmemWriteVgprInOrder : SubtargetFeature<"vmem-write-vgpr-in-order",
978978
"VMEM instructions of the same type write VGPR results in order"
979979
>;
980980

981+
def FeaturePrngInst : SubtargetFeature<"prng-inst",
982+
"HasPrngInst",
983+
"true",
984+
"Has v_prng_b32 instruction"
985+
>;
986+
981987
//===------------------------------------------------------------===//
982988
// Subtarget Features (options and debugging)
983989
//===------------------------------------------------------------===//
@@ -1495,7 +1501,8 @@ def FeatureISAVersion9_5_Common : FeatureSet<
14951501
FeatureFP8ConversionInsts,
14961502
FeatureCvtFP8VOP1Bug,
14971503
FeatureGFX950Insts,
1498-
FeatureAddressableLocalMemorySize163840
1504+
FeatureAddressableLocalMemorySize163840,
1505+
FeaturePrngInst
14991506
])>;
15001507

15011508
def FeatureISAVersion9_4_0 : FeatureSet<
@@ -2344,6 +2351,9 @@ def HasSALUFloatInsts : Predicate<"Subtarget->hasSALUFloatInsts()">,
23442351
def HasPseudoScalarTrans : Predicate<"Subtarget->hasPseudoScalarTrans()">,
23452352
AssemblerPredicate<(all_of FeaturePseudoScalarTrans)>;
23462353

2354+
def HasPrngInst : Predicate<"Subtarget->hasPrngInst()">,
2355+
AssemblerPredicate<(all_of FeaturePrngInst)>;
2356+
23472357
def HasGDS : Predicate<"Subtarget->hasGDS()">;
23482358

23492359
def HasGWS : Predicate<"Subtarget->hasGWS()">;

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1253,6 +1253,12 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
12531253

12541254
break;
12551255
}
1256+
case Intrinsic::amdgcn_prng_b32: {
1257+
auto *Src = II.getArgOperand(0);
1258+
if (isa<UndefValue>(Src)) {
1259+
return IC.replaceInstUsesWith(II, Src);
1260+
}
1261+
}
12561262
}
12571263
if (const AMDGPU::ImageDimIntrinsicInfo *ImageDimIntr =
12581264
AMDGPU::getImageDimIntrinsicInfo(II.getIntrinsicID())) {

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4515,6 +4515,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
45154515
case Intrinsic::amdgcn_cvt_pk_u8_f32:
45164516
case Intrinsic::amdgcn_alignbyte:
45174517
case Intrinsic::amdgcn_perm:
4518+
case Intrinsic::amdgcn_prng_b32:
45184519
case Intrinsic::amdgcn_fdot2:
45194520
case Intrinsic::amdgcn_sdot2:
45204521
case Intrinsic::amdgcn_udot2:

‎llvm/lib/Target/AMDGPU/GCNSubtarget.h

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -220,7 +220,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
220220
bool HasSALUFloatInsts = false;
221221
bool HasPseudoScalarTrans = false;
222222
bool HasRestrictedSOffset = false;
223-
223+
bool HasPrngInst = false;
224224
bool HasVcmpxPermlaneHazard = false;
225225
bool HasVMEMtoScalarWriteHazard = false;
226226
bool HasSMEMtoVectorWriteHazard = false;
@@ -1321,6 +1321,8 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
13211321
/// instruction.
13221322
unsigned maxHardClauseLength() const { return MaxHardClauseLength; }
13231323

1324+
bool hasPrngInst() const { return HasPrngInst; }
1325+
13241326
/// Return the maximum number of waves per SIMD for kernels using \p SGPRs
13251327
/// SGPRs
13261328
unsigned getOccupancyWithNumSGPRs(unsigned SGPRs) const;

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -761,6 +761,9 @@ let SubtargetPredicate = isGFX11Plus in {
761761
defm V_CVT_U32_U16 : VOP1Inst_t16<"v_cvt_u32_u16", VOP_I32_I16>;
762762
} // End SubtargetPredicate = isGFX11Plus
763763

764+
let SubtargetPredicate = HasPrngInst in
765+
defm V_PRNG_B32 : VOP1Inst <"v_prng_b32", VOP_I32_I32, int_amdgcn_prng_b32>;
766+
764767
foreach vt = Reg32Types.types in {
765768
def : GCNPat<(int_amdgcn_permlane64 (vt VRegSrc_32:$src0)),
766769
(vt (V_PERMLANE64_B32 (vt VRegSrc_32:$src0)))
@@ -1516,6 +1519,8 @@ defm V_CVT_F32_BF8 : VOP1_Real_NoDstSel_SDWA_gfx9<0x55>;
15161519
defm V_CVT_PK_F32_FP8 : VOP1_Real_NoDstSel_SDWA_gfx9<0x56>;
15171520
defm V_CVT_PK_F32_BF8 : VOP1_Real_NoDstSel_SDWA_gfx9<0x57>;
15181521

1522+
defm V_PRNG_B32 : VOP1_Real_gfx9 <0x58>;
1523+
15191524
class MovDPP8Pattern<Predicate Pred, Instruction Inst, ValueType vt> : GCNPat <
15201525
(vt (int_amdgcn_mov_dpp8 vt:$src, timm:$dpp8)),
15211526
(Inst VGPR_32:$src, VGPR_32:$src, (as_i32timm $dpp8), (i32 DPP8Mode.FI_0))> {

‎llvm/lib/TargetParser/TargetParser.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -470,6 +470,7 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const Triple &T,
470470
Features["gws"] = true;
471471
break;
472472
case GK_GFX950:
473+
Features["prng-inst"] = true;
473474
Features["gfx950-insts"] = true;
474475
[[fallthrough]];
475476
case GK_GFX942:

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

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
; RUN: llc -global-isel=0 -march=amdgcn -mcpu=gfx950 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN %s
2+
; RUN: llc -global-isel=1 -march=amdgcn -mcpu=gfx950 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
3+
4+
declare i32 @llvm.amdgcn.prng.b32(i32) #0
5+
6+
; GCN-LABEL: {{^}}prng_b32:
7+
; GCN: v_prng_b32_e32 {{v[0-9]+}}, {{s[0-9]+}}
8+
define amdgpu_kernel void @prng_b32(ptr addrspace(1) %out, i32 %src) #1 {
9+
%prng = call i32 @llvm.amdgcn.prng.b32(i32 %src) #0
10+
store i32 %prng, ptr addrspace(1) %out, align 4
11+
ret void
12+
}
13+
14+
; GCN-LABEL: {{^}}prng_b32_constant_4
15+
; GCN: v_prng_b32_e32 {{v[0-9]+}}, 4
16+
define amdgpu_kernel void @prng_b32_constant_4(ptr addrspace(1) %out) #1 {
17+
%prng = call i32 @llvm.amdgcn.prng.b32(i32 4) #0
18+
store i32 %prng, ptr addrspace(1) %out, align 4
19+
ret void
20+
}
21+
22+
; GCN-LABEL: {{^}}prng_b32_constant_100
23+
; GCN: v_prng_b32_e32 {{v[0-9]+}}, 0x64
24+
define amdgpu_kernel void @prng_b32_constant_100(ptr addrspace(1) %out) #1 {
25+
%prng = call i32 @llvm.amdgcn.prng.b32(i32 100) #0
26+
store i32 %prng, ptr addrspace(1) %out, align 4
27+
ret void
28+
}
29+
30+
31+
attributes #0 = { nounwind readnone }
32+
attributes #1 = { nounwind }

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

Lines changed: 57 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,57 @@
1+
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx950 -show-encoding %s | FileCheck --check-prefix=GFX950 %s
2+
3+
v_prng_b32 v5, v1
4+
// GFX950: v_prng_b32_e32 v5, v1 ; encoding: [0x01,0xb1,0x0a,0x7e]
5+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
6+
7+
v_prng_b32 v5, v255
8+
// GFX950: v_prng_b32_e32 v5, v255 ; encoding: [0xff,0xb1,0x0a,0x7e]
9+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
10+
11+
v_prng_b32 v5, s1
12+
// GFX950: v_prng_b32_e32 v5, s1 ; encoding: [0x01,0xb0,0x0a,0x7e]
13+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
14+
15+
v_prng_b32 v5, s101
16+
// GFX950: v_prng_b32_e32 v5, s101 ; encoding: [0x65,0xb0,0x0a,0x7e]
17+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
18+
19+
v_prng_b32 v5, vcc_lo
20+
// GFX950: v_prng_b32_e32 v5, vcc_lo ; encoding: [0x6a,0xb0,0x0a,0x7e]
21+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
22+
23+
v_prng_b32 v5, vcc_hi
24+
// GFX950: v_prng_b32_e32 v5, vcc_hi ; encoding: [0x6b,0xb0,0x0a,0x7e]
25+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
26+
27+
v_prng_b32 v5, ttmp15
28+
// GFX950: v_prng_b32_e32 v5, ttmp15 ; encoding: [0x7b,0xb0,0x0a,0x7e]
29+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
30+
31+
v_prng_b32 v5, m0
32+
// GFX950: v_prng_b32_e32 v5, m0 ; encoding: [0x7c,0xb0,0x0a,0x7e]
33+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
34+
35+
v_prng_b32 v5, exec_lo
36+
// GFX950: v_prng_b32_e32 v5, exec_lo ; encoding: [0x7e,0xb0,0x0a,0x7e]
37+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
38+
39+
v_prng_b32 v5, exec_hi
40+
// GFX950: v_prng_b32_e32 v5, exec_hi ; encoding: [0x7f,0xb0,0x0a,0x7e]
41+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
42+
43+
v_prng_b32 v5, -1
44+
// GFX950: v_prng_b32_e32 v5, -1 ; encoding: [0xc1,0xb0,0x0a,0x7e]
45+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
46+
47+
v_prng_b32 v5, 0.5
48+
// GFX950: v_prng_b32_e32 v5, 0.5 ; encoding: [0xf0,0xb0,0x0a,0x7e]
49+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
50+
51+
v_prng_b32 v5, src_scc
52+
// GFX950: v_prng_b32_e32 v5, src_scc ; encoding: [0xfd,0xb0,0x0a,0x7e]
53+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU
54+
55+
v_prng_b32 v255, 0xaf123456
56+
// GFX950: v_prng_b32_e32 v255, 0xaf123456 ; encoding: [0xff,0xb0,0xfe,0x7f,0x56,0x34,0x12,0xaf]
57+
// GFX940-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU

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

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx950 -show-encoding %s | FileCheck --check-prefixes=GFX950 %s
2+
3+
v_prng_b32 v5, v1 quad_perm:[3,2,1,0]
4+
// GFX950: v_prng_b32_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x1b,0x00,0xff]
5+
6+
v_prng_b32 v5, v1 quad_perm:[0,1,2,3]
7+
// GFX950: v_prng_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0xe4,0x00,0xff]
8+
9+
v_prng_b32 v5, v1 row_mirror
10+
// GFX950: v_prng_b32_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x40,0x01,0xff]
11+
12+
v_prng_b32 v5, v1 row_half_mirror
13+
// GFX950: v_prng_b32_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x41,0x01,0xff]
14+
15+
v_prng_b32 v5, v1 row_shl:1
16+
// GFX950: v_prng_b32_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x01,0x01,0xff]
17+
18+
v_prng_b32 v5, v1 row_shl:15
19+
// GFX950: v_prng_b32_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x0f,0x01,0xff]
20+
21+
v_prng_b32 v5, v1 row_shr:1
22+
// GFX950: v_prng_b32_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x11,0x01,0xff]
23+
24+
v_prng_b32 v5, v1 row_shr:15
25+
// GFX950: v_prng_b32_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x1f,0x01,0xff]
26+
27+
v_prng_b32 v5, v1 row_ror:1
28+
// GFX950: v_prng_b32_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x21,0x01,0xff]
29+
30+
v_prng_b32 v5, v1 row_ror:15
31+
// GFX950: v_prng_b32_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xb0,0x0a,0x7e,0x01,0x2f,0x01,0xff]

‎llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_vop1.txt

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
# RUN: llvm-mc -arch=amdgcn -mcpu=gfx950 -disassemble -show-encoding < %s | FileCheck -check-prefix=GFX950 %s
2+
3+
# GFX950: v_prng_b32_e32 v5, v1 ; encoding: [0x01,0xb1,0x0a,0x7e]
4+
0x01,0xb1,0x0a,0x7e
5+
6+
# GFX950: v_prng_b32_e32 v5, v255 ; encoding: [0xff,0xb1,0x0a,0x7e]
7+
0xff,0xb1,0x0a,0x7e
8+
9+
# GFX950: v_prng_b32_e32 v5, s1 ; encoding: [0x01,0xb0,0x0a,0x7e]
10+
0x01,0xb0,0x0a,0x7e
11+
12+
# GFX950: v_prng_b32_e32 v5, s101 ; encoding: [0x65,0xb0,0x0a,0x7e]
13+
0x65,0xb0,0x0a,0x7e
14+
15+
# GFX950: v_prng_b32_e32 v5, vcc_lo ; encoding: [0x6a,0xb0,0x0a,0x7e]
16+
0x6a,0xb0,0x0a,0x7e
17+
18+
# GFX950: v_prng_b32_e32 v5, vcc_hi ; encoding: [0x6b,0xb0,0x0a,0x7e]
19+
0x6b,0xb0,0x0a,0x7e
20+
21+
# GFX950: v_prng_b32_e32 v5, ttmp15 ; encoding: [0x7b,0xb0,0x0a,0x7e]
22+
0x7b,0xb0,0x0a,0x7e
23+
24+
# GFX950: v_prng_b32_e32 v5, m0 ; encoding: [0x7c,0xb0,0x0a,0x7e]
25+
0x7c,0xb0,0x0a,0x7e
26+
27+
# GFX950: v_prng_b32_e32 v5, exec_lo ; encoding: [0x7e,0xb0,0x0a,0x7e]
28+
0x7e,0xb0,0x0a,0x7e
29+
30+
# GFX950: v_prng_b32_e32 v5, exec_hi ; encoding: [0x7f,0xb0,0x0a,0x7e]
31+
0x7f,0xb0,0x0a,0x7e
32+
33+
# GFX950: v_prng_b32_e32 v5, -1 ; encoding: [0xc1,0xb0,0x0a,0x7e]
34+
0xc1,0xb0,0x0a,0x7e
35+
36+
# GFX950: v_prng_b32_e32 v5, 0.5 ; encoding: [0xf0,0xb0,0x0a,0x7e]
37+
0xf0,0xb0,0x0a,0x7e
38+
39+
# GFX950: v_prng_b32_e32 v5, src_scc ; encoding: [0xfd,0xb0,0x0a,0x7e]
40+
0xfd,0xb0,0x0a,0x7e
41+
42+
# GFX950: v_prng_b32_e32 v255, 0xaf123456 ; encoding: [0xff,0xb0,0xfe,0x7f,0x56,0x34,0x12,0xaf]
43+
0xff,0xb0,0xfe,0x7f,0x56,0x34,0x12,0xaf

‎llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6547,3 +6547,21 @@ define half @test_constant_fold_exp2_f16_neg_denorm() {
65476547
%val = call half @llvm.amdgcn.exp2.f16(half 0xH83ff)
65486548
ret half %val
65496549
}
6550+
6551+
; --------------------------------------------------------------------
6552+
; llvm.amdgcn.prng
6553+
; --------------------------------------------------------------------
6554+
declare i32 @llvm.amdgcn.prng.b32(i32)
6555+
define i32 @prng_undef_i32() {
6556+
; CHECK-LABEL: @prng_undef_i32(
6557+
; CHECK-NEXT: ret i32 undef
6558+
%prng = call i32 @llvm.amdgcn.prng.b32(i32 undef)
6559+
ret i32 %prng
6560+
}
6561+
6562+
define i32 @prng_poison_i32() {
6563+
; CHECK-LABEL: @prng_poison_i32(
6564+
; CHECK-NEXT: ret i32 poison
6565+
%prng = call i32 @llvm.amdgcn.prng.b32(i32 poison)
6566+
ret i32 %prng
6567+
}

0 commit comments

Comments
 (0)