llvm.org GIT mirror llvm / 224ee47
AMDGPU/SI: Add llvm.amdgcn.mov.dpp intrinsic This intrinsic will be used to expose dpp functionality to higher-level languages. It will map to the dpp version of v_mov_b32. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@260792 91177308-0d34-0410-b5e6-96231b3b80d8 Tom Stellard 4 years ago
3 changed file(s) with 30 addition(s) and 0 deletion(s). Raw diff Collapse all Expand all
202202 // VI Intrinsics
203203 //===----------------------------------------------------------------------===//
204204
205 // llvm.amdgcn.mov.dpp.i32
206 def int_amdgcn_mov_dpp :
207 Intrinsic<[llvm_anyint_ty],
208 [LLVMMatchType<0>, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty,
209 llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
210
205211 def int_amdgcn_s_dcache_wb :
206212 GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
207213 Intrinsic<[], [], []>;
108108 (S_BUFFER_LOAD_DWORD_IMM $sbase, (as_i32imm $offset))
109109 >;
110110
111 //===----------------------------------------------------------------------===//
112 // DPP Paterns
113 //===----------------------------------------------------------------------===//
114
115 def : Pat <
116 (int_amdgcn_mov_dpp i32:$src, imm:$dpp_ctrl, imm:$bound_ctrl,
117 imm:$bank_mask, imm:$row_mask),
118 (V_MOV_B32_dpp $src, (as_i32imm $dpp_ctrl), (as_i1imm $bound_ctrl),
119 (as_i32imm $bank_mask), (as_i32imm $row_mask))
120 >;
121
111122 } // End Predicates = [isVI]
0 ; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs -show-mc-encoding < %s | FileCheck -check-prefix=VI %s
1
2 ; VI-LABEL: {{^}}dpp_test:
3 ; VI: v_mov_b32 v0, v0, 1, -1, 1, 1 ; encoding: [0xfa,0x02,0x00,0x7e,0x00,0x01,0x08,0x11]
4 define void @dpp_test(i32 addrspace(1)* %out, i32 %in) {
5 %tmp0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %in, i32 1, i1 1, i32 1, i32 1) #0
6 store i32 %tmp0, i32 addrspace(1)* %out
7 ret void
8 }
9
10 declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32, i1, i32, i32) #0
11
12 attributes #0 = { nounwind readnone convergent }