RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-24 19:52:54 +01:00
Code Issues Projects Releases Wiki Activity

AMDGPU/SI: Implement DS_PERMUTE/DS_BPERMUTE Instruction Definitions and Intrinsics

Browse Source 
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

llvm-svn: 262356
This commit is contained in:
Changpeng Fang 2016-03-01 17:51:23 +00:00
parent 306353eaf8
commit 929a348e60
6 changed files with 67 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -258,4 +258,13 @@ def int_amdgcn_s_dcache_wb_vol :
def int_amdgcn_s_memrealtime :
GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
Intrinsic<[llvm_i64_ty], [], []>;
// llvm.amdgcn.ds.permute <index> <src>
def int_amdgcn_ds_permute :
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
// llvm.amdgcn.ds.bpermute <index> <src>
def int_amdgcn_ds_bpermute :
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
}

4
lib/Target/AMDGPU/SIInstrInfo.cpp
View File

@ -224,6 +224,10 @@ bool SIInstrInfo::getMemOpBaseRegImmOfs(MachineInstr *LdSt, unsigned &BaseReg,
// will use this for some partially aligned loads.
const MachineOperand *Offset0Imm = getNamedOperand(*LdSt,
AMDGPU::OpName::offset0);
// DS_PERMUTE does not have Offset0Imm (and Offset1Imm).
if (!Offset0Imm)
return false;
const MachineOperand *Offset1Imm = getNamedOperand(*LdSt,
AMDGPU::OpName::offset1);

17
lib/Target/AMDGPU/SIInstrInfo.td
View File

@ -2409,6 +2409,23 @@ multiclass DS_1A1D_RET <bits<8> op, string opName, RegisterClass rc,
}
}
multiclass DS_1A1D_PERMUTE <bits<8> op, string opName, RegisterClass rc,
SDPatternOperator node = null_frag,
dag outs = (outs rc:$vdst),
dag ins = (ins VGPR_32:$addr, rc:$data0),
string asm = opName#" $vdst, $addr, $data0"> {
let mayLoad = 0, mayStore = 0, isConvergent = 1 in {
def "" : DS_Pseudo <opName, outs, ins,
[(set (i32 rc:$vdst),
(node (i32 VGPR_32:$addr), (i32 rc:$data0)))]>;
let data1 = 0, offset0 = 0, offset1 = 0, gds = 0 in {
def "_vi" : DS_Real_vi <op, opName, outs, ins, asm>;
}
}
}
multiclass DS_1A2D_RET_m <bits<8> op, string opName, RegisterClass rc,
string noRetOp = "", dag ins,
dag outs = (outs rc:$vdst),

11
lib/Target/AMDGPU/VIInstructions.td
View File

@ -136,4 +136,15 @@ def : Pat <
(S_MEMREALTIME)
>;
//===----------------------------------------------------------------------===//
// DS_PERMUTE/DS_BPERMUTE Instructions.
//===----------------------------------------------------------------------===//
let Uses = [EXEC] in {
defm DS_PERMUTE_B32 : DS_1A1D_PERMUTE <0x3e, "ds_permute_b32", VGPR_32,
int_amdgcn_ds_permute>;
defm DS_BPERMUTE_B32 : DS_1A1D_PERMUTE <0x3f, "ds_bpermute_b32", VGPR_32,
int_amdgcn_ds_bpermute>;
}
} // End Predicates = [isVI]

13
test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll Normal file
View File

@ -0,0 +1,13 @@
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
declare i32 @llvm.amdgcn.ds.bpermute(i32, i32) #0
; FUNC-LABEL: {{^}}ds_bpermute:
; CHECK: ds_bpermute_b32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
define void @ds_bpermute(i32 addrspace(1)* %out, i32 %index, i32 %src) nounwind {
%bpermute = call i32 @llvm.amdgcn.ds.bpermute(i32 %index, i32 %src) #0
store i32 %bpermute, i32 addrspace(1)* %out, align 4
ret void
}
attributes #0 = { nounwind readnone convergent }

13
test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll Normal file
View File

@ -0,0 +1,13 @@
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
declare i32 @llvm.amdgcn.ds.permute(i32, i32) #0
; FUNC-LABEL: {{^}}ds_permute:
; CHECK: ds_permute_b32 v{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
define void @ds_permute(i32 addrspace(1)* %out, i32 %index, i32 %src) nounwind {
%bpermute = call i32 @llvm.amdgcn.ds.permute(i32 %index, i32 %src) #0
store i32 %bpermute, i32 addrspace(1)* %out, align 4
ret void
}
attributes #0 = { nounwind readnone convergent }