diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 07d5b5ea40dc..ddab44c30658 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -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<[], [], []>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td b/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
index f079c8d0c70c..c65db5779d79 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
@@ -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,
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index f17af6dc7619..e3143451b009 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -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() ?
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index ebaefae3bfef..ff036a1176cf 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -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;
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 9ce717f43c03..80d1da902abc 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -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
 //===----------------------------------------------------------------------===//
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.compr.ll
new file mode 100644
index 000000000000..0b5b20be334a
--- /dev/null
+++ b/llvm/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> <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 }
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.exp.ll
new file mode 100644
index 000000000000..9b5836fa56a6
--- /dev/null
+++ b/llvm/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 }