This is an archive of the discontinued LLVM Phabricator instance.

Differential D27682

AMDGPU: Add replacement export intrinsics
ClosedPublic

Authored by arsenm on Dec 12 2016, 11:39 AM.

Details

Reviewers
• tstellarAMD
arsenm
Summary

Replace llvm.SI.export with llvm.amdgcn.exp and llvm.amdgcn.exp.compr.. The compr bit is separate because it changes the number and type of the operands.

This is a WIP patch because I'm not sure whether the en mask should be an explicit operand, or whether it should be computed from which inputs are undef. Right now dead code is emitted for disabled inputs.

Diff Detail

Event Timeline

arsenm updated this revision to Diff 81117.Dec 12 2016, 11:39 AM
arsenm retitled this revision from to AMDGPU: Add replacement export intrinsics.
arsenm updated this object.
arsenm added a subscriber: llvm-commits.
Herald added a reviewer: • tstellarAMD. · View Herald TranscriptDec 12 2016, 11:39 AM
Herald added subscribers: tony-tye, yaxunl, nhaehnle and 2 others. · View Herald Transcript
mareko added a subscriber: mareko.Dec 13 2016, 6:55 AM
mareko added inline comments.
include/llvm/IR/IntrinsicsAMDGPU.td
458

// src3

470

v2f16 isn't the best choice here. The compressed export can be used with one of these types:

  • v2f16
  • v2i16
  • v2u16

Using i32 would be better, because the last 2 are packed as i32 anyway. It really depends on the output type of packing instructions. This is the complete list of instructions we should be using for compressed exports:

  • v_cvt_pkrtz_f16_f32
  • v_cvt_pknorm_u16_f32
  • v_cvt_pknorm_i16_f32
  • v_cvt_pk_u16_u32
  • v_cvt_pk_i16_i32
473

While IntrInaccessibleMemOnly makes sense for EXP in theory, in practice we might need something more limiting, because the first executed EXP instruction limits parallelism and therefore reduces the ability to hide latencies (the first EXP triggers EXP_ALLOC and if there is not enough EXP memory, the wave has to wait), so we don't want to move the first EXP across any load or store that's above it.

Depending on the chip and other parameters, EXP_ALLOC is sometimes done at wave launch, in which case the EXP scheduling doesn't matter. These are the only cases where EXP_ALLOC is done at wave launch:

  • SI: all vertex shaders (not configurable)
  • CIK-VI: all vertex shaders if the number of good CUs is <= 4 (e.g. Kabini, Mullins, Stoney, some Kaveri chips, Carrizo B4), it's configurable via a context register.
arsenm updated this revision to Diff 83775.Jan 9 2017, 9:28 PM
arsenm edited edge metadata.
arsenm marked 3 inline comments as done.

Allow mangling to v2i16 or v2f16

Herald edited edge metadata. · View Herald TranscriptJan 9 2017, 9:28 PM
mareko added a comment.Jan 11 2017, 6:02 AM

LGTM if I can just bitcast from i32 to v2i16.

arsenm updated this revision to Diff 84011.Jan 11 2017, 1:06 PM
arsenm edited edge metadata.

Also mangle for llvm.amdgcn.exp

Herald edited edge metadata. · View Herald TranscriptJan 11 2017, 1:06 PM
mareko added a comment.Jan 11 2017, 1:42 PM

LGTM.

arsenm accepted this revision.Jan 16 2017, 11:47 PM

r292205

This revision is now accepted and ready to land.Jan 16 2017, 11:47 PM
arsenm closed this revision.Jan 16 2017, 11:47 PM

Revision Contents

PathSize
include/
llvm/
IR/
26 lines
lib/
Target/
AMDGPU/
17 lines
67 lines
7 lines
9 lines
test/
CodeGen/
AMDGPU/
143 lines
484 lines

Diff 84011

include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 441 Lines • ▼ Show 20 Linesdef int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
[llvm_i32_ty, // src(VGPR) [llvm_i32_ty, // src(VGPR)
llvm_i32_ty, // cmp(VGPR) llvm_i32_ty, // cmp(VGPR)
llvm_v4i32_ty, // rsrc(SGPR) llvm_v4i32_ty, // rsrc(SGPR)
llvm_i32_ty, // vindex(VGPR) llvm_i32_ty, // vindex(VGPR)
llvm_i32_ty, // offset(SGPR/VGPR/imm) llvm_i32_ty, // offset(SGPR/VGPR/imm)
llvm_i1_ty], // slc(imm) 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
marekoUnsubmitted
Done
ReplyInline Actions

// src3

mareko: // 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
marekoUnsubmitted
Done
ReplyInline Actions

v2f16 isn't the best choice here. The compressed export can be used with one of these types:

  • v2f16
  • v2i16
  • v2u16

Using i32 would be better, because the last 2 are packed as i32 anyway. It really depends on the output type of packing instructions. This is the complete list of instructions we should be using for compressed exports:

  • v_cvt_pkrtz_f16_f32
  • v_cvt_pknorm_u16_f32
  • v_cvt_pknorm_i16_f32
  • v_cvt_pk_u16_u32
  • v_cvt_pk_i16_i32
mareko: v2f16 isn't the best choice here. The compressed export can be used with one of these types…
llvm_i1_ty, // done
llvm_i1_ty], // vm
[]
marekoUnsubmitted
Done
ReplyInline Actions

While IntrInaccessibleMemOnly makes sense for EXP in theory, in practice we might need something more limiting, because the first executed EXP instruction limits parallelism and therefore reduces the ability to hide latencies (the first EXP triggers EXP_ALLOC and if there is not enough EXP memory, the wave has to wait), so we don't want to move the first EXP across any load or store that's above it.

Depending on the chip and other parameters, EXP_ALLOC is sometimes done at wave launch, in which case the EXP scheduling doesn't matter. These are the only cases where EXP_ALLOC is done at wave launch:

  • SI: all vertex shaders (not configurable)
  • CIK-VI: all vertex shaders if the number of good CUs is <= 4 (e.g. Kabini, Mullins, Stoney, some Kaveri chips, Carrizo B4), it's configurable via a context register.
mareko: While IntrInaccessibleMemOnly makes sense for EXP in theory, in practice we might need…
>;
def int_amdgcn_buffer_wbinvl1_sc : def int_amdgcn_buffer_wbinvl1_sc :
GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">, GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
def int_amdgcn_buffer_wbinvl1 : def int_amdgcn_buffer_wbinvl1 :
GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">, GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
▲ Show 20 Lines • Show All 167 Lines • Show Last 20 Lines

lib/Target/AMDGPU/AMDGPUInstrInfo.td

Show First 20 Lines • Show All 282 Lines • ▼ Show 20 Linesdef AMDGPUinterp_p2 : SDNode<"AMDGPUISD::INTERP_P2",
[SDNPInGlue]>; [SDNPInGlue]>;
def AMDGPUkill : SDNode<"AMDGPUISD::KILL", AMDGPUKillSDT, def AMDGPUkill : SDNode<"AMDGPUISD::KILL", AMDGPUKillSDT,
[SDNPHasChain, SDNPSideEffect]>; [SDNPHasChain, SDNPSideEffect]>;
// SI+ export // SI+ export
def AMDGPUExportOp : SDTypeProfile<0, 8, [ def AMDGPUExportOp : SDTypeProfile<0, 8, [
SDTCisInt<0>, // i8 en SDTCisInt<0>, // i8 tgt
SDTCisInt<1>, // i1 vm 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 // skip done
SDTCisInt<2>, // i8 tgt SDTCisInt<1> // i1 vm
SDTCisSameAs<3, 1>, // i1 compr
SDTCisFP<4>, // f32 src0
SDTCisSameAs<5, 4>, // f32 src1
SDTCisSameAs<6, 4>, // f32 src2
SDTCisSameAs<7, 4> // f32 src3
]>; ]>;
def AMDGPUexport: SDNode<"AMDGPUISD::EXPORT", AMDGPUExportOp, def AMDGPUexport: SDNode<"AMDGPUISD::EXPORT", AMDGPUExportOp,
[SDNPHasChain, SDNPMayStore]>; [SDNPHasChain, SDNPMayStore]>;
def AMDGPUexport_done: SDNode<"AMDGPUISD::EXPORT_DONE", AMDGPUExportOp, def AMDGPUexport_done: SDNode<"AMDGPUISD::EXPORT_DONE", AMDGPUExportOp,
[SDNPHasChain, SDNPMayLoad, SDNPMayStore]>; [SDNPHasChain, SDNPMayLoad, SDNPMayStore]>;
Show All 27 Lines

lib/Target/AMDGPU/SIISelLowering.cpp

Show First 20 Lines • Show All 139 Lines • ▼ Show 20 LinesSITargetLowering::SITargetLowering(const TargetMachine &TM,
setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::v4i8, Custom); setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::v4i8, Custom);
setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::v2i16, Custom); setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::v2i16, Custom);
setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::v4i16, Custom); setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::v4i16, Custom);
setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::Other, Custom); setOperationAction(ISD::SIGN_EXTEND_INREG, MVT::Other, Custom);
setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::f32, Custom); setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::f32, Custom);
setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v4f32, Custom); setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::v4f32, Custom);
setOperationAction(ISD::INTRINSIC_W_CHAIN, MVT::Other, 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::BRCOND, MVT::Other, Custom);
setOperationAction(ISD::BR_CC, MVT::i1, Expand); setOperationAction(ISD::BR_CC, MVT::i1, Expand);
setOperationAction(ISD::BR_CC, MVT::i32, Expand); setOperationAction(ISD::BR_CC, MVT::i32, Expand);
setOperationAction(ISD::BR_CC, MVT::i64, Expand); setOperationAction(ISD::BR_CC, MVT::i64, Expand);
setOperationAction(ISD::BR_CC, MVT::f32, Expand); setOperationAction(ISD::BR_CC, MVT::f32, Expand);
setOperationAction(ISD::BR_CC, MVT::f64, Expand); setOperationAction(ISD::BR_CC, MVT::f64, Expand);
▲ Show 20 Lines • Show All 2,560 Lines • ▼ Show 20 Lines
SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
SelectionDAG &DAG) const { SelectionDAG &DAG) const {
MachineFunction &MF = DAG.getMachineFunction(); MachineFunction &MF = DAG.getMachineFunction();
SDLoc DL(Op); SDLoc DL(Op);
SDValue Chain = Op.getOperand(0); SDValue Chain = Op.getOperand(0);
unsigned IntrinsicID = cast<ConstantSDNode>(Op.getOperand(1))->getZExtValue(); unsigned IntrinsicID = cast<ConstantSDNode>(Op.getOperand(1))->getZExtValue();
switch (IntrinsicID) { switch (IntrinsicID) {
case AMDGPUIntrinsic::SI_sendmsg: case Intrinsic::amdgcn_exp: {
case Intrinsic::amdgcn_s_sendmsg: { 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)); Chain = copyToM0(DAG, Chain, DL, Op.getOperand(3));
SDValue Glue = Chain.getValue(1); SDValue Glue = Chain.getValue(1);
return DAG.getNode(AMDGPUISD::SENDMSG, DL, MVT::Other, Chain, return DAG.getNode(AMDGPUISD::SENDMSG, DL, MVT::Other, Chain,
Op.getOperand(2), Glue); Op.getOperand(2), Glue);
} }
case Intrinsic::amdgcn_s_sendmsghalt: { case Intrinsic::amdgcn_s_sendmsghalt: {
Chain = copyToM0(DAG, Chain, DL, Op.getOperand(3)); Chain = copyToM0(DAG, Chain, DL, Op.getOperand(3));
SDValue Glue = Chain.getValue(1); SDValue Glue = Chain.getValue(1);
Show All 35 Lines if (const ConstantFPSDNode *K = dyn_cast<ConstantFPSDNode>(Src)) {
SDValue NegOne = DAG.getTargetConstant(FloatToBits(-1.0f), DL, MVT::i32); SDValue NegOne = DAG.getTargetConstant(FloatToBits(-1.0f), DL, MVT::i32);
return DAG.getNode(AMDGPUISD::KILL, DL, MVT::Other, Chain, NegOne); return DAG.getNode(AMDGPUISD::KILL, DL, MVT::Other, Chain, NegOne);
} }
SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::i32, Src); SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::i32, Src);
return DAG.getNode(AMDGPUISD::KILL, DL, MVT::Other, Chain, Cast); 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 *En = cast<ConstantSDNode>(Op.getOperand(2));
const ConstantSDNode *VM = cast<ConstantSDNode>(Op.getOperand(3)); const ConstantSDNode *VM = cast<ConstantSDNode>(Op.getOperand(3));
const ConstantSDNode *Done = cast<ConstantSDNode>(Op.getOperand(4)); const ConstantSDNode *Done = cast<ConstantSDNode>(Op.getOperand(4));
const ConstantSDNode *Tgt = cast<ConstantSDNode>(Op.getOperand(5)); const ConstantSDNode *Tgt = cast<ConstantSDNode>(Op.getOperand(5));
const ConstantSDNode *Compr = cast<ConstantSDNode>(Op.getOperand(6)); const ConstantSDNode *Compr = cast<ConstantSDNode>(Op.getOperand(6));
const SDValue Ops[] = { const SDValue Ops[] = {
Chain, Chain,
DAG.getTargetConstant(En->getZExtValue(), DL, MVT::i8),
DAG.getTargetConstant(VM->getZExtValue(), DL, MVT::i1),
DAG.getTargetConstant(Tgt->getZExtValue(), DL, MVT::i8), DAG.getTargetConstant(Tgt->getZExtValue(), DL, MVT::i8),
DAG.getTargetConstant(Compr->getZExtValue(), DL, MVT::i1), DAG.getTargetConstant(En->getZExtValue(), DL, MVT::i8),
Op.getOperand(7), // src0 Op.getOperand(7), // src0
Op.getOperand(8), // src1 Op.getOperand(8), // src1
Op.getOperand(9), // src2 Op.getOperand(9), // src2
Op.getOperand(10) // src3 Op.getOperand(10), // src3
DAG.getTargetConstant(Compr->getZExtValue(), DL, MVT::i1),
DAG.getTargetConstant(VM->getZExtValue(), DL, MVT::i1)
}; };
unsigned Opc = Done->isNullValue() ? unsigned Opc = Done->isNullValue() ?
AMDGPUISD::EXPORT : AMDGPUISD::EXPORT_DONE; AMDGPUISD::EXPORT : AMDGPUISD::EXPORT_DONE;
return DAG.getNode(Opc, DL, Op->getVTList(), Ops); return DAG.getNode(Opc, DL, Op->getVTList(), Ops);
} }
default: default:
return SDValue(); return SDValue();
▲ Show 20 Lines • Show All 1,740 Lines • Show Last 20 Lines

lib/Target/AMDGPU/SIInstrInfo.td

Show First 20 Lines • Show All 642 Lines • ▼ Show 20 Lines
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
class EXP_Helper<bit done, SDPatternOperator node = null_frag> : EXPCommon< class EXP_Helper<bit done, SDPatternOperator node = null_frag> : EXPCommon<
(outs), (outs),
(ins exp_tgt:$tgt, (ins exp_tgt:$tgt,
ExpSrc0:$src0, ExpSrc1:$src1, ExpSrc2:$src2, ExpSrc3:$src3, ExpSrc0:$src0, ExpSrc1:$src1, ExpSrc2:$src2, ExpSrc3:$src3,
exp_vm:$vm, exp_compr:$compr, i8imm:$en), exp_vm:$vm, exp_compr:$compr, i8imm:$en),
"exp$tgt $src0, $src1, $src2, $src3"#!if(done, " done", "")#"$compr$vm", "exp$tgt $src0, $src1, $src2, $src3"#!if(done, " done", "")#"$compr$vm",
[(node (i8 timm:$en), (i1 timm:$vm), (i8 timm:$tgt), (i1 timm:$compr), [(node (i8 timm:$tgt), (i8 timm:$en),
f32:$src0, f32:$src1, f32:$src2, f32:$src3)]> { f32:$src0, f32:$src1, f32:$src2, f32:$src3,
(i1 timm:$compr), (i1 timm:$vm))]> {
let AsmMatchConverter = "cvtExp"; let AsmMatchConverter = "cvtExp";
} }
// Split EXP instruction into EXP and EXP_DONE so we can set // Split EXP instruction into EXP and EXP_DONE so we can set
// mayLoad for done=1. // mayLoad for done=1.
multiclass EXP_m<bit done, SDPatternOperator node> { multiclass EXP_m<bit done, SDPatternOperator node> {
let mayLoad = done in { let mayLoad = done in {
let isPseudo = 1, isCodeGenOnly = 1 in { let isPseudo = 1, isCodeGenOnly = 1 in {
▲ Show 20 Lines • Show All 369 Lines • ▼ Show 20 Linesclass VOPProfile <list<ValueType> _ArgVT> {
field RegisterClass Src1SDWA = getVregSrcForVT<Src1VT>.ret; field RegisterClass Src1SDWA = getVregSrcForVT<Src1VT>.ret;
field Operand Src0Mod = getSrcMod<Src0VT>.ret; field Operand Src0Mod = getSrcMod<Src0VT>.ret;
field Operand Src1Mod = getSrcMod<Src1VT>.ret; field Operand Src1Mod = getSrcMod<Src1VT>.ret;
field Operand Src2Mod = getSrcMod<Src2VT>.ret; field Operand Src2Mod = getSrcMod<Src2VT>.ret;
field Operand Src0ModDPP = getSrcModExt<Src0VT>.ret; field Operand Src0ModDPP = getSrcModExt<Src0VT>.ret;
field Operand Src1ModDPP = getSrcModExt<Src1VT>.ret; field Operand Src1ModDPP = getSrcModExt<Src1VT>.ret;
field Operand Src0ModSDWA = getSrcModExt<Src0VT>.ret; field Operand Src0ModSDWA = getSrcModExt<Src0VT>.ret;
field Operand Src1ModSDWA = getSrcModExt<Src1VT>.ret; field Operand Src1ModSDWA = getSrcModExt<Src1VT>.ret;
field bit HasDst = !if(!eq(DstVT.Value, untyped.Value), 0, 1); field bit HasDst = !if(!eq(DstVT.Value, untyped.Value), 0, 1);
field bit HasDst32 = HasDst; field bit HasDst32 = HasDst;
field bit EmitDst = HasDst; // force dst encoding, see v_movreld_b32 special case field bit EmitDst = HasDst; // force dst encoding, see v_movreld_b32 special case
field int NumSrcArgs = getNumSrcArgs<Src0VT, Src1VT, Src2VT>.ret; field int NumSrcArgs = getNumSrcArgs<Src0VT, Src1VT, Src2VT>.ret;
field bit HasSrc0 = !if(!eq(Src0VT.Value, untyped.Value), 0, 1); field bit HasSrc0 = !if(!eq(Src0VT.Value, untyped.Value), 0, 1);
field bit HasSrc1 = !if(!eq(Src1VT.Value, untyped.Value), 0, 1); field bit HasSrc1 = !if(!eq(Src1VT.Value, untyped.Value), 0, 1);
field bit HasSrc2 = !if(!eq(Src2VT.Value, untyped.Value), 0, 1); field bit HasSrc2 = !if(!eq(Src2VT.Value, untyped.Value), 0, 1);
▲ Show 20 Lines • Show All 244 Lines • Show Last 20 Lines

lib/Target/AMDGPU/SIInstructions.td

Show First 20 Lines • Show All 1,030 Lines • ▼ Show 20 Linesdef : Pat<
(V_MUL_F32_e64 0, (i32 CONST.FP32_ONE), 0, $src, 0, 0) (V_MUL_F32_e64 0, (i32 CONST.FP32_ONE), 0, $src, 0, 0)
>; >;
def : Pat< def : Pat<
(fcanonicalize f64:$src), (fcanonicalize f64:$src),
(V_MUL_F64 0, CONST.FP64_ONE, 0, $src, 0, 0) (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 // Fract Patterns
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
let Predicates = [isSI] in { let Predicates = [isSI] in {
// V_FRACT is buggy on SI, so the F32 version is never used and (x-floor(x)) is // V_FRACT is buggy on SI, so the F32 version is never used and (x-floor(x)) is
// used instead. However, SI doesn't have V_FLOOR_F64, so the most efficient // used instead. However, SI doesn't have V_FLOOR_F64, so the most efficient
▲ Show 20 Lines • Show All 43 Lines • Show Last 20 Lines

test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll

  • This file was added.
; 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 }

test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll

  • This file was added.
; 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 }