Skip to content

Commit 86976f2

Browse files
Acim-MaravicAcim-Maravic
committed
[LLVM] Add intrinsics for v_cvt_pk_norm_{i16,u16}_f16
Added builtin and intrinsic for v_cvt_pk_norm_i16_f16 and v_cvt_pk_norm_u16_f16
1 parent a947419 commit 86976f2

File tree

8 files changed

+287
-2
lines changed

8 files changed

+287
-2
lines changed

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -259,6 +259,9 @@ TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atom
259259
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts")
260260
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "vmem-to-lds-load-insts")
261261

262+
TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_norm_i16_f16, "V2sxx", "nc", "gfx9-insts")
263+
TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_norm_u16_f16, "V2Usxx", "nc", "gfx9-insts")
264+
262265
//===----------------------------------------------------------------------===//
263266
// Deep learning builtins.
264267
//===----------------------------------------------------------------------===//

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

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,8 @@
55
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
66
typedef unsigned int uint;
77
typedef unsigned long ulong;
8+
typedef short __attribute__((ext_vector_type(2))) short2;
9+
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
810

911
// CHECK-LABEL: @test_fmed3_f16
1012
// CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c)
@@ -26,3 +28,17 @@ void test_groupstaticsize(global uint* out)
2628
{
2729
*out = __builtin_amdgcn_groupstaticsize();
2830
}
31+
32+
// CHECK-LABEL: define dso_local void @test_cvt_pk_norm_i16_f16(
33+
// CHECK: call <2 x i16> @llvm.amdgcn.cvt.pk.norm.i16.f16(half %src0, half %src1)
34+
void test_cvt_pk_norm_i16_f16(global short2* out, half src0, half src1)
35+
{
36+
*out = __builtin_amdgcn_cvt_pk_norm_i16_f16(src0, src1);
37+
}
38+
39+
// CHECK-LABEL: define dso_local void @test_cvt_pk_norm_u16_f16(
40+
// CHECK: call <2 x i16> @llvm.amdgcn.cvt.pk.norm.u16.f16(half %src0, half %src1)
41+
void test_cvt_pk_norm_u16_f16(global ushort2* out, half src0, half src1)
42+
{
43+
*out = __builtin_amdgcn_cvt_pk_norm_u16_f16(src0, src1);
44+
}

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2644,6 +2644,18 @@ def int_amdgcn_global_load_lds : AMDGPUGlobalLoadLDS;
26442644
def int_amdgcn_pops_exiting_wave_id :
26452645
DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrHasSideEffects]>;
26462646

2647+
def int_amdgcn_cvt_pk_norm_i16_f16 :
2648+
ClangBuiltin<"__builtin_amdgcn_cvt_pk_norm_i16_f16">,
2649+
DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_half_ty, llvm_half_ty],
2650+
[IntrNoMem, IntrSpeculatable]
2651+
>;
2652+
2653+
def int_amdgcn_cvt_pk_norm_u16_f16 :
2654+
ClangBuiltin<"__builtin_amdgcn_cvt_pk_norm_u16_f16">,
2655+
DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_half_ty, llvm_half_ty],
2656+
[IntrNoMem, IntrSpeculatable]
2657+
>;
2658+
26472659
//===----------------------------------------------------------------------===//
26482660
// GFX10 Intrinsics
26492661
//===----------------------------------------------------------------------===//

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4669,6 +4669,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
46694669
case Intrinsic::amdgcn_swmmac_f32_16x16x32_fp8_bf8:
46704670
case Intrinsic::amdgcn_swmmac_f32_16x16x32_bf8_fp8:
46714671
case Intrinsic::amdgcn_swmmac_f32_16x16x32_bf8_bf8:
4672+
case Intrinsic::amdgcn_cvt_pk_norm_i16_f16:
4673+
case Intrinsic::amdgcn_cvt_pk_norm_u16_f16:
46724674
return getDefaultMappingVOP(MI);
46734675
case Intrinsic::amdgcn_log:
46744676
case Intrinsic::amdgcn_exp2:

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2856,6 +2856,7 @@ def VOP_V32BF16_V6I32_F32 : VOPProfile <[v32bf16, v6i32, f32, untyped]>;
28562856
def VOP_V6I32_V32F16_F32 : VOPProfile<[v6i32, v32f16, f32, untyped]>;
28572857
def VOP_V6I32_V32BF16_F32 : VOPProfile<[v6i32, v32bf16, f32, untyped]>;
28582858
def VOP_V6I32_V16F32_V16F32_F32 : VOPProfile<[v6i32, v16f32, v16f32, f32]>;
2859+
def VOP_V2I16_F16_F16 : VOPProfile<[v2i16, f16, f16, untyped]>;
28592860
def VOP_V2F16_I32_F32 : VOPProfile<[v2f16, i32, f32, untyped]>;
28602861
def VOP_V2I16_F32_F32_F32 : VOPProfile<[v2i16, f32, f32, f32]>;
28612862
def VOP_V2I16_V2F16_F32 : VOPProfile<[v2i16, v2f16, f32, untyped]>;

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

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -666,8 +666,11 @@ let isCommutable = 1 in {
666666
defm V_MAD_I32_I16 : VOP3Inst_t16 <"v_mad_i32_i16", VOP_I32_I16_I16_I32>;
667667
} // End isCommutable = 1
668668

669-
defm V_CVT_PKNORM_I16_F16 : VOP3Inst_t16 <"v_cvt_pknorm_i16_f16", VOP_B32_F16_F16>;
670-
defm V_CVT_PKNORM_U16_F16 : VOP3Inst_t16 <"v_cvt_pknorm_u16_f16", VOP_B32_F16_F16>;
669+
defm V_CVT_PKNORM_I16_F16 : VOP3Inst_t16 <"v_cvt_pknorm_i16_f16", VOP_V2I16_F16_F16, int_amdgcn_cvt_pk_norm_i16_f16>;
670+
defm V_CVT_PKNORM_U16_F16 : VOP3Inst_t16 <"v_cvt_pknorm_u16_f16", VOP_V2I16_F16_F16, int_amdgcn_cvt_pk_norm_u16_f16>;
671+
// multiclass VOP3Inst<string OpName, VOPProfile P, SDPatternOperator node = null_frag> {
672+
673+
// multiclass VOP3Inst_t16<string OpName, VOPProfile P, SDPatternOperator node = null_frag, SDPatternOperator node_t16 = node>
671674

672675
defm V_PACK_B32_F16 : VOP3Inst_t16 <"v_pack_b32_f16", VOP_B32_F16_F16>;
673676

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.pk.norm.i16.f16.ll‎

Lines changed: 124 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,124 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GFX9-SDAG %s
3+
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GFX9-GISEL %s
4+
5+
define amdgpu_kernel void @cvt_pk_norm_i16_f16_sgpr(ptr addrspace(1) %out, half %a, half %b) {
6+
; GFX9-SDAG-LABEL: cvt_pk_norm_i16_f16_sgpr:
7+
; GFX9-SDAG: ; %bb.0:
8+
; GFX9-SDAG-NEXT: s_load_dword s2, s[4:5], 0x2c
9+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
10+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v0, 0
11+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
12+
; GFX9-SDAG-NEXT: s_lshr_b32 s3, s2, 16
13+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v1, s3
14+
; GFX9-SDAG-NEXT: v_cvt_pknorm_i16_f16 v1, s2, v1
15+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
16+
; GFX9-SDAG-NEXT: s_endpgm
17+
;
18+
; GFX9-GISEL-LABEL: cvt_pk_norm_i16_f16_sgpr:
19+
; GFX9-GISEL: ; %bb.0:
20+
; GFX9-GISEL-NEXT: s_load_dword s2, s[4:5], 0x2c
21+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
22+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v1, 0
23+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
24+
; GFX9-GISEL-NEXT: s_lshr_b32 s3, s2, 16
25+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v0, s3
26+
; GFX9-GISEL-NEXT: v_cvt_pknorm_i16_f16 v0, s2, v0
27+
; GFX9-GISEL-NEXT: global_store_dword v1, v0, s[0:1]
28+
; GFX9-GISEL-NEXT: s_endpgm
29+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.i16.f16(half %a, half %b)
30+
%res = bitcast <2 x i16> %r to i32
31+
store i32 %res, ptr addrspace(1) %out
32+
ret void
33+
}
34+
35+
define amdgpu_kernel void @cvt_pk_norm_i16_f16_const(ptr addrspace(1) %out) {
36+
; GFX9-SDAG-LABEL: cvt_pk_norm_i16_f16_const:
37+
; GFX9-SDAG: ; %bb.0:
38+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
39+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v0, 0
40+
; GFX9-SDAG-NEXT: v_cvt_pknorm_i16_f16 v1, 1.0, 2.0
41+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
42+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
43+
; GFX9-SDAG-NEXT: s_endpgm
44+
;
45+
; GFX9-GISEL-LABEL: cvt_pk_norm_i16_f16_const:
46+
; GFX9-GISEL: ; %bb.0:
47+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
48+
; GFX9-GISEL-NEXT: v_cvt_pknorm_i16_f16 v0, 1.0, 2.0
49+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v1, 0
50+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
51+
; GFX9-GISEL-NEXT: global_store_dword v1, v0, s[0:1]
52+
; GFX9-GISEL-NEXT: s_endpgm
53+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.i16.f16(half 1.0, half 2.0)
54+
%res = bitcast <2 x i16> %r to i32
55+
store i32 %res, ptr addrspace(1) %out
56+
ret void
57+
}
58+
59+
define amdgpu_kernel void @cvt_pk_norm_i16_f16_undef(ptr addrspace(1) %out) {
60+
; GFX9-SDAG-LABEL: cvt_pk_norm_i16_f16_undef:
61+
; GFX9-SDAG: ; %bb.0:
62+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
63+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v0, 0
64+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
65+
; GFX9-SDAG-NEXT: v_cvt_pknorm_i16_f16 v1, s0, v0
66+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
67+
; GFX9-SDAG-NEXT: s_endpgm
68+
;
69+
; GFX9-GISEL-LABEL: cvt_pk_norm_i16_f16_undef:
70+
; GFX9-GISEL: ; %bb.0:
71+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
72+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v1, 0
73+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
74+
; GFX9-GISEL-NEXT: v_cvt_pknorm_i16_f16 v0, s0, s0
75+
; GFX9-GISEL-NEXT: global_store_dword v1, v0, s[0:1]
76+
; GFX9-GISEL-NEXT: s_endpgm
77+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.i16.f16(half undef, half undef)
78+
%res = bitcast <2 x i16> %r to i32
79+
store i32 %res, ptr addrspace(1) %out
80+
ret void
81+
}
82+
83+
define amdgpu_kernel void @cvt_pk_norm_i16_f16_vgpr(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr) {
84+
; GFX9-SDAG-LABEL: cvt_pk_norm_i16_f16_vgpr:
85+
; GFX9-SDAG: ; %bb.0:
86+
; GFX9-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
87+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
88+
; GFX9-SDAG-NEXT: v_lshlrev_b32_e32 v1, 1, v0
89+
; GFX9-SDAG-NEXT: v_lshlrev_b32_e32 v0, 2, v0
90+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
91+
; GFX9-SDAG-NEXT: global_load_ushort v2, v1, s[2:3] glc
92+
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0)
93+
; GFX9-SDAG-NEXT: global_load_ushort v3, v1, s[6:7] glc
94+
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0)
95+
; GFX9-SDAG-NEXT: v_cvt_pknorm_i16_f16 v1, v2, v3
96+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
97+
; GFX9-SDAG-NEXT: s_endpgm
98+
;
99+
; GFX9-GISEL-LABEL: cvt_pk_norm_i16_f16_vgpr:
100+
; GFX9-GISEL: ; %bb.0:
101+
; GFX9-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
102+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
103+
; GFX9-GISEL-NEXT: v_lshlrev_b32_e32 v1, 1, v0
104+
; GFX9-GISEL-NEXT: v_lshlrev_b32_e32 v0, 2, v0
105+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
106+
; GFX9-GISEL-NEXT: global_load_ushort v2, v1, s[2:3] glc
107+
; GFX9-GISEL-NEXT: s_waitcnt vmcnt(0)
108+
; GFX9-GISEL-NEXT: global_load_ushort v3, v1, s[6:7] glc
109+
; GFX9-GISEL-NEXT: s_waitcnt vmcnt(0)
110+
; GFX9-GISEL-NEXT: v_cvt_pknorm_i16_f16 v1, v2, v3
111+
; GFX9-GISEL-NEXT: global_store_dword v0, v1, s[0:1]
112+
; GFX9-GISEL-NEXT: s_endpgm
113+
%tid = call i32 @llvm.amdgcn.workitem.id.x()
114+
%tid.ext = sext i32 %tid to i64
115+
%a.gep = getelementptr inbounds half, ptr addrspace(1) %a.ptr, i64 %tid.ext
116+
%b.gep = getelementptr inbounds half, ptr addrspace(1) %b.ptr, i64 %tid.ext
117+
%out.gep = getelementptr inbounds i32, ptr addrspace(1) %out, i64 %tid.ext
118+
%a = load volatile half, ptr addrspace(1) %a.gep
119+
%b = load volatile half, ptr addrspace(1) %b.gep
120+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.i16.f16(half %a, half %b)
121+
%res = bitcast <2 x i16> %r to i32
122+
store i32 %res, ptr addrspace(1) %out.gep
123+
ret void
124+
}

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.pk.norm.u16.f16.ll‎

Lines changed: 124 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,124 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GFX9-SDAG %s
3+
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GFX9-GISEL %s
4+
5+
define amdgpu_kernel void @cvt_pk_norm_u16_f16_sgpr(ptr addrspace(1) %out, half %a, half %b) {
6+
; GFX9-SDAG-LABEL: cvt_pk_norm_u16_f16_sgpr:
7+
; GFX9-SDAG: ; %bb.0:
8+
; GFX9-SDAG-NEXT: s_load_dword s2, s[4:5], 0x2c
9+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
10+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v0, 0
11+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
12+
; GFX9-SDAG-NEXT: s_lshr_b32 s3, s2, 16
13+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v1, s3
14+
; GFX9-SDAG-NEXT: v_cvt_pknorm_u16_f16 v1, s2, v1
15+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
16+
; GFX9-SDAG-NEXT: s_endpgm
17+
;
18+
; GFX9-GISEL-LABEL: cvt_pk_norm_u16_f16_sgpr:
19+
; GFX9-GISEL: ; %bb.0:
20+
; GFX9-GISEL-NEXT: s_load_dword s2, s[4:5], 0x2c
21+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
22+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v1, 0
23+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
24+
; GFX9-GISEL-NEXT: s_lshr_b32 s3, s2, 16
25+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v0, s3
26+
; GFX9-GISEL-NEXT: v_cvt_pknorm_u16_f16 v0, s2, v0
27+
; GFX9-GISEL-NEXT: global_store_dword v1, v0, s[0:1]
28+
; GFX9-GISEL-NEXT: s_endpgm
29+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.u16.f16(half %a, half %b)
30+
%res = bitcast <2 x i16> %r to i32
31+
store i32 %res, ptr addrspace(1) %out
32+
ret void
33+
}
34+
35+
define amdgpu_kernel void @cvt_pk_norm_u16_f16_const(ptr addrspace(1) %out) {
36+
; GFX9-SDAG-LABEL: cvt_pk_norm_u16_f16_const:
37+
; GFX9-SDAG: ; %bb.0:
38+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
39+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v0, 0
40+
; GFX9-SDAG-NEXT: v_cvt_pknorm_u16_f16 v1, 1.0, 2.0
41+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
42+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
43+
; GFX9-SDAG-NEXT: s_endpgm
44+
;
45+
; GFX9-GISEL-LABEL: cvt_pk_norm_u16_f16_const:
46+
; GFX9-GISEL: ; %bb.0:
47+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
48+
; GFX9-GISEL-NEXT: v_cvt_pknorm_u16_f16 v0, 1.0, 2.0
49+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v1, 0
50+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
51+
; GFX9-GISEL-NEXT: global_store_dword v1, v0, s[0:1]
52+
; GFX9-GISEL-NEXT: s_endpgm
53+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.u16.f16(half 1.0, half 2.0)
54+
%res = bitcast <2 x i16> %r to i32
55+
store i32 %res, ptr addrspace(1) %out
56+
ret void
57+
}
58+
59+
define amdgpu_kernel void @cvt_pk_norm_u16_f16_undef(ptr addrspace(1) %out) {
60+
; GFX9-SDAG-LABEL: cvt_pk_norm_u16_f16_undef:
61+
; GFX9-SDAG: ; %bb.0:
62+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
63+
; GFX9-SDAG-NEXT: v_mov_b32_e32 v0, 0
64+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
65+
; GFX9-SDAG-NEXT: v_cvt_pknorm_u16_f16 v1, s0, v0
66+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
67+
; GFX9-SDAG-NEXT: s_endpgm
68+
;
69+
; GFX9-GISEL-LABEL: cvt_pk_norm_u16_f16_undef:
70+
; GFX9-GISEL: ; %bb.0:
71+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
72+
; GFX9-GISEL-NEXT: v_mov_b32_e32 v1, 0
73+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
74+
; GFX9-GISEL-NEXT: v_cvt_pknorm_u16_f16 v0, s0, s0
75+
; GFX9-GISEL-NEXT: global_store_dword v1, v0, s[0:1]
76+
; GFX9-GISEL-NEXT: s_endpgm
77+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.u16.f16(half undef, half undef)
78+
%res = bitcast <2 x i16> %r to i32
79+
store i32 %res, ptr addrspace(1) %out
80+
ret void
81+
}
82+
83+
define amdgpu_kernel void @cvt_pk_norm_u16_f16_vgpr(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr) {
84+
; GFX9-SDAG-LABEL: cvt_pk_norm_u16_f16_vgpr:
85+
; GFX9-SDAG: ; %bb.0:
86+
; GFX9-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
87+
; GFX9-SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
88+
; GFX9-SDAG-NEXT: v_lshlrev_b32_e32 v1, 1, v0
89+
; GFX9-SDAG-NEXT: v_lshlrev_b32_e32 v0, 2, v0
90+
; GFX9-SDAG-NEXT: s_waitcnt lgkmcnt(0)
91+
; GFX9-SDAG-NEXT: global_load_ushort v2, v1, s[2:3] glc
92+
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0)
93+
; GFX9-SDAG-NEXT: global_load_ushort v3, v1, s[6:7] glc
94+
; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0)
95+
; GFX9-SDAG-NEXT: v_cvt_pknorm_u16_f16 v1, v2, v3
96+
; GFX9-SDAG-NEXT: global_store_dword v0, v1, s[0:1]
97+
; GFX9-SDAG-NEXT: s_endpgm
98+
;
99+
; GFX9-GISEL-LABEL: cvt_pk_norm_u16_f16_vgpr:
100+
; GFX9-GISEL: ; %bb.0:
101+
; GFX9-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
102+
; GFX9-GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x34
103+
; GFX9-GISEL-NEXT: v_lshlrev_b32_e32 v1, 1, v0
104+
; GFX9-GISEL-NEXT: v_lshlrev_b32_e32 v0, 2, v0
105+
; GFX9-GISEL-NEXT: s_waitcnt lgkmcnt(0)
106+
; GFX9-GISEL-NEXT: global_load_ushort v2, v1, s[2:3] glc
107+
; GFX9-GISEL-NEXT: s_waitcnt vmcnt(0)
108+
; GFX9-GISEL-NEXT: global_load_ushort v3, v1, s[6:7] glc
109+
; GFX9-GISEL-NEXT: s_waitcnt vmcnt(0)
110+
; GFX9-GISEL-NEXT: v_cvt_pknorm_u16_f16 v1, v2, v3
111+
; GFX9-GISEL-NEXT: global_store_dword v0, v1, s[0:1]
112+
; GFX9-GISEL-NEXT: s_endpgm
113+
%tid = call i32 @llvm.amdgcn.workitem.id.x()
114+
%tid.ext = sext i32 %tid to i64
115+
%a.gep = getelementptr inbounds half, ptr addrspace(1) %a.ptr, i64 %tid.ext
116+
%b.gep = getelementptr inbounds half, ptr addrspace(1) %b.ptr, i64 %tid.ext
117+
%out.gep = getelementptr inbounds i32, ptr addrspace(1) %out, i64 %tid.ext
118+
%a = load volatile half, ptr addrspace(1) %a.gep
119+
%b = load volatile half, ptr addrspace(1) %b.gep
120+
%r = call <2 x i16> @llvm.amdgcn.cvt.pk.norm.u16.f16(half %a, half %b)
121+
%res = bitcast <2 x i16> %r to i32
122+
store i32 %res, ptr addrspace(1) %out.gep
123+
ret void
124+
}

0 commit comments

Comments
 (0)