forked from OSchip/llvm-project
AMDGPU: Add replacement export intrinsics
llvm-svn: 292205
This commit is contained in:
Matt Arsenault
parent
e4975487f5
commit
4165efdc58
|
|
@ -447,6 +447,32 @@ def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
|
|||
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<[], [], []>;
|
||||
|
|
|
|||
|
|
@ -288,15 +288,16 @@ def AMDGPUkill : SDNode<"AMDGPUISD::KILL", AMDGPUKillSDT,
|
|||
|
||||
// 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,
|
||||
|
|
|
|||
|
|
@ -145,6 +145,8 @@ SITargetLowering::SITargetLowering(const TargetMachine &TM,
|
|||
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_VOID(SDValue Op,
|
|||
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_VOID(SDValue Op,
|
|||
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_VOID(SDValue Op,
|
|||
|
||||
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() ?
|
||||
|
|
|
|||
|
|
@ -648,8 +648,9 @@ class EXP_Helper<bit done, SDPatternOperator node = null_frag> : EXPCommon<
|
|||
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;
|
||||
|
|
|
|||
|
|
@ -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
|
||||
//===----------------------------------------------------------------------===//
|
||||
|
|
|
|||
|
|
@ -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 }
|
||||
|
|
@ -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 }
|
||||
Loading…
Reference in New Issue