[llvm] r262356 - AMDGPU/SI: Implement DS_PERMUTE/DS_BPERMUTE Instruction Definitions and Intrinsics
Changpeng Fang via llvm-commits
llvm-commits at lists.llvm.org
Tue Mar 1 09:51:24 PST 2016
Author: chfang
Date: Tue Mar 1 11:51:23 2016
New Revision: 262356
URL: http://llvm.org/viewvc/llvm-project?rev=262356&view=rev
Log:
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
Added:
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp
llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
llvm/trunk/lib/Target/AMDGPU/VIInstructions.td
Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=262356&r1=262355&r2=262356&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Tue Mar 1 11:51:23 2016
@@ -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]>;
+
}
Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp?rev=262356&r1=262355&r2=262356&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp Tue Mar 1 11:51:23 2016
@@ -224,6 +224,10 @@ bool SIInstrInfo::getMemOpBaseRegImmOfs(
// 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);
Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td?rev=262356&r1=262355&r2=262356&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td Tue Mar 1 11:51:23 2016
@@ -2409,6 +2409,23 @@ multiclass DS_1A1D_RET <bits<8> op, stri
}
}
+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),
Modified: llvm/trunk/lib/Target/AMDGPU/VIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/VIInstructions.td?rev=262356&r1=262355&r2=262356&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/VIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/VIInstructions.td Tue Mar 1 11:51:23 2016
@@ -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]
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll?rev=262356&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.ll Tue Mar 1 11:51:23 2016
@@ -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 }
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll?rev=262356&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.ds.permute.ll Tue Mar 1 11:51:23 2016
@@ -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 }
More information about the llvm-commits
mailing list