llvm.org GIT mirror llvm / 5bbcde0
AMDGPU/SI: Implement DS_PERMUTE/DS_BPERMUTE Instruction Definitions and Intrinsics Summary: This patch impleemnts DS_PERMUTE/DS_BPERMUTE instruction definitions and intrinsics, which are new since VI. Reviewers: tstellarAMD, arsenm Subscribers: llvm-commits, arsenm Differential Revision: http://reviews.llvm.org/D17614 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@262356 91177308-0d34-0410-b5e6-96231b3b80d8 Changpeng Fang 4 years ago
6 changed file(s) with 68 addition(s) and 1 deletion(s). Raw diff Collapse all Expand all
257257 def int_amdgcn_s_memrealtime :
258258 GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
259259 Intrinsic<[llvm_i64_ty], [], []>;
260 }
260
261 // llvm.amdgcn.ds.permute
262 def int_amdgcn_ds_permute :
263 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
264
265 // llvm.amdgcn.ds.bpermute
266 def int_amdgcn_ds_bpermute :
267 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
268
269 }
223223 // will use this for some partially aligned loads.
224224 const MachineOperand *Offset0Imm = getNamedOperand(*LdSt,
225225 AMDGPU::OpName::offset0);
226 // DS_PERMUTE does not have Offset0Imm (and Offset1Imm).
227 if (!Offset0Imm)
228 return false;
229
226230 const MachineOperand *Offset1Imm = getNamedOperand(*LdSt,
227231 AMDGPU::OpName::offset1);
228232
24082408 }
24092409 }
24102410
2411 multiclass DS_1A1D_PERMUTE op, string opName, RegisterClass rc,
2412 SDPatternOperator node = null_frag,
2413 dag outs = (outs rc:$vdst),
2414 dag ins = (ins VGPR_32:$addr, rc:$data0),
2415 string asm = opName#" $vdst, $addr, $data0"> {
2416
2417 let mayLoad = 0, mayStore = 0, isConvergent = 1 in {
2418 def "" : DS_Pseudo
2419 [(set (i32 rc:$vdst),
2420 (node (i32 VGPR_32:$addr), (i32 rc:$data0)))]>;
2421
2422 let data1 = 0, offset0 = 0, offset1 = 0, gds = 0 in {
2423 def "_vi" : DS_Real_vi ;
2424 }
2425 }
2426 }
2427
24112428 multiclass DS_1A2D_RET_m op, string opName, RegisterClass rc,
24122429 string noRetOp = "", dag ins,
24132430 dag outs = (outs rc:$vdst),
135135 (S_MEMREALTIME)
136136 >;
137137
138 //===----------------------------------------------------------------------===//
139 // DS_PERMUTE/DS_BPERMUTE Instructions.
140 //===----------------------------------------------------------------------===//
141
142 let Uses = [EXEC] in {
143 defm DS_PERMUTE_B32 : DS_1A1D_PERMUTE <0x3e, "ds_permute_b32", VGPR_32,
144 int_amdgcn_ds_permute>;
145 defm DS_BPERMUTE_B32 : DS_1A1D_PERMUTE <0x3f, "ds_bpermute_b32", VGPR_32,
146 int_amdgcn_ds_bpermute>;
147 }
148
138149 } // End Predicates = [isVI]
0 ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
1
2 declare i32 @llvm.amdgcn.ds.bpermute(i32, i32) #0
3
4 ; FUNC-LABEL: {{^}}ds_bpermute:
5 ; CHECK: ds_bpermute_b32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
6 define void @ds_bpermute(i32 addrspace(1)* %out, i32 %index, i32 %src) nounwind {
7 %bpermute = call i32 @llvm.amdgcn.ds.bpermute(i32 %index, i32 %src) #0
8 store i32 %bpermute, i32 addrspace(1)* %out, align 4
9 ret void
10 }
11
12 attributes #0 = { nounwind readnone convergent }
0 ; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
1
2 declare i32 @llvm.amdgcn.ds.permute(i32, i32) #0
3
4 ; FUNC-LABEL: {{^}}ds_permute:
5 ; CHECK: ds_permute_b32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
6 define void @ds_permute(i32 addrspace(1)* %out, i32 %index, i32 %src) nounwind {
7 %bpermute = call i32 @llvm.amdgcn.ds.permute(i32 %index, i32 %src) #0
8 store i32 %bpermute, i32 addrspace(1)* %out, align 4
9 ret void
10 }
11
12 attributes #0 = { nounwind readnone convergent }