[llvm] r292205 - AMDGPU: Add replacement export intrinsics
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Mon Jan 16 23:26:53 PST 2017
Author: arsenm
Date: Tue Jan 17 01:26:53 2017
New Revision: 292205
URL: http://llvm.org/viewvc/llvm-project?rev=292205&view=rev
Log:
AMDGPU: Add replacement export intrinsics
Added:
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll
llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.td
llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=292205&r1=292204&r2=292205&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Tue Jan 17 01:26:53 2017
@@ -447,6 +447,32 @@ def int_amdgcn_buffer_atomic_cmpswap : I
llvm_i1_ty], // slc(imm)
[]>;
+// Uses that do not set the done bit should set IntrWriteMem on the
+// call site.
+def int_amdgcn_exp : Intrinsic <[], [
+ llvm_i32_ty, // tgt,
+ llvm_i32_ty, // en
+ llvm_any_ty, // src0 (f32 or i32)
+ LLVMMatchType<0>, // src1
+ LLVMMatchType<0>, // src2
+ LLVMMatchType<0>, // src3
+ llvm_i1_ty, // done
+ llvm_i1_ty // vm
+ ],
+ []
+>;
+
+// exp with compr bit set.
+def int_amdgcn_exp_compr : Intrinsic <[], [
+ llvm_i32_ty, // tgt,
+ llvm_i32_ty, // en
+ llvm_anyvector_ty, // src0 (v2f16 or v2i16)
+ LLVMMatchType<0>, // src1
+ llvm_i1_ty, // done
+ llvm_i1_ty], // vm
+ []
+>;
+
def int_amdgcn_buffer_wbinvl1_sc :
GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
Intrinsic<[], [], []>;
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.td?rev=292205&r1=292204&r2=292205&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUInstrInfo.td Tue Jan 17 01:26:53 2017
@@ -288,15 +288,16 @@ def AMDGPUkill : SDNode<"AMDGPUISD::KILL
// SI+ export
def AMDGPUExportOp : SDTypeProfile<0, 8, [
- SDTCisInt<0>, // i8 en
- SDTCisInt<1>, // i1 vm
+ SDTCisInt<0>, // i8 tgt
+ SDTCisInt<1>, // i8 en
+ // i32 or f32 src0
+ SDTCisSameAs<3, 2>, // f32 src1
+ SDTCisSameAs<4, 2>, // f32 src2
+ SDTCisSameAs<5, 2>, // f32 src3
+ SDTCisInt<6>, // i1 compr
// skip done
- SDTCisInt<2>, // i8 tgt
- SDTCisSameAs<3, 1>, // i1 compr
- SDTCisFP<4>, // f32 src0
- SDTCisSameAs<5, 4>, // f32 src1
- SDTCisSameAs<6, 4>, // f32 src2
- SDTCisSameAs<7, 4> // f32 src3
+ SDTCisInt<1> // i1 vm
+
]>;
def AMDGPUexport: SDNode<"AMDGPUISD::EXPORT", AMDGPUExportOp,
Modified: llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp?rev=292205&r1=292204&r2=292205&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp Tue Jan 17 01:26:53 2017
@@ -145,6 +145,8 @@ SITargetLowering::SITargetLowering(const
setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::f32, Custom);
setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v4f32, Custom);
setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::Other, Custom);
+ setOperationAction(ISD::INTRINSIC_VOID, MVT::v2i16, Custom);
+ setOperationAction(ISD::INTRINSIC_VOID, MVT::v2f16, Custom);
setOperationAction(ISD::BRCOND, MVT::Other, Custom);
setOperationAction(ISD::BR_CC, MVT::i1, Expand);
@@ -2723,8 +2725,55 @@ SDValue SITargetLowering::LowerINTRINSIC
unsigned IntrinsicID = cast<ConstantSDNode>(Op.getOperand(1))->getZExtValue();
switch (IntrinsicID) {
- case AMDGPUIntrinsic::SI_sendmsg:
- case Intrinsic::amdgcn_s_sendmsg: {
+ case Intrinsic::amdgcn_exp: {
+ const ConstantSDNode *Tgt = cast<ConstantSDNode>(Op.getOperand(2));
+ const ConstantSDNode *En = cast<ConstantSDNode>(Op.getOperand(3));
+ const ConstantSDNode *Done = cast<ConstantSDNode>(Op.getOperand(8));
+ const ConstantSDNode *VM = cast<ConstantSDNode>(Op.getOperand(9));
+
+ const SDValue Ops[] = {
+ Chain,
+ DAG.getTargetConstant(Tgt->getZExtValue(), DL, MVT::i8), // tgt
+ DAG.getTargetConstant(En->getZExtValue(), DL, MVT::i8), // en
+ Op.getOperand(4), // src0
+ Op.getOperand(5), // src1
+ Op.getOperand(6), // src2
+ Op.getOperand(7), // src3
+ DAG.getTargetConstant(0, DL, MVT::i1), // compr
+ DAG.getTargetConstant(VM->getZExtValue(), DL, MVT::i1)
+ };
+
+ unsigned Opc = Done->isNullValue() ?
+ AMDGPUISD::EXPORT : AMDGPUISD::EXPORT_DONE;
+ return DAG.getNode(Opc, DL, Op->getVTList(), Ops);
+ }
+ case Intrinsic::amdgcn_exp_compr: {
+ const ConstantSDNode *Tgt = cast<ConstantSDNode>(Op.getOperand(2));
+ const ConstantSDNode *En = cast<ConstantSDNode>(Op.getOperand(3));
+ SDValue Src0 = Op.getOperand(4);
+ SDValue Src1 = Op.getOperand(5);
+ const ConstantSDNode *Done = cast<ConstantSDNode>(Op.getOperand(6));
+ const ConstantSDNode *VM = cast<ConstantSDNode>(Op.getOperand(7));
+
+ SDValue Undef = DAG.getUNDEF(MVT::f32);
+ const SDValue Ops[] = {
+ Chain,
+ DAG.getTargetConstant(Tgt->getZExtValue(), DL, MVT::i8), // tgt
+ DAG.getTargetConstant(En->getZExtValue(), DL, MVT::i8), // en
+ DAG.getNode(ISD::BITCAST, DL, MVT::f32, Src0),
+ DAG.getNode(ISD::BITCAST, DL, MVT::f32, Src1),
+ Undef, // src2
+ Undef, // src3
+ DAG.getTargetConstant(1, DL, MVT::i1), // compr
+ DAG.getTargetConstant(VM->getZExtValue(), DL, MVT::i1)
+ };
+
+ unsigned Opc = Done->isNullValue() ?
+ AMDGPUISD::EXPORT : AMDGPUISD::EXPORT_DONE;
+ return DAG.getNode(Opc, DL, Op->getVTList(), Ops);
+ }
+ case Intrinsic::amdgcn_s_sendmsg:
+ case AMDGPUIntrinsic::SI_sendmsg: {
Chain = copyToM0(DAG, Chain, DL, Op.getOperand(3));
SDValue Glue = Chain.getValue(1);
return DAG.getNode(AMDGPUISD::SENDMSG, DL, MVT::Other, Chain,
@@ -2776,7 +2825,7 @@ SDValue SITargetLowering::LowerINTRINSIC
SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::i32, Src);
return DAG.getNode(AMDGPUISD::KILL, DL, MVT::Other, Chain, Cast);
}
- case AMDGPUIntrinsic::SI_export: {
+ case AMDGPUIntrinsic::SI_export: { // Legacy intrinsic.
const ConstantSDNode *En = cast<ConstantSDNode>(Op.getOperand(2));
const ConstantSDNode *VM = cast<ConstantSDNode>(Op.getOperand(3));
const ConstantSDNode *Done = cast<ConstantSDNode>(Op.getOperand(4));
@@ -2785,14 +2834,14 @@ SDValue SITargetLowering::LowerINTRINSIC
const SDValue Ops[] = {
Chain,
- DAG.getTargetConstant(En->getZExtValue(), DL, MVT::i8),
- DAG.getTargetConstant(VM->getZExtValue(), DL, MVT::i1),
DAG.getTargetConstant(Tgt->getZExtValue(), DL, MVT::i8),
+ DAG.getTargetConstant(En->getZExtValue(), DL, MVT::i8),
+ Op.getOperand(7), // src0
+ Op.getOperand(8), // src1
+ Op.getOperand(9), // src2
+ Op.getOperand(10), // src3
DAG.getTargetConstant(Compr->getZExtValue(), DL, MVT::i1),
- Op.getOperand(7), // src0
- Op.getOperand(8), // src1
- Op.getOperand(9), // src2
- Op.getOperand(10) // src3
+ DAG.getTargetConstant(VM->getZExtValue(), DL, MVT::i1)
};
unsigned Opc = Done->isNullValue() ?
Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td?rev=292205&r1=292204&r2=292205&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td Tue Jan 17 01:26:53 2017
@@ -648,8 +648,9 @@ class EXP_Helper<bit done, SDPatternOper
ExpSrc0:$src0, ExpSrc1:$src1, ExpSrc2:$src2, ExpSrc3:$src3,
exp_vm:$vm, exp_compr:$compr, i8imm:$en),
"exp$tgt $src0, $src1, $src2, $src3"#!if(done, " done", "")#"$compr$vm",
- [(node (i8 timm:$en), (i1 timm:$vm), (i8 timm:$tgt), (i1 timm:$compr),
- f32:$src0, f32:$src1, f32:$src2, f32:$src3)]> {
+ [(node (i8 timm:$tgt), (i8 timm:$en),
+ f32:$src0, f32:$src1, f32:$src2, f32:$src3,
+ (i1 timm:$compr), (i1 timm:$vm))]> {
let AsmMatchConverter = "cvtExp";
}
@@ -1035,7 +1036,7 @@ class VOPProfile <list<ValueType> _ArgVT
field Operand Src1ModDPP = getSrcModExt<Src1VT>.ret;
field Operand Src0ModSDWA = getSrcModExt<Src0VT>.ret;
field Operand Src1ModSDWA = getSrcModExt<Src1VT>.ret;
-
+
field bit HasDst = !if(!eq(DstVT.Value, untyped.Value), 0, 1);
field bit HasDst32 = HasDst;
Modified: llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstructions.td?rev=292205&r1=292204&r2=292205&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstructions.td Tue Jan 17 01:26:53 2017
@@ -1067,6 +1067,15 @@ def : Pat<
(V_MUL_F64 0, CONST.FP64_ONE, 0, $src, 0, 0)
>;
+// Allow integer inputs
+class ExpPattern<SDPatternOperator node, ValueType vt, Instruction Inst> : Pat<
+ (node (i8 timm:$tgt), (i8 timm:$en), vt:$src0, vt:$src1, vt:$src2, vt:$src3, (i1 timm:$compr), (i1 timm:$vm)),
+ (Inst i8:$tgt, vt:$src0, vt:$src1, vt:$src2, vt:$src3, i1:$vm, i1:$compr, i8:$en)
+>;
+
+def : ExpPattern<AMDGPUexport, i32, EXP>;
+def : ExpPattern<AMDGPUexport_done, i32, EXP_DONE>;
+
//===----------------------------------------------------------------------===//
// Fract Patterns
//===----------------------------------------------------------------------===//
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll?rev=292205&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll Tue Jan 17 01:26:53 2017
@@ -0,0 +1,143 @@
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -strict-whitespace -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -strict-whitespace -check-prefix=GCN %s
+
+declare void @llvm.amdgcn.exp.compr.v2f16(i32, i32, <2 x half>, <2 x half>, i1, i1) #0
+declare void @llvm.amdgcn.exp.compr.v2i16(i32, i32, <2 x i16>, <2 x i16>, i1, i1) #0
+
+; GCN-LABEL: {{^}}test_export_compr_zeroes_v2f16:
+; GCN: exp mrt0 off, off, off, off compr{{$}}
+; GCN: exp mrt0 off, off, off, off done compr{{$}}
+define void @test_export_compr_zeroes_v2f16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 0, <2 x half> zeroinitializer, <2 x half> zeroinitializer, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 0, <2 x half> zeroinitializer, <2 x half> zeroinitializer, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_en_src0_v2f16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x40003c00
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x44003800
+; GCN: exp mrt0 [[SRC0]], off, off, off done compr{{$}}
+define void @test_export_compr_en_src0_v2f16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 1, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_en_src1_v2f16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x40003c00
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x44003800
+; GCN: exp mrt0 off, [[SRC1]], off, off done compr{{$}}
+define void @test_export_compr_en_src1_v2f16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 2, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_en_src0_src1_v2f16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x40003c00
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x44003800
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off done compr{{$}}
+define void @test_export_compr_en_src0_src1_v2f16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 3, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_mrt7_v2f16:
+; GCN-DAG: v_mov_b32_e32 [[VHALF:v[0-9]+]], 0x38003800
+; GCN: exp mrt7 [[VHALF]], [[VHALF]], [[VHALF]], [[VHALF]] compr{{$}}
+; GCN: exp mrt7 [[VHALF]], [[VHALF]], [[VHALF]], [[VHALF]] done compr{{$}}
+define void @test_export_compr_mrt7_v2f16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 7, i32 15, <2 x half> <half 0.5, half 0.5>, <2 x half> <half 0.5, half 0.5>, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 7, i32 15, <2 x half> <half 0.5, half 0.5>, <2 x half> <half 0.5, half 0.5>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_z_v2f16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x40003c00
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x44003800
+; GCN: exp mrtz [[SRC0]], [[SRC1]], off, off compr{{$}}
+; GCN: exp mrtz [[SRC0]], [[SRC1]], off, off done compr{{$}}
+define void @test_export_compr_z_v2f16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 8, i32 3, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 8, i32 3, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_vm_v2f16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x40003c00
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x44003800
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off compr vm{{$}}
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off done compr vm{{$}}
+define void @test_export_compr_vm_v2f16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 3, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 false, i1 true)
+ call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 3, <2 x half> <half 1.0, half 2.0>, <2 x half> <half 0.5, half 4.0>, i1 true, i1 true)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_zeroes_v2i16:
+; GCN: exp mrt0 off, off, off, off compr{{$}}
+; GCN: exp mrt0 off, off, off, off done compr{{$}}
+define void @test_export_compr_zeroes_v2i16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 0, <2 x i16> zeroinitializer, <2 x i16> zeroinitializer, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 0, <2 x i16> zeroinitializer, <2 x i16> zeroinitializer, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_en_src0_v2i16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x20001
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x40005
+; GCN: exp mrt0 [[SRC0]], off, off, off done compr{{$}}
+define void @test_export_compr_en_src0_v2i16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 1, <2 x i16> <i16 1, i16 2>, <2 x i16> <i16 5, i16 4>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_en_src1_v2i16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x20001
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x40005
+; GCN: exp mrt0 off, [[SRC1]], off, off done compr{{$}}
+define void @test_export_compr_en_src1_v2i16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 2, <2 x i16> <i16 1, i16 2>, <2 x i16> <i16 5, i16 4>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_en_src0_src1_v2i16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x20001
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x40005
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off done compr{{$}}
+define void @test_export_compr_en_src0_src1_v2i16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 3, <2 x i16> <i16 1, i16 2>, <2 x i16> <i16 5, i16 4>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_mrt7_v2i16:
+; GCN-DAG: v_mov_b32_e32 [[VI16:v[0-9]+]], 0x50005
+; GCN: exp mrt7 [[VI16]], [[VI16]], [[VI16]], [[VI16]] compr{{$}}
+; GCN: exp mrt7 [[VI16]], [[VI16]], [[VI16]], [[VI16]] done compr{{$}}
+define void @test_export_compr_mrt7_v2i16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 7, i32 15, <2 x i16> <i16 5, i16 5>, <2 x i16> <i16 5, i16 5>, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 7, i32 15, <2 x i16> <i16 5, i16 5>, <2 x i16> <i16 5, i16 5>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_z_v2i16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x20001
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x40005
+; GCN: exp mrtz [[SRC0]], [[SRC1]], off, off compr{{$}}
+; GCN: exp mrtz [[SRC0]], [[SRC1]], off, off done compr{{$}}
+define void @test_export_compr_z_v2i16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 8, i32 3, <2 x i16> <i16 1, i16 2>, <2 x i16> <i16 5, i16 4>, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 8, i32 3, <2 x i16> <i16 1, i16 2>, <2 x i16> <i16 5, i16 4>, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_compr_vm_v2i16:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 0x20001
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 0x40005
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off compr vm{{$}}
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off done compr vm{{$}}
+define void @test_export_compr_vm_v2i16() #0 {
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 3, <2 x i16> <i16 1, i16 2>, <2 x i16> <i16 5, i16 4>, i1 false, i1 true)
+ call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 3, <2 x i16> <i16 1, i16 2>, <2 x i16> <i16 5, i16 4>, i1 true, i1 true)
+ ret void
+}
+
+attributes #0 = { nounwind }
Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll?rev=292205&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll Tue Jan 17 01:26:53 2017
@@ -0,0 +1,484 @@
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -strict-whitespace -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -strict-whitespace -check-prefix=GCN %s
+
+declare void @llvm.amdgcn.exp.f32(i32, i32, float, float, float, float, i1, i1) #1
+declare void @llvm.amdgcn.exp.i32(i32, i32, i32, i32, i32, i32, i1, i1) #1
+
+; GCN-LABEL: {{^}}test_export_zeroes_f32:
+; GCN: exp mrt0 off, off, off, off{{$}}
+; GCN: exp mrt0 off, off, off, off done{{$}}
+define void @test_export_zeroes_f32() #0 {
+
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 0, float 0.0, float 0.0, float 0.0, float 0.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 0, float 0.0, float 0.0, float 0.0, float 0.0, i1 true, i1 false)
+ ret void
+}
+
+; FIXME: Should not set up registers for the unused source registers.
+
+; GCN-LABEL: {{^}}test_export_en_src0_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 [[SRC0]], off, off, off done{{$}}
+define void @test_export_en_src0_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 1, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src1_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 off, [[SRC1]], off, off done{{$}}
+define void @test_export_en_src1_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 2, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src2_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 off, off, [[SRC2]], off done{{$}}
+define void @test_export_en_src2_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 4, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src3_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 off, off, off, [[SRC3]] done{{$}}
+define void @test_export_en_src3_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 8, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src1_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off done{{$}}
+define void @test_export_en_src0_src1_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 3, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src2_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 [[SRC0]], off, [[SRC2]], off done{{$}}
+define void @test_export_en_src0_src2_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 5, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src3_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 [[SRC0]], off, off, [[SRC3]]{{$}}
+; GCN: exp mrt0 [[SRC0]], off, off, [[SRC3]] done{{$}}
+define void @test_export_en_src0_src3_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 9, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 9, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src1_src2_src3_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_en_src0_src1_src2_src3_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_mrt7_f32:
+; GCN-DAG: v_mov_b32_e32 [[VHALF:v[0-9]+]], 0.5
+; GCN: exp mrt7 [[VHALF]], [[VHALF]], [[VHALF]], [[VHALF]]{{$}}
+; GCN: exp mrt7 [[VHALF]], [[VHALF]], [[VHALF]], [[VHALF]] done{{$}}
+define void @test_export_mrt7_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 7, i32 15, float 0.5, float 0.5, float 0.5, float 0.5, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 7, i32 15, float 0.5, float 0.5, float 0.5, float 0.5, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_z_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrtz [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp mrtz [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_z_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 8, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 8, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_null_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp null [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp null [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_null_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 9, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 9, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_reserved10_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp invalid_target_10 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp invalid_target_10 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_reserved10_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 10, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 10, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_reserved11_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp invalid_target_11 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp invalid_target_11 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_reserved11_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 11, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 11, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_pos0_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp pos0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp pos0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_pos0_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 12, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 12, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_pos3_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp pos3 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp pos3 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_pos3_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 15, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 15, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_param0_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp param0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp param0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_param0_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 32, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 32, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_param31_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp param31 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp param31 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_param31_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 63, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.f32(i32 63, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_vm_f32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1.0
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2.0
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 0.5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4.0
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] vm{{$}}
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done vm{{$}}
+define void @test_export_vm_f32() #0 {
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 false, i1 true)
+ call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float 1.0, float 2.0, float 0.5, float 4.0, i1 true, i1 true)
+ ret void
+}
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+; GCN-LABEL: {{^}}test_export_zeroes_i32:
+; GCN: exp mrt0 off, off, off, off{{$}}
+; GCN: exp mrt0 off, off, off, off done{{$}}
+define void @test_export_zeroes_i32() #0 {
+
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i1 true, i1 false)
+ ret void
+}
+
+; FIXME: Should not set up registers for the unused source registers.
+
+; GCN-LABEL: {{^}}test_export_en_src0_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 [[SRC0]], off, off, off done{{$}}
+define void @test_export_en_src0_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 1, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src1_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 off, [[SRC1]], off, off done{{$}}
+define void @test_export_en_src1_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 2, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src2_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 off, off, [[SRC2]], off done{{$}}
+define void @test_export_en_src2_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 4, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src3_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 off, off, off, [[SRC3]] done{{$}}
+define void @test_export_en_src3_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 8, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src1_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], off, off done{{$}}
+define void @test_export_en_src0_src1_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 3, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src2_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 [[SRC0]], off, [[SRC2]], off done{{$}}
+define void @test_export_en_src0_src2_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 5, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src3_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 [[SRC0]], off, off, [[SRC3]]{{$}}
+; GCN: exp mrt0 [[SRC0]], off, off, [[SRC3]] done{{$}}
+define void @test_export_en_src0_src3_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 9, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 9, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_en_src0_src1_src2_src3_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_en_src0_src1_src2_src3_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_mrt7_i32:
+; GCN-DAG: v_mov_b32_e32 [[VHALF:v[0-9]+]], 5
+; GCN: exp mrt7 [[VHALF]], [[VHALF]], [[VHALF]], [[VHALF]]{{$}}
+; GCN: exp mrt7 [[VHALF]], [[VHALF]], [[VHALF]], [[VHALF]] done{{$}}
+define void @test_export_mrt7_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 7, i32 15, i32 5, i32 5, i32 5, i32 5, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 7, i32 15, i32 5, i32 5, i32 5, i32 5, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_z_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrtz [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp mrtz [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_z_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 8, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 8, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_null_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp null [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp null [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_null_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 9, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 9, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_reserved10_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp invalid_target_10 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp invalid_target_10 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_reserved10_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 10, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 10, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_reserved11_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp invalid_target_11 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp invalid_target_11 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_reserved11_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 11, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 11, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_pos0_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp pos0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp pos0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_pos0_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 12, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 12, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_pos3_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp pos3 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp pos3 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_pos3_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 15, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 15, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_param0_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp param0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp param0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_param0_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 32, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 32, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_param31_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp param31 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]]{{$}}
+; GCN: exp param31 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done{{$}}
+define void @test_export_param31_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 63, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 false)
+ call void @llvm.amdgcn.exp.i32(i32 63, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 false)
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_export_vm_i32:
+; GCN-DAG: v_mov_b32_e32 [[SRC0:v[0-9]+]], 1
+; GCN-DAG: v_mov_b32_e32 [[SRC1:v[0-9]+]], 2
+; GCN-DAG: v_mov_b32_e32 [[SRC2:v[0-9]+]], 5
+; GCN-DAG: v_mov_b32_e32 [[SRC3:v[0-9]+]], 4
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] vm{{$}}
+; GCN: exp mrt0 [[SRC0]], [[SRC1]], [[SRC2]], [[SRC3]] done vm{{$}}
+define void @test_export_vm_i32() #0 {
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 15, i32 1, i32 2, i32 5, i32 4, i1 false, i1 true)
+ call void @llvm.amdgcn.exp.i32(i32 0, i32 15, i32 1, i32 2, i32 5, i32 4, i1 true, i1 true)
+ ret void
+}
+
+attributes #0 = { nounwind }
+attributes #1 = { nounwind inaccessiblememonly }
More information about the llvm-commits
mailing list