Index: include/llvm/IR/IntrinsicsAMDGPU.td =================================================================== --- include/llvm/IR/IntrinsicsAMDGPU.td +++ include/llvm/IR/IntrinsicsAMDGPU.td @@ -447,6 +447,32 @@ 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<[], [], []>; Index: lib/Target/AMDGPU/AMDGPUInstrInfo.td =================================================================== --- lib/Target/AMDGPU/AMDGPUInstrInfo.td +++ lib/Target/AMDGPU/AMDGPUInstrInfo.td @@ -288,15 +288,16 @@ // 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, Index: lib/Target/AMDGPU/SIISelLowering.cpp =================================================================== --- lib/Target/AMDGPU/SIISelLowering.cpp +++ lib/Target/AMDGPU/SIISelLowering.cpp @@ -145,6 +145,8 @@ 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); @@ -2721,8 +2723,55 @@ unsigned IntrinsicID = cast(Op.getOperand(1))->getZExtValue(); switch (IntrinsicID) { - case AMDGPUIntrinsic::SI_sendmsg: - case Intrinsic::amdgcn_s_sendmsg: { + case Intrinsic::amdgcn_exp: { + const ConstantSDNode *Tgt = cast(Op.getOperand(2)); + const ConstantSDNode *En = cast(Op.getOperand(3)); + const ConstantSDNode *Done = cast(Op.getOperand(8)); + const ConstantSDNode *VM = cast(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(Op.getOperand(2)); + const ConstantSDNode *En = cast(Op.getOperand(3)); + SDValue Src0 = Op.getOperand(4); + SDValue Src1 = Op.getOperand(5); + const ConstantSDNode *Done = cast(Op.getOperand(6)); + const ConstantSDNode *VM = cast(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, @@ -2774,7 +2823,7 @@ 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(Op.getOperand(2)); const ConstantSDNode *VM = cast(Op.getOperand(3)); const ConstantSDNode *Done = cast(Op.getOperand(4)); @@ -2783,14 +2832,14 @@ 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() ? Index: lib/Target/AMDGPU/SIInstrInfo.td =================================================================== --- lib/Target/AMDGPU/SIInstrInfo.td +++ lib/Target/AMDGPU/SIInstrInfo.td @@ -648,8 +648,9 @@ 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 @@ field Operand Src1ModDPP = getSrcModExt.ret; field Operand Src0ModSDWA = getSrcModExt.ret; field Operand Src1ModSDWA = getSrcModExt.ret; - + field bit HasDst = !if(!eq(DstVT.Value, untyped.Value), 0, 1); field bit HasDst32 = HasDst; Index: lib/Target/AMDGPU/SIInstructions.td =================================================================== --- lib/Target/AMDGPU/SIInstructions.td +++ lib/Target/AMDGPU/SIInstructions.td @@ -1036,6 +1036,15 @@ (V_MUL_F64 0, CONST.FP64_ONE, 0, $src, 0, 0) >; +// Allow integer inputs +class ExpPattern : 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; +def : ExpPattern; + //===----------------------------------------------------------------------===// // Fract Patterns //===----------------------------------------------------------------------===// Index: test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll =================================================================== --- /dev/null +++ test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll @@ -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> , <2 x half> , 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> , <2 x half> , 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> , <2 x half> , 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> , <2 x half> , i1 false, i1 false) + call void @llvm.amdgcn.exp.compr.v2f16(i32 7, i32 15, <2 x half> , <2 x half> , 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> , <2 x half> , i1 false, i1 false) + call void @llvm.amdgcn.exp.compr.v2f16(i32 8, i32 3, <2 x half> , <2 x half> , 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> , <2 x half> , i1 false, i1 true) + call void @llvm.amdgcn.exp.compr.v2f16(i32 0, i32 3, <2 x half> , <2 x half> , 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> , <2 x i16> , 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> , <2 x i16> , 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> , <2 x i16> , 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> , <2 x i16> , i1 false, i1 false) + call void @llvm.amdgcn.exp.compr.v2i16(i32 7, i32 15, <2 x i16> , <2 x i16> , 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> , <2 x i16> , i1 false, i1 false) + call void @llvm.amdgcn.exp.compr.v2i16(i32 8, i32 3, <2 x i16> , <2 x i16> , 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> , <2 x i16> , i1 false, i1 true) + call void @llvm.amdgcn.exp.compr.v2i16(i32 0, i32 3, <2 x i16> , <2 x i16> , i1 true, i1 true) + ret void +} + +attributes #0 = { nounwind } Index: test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll =================================================================== --- /dev/null +++ test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll @@ -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 }