diff options
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 244 | ||||
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 23 | ||||
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 18 | ||||
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 2630 | ||||
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/prmt-const-folding.ll | 171 |
5 files changed, 2132 insertions, 954 deletions
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index d017c65..7aa06f9 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -1048,9 +1048,12 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, MVT::v32i32, MVT::v64i32, MVT::v128i32}, Custom); - setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::Other, Custom); - // Enable custom lowering for the i128 bit operand with clusterlaunchcontrol - setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::i128, Custom); + // Enable custom lowering for the following: + // * MVT::i128 - clusterlaunchcontrol + // * MVT::i32 - prmt + // * MVT::Other - internal.addrspace.wrap + setOperationAction(ISD::INTRINSIC_WO_CHAIN, {MVT::i32, MVT::i128, MVT::Other}, + Custom); } const char *NVPTXTargetLowering::getTargetNodeName(unsigned Opcode) const { @@ -2060,6 +2063,19 @@ NVPTXTargetLowering::LowerCONCAT_VECTORS(SDValue Op, SelectionDAG &DAG) const { return DAG.getBuildVector(Node->getValueType(0), dl, Ops); } +static SDValue getPRMT(SDValue A, SDValue B, SDValue Selector, SDLoc DL, + SelectionDAG &DAG, + unsigned Mode = NVPTX::PTXPrmtMode::NONE) { + return DAG.getNode(NVPTXISD::PRMT, DL, MVT::i32, + {A, B, Selector, DAG.getConstant(Mode, DL, MVT::i32)}); +} + +static SDValue getPRMT(SDValue A, SDValue B, uint64_t Selector, SDLoc DL, + SelectionDAG &DAG, + unsigned Mode = NVPTX::PTXPrmtMode::NONE) { + return getPRMT(A, B, DAG.getConstant(Selector, DL, MVT::i32), DL, DAG, Mode); +} + SDValue NVPTXTargetLowering::LowerBITCAST(SDValue Op, SelectionDAG &DAG) const { // Handle bitcasting from v2i8 without hitting the default promotion // strategy which goes through stack memory. @@ -2111,15 +2127,12 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, L = DAG.getAnyExtOrTrunc(L, DL, MVT::i32); R = DAG.getAnyExtOrTrunc(R, DL, MVT::i32); } - return DAG.getNode( - NVPTXISD::PRMT, DL, MVT::v4i8, - {L, R, DAG.getConstant(SelectionValue, DL, MVT::i32), - DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); + return getPRMT(L, R, SelectionValue, DL, DAG); }; auto PRMT__10 = GetPRMT(Op->getOperand(0), Op->getOperand(1), true, 0x3340); auto PRMT__32 = GetPRMT(Op->getOperand(2), Op->getOperand(3), true, 0x3340); auto PRMT3210 = GetPRMT(PRMT__10, PRMT__32, false, 0x5410); - return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210); + return DAG.getBitcast(VT, PRMT3210); } // Get value or the Nth operand as an APInt(32). Undef values treated as 0. @@ -2176,11 +2189,14 @@ SDValue NVPTXTargetLowering::LowerEXTRACT_VECTOR_ELT(SDValue Op, SDValue Selector = DAG.getNode(ISD::OR, DL, MVT::i32, DAG.getZExtOrTrunc(Index, DL, MVT::i32), DAG.getConstant(0x7770, DL, MVT::i32)); - SDValue PRMT = DAG.getNode( - NVPTXISD::PRMT, DL, MVT::i32, - {DAG.getBitcast(MVT::i32, Vector), DAG.getConstant(0, DL, MVT::i32), - Selector, DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); - return DAG.getAnyExtOrTrunc(PRMT, DL, Op->getValueType(0)); + SDValue PRMT = getPRMT(DAG.getBitcast(MVT::i32, Vector), + DAG.getConstant(0, DL, MVT::i32), Selector, DL, DAG); + SDValue Ext = DAG.getAnyExtOrTrunc(PRMT, DL, Op->getValueType(0)); + SDNodeFlags Flags; + Flags.setNoSignedWrap(Ext.getScalarValueSizeInBits() > 8); + Flags.setNoUnsignedWrap(Ext.getScalarValueSizeInBits() >= 8); + Ext->setFlags(Flags); + return Ext; } // Constant index will be matched by tablegen. @@ -2242,9 +2258,9 @@ SDValue NVPTXTargetLowering::LowerVECTOR_SHUFFLE(SDValue Op, } SDLoc DL(Op); - return DAG.getNode(NVPTXISD::PRMT, DL, MVT::v4i8, V1, V2, - DAG.getConstant(Selector, DL, MVT::i32), - DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)); + SDValue PRMT = getPRMT(DAG.getBitcast(MVT::i32, V1), + DAG.getBitcast(MVT::i32, V2), Selector, DL, DAG); + return DAG.getBitcast(Op.getValueType(), PRMT); } /// LowerShiftRightParts - Lower SRL_PARTS, SRA_PARTS, which /// 1) returns two i32 values and take a 2 x i32 value to shift plus a shift @@ -2729,10 +2745,46 @@ static SDValue LowerClusterLaunchControlQueryCancel(SDValue Op, {TryCancelResponse0, TryCancelResponse1}); } +static SDValue lowerPrmtIntrinsic(SDValue Op, SelectionDAG &DAG) { + const unsigned Mode = [&]() { + switch (Op->getConstantOperandVal(0)) { + case Intrinsic::nvvm_prmt: + return NVPTX::PTXPrmtMode::NONE; + case Intrinsic::nvvm_prmt_b4e: + return NVPTX::PTXPrmtMode::B4E; + case Intrinsic::nvvm_prmt_ecl: + return NVPTX::PTXPrmtMode::ECL; + case Intrinsic::nvvm_prmt_ecr: + return NVPTX::PTXPrmtMode::ECR; + case Intrinsic::nvvm_prmt_f4e: + return NVPTX::PTXPrmtMode::F4E; + case Intrinsic::nvvm_prmt_rc16: + return NVPTX::PTXPrmtMode::RC16; + case Intrinsic::nvvm_prmt_rc8: + return NVPTX::PTXPrmtMode::RC8; + default: + llvm_unreachable("unsupported/unhandled intrinsic"); + } + }(); + SDLoc DL(Op); + SDValue A = Op->getOperand(1); + SDValue B = Op.getNumOperands() == 4 ? Op.getOperand(2) + : DAG.getConstant(0, DL, MVT::i32); + SDValue Selector = (Op->op_end() - 1)->get(); + return getPRMT(A, B, Selector, DL, DAG, Mode); +} static SDValue lowerIntrinsicWOChain(SDValue Op, SelectionDAG &DAG) { switch (Op->getConstantOperandVal(0)) { default: return Op; + case Intrinsic::nvvm_prmt: + case Intrinsic::nvvm_prmt_b4e: + case Intrinsic::nvvm_prmt_ecl: + case Intrinsic::nvvm_prmt_ecr: + case Intrinsic::nvvm_prmt_f4e: + case Intrinsic::nvvm_prmt_rc16: + case Intrinsic::nvvm_prmt_rc8: + return lowerPrmtIntrinsic(Op, DAG); case Intrinsic::nvvm_internal_addrspace_wrap: return Op.getOperand(1); case Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_is_canceled: @@ -5775,11 +5827,10 @@ PerformBUILD_VECTORCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) { SDLoc DL(N); auto &DAG = DCI.DAG; - auto PRMT = DAG.getNode( - NVPTXISD::PRMT, DL, MVT::v4i8, - {Op0, Op1, DAG.getConstant((Op1Bytes << 8) | Op0Bytes, DL, MVT::i32), - DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); - return DAG.getNode(ISD::BITCAST, DL, VT, PRMT); + auto PRMT = + getPRMT(DAG.getBitcast(MVT::i32, Op0), DAG.getBitcast(MVT::i32, Op1), + (Op1Bytes << 8) | Op0Bytes, DL, DAG); + return DAG.getBitcast(VT, PRMT); } static SDValue combineADDRSPACECAST(SDNode *N, @@ -5797,47 +5848,120 @@ static SDValue combineADDRSPACECAST(SDNode *N, return SDValue(); } +// Given a constant selector value and a prmt mode, return the selector value +// normalized to the generic prmt mode. See the PTX ISA documentation for more +// details: +// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-prmt +static APInt getPRMTSelector(const APInt &Selector, unsigned Mode) { + if (Mode == NVPTX::PTXPrmtMode::NONE) + return Selector; + + const unsigned V = Selector.trunc(2).getZExtValue(); + + const auto GetSelector = [](unsigned S0, unsigned S1, unsigned S2, + unsigned S3) { + return APInt(32, S0 | (S1 << 4) | (S2 << 8) | (S3 << 12)); + }; + + switch (Mode) { + case NVPTX::PTXPrmtMode::F4E: + return GetSelector(V, V + 1, V + 2, V + 3); + case NVPTX::PTXPrmtMode::B4E: + return GetSelector(V, (V - 1) & 7, (V - 2) & 7, (V - 3) & 7); + case NVPTX::PTXPrmtMode::RC8: + return GetSelector(V, V, V, V); + case NVPTX::PTXPrmtMode::ECL: + return GetSelector(V, std::max(V, 1U), std::max(V, 2U), 3U); + case NVPTX::PTXPrmtMode::ECR: + return GetSelector(0, std::min(V, 1U), std::min(V, 2U), V); + case NVPTX::PTXPrmtMode::RC16: { + unsigned V1 = (V & 1) << 1; + return GetSelector(V1, V1 + 1, V1, V1 + 1); + } + default: + llvm_unreachable("Invalid PRMT mode"); + } +} + +static APInt computePRMT(APInt A, APInt B, APInt Selector, unsigned Mode) { + // {b, a} = {{b7, b6, b5, b4}, {b3, b2, b1, b0}} + APInt BitField = B.concat(A); + APInt SelectorVal = getPRMTSelector(Selector, Mode); + APInt Result(32, 0); + for (unsigned I : llvm::seq(4U)) { + APInt Sel = SelectorVal.extractBits(4, I * 4); + unsigned Idx = Sel.getLoBits(3).getZExtValue(); + unsigned Sign = Sel.getHiBits(1).getZExtValue(); + APInt Byte = BitField.extractBits(8, Idx * 8); + if (Sign) + Byte = Byte.ashr(8); + Result.insertBits(Byte, I * 8); + } + return Result; +} + +static SDValue combinePRMT(SDNode *N, TargetLowering::DAGCombinerInfo &DCI, + CodeGenOptLevel OptLevel) { + if (OptLevel == CodeGenOptLevel::None) + return SDValue(); + + // Constant fold PRMT + if (isa<ConstantSDNode>(N->getOperand(0)) && + isa<ConstantSDNode>(N->getOperand(1)) && + isa<ConstantSDNode>(N->getOperand(2))) + return DCI.DAG.getConstant(computePRMT(N->getConstantOperandAPInt(0), + N->getConstantOperandAPInt(1), + N->getConstantOperandAPInt(2), + N->getConstantOperandVal(3)), + SDLoc(N), N->getValueType(0)); + + return SDValue(); +} + SDValue NVPTXTargetLowering::PerformDAGCombine(SDNode *N, DAGCombinerInfo &DCI) const { CodeGenOptLevel OptLevel = getTargetMachine().getOptLevel(); switch (N->getOpcode()) { - default: break; - case ISD::ADD: - return PerformADDCombine(N, DCI, OptLevel); - case ISD::FADD: - return PerformFADDCombine(N, DCI, OptLevel); - case ISD::MUL: - return PerformMULCombine(N, DCI, OptLevel); - case ISD::SHL: - return PerformSHLCombine(N, DCI, OptLevel); - case ISD::AND: - return PerformANDCombine(N, DCI); - case ISD::UREM: - case ISD::SREM: - return PerformREMCombine(N, DCI, OptLevel); - case ISD::SETCC: - return PerformSETCCCombine(N, DCI, STI.getSmVersion()); - case ISD::LOAD: - case NVPTXISD::LoadParamV2: - case NVPTXISD::LoadV2: - case NVPTXISD::LoadV4: - return combineUnpackingMovIntoLoad(N, DCI); - case NVPTXISD::StoreParam: - case NVPTXISD::StoreParamV2: - case NVPTXISD::StoreParamV4: - return PerformStoreParamCombine(N, DCI); - case ISD::STORE: - case NVPTXISD::StoreV2: - case NVPTXISD::StoreV4: - return PerformStoreCombine(N, DCI); - case ISD::EXTRACT_VECTOR_ELT: - return PerformEXTRACTCombine(N, DCI); - case ISD::VSELECT: - return PerformVSELECTCombine(N, DCI); - case ISD::BUILD_VECTOR: - return PerformBUILD_VECTORCombine(N, DCI); - case ISD::ADDRSPACECAST: - return combineADDRSPACECAST(N, DCI); + default: + break; + case ISD::ADD: + return PerformADDCombine(N, DCI, OptLevel); + case ISD::ADDRSPACECAST: + return combineADDRSPACECAST(N, DCI); + case ISD::AND: + return PerformANDCombine(N, DCI); + case ISD::BUILD_VECTOR: + return PerformBUILD_VECTORCombine(N, DCI); + case ISD::EXTRACT_VECTOR_ELT: + return PerformEXTRACTCombine(N, DCI); + case ISD::FADD: + return PerformFADDCombine(N, DCI, OptLevel); + case ISD::LOAD: + case NVPTXISD::LoadParamV2: + case NVPTXISD::LoadV2: + case NVPTXISD::LoadV4: + return combineUnpackingMovIntoLoad(N, DCI); + case ISD::MUL: + return PerformMULCombine(N, DCI, OptLevel); + case NVPTXISD::PRMT: + return combinePRMT(N, DCI, OptLevel); + case ISD::SETCC: + return PerformSETCCCombine(N, DCI, STI.getSmVersion()); + case ISD::SHL: + return PerformSHLCombine(N, DCI, OptLevel); + case ISD::SREM: + case ISD::UREM: + return PerformREMCombine(N, DCI, OptLevel); + case NVPTXISD::StoreParam: + case NVPTXISD::StoreParamV2: + case NVPTXISD::StoreParamV4: + return PerformStoreParamCombine(N, DCI); + case ISD::STORE: + case NVPTXISD::StoreV2: + case NVPTXISD::StoreV4: + return PerformStoreCombine(N, DCI); + case ISD::VSELECT: + return PerformVSELECTCombine(N, DCI); } return SDValue(); } @@ -6387,7 +6511,7 @@ static void computeKnownBitsForPRMT(const SDValue Op, KnownBits &Known, ConstantSDNode *Selector = dyn_cast<ConstantSDNode>(Op.getOperand(2)); unsigned Mode = Op.getConstantOperandVal(3); - if (Mode != NVPTX::PTXPrmtMode::NONE || !Selector) + if (!Selector) return; KnownBits AKnown = DAG.computeKnownBits(A, Depth); @@ -6396,7 +6520,7 @@ static void computeKnownBitsForPRMT(const SDValue Op, KnownBits &Known, // {b, a} = {{b7, b6, b5, b4}, {b3, b2, b1, b0}} KnownBits BitField = BKnown.concat(AKnown); - APInt SelectorVal = Selector->getAPIntValue(); + APInt SelectorVal = getPRMTSelector(Selector->getAPIntValue(), Mode); for (unsigned I : llvm::seq(std::min(4U, Known.getBitWidth() / 8))) { APInt Sel = SelectorVal.extractBits(4, I * 4); unsigned Idx = Sel.getLoBits(3).getZExtValue(); diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td index 4eef6c9..a5bb83d 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -1453,18 +1453,33 @@ let hasSideEffects = false in { (ins PrmtMode:$mode), "prmt.b32$mode", [(set i32:$d, (prmt i32:$a, i32:$b, imm:$c, imm:$mode))]>; + def PRMT_B32rir + : BasicFlagsNVPTXInst<(outs B32:$d), + (ins B32:$a, i32imm:$b, B32:$c), + (ins PrmtMode:$mode), + "prmt.b32$mode", + [(set i32:$d, (prmt i32:$a, imm:$b, i32:$c, imm:$mode))]>; def PRMT_B32rii : BasicFlagsNVPTXInst<(outs B32:$d), (ins B32:$a, i32imm:$b, Hexu32imm:$c), (ins PrmtMode:$mode), "prmt.b32$mode", [(set i32:$d, (prmt i32:$a, imm:$b, imm:$c, imm:$mode))]>; - def PRMT_B32rir + def PRMT_B32irr : BasicFlagsNVPTXInst<(outs B32:$d), - (ins B32:$a, i32imm:$b, B32:$c), - (ins PrmtMode:$mode), + (ins i32imm:$a, B32:$b, B32:$c), (ins PrmtMode:$mode), + "prmt.b32$mode", + [(set i32:$d, (prmt imm:$a, i32:$b, i32:$c, imm:$mode))]>; + def PRMT_B32iri + : BasicFlagsNVPTXInst<(outs B32:$d), + (ins i32imm:$a, B32:$b, Hexu32imm:$c), (ins PrmtMode:$mode), + "prmt.b32$mode", + [(set i32:$d, (prmt imm:$a, i32:$b, imm:$c, imm:$mode))]>; + def PRMT_B32iir + : BasicFlagsNVPTXInst<(outs B32:$d), + (ins i32imm:$a, i32imm:$b, B32:$c), (ins PrmtMode:$mode), "prmt.b32$mode", - [(set i32:$d, (prmt i32:$a, imm:$b, i32:$c, imm:$mode))]>; + [(set i32:$d, (prmt imm:$a, imm:$b, i32:$c, imm:$mode))]>; } diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index bad4c3c..70150bd 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -1047,24 +1047,6 @@ class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass, // MISC // -class PRMT3Pat<Intrinsic prmt_intrinsic, PatLeaf prmt_mode> - : Pat<(prmt_intrinsic i32:$a, i32:$b, i32:$c), - (PRMT_B32rrr $a, $b, $c, prmt_mode)>; - -class PRMT2Pat<Intrinsic prmt_intrinsic, PatLeaf prmt_mode> - : Pat<(prmt_intrinsic i32:$a, i32:$c), - (PRMT_B32rir $a, (i32 0), $c, prmt_mode)>; - -def : PRMT3Pat<int_nvvm_prmt, PrmtNONE>; -def : PRMT3Pat<int_nvvm_prmt_f4e, PrmtF4E>; -def : PRMT3Pat<int_nvvm_prmt_b4e, PrmtB4E>; - -def : PRMT2Pat<int_nvvm_prmt_rc8, PrmtRC8>; -def : PRMT2Pat<int_nvvm_prmt_ecl, PrmtECL>; -def : PRMT2Pat<int_nvvm_prmt_ecr, PrmtECR>; -def : PRMT2Pat<int_nvvm_prmt_rc16, PrmtRC16>; - - def INT_NVVM_NANOSLEEP_I : BasicNVPTXInst<(outs), (ins i32imm:$i), "nanosleep.u32", [(int_nvvm_nanosleep imm:$i)]>, Requires<[hasPTX<63>, hasSM<70>]>; diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll index 410c001..cbc9f70 100644 --- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll @@ -1,14 +1,19 @@ ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 3 ; ## Support i16x2 instructions -; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx80 \ -; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ -; RUN: | FileCheck -allow-deprecated-dag-overlap %s -; RUN: %if ptxas %{ \ -; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 \ -; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ -; RUN: | %ptxas-verify -arch=sm_90 \ +; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ +; RUN: | FileCheck %s --check-prefixes=CHECK,O0 +; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -verify-machineinstrs \ +; RUN: | FileCheck %s --check-prefixes=CHECK,O3 +; RUN: %if ptxas %{ \ +; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \ +; RUN: | %ptxas-verify -arch=sm_90 \ +; RUN: %} +; RUN: %if ptxas %{ \ +; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -verify-machineinstrs \ +; RUN: | %ptxas-verify -arch=sm_90 \ ; RUN: %} +target triple = "nvptx64-nvidia-cuda" target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128" define <4 x i8> @test_ret_const() #0 { @@ -79,61 +84,111 @@ define i8 @test_extract_3(<4 x i8> %a) #0 { } define i8 @test_extract_i(<4 x i8> %a, i64 %idx) #0 { -; CHECK-LABEL: test_extract_i( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-NEXT: .reg .b64 %rd<2>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_extract_i_param_0]; -; CHECK-NEXT: cvt.u32.u64 %r2, %rd1; -; CHECK-NEXT: or.b32 %r3, %r2, 30576; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, %r3; -; CHECK-NEXT: st.param.b32 [func_retval0], %r4; -; CHECK-NEXT: ret; +; O0-LABEL: test_extract_i( +; O0: { +; O0-NEXT: .reg .b32 %r<5>; +; O0-NEXT: .reg .b64 %rd<2>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_extract_i_param_0]; +; O0-NEXT: cvt.u32.u64 %r2, %rd1; +; O0-NEXT: or.b32 %r3, %r2, 30576; +; O0-NEXT: prmt.b32 %r4, %r1, 0, %r3; +; O0-NEXT: st.param.b32 [func_retval0], %r4; +; O0-NEXT: ret; +; +; O3-LABEL: test_extract_i( +; O3: { +; O3-NEXT: .reg .b32 %r<5>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_extract_i_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_extract_i_param_1]; +; O3-NEXT: or.b32 %r3, %r2, 30576; +; O3-NEXT: prmt.b32 %r4, %r1, 0, %r3; +; O3-NEXT: st.param.b32 [func_retval0], %r4; +; O3-NEXT: ret; %e = extractelement <4 x i8> %a, i64 %idx ret i8 %e } define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_add( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_add_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_add_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; -; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; -; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; -; CHECK-NEXT: add.s16 %rs3, %rs2, %rs1; -; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; -; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; -; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; -; CHECK-NEXT: add.s16 %rs6, %rs5, %rs4; -; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; -; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; -; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; -; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; -; CHECK-NEXT: add.s16 %rs9, %rs8, %rs7; -; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; -; CHECK-NEXT: cvt.u16.u32 %rs10, %r13; -; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r14; -; CHECK-NEXT: add.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_add( +; O0: { +; O0-NEXT: .reg .b16 %rs<13>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_add_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_add_param_0]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; +; O0-NEXT: cvt.u16.u32 %rs1, %r3; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; +; O0-NEXT: cvt.u16.u32 %rs2, %r4; +; O0-NEXT: add.s16 %rs3, %rs2, %rs1; +; O0-NEXT: cvt.u32.u16 %r5, %rs3; +; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; +; O0-NEXT: cvt.u16.u32 %rs4, %r6; +; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; +; O0-NEXT: cvt.u16.u32 %rs5, %r7; +; O0-NEXT: add.s16 %rs6, %rs5, %rs4; +; O0-NEXT: cvt.u32.u16 %r8, %rs6; +; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; +; O0-NEXT: cvt.u16.u32 %rs7, %r10; +; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; +; O0-NEXT: cvt.u16.u32 %rs8, %r11; +; O0-NEXT: add.s16 %rs9, %rs8, %rs7; +; O0-NEXT: cvt.u32.u16 %r12, %rs9; +; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; +; O0-NEXT: cvt.u16.u32 %rs10, %r13; +; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; +; O0-NEXT: cvt.u16.u32 %rs11, %r14; +; O0-NEXT: add.s16 %rs12, %rs11, %rs10; +; O0-NEXT: cvt.u32.u16 %r15, %rs12; +; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_add( +; O3: { +; O3-NEXT: .reg .b16 %rs<13>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_add_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_add_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; +; O3-NEXT: cvt.u16.u32 %rs1, %r3; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; +; O3-NEXT: cvt.u16.u32 %rs2, %r4; +; O3-NEXT: add.s16 %rs3, %rs2, %rs1; +; O3-NEXT: cvt.u32.u16 %r5, %rs3; +; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; +; O3-NEXT: cvt.u16.u32 %rs4, %r6; +; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; +; O3-NEXT: cvt.u16.u32 %rs5, %r7; +; O3-NEXT: add.s16 %rs6, %rs5, %rs4; +; O3-NEXT: cvt.u32.u16 %r8, %rs6; +; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; +; O3-NEXT: cvt.u16.u32 %rs7, %r10; +; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; +; O3-NEXT: cvt.u16.u32 %rs8, %r11; +; O3-NEXT: add.s16 %rs9, %rs8, %rs7; +; O3-NEXT: cvt.u32.u16 %r12, %rs9; +; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; +; O3-NEXT: cvt.u16.u32 %rs10, %r13; +; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; +; O3-NEXT: cvt.u16.u32 %rs11, %r14; +; O3-NEXT: add.s16 %rs12, %rs11, %rs10; +; O3-NEXT: cvt.u32.u16 %r15, %rs12; +; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r17; +; O3-NEXT: ret; %r = add <4 x i8> %a, %b ret <4 x i8> %r } @@ -205,341 +260,631 @@ define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 { } define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_sub( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_sub_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_sub_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; -; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; -; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; -; CHECK-NEXT: sub.s16 %rs3, %rs2, %rs1; -; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; -; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; -; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; -; CHECK-NEXT: sub.s16 %rs6, %rs5, %rs4; -; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; -; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; -; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; -; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; -; CHECK-NEXT: sub.s16 %rs9, %rs8, %rs7; -; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; -; CHECK-NEXT: cvt.u16.u32 %rs10, %r13; -; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r14; -; CHECK-NEXT: sub.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_sub( +; O0: { +; O0-NEXT: .reg .b16 %rs<13>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_sub_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_sub_param_0]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; +; O0-NEXT: cvt.u16.u32 %rs1, %r3; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; +; O0-NEXT: cvt.u16.u32 %rs2, %r4; +; O0-NEXT: sub.s16 %rs3, %rs2, %rs1; +; O0-NEXT: cvt.u32.u16 %r5, %rs3; +; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; +; O0-NEXT: cvt.u16.u32 %rs4, %r6; +; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; +; O0-NEXT: cvt.u16.u32 %rs5, %r7; +; O0-NEXT: sub.s16 %rs6, %rs5, %rs4; +; O0-NEXT: cvt.u32.u16 %r8, %rs6; +; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; +; O0-NEXT: cvt.u16.u32 %rs7, %r10; +; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; +; O0-NEXT: cvt.u16.u32 %rs8, %r11; +; O0-NEXT: sub.s16 %rs9, %rs8, %rs7; +; O0-NEXT: cvt.u32.u16 %r12, %rs9; +; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; +; O0-NEXT: cvt.u16.u32 %rs10, %r13; +; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; +; O0-NEXT: cvt.u16.u32 %rs11, %r14; +; O0-NEXT: sub.s16 %rs12, %rs11, %rs10; +; O0-NEXT: cvt.u32.u16 %r15, %rs12; +; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_sub( +; O3: { +; O3-NEXT: .reg .b16 %rs<13>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_sub_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_sub_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; +; O3-NEXT: cvt.u16.u32 %rs1, %r3; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; +; O3-NEXT: cvt.u16.u32 %rs2, %r4; +; O3-NEXT: sub.s16 %rs3, %rs2, %rs1; +; O3-NEXT: cvt.u32.u16 %r5, %rs3; +; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; +; O3-NEXT: cvt.u16.u32 %rs4, %r6; +; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; +; O3-NEXT: cvt.u16.u32 %rs5, %r7; +; O3-NEXT: sub.s16 %rs6, %rs5, %rs4; +; O3-NEXT: cvt.u32.u16 %r8, %rs6; +; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; +; O3-NEXT: cvt.u16.u32 %rs7, %r10; +; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; +; O3-NEXT: cvt.u16.u32 %rs8, %r11; +; O3-NEXT: sub.s16 %rs9, %rs8, %rs7; +; O3-NEXT: cvt.u32.u16 %r12, %rs9; +; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; +; O3-NEXT: cvt.u16.u32 %rs10, %r13; +; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; +; O3-NEXT: cvt.u16.u32 %rs11, %r14; +; O3-NEXT: sub.s16 %rs12, %rs11, %rs10; +; O3-NEXT: cvt.u32.u16 %r15, %rs12; +; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r17; +; O3-NEXT: ret; %r = sub <4 x i8> %a, %b ret <4 x i8> %r } define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_smax( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_smax_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_smax_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; -; CHECK-NEXT: setp.gt.s32 %p1, %r4, %r3; -; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; -; CHECK-NEXT: setp.gt.s32 %p2, %r6, %r5; -; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; -; CHECK-NEXT: setp.gt.s32 %p3, %r8, %r7; -; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; -; CHECK-NEXT: setp.gt.s32 %p4, %r10, %r9; -; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; -; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; -; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_smax( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_smax_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_smax_param_0]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O0-NEXT: setp.gt.s32 %p1, %r4, %r3; +; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O0-NEXT: setp.gt.s32 %p2, %r6, %r5; +; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O0-NEXT: setp.gt.s32 %p3, %r8, %r7; +; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O0-NEXT: setp.gt.s32 %p4, %r10, %r9; +; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_smax( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_smax_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_smax_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O3-NEXT: setp.gt.s32 %p1, %r4, %r3; +; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O3-NEXT: setp.gt.s32 %p2, %r6, %r5; +; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O3-NEXT: setp.gt.s32 %p3, %r8, %r7; +; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O3-NEXT: setp.gt.s32 %p4, %r10, %r9; +; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r17; +; O3-NEXT: ret; %cmp = icmp sgt <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r } define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_umax( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_umax_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_umax_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; -; CHECK-NEXT: setp.gt.u32 %p1, %r4, %r3; -; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; -; CHECK-NEXT: setp.gt.u32 %p2, %r6, %r5; -; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; -; CHECK-NEXT: setp.gt.u32 %p3, %r8, %r7; -; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; -; CHECK-NEXT: setp.gt.u32 %p4, %r10, %r9; -; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; -; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; -; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_umax( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_umax_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_umax_param_0]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O0-NEXT: setp.gt.u32 %p1, %r4, %r3; +; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O0-NEXT: setp.gt.u32 %p2, %r6, %r5; +; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O0-NEXT: setp.gt.u32 %p3, %r8, %r7; +; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O0-NEXT: setp.gt.u32 %p4, %r10, %r9; +; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_umax( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_umax_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_umax_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O3-NEXT: setp.gt.u32 %p1, %r4, %r3; +; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O3-NEXT: setp.gt.u32 %p2, %r6, %r5; +; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O3-NEXT: setp.gt.u32 %p3, %r8, %r7; +; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O3-NEXT: setp.gt.u32 %p4, %r10, %r9; +; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r17; +; O3-NEXT: ret; %cmp = icmp ugt <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r } define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_smin( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_smin_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_smin_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; -; CHECK-NEXT: setp.le.s32 %p1, %r4, %r3; -; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; -; CHECK-NEXT: setp.le.s32 %p2, %r6, %r5; -; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; -; CHECK-NEXT: setp.le.s32 %p3, %r8, %r7; -; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; -; CHECK-NEXT: setp.le.s32 %p4, %r10, %r9; -; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; -; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; -; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_smin( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_smin_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_smin_param_0]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O0-NEXT: setp.le.s32 %p1, %r4, %r3; +; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O0-NEXT: setp.le.s32 %p2, %r6, %r5; +; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O0-NEXT: setp.le.s32 %p3, %r8, %r7; +; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O0-NEXT: setp.le.s32 %p4, %r10, %r9; +; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_smin( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_smin_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_smin_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O3-NEXT: setp.le.s32 %p1, %r4, %r3; +; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O3-NEXT: setp.le.s32 %p2, %r6, %r5; +; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O3-NEXT: setp.le.s32 %p3, %r8, %r7; +; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O3-NEXT: setp.le.s32 %p4, %r10, %r9; +; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r17; +; O3-NEXT: ret; %cmp = icmp sle <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r } define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_umin( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_umin_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_umin_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; -; CHECK-NEXT: setp.le.u32 %p1, %r4, %r3; -; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; -; CHECK-NEXT: setp.le.u32 %p2, %r6, %r5; -; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; -; CHECK-NEXT: setp.le.u32 %p3, %r8, %r7; -; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; -; CHECK-NEXT: setp.le.u32 %p4, %r10, %r9; -; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; -; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; -; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_umin( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_umin_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_umin_param_0]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O0-NEXT: setp.le.u32 %p1, %r4, %r3; +; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O0-NEXT: setp.le.u32 %p2, %r6, %r5; +; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O0-NEXT: setp.le.u32 %p3, %r8, %r7; +; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O0-NEXT: setp.le.u32 %p4, %r10, %r9; +; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_umin( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_umin_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_umin_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O3-NEXT: setp.le.u32 %p1, %r4, %r3; +; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O3-NEXT: setp.le.u32 %p2, %r6, %r5; +; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O3-NEXT: setp.le.u32 %p3, %r8, %r7; +; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O3-NEXT: setp.le.u32 %p4, %r10, %r9; +; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4; +; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3; +; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2; +; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r17; +; O3-NEXT: ret; %cmp = icmp ule <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r } define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { -; CHECK-LABEL: test_eq( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<23>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r3, [test_eq_param_2]; -; CHECK-NEXT: ld.param.b32 %r2, [test_eq_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_eq_param_0]; -; CHECK-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U; -; CHECK-NEXT: setp.eq.b32 %p1, %r5, %r4; -; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U; -; CHECK-NEXT: setp.eq.b32 %p2, %r7, %r6; -; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U; -; CHECK-NEXT: setp.eq.b32 %p3, %r9, %r8; -; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U; -; CHECK-NEXT: setp.eq.b32 %p4, %r11, %r10; -; CHECK-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U; -; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; -; CHECK-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U; -; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U; -; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; -; CHECK-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U; -; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1; -; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U; -; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r22; -; CHECK-NEXT: ret; +; O0-LABEL: test_eq( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<23>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r3, [test_eq_param_2]; +; O0-NEXT: ld.param.b32 %r2, [test_eq_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_eq_param_0]; +; O0-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U; +; O0-NEXT: setp.eq.b32 %p1, %r5, %r4; +; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U; +; O0-NEXT: setp.eq.b32 %p2, %r7, %r6; +; O0-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U; +; O0-NEXT: setp.eq.b32 %p3, %r9, %r8; +; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U; +; O0-NEXT: setp.eq.b32 %p4, %r11, %r10; +; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U; +; O0-NEXT: selp.b32 %r13, %r11, %r12, %p4; +; O0-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U; +; O0-NEXT: selp.b32 %r15, %r9, %r14, %p3; +; O0-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U; +; O0-NEXT: selp.b32 %r18, %r7, %r17, %p2; +; O0-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U; +; O0-NEXT: selp.b32 %r20, %r5, %r19, %p1; +; O0-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U; +; O0-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r22; +; O0-NEXT: ret; +; +; O3-LABEL: test_eq( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<23>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_eq_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_eq_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O3-NEXT: setp.eq.b32 %p1, %r4, %r3; +; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O3-NEXT: setp.eq.b32 %p2, %r6, %r5; +; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O3-NEXT: setp.eq.b32 %p3, %r8, %r7; +; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O3-NEXT: setp.eq.b32 %p4, %r10, %r9; +; O3-NEXT: ld.param.b32 %r11, [test_eq_param_2]; +; O3-NEXT: prmt.b32 %r12, %r11, 0, 0x7773U; +; O3-NEXT: selp.b32 %r13, %r10, %r12, %p4; +; O3-NEXT: prmt.b32 %r14, %r11, 0, 0x7772U; +; O3-NEXT: selp.b32 %r15, %r8, %r14, %p3; +; O3-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r11, 0, 0x7771U; +; O3-NEXT: selp.b32 %r18, %r6, %r17, %p2; +; O3-NEXT: prmt.b32 %r19, %r11, 0, 0x7770U; +; O3-NEXT: selp.b32 %r20, %r4, %r19, %p1; +; O3-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U; +; O3-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r22; +; O3-NEXT: ret; %cmp = icmp eq <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c ret <4 x i8> %r } define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { -; CHECK-LABEL: test_ne( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<23>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r3, [test_ne_param_2]; -; CHECK-NEXT: ld.param.b32 %r2, [test_ne_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_ne_param_0]; -; CHECK-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U; -; CHECK-NEXT: setp.ne.b32 %p1, %r5, %r4; -; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U; -; CHECK-NEXT: setp.ne.b32 %p2, %r7, %r6; -; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U; -; CHECK-NEXT: setp.ne.b32 %p3, %r9, %r8; -; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U; -; CHECK-NEXT: setp.ne.b32 %p4, %r11, %r10; -; CHECK-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U; -; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; -; CHECK-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U; -; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U; -; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; -; CHECK-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U; -; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1; -; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U; -; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r22; -; CHECK-NEXT: ret; +; O0-LABEL: test_ne( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<23>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r3, [test_ne_param_2]; +; O0-NEXT: ld.param.b32 %r2, [test_ne_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_ne_param_0]; +; O0-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U; +; O0-NEXT: setp.ne.b32 %p1, %r5, %r4; +; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U; +; O0-NEXT: setp.ne.b32 %p2, %r7, %r6; +; O0-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U; +; O0-NEXT: setp.ne.b32 %p3, %r9, %r8; +; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U; +; O0-NEXT: setp.ne.b32 %p4, %r11, %r10; +; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U; +; O0-NEXT: selp.b32 %r13, %r11, %r12, %p4; +; O0-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U; +; O0-NEXT: selp.b32 %r15, %r9, %r14, %p3; +; O0-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U; +; O0-NEXT: selp.b32 %r18, %r7, %r17, %p2; +; O0-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U; +; O0-NEXT: selp.b32 %r20, %r5, %r19, %p1; +; O0-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U; +; O0-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r22; +; O0-NEXT: ret; +; +; O3-LABEL: test_ne( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<23>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_ne_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_ne_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O3-NEXT: setp.ne.b32 %p1, %r4, %r3; +; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O3-NEXT: setp.ne.b32 %p2, %r6, %r5; +; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O3-NEXT: setp.ne.b32 %p3, %r8, %r7; +; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O3-NEXT: setp.ne.b32 %p4, %r10, %r9; +; O3-NEXT: ld.param.b32 %r11, [test_ne_param_2]; +; O3-NEXT: prmt.b32 %r12, %r11, 0, 0x7773U; +; O3-NEXT: selp.b32 %r13, %r10, %r12, %p4; +; O3-NEXT: prmt.b32 %r14, %r11, 0, 0x7772U; +; O3-NEXT: selp.b32 %r15, %r8, %r14, %p3; +; O3-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r11, 0, 0x7771U; +; O3-NEXT: selp.b32 %r18, %r6, %r17, %p2; +; O3-NEXT: prmt.b32 %r19, %r11, 0, 0x7770U; +; O3-NEXT: selp.b32 %r20, %r4, %r19, %p1; +; O3-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U; +; O3-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r22; +; O3-NEXT: ret; %cmp = icmp ne <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c ret <4 x i8> %r } define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_mul( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_mul_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_mul_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; -; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; -; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; -; CHECK-NEXT: mul.lo.s16 %rs3, %rs2, %rs1; -; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; -; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; -; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; -; CHECK-NEXT: mul.lo.s16 %rs6, %rs5, %rs4; -; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; -; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; -; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; -; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; -; CHECK-NEXT: mul.lo.s16 %rs9, %rs8, %rs7; -; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; -; CHECK-NEXT: cvt.u16.u32 %rs10, %r13; -; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r14; -; CHECK-NEXT: mul.lo.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_mul( +; O0: { +; O0-NEXT: .reg .b16 %rs<13>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_mul_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_mul_param_0]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; +; O0-NEXT: cvt.u16.u32 %rs1, %r3; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; +; O0-NEXT: cvt.u16.u32 %rs2, %r4; +; O0-NEXT: mul.lo.s16 %rs3, %rs2, %rs1; +; O0-NEXT: cvt.u32.u16 %r5, %rs3; +; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; +; O0-NEXT: cvt.u16.u32 %rs4, %r6; +; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; +; O0-NEXT: cvt.u16.u32 %rs5, %r7; +; O0-NEXT: mul.lo.s16 %rs6, %rs5, %rs4; +; O0-NEXT: cvt.u32.u16 %r8, %rs6; +; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; +; O0-NEXT: cvt.u16.u32 %rs7, %r10; +; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; +; O0-NEXT: cvt.u16.u32 %rs8, %r11; +; O0-NEXT: mul.lo.s16 %rs9, %rs8, %rs7; +; O0-NEXT: cvt.u32.u16 %r12, %rs9; +; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; +; O0-NEXT: cvt.u16.u32 %rs10, %r13; +; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; +; O0-NEXT: cvt.u16.u32 %rs11, %r14; +; O0-NEXT: mul.lo.s16 %rs12, %rs11, %rs10; +; O0-NEXT: cvt.u32.u16 %r15, %rs12; +; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_mul( +; O3: { +; O3-NEXT: .reg .b16 %rs<13>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_mul_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_mul_param_1]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U; +; O3-NEXT: cvt.u16.u32 %rs1, %r3; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; +; O3-NEXT: cvt.u16.u32 %rs2, %r4; +; O3-NEXT: mul.lo.s16 %rs3, %rs2, %rs1; +; O3-NEXT: cvt.u32.u16 %r5, %rs3; +; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U; +; O3-NEXT: cvt.u16.u32 %rs4, %r6; +; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U; +; O3-NEXT: cvt.u16.u32 %rs5, %r7; +; O3-NEXT: mul.lo.s16 %rs6, %rs5, %rs4; +; O3-NEXT: cvt.u32.u16 %r8, %rs6; +; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U; +; O3-NEXT: cvt.u16.u32 %rs7, %r10; +; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U; +; O3-NEXT: cvt.u16.u32 %rs8, %r11; +; O3-NEXT: mul.lo.s16 %rs9, %rs8, %rs7; +; O3-NEXT: cvt.u32.u16 %r12, %rs9; +; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U; +; O3-NEXT: cvt.u16.u32 %rs10, %r13; +; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U; +; O3-NEXT: cvt.u16.u32 %rs11, %r14; +; O3-NEXT: mul.lo.s16 %rs12, %rs11, %rs10; +; O3-NEXT: cvt.u32.u16 %r15, %rs12; +; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r17; +; O3-NEXT: ret; %r = mul <4 x i8> %a, %b ret <4 x i8> %r } define <4 x i8> @test_or(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_or( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_or_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_or_param_0]; -; CHECK-NEXT: or.b32 %r3, %r1, %r2; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_or( +; O0: { +; O0-NEXT: .reg .b32 %r<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_or_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_or_param_0]; +; O0-NEXT: or.b32 %r3, %r1, %r2; +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_or( +; O3: { +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_or_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_or_param_1]; +; O3-NEXT: or.b32 %r3, %r1, %r2; +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %r = or <4 x i8> %a, %b ret <4 x i8> %r } define <4 x i8> @test_or_computed(i8 %a) { -; CHECK-LABEL: test_or_computed( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<8>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [test_or_computed_param_0]; -; CHECK-NEXT: mov.b32 %r1, 0; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; -; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; -; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U; -; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; -; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; -; CHECK-NEXT: or.b32 %r7, %r6, %r5; -; CHECK-NEXT: st.param.b32 [func_retval0], %r7; -; CHECK-NEXT: ret; +; O0-LABEL: test_or_computed( +; O0: { +; O0-NEXT: .reg .b16 %rs<2>; +; O0-NEXT: .reg .b32 %r<8>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b8 %rs1, [test_or_computed_param_0]; +; O0-NEXT: mov.b32 %r1, 0; +; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; +; O0-NEXT: cvt.u32.u16 %r3, %rs1; +; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U; +; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; +; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; +; O0-NEXT: or.b32 %r7, %r6, %r5; +; O0-NEXT: st.param.b32 [func_retval0], %r7; +; O0-NEXT: ret; +; +; O3-LABEL: test_or_computed( +; O3: { +; O3-NEXT: .reg .b32 %r<6>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b8 %r1, [test_or_computed_param_0]; +; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U; +; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8; +; O3-NEXT: or.b32 %r5, %r4, %r3; +; O3-NEXT: st.param.b32 [func_retval0], %r5; +; O3-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 %r = or <4 x i8> %ins.1, %ins.0 @@ -575,37 +920,61 @@ define <4 x i8> @test_or_imm_1(<4 x i8> %a) #0 { } define <4 x i8> @test_xor(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_xor( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_xor_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_xor_param_0]; -; CHECK-NEXT: xor.b32 %r3, %r1, %r2; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_xor( +; O0: { +; O0-NEXT: .reg .b32 %r<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_xor_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_xor_param_0]; +; O0-NEXT: xor.b32 %r3, %r1, %r2; +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_xor( +; O3: { +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_xor_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_xor_param_1]; +; O3-NEXT: xor.b32 %r3, %r1, %r2; +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %r = xor <4 x i8> %a, %b ret <4 x i8> %r } define <4 x i8> @test_xor_computed(i8 %a) { -; CHECK-LABEL: test_xor_computed( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<8>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [test_xor_computed_param_0]; -; CHECK-NEXT: mov.b32 %r1, 0; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; -; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; -; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U; -; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; -; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; -; CHECK-NEXT: xor.b32 %r7, %r6, %r5; -; CHECK-NEXT: st.param.b32 [func_retval0], %r7; -; CHECK-NEXT: ret; +; O0-LABEL: test_xor_computed( +; O0: { +; O0-NEXT: .reg .b16 %rs<2>; +; O0-NEXT: .reg .b32 %r<8>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b8 %rs1, [test_xor_computed_param_0]; +; O0-NEXT: mov.b32 %r1, 0; +; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; +; O0-NEXT: cvt.u32.u16 %r3, %rs1; +; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U; +; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; +; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; +; O0-NEXT: xor.b32 %r7, %r6, %r5; +; O0-NEXT: st.param.b32 [func_retval0], %r7; +; O0-NEXT: ret; +; +; O3-LABEL: test_xor_computed( +; O3: { +; O3-NEXT: .reg .b32 %r<6>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b8 %r1, [test_xor_computed_param_0]; +; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U; +; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8; +; O3-NEXT: xor.b32 %r5, %r4, %r3; +; O3-NEXT: st.param.b32 [func_retval0], %r5; +; O3-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 %r = xor <4 x i8> %ins.1, %ins.0 @@ -641,37 +1010,61 @@ define <4 x i8> @test_xor_imm_1(<4 x i8> %a) #0 { } define <4 x i8> @test_and(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_and( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_and_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_and_param_0]; -; CHECK-NEXT: and.b32 %r3, %r1, %r2; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_and( +; O0: { +; O0-NEXT: .reg .b32 %r<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_and_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_and_param_0]; +; O0-NEXT: and.b32 %r3, %r1, %r2; +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_and( +; O3: { +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_and_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_and_param_1]; +; O3-NEXT: and.b32 %r3, %r1, %r2; +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %r = and <4 x i8> %a, %b ret <4 x i8> %r } define <4 x i8> @test_and_computed(i8 %a) { -; CHECK-LABEL: test_and_computed( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<8>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [test_and_computed_param_0]; -; CHECK-NEXT: mov.b32 %r1, 0; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; -; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; -; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U; -; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; -; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; -; CHECK-NEXT: and.b32 %r7, %r6, %r5; -; CHECK-NEXT: st.param.b32 [func_retval0], %r7; -; CHECK-NEXT: ret; +; O0-LABEL: test_and_computed( +; O0: { +; O0-NEXT: .reg .b16 %rs<2>; +; O0-NEXT: .reg .b32 %r<8>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b8 %rs1, [test_and_computed_param_0]; +; O0-NEXT: mov.b32 %r1, 0; +; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; +; O0-NEXT: cvt.u32.u16 %r3, %rs1; +; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U; +; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; +; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; +; O0-NEXT: and.b32 %r7, %r6, %r5; +; O0-NEXT: st.param.b32 [func_retval0], %r7; +; O0-NEXT: ret; +; +; O3-LABEL: test_and_computed( +; O3: { +; O3-NEXT: .reg .b32 %r<6>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b8 %r1, [test_and_computed_param_0]; +; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U; +; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8; +; O3-NEXT: and.b32 %r5, %r4, %r3; +; O3-NEXT: st.param.b32 [func_retval0], %r5; +; O3-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 %r = and <4 x i8> %ins.1, %ins.0 @@ -707,76 +1100,132 @@ define <4 x i8> @test_and_imm_1(<4 x i8> %a) #0 { } define void @test_ldst_v2i8(ptr %a, ptr %b) { -; CHECK-LABEL: test_ldst_v2i8( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<2>; -; CHECK-NEXT: .reg .b64 %rd<3>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0]; -; CHECK-NEXT: ld.b32 %r1, [%rd1]; -; CHECK-NEXT: st.b32 [%rd2], %r1; -; CHECK-NEXT: ret; +; O0-LABEL: test_ldst_v2i8( +; O0: { +; O0-NEXT: .reg .b32 %r<2>; +; O0-NEXT: .reg .b64 %rd<3>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0]; +; O0-NEXT: ld.b32 %r1, [%rd1]; +; O0-NEXT: st.b32 [%rd2], %r1; +; O0-NEXT: ret; +; +; O3-LABEL: test_ldst_v2i8( +; O3: { +; O3-NEXT: .reg .b32 %r<2>; +; O3-NEXT: .reg .b64 %rd<3>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0]; +; O3-NEXT: ld.b32 %r1, [%rd1]; +; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1]; +; O3-NEXT: st.b32 [%rd2], %r1; +; O3-NEXT: ret; %t1 = load <4 x i8>, ptr %a store <4 x i8> %t1, ptr %b, align 16 ret void } define void @test_ldst_v3i8(ptr %a, ptr %b) { -; CHECK-LABEL: test_ldst_v3i8( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<3>; -; CHECK-NEXT: .reg .b64 %rd<3>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0]; -; CHECK-NEXT: ld.b32 %r1, [%rd1]; -; CHECK-NEXT: st.b16 [%rd2], %r1; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U; -; CHECK-NEXT: st.b8 [%rd2+2], %r2; -; CHECK-NEXT: ret; +; O0-LABEL: test_ldst_v3i8( +; O0: { +; O0-NEXT: .reg .b32 %r<3>; +; O0-NEXT: .reg .b64 %rd<3>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0]; +; O0-NEXT: ld.b32 %r1, [%rd1]; +; O0-NEXT: st.b16 [%rd2], %r1; +; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U; +; O0-NEXT: st.b8 [%rd2+2], %r2; +; O0-NEXT: ret; +; +; O3-LABEL: test_ldst_v3i8( +; O3: { +; O3-NEXT: .reg .b32 %r<3>; +; O3-NEXT: .reg .b64 %rd<3>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0]; +; O3-NEXT: ld.b32 %r1, [%rd1]; +; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1]; +; O3-NEXT: st.b16 [%rd2], %r1; +; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U; +; O3-NEXT: st.b8 [%rd2+2], %r2; +; O3-NEXT: ret; %t1 = load <3 x i8>, ptr %a store <3 x i8> %t1, ptr %b, align 16 ret void } define void @test_ldst_v4i8(ptr %a, ptr %b) { -; CHECK-LABEL: test_ldst_v4i8( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<2>; -; CHECK-NEXT: .reg .b64 %rd<3>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0]; -; CHECK-NEXT: ld.b32 %r1, [%rd1]; -; CHECK-NEXT: st.b32 [%rd2], %r1; -; CHECK-NEXT: ret; +; O0-LABEL: test_ldst_v4i8( +; O0: { +; O0-NEXT: .reg .b32 %r<2>; +; O0-NEXT: .reg .b64 %rd<3>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0]; +; O0-NEXT: ld.b32 %r1, [%rd1]; +; O0-NEXT: st.b32 [%rd2], %r1; +; O0-NEXT: ret; +; +; O3-LABEL: test_ldst_v4i8( +; O3: { +; O3-NEXT: .reg .b32 %r<2>; +; O3-NEXT: .reg .b64 %rd<3>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0]; +; O3-NEXT: ld.b32 %r1, [%rd1]; +; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1]; +; O3-NEXT: st.b32 [%rd2], %r1; +; O3-NEXT: ret; %t1 = load <4 x i8>, ptr %a store <4 x i8> %t1, ptr %b, align 16 ret void } define void @test_ldst_v4i8_unaligned(ptr %a, ptr %b) { -; CHECK-LABEL: test_ldst_v4i8_unaligned( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-NEXT: .reg .b64 %rd<3>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0]; -; CHECK-NEXT: ld.b8 %r1, [%rd1]; -; CHECK-NEXT: ld.b8 %r2, [%rd1+1]; -; CHECK-NEXT: ld.b8 %r3, [%rd1+2]; -; CHECK-NEXT: ld.b8 %r4, [%rd1+3]; -; CHECK-NEXT: st.b8 [%rd2+3], %r4; -; CHECK-NEXT: st.b8 [%rd2+2], %r3; -; CHECK-NEXT: st.b8 [%rd2+1], %r2; -; CHECK-NEXT: st.b8 [%rd2], %r1; -; CHECK-NEXT: ret; +; O0-LABEL: test_ldst_v4i8_unaligned( +; O0: { +; O0-NEXT: .reg .b32 %r<5>; +; O0-NEXT: .reg .b64 %rd<3>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0]; +; O0-NEXT: ld.b8 %r1, [%rd1]; +; O0-NEXT: ld.b8 %r2, [%rd1+1]; +; O0-NEXT: ld.b8 %r3, [%rd1+2]; +; O0-NEXT: ld.b8 %r4, [%rd1+3]; +; O0-NEXT: st.b8 [%rd2+3], %r4; +; O0-NEXT: st.b8 [%rd2+2], %r3; +; O0-NEXT: st.b8 [%rd2+1], %r2; +; O0-NEXT: st.b8 [%rd2], %r1; +; O0-NEXT: ret; +; +; O3-LABEL: test_ldst_v4i8_unaligned( +; O3: { +; O3-NEXT: .reg .b32 %r<5>; +; O3-NEXT: .reg .b64 %rd<3>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0]; +; O3-NEXT: ld.b8 %r1, [%rd1+1]; +; O3-NEXT: ld.b8 %r2, [%rd1]; +; O3-NEXT: ld.b8 %r3, [%rd1+3]; +; O3-NEXT: ld.b8 %r4, [%rd1+2]; +; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1]; +; O3-NEXT: st.b8 [%rd2+2], %r4; +; O3-NEXT: st.b8 [%rd2+3], %r3; +; O3-NEXT: st.b8 [%rd2], %r2; +; O3-NEXT: st.b8 [%rd2+1], %r1; +; O3-NEXT: ret; %t1 = load <4 x i8>, ptr %a, align 1 store <4 x i8> %t1, ptr %b, align 1 ret void @@ -784,17 +1233,29 @@ define void @test_ldst_v4i8_unaligned(ptr %a, ptr %b) { define void @test_ldst_v8i8(ptr %a, ptr %b) { -; CHECK-LABEL: test_ldst_v8i8( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<3>; -; CHECK-NEXT: .reg .b64 %rd<3>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0]; -; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1]; -; CHECK-NEXT: st.v2.b32 [%rd2], {%r1, %r2}; -; CHECK-NEXT: ret; +; O0-LABEL: test_ldst_v8i8( +; O0: { +; O0-NEXT: .reg .b32 %r<3>; +; O0-NEXT: .reg .b64 %rd<3>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0]; +; O0-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1]; +; O0-NEXT: st.v2.b32 [%rd2], {%r1, %r2}; +; O0-NEXT: ret; +; +; O3-LABEL: test_ldst_v8i8( +; O3: { +; O3-NEXT: .reg .b32 %r<3>; +; O3-NEXT: .reg .b64 %rd<3>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0]; +; O3-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1]; +; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1]; +; O3-NEXT: st.v2.b32 [%rd2], {%r1, %r2}; +; O3-NEXT: ret; %t1 = load <8 x i8>, ptr %a store <8 x i8> %t1, ptr %b, align 16 ret void @@ -803,168 +1264,310 @@ define void @test_ldst_v8i8(ptr %a, ptr %b) { declare <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) #0 define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_call( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_call_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_call_param_0]; -; CHECK-NEXT: { // callseq 0, 0 -; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0], %r1; -; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1], %r2; -; CHECK-NEXT: .param .align 4 .b8 retval0[4]; -; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1); -; CHECK-NEXT: ld.param.b32 %r3, [retval0]; -; CHECK-NEXT: } // callseq 0 -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_call( +; O0: { +; O0-NEXT: .reg .b32 %r<5>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_call_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_call_param_0]; +; O0-NEXT: { // callseq 0, 0 +; O0-NEXT: .param .align 4 .b8 param0[4]; +; O0-NEXT: st.param.b32 [param0], %r1; +; O0-NEXT: .param .align 4 .b8 param1[4]; +; O0-NEXT: st.param.b32 [param1], %r2; +; O0-NEXT: .param .align 4 .b8 retval0[4]; +; O0-NEXT: call.uni (retval0), test_callee, (param0, param1); +; O0-NEXT: ld.param.b32 %r3, [retval0]; +; O0-NEXT: } // callseq 0 +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_call( +; O3: { +; O3-NEXT: .reg .b32 %r<5>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_call_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_call_param_1]; +; O3-NEXT: { // callseq 0, 0 +; O3-NEXT: .param .align 4 .b8 param0[4]; +; O3-NEXT: st.param.b32 [param0], %r1; +; O3-NEXT: .param .align 4 .b8 param1[4]; +; O3-NEXT: st.param.b32 [param1], %r2; +; O3-NEXT: .param .align 4 .b8 retval0[4]; +; O3-NEXT: call.uni (retval0), test_callee, (param0, param1); +; O3-NEXT: ld.param.b32 %r3, [retval0]; +; O3-NEXT: } // callseq 0 +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %r = call <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) ret <4 x i8> %r } define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_call_flipped( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0]; -; CHECK-NEXT: { // callseq 1, 0 -; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0], %r2; -; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1], %r1; -; CHECK-NEXT: .param .align 4 .b8 retval0[4]; -; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1); -; CHECK-NEXT: ld.param.b32 %r3, [retval0]; -; CHECK-NEXT: } // callseq 1 -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_call_flipped( +; O0: { +; O0-NEXT: .reg .b32 %r<5>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0]; +; O0-NEXT: { // callseq 1, 0 +; O0-NEXT: .param .align 4 .b8 param0[4]; +; O0-NEXT: st.param.b32 [param0], %r2; +; O0-NEXT: .param .align 4 .b8 param1[4]; +; O0-NEXT: st.param.b32 [param1], %r1; +; O0-NEXT: .param .align 4 .b8 retval0[4]; +; O0-NEXT: call.uni (retval0), test_callee, (param0, param1); +; O0-NEXT: ld.param.b32 %r3, [retval0]; +; O0-NEXT: } // callseq 1 +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_call_flipped( +; O3: { +; O3-NEXT: .reg .b32 %r<5>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1]; +; O3-NEXT: { // callseq 1, 0 +; O3-NEXT: .param .align 4 .b8 param0[4]; +; O3-NEXT: st.param.b32 [param0], %r2; +; O3-NEXT: .param .align 4 .b8 param1[4]; +; O3-NEXT: st.param.b32 [param1], %r1; +; O3-NEXT: .param .align 4 .b8 retval0[4]; +; O3-NEXT: call.uni (retval0), test_callee, (param0, param1); +; O3-NEXT: ld.param.b32 %r3, [retval0]; +; O3-NEXT: } // callseq 1 +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %r = call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) ret <4 x i8> %r } define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_tailcall_flipped( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0]; -; CHECK-NEXT: { // callseq 2, 0 -; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0], %r2; -; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1], %r1; -; CHECK-NEXT: .param .align 4 .b8 retval0[4]; -; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1); -; CHECK-NEXT: ld.param.b32 %r3, [retval0]; -; CHECK-NEXT: } // callseq 2 -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_tailcall_flipped( +; O0: { +; O0-NEXT: .reg .b32 %r<5>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0]; +; O0-NEXT: { // callseq 2, 0 +; O0-NEXT: .param .align 4 .b8 param0[4]; +; O0-NEXT: st.param.b32 [param0], %r2; +; O0-NEXT: .param .align 4 .b8 param1[4]; +; O0-NEXT: st.param.b32 [param1], %r1; +; O0-NEXT: .param .align 4 .b8 retval0[4]; +; O0-NEXT: call.uni (retval0), test_callee, (param0, param1); +; O0-NEXT: ld.param.b32 %r3, [retval0]; +; O0-NEXT: } // callseq 2 +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_tailcall_flipped( +; O3: { +; O3-NEXT: .reg .b32 %r<5>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1]; +; O3-NEXT: { // callseq 2, 0 +; O3-NEXT: .param .align 4 .b8 param0[4]; +; O3-NEXT: st.param.b32 [param0], %r2; +; O3-NEXT: .param .align 4 .b8 param1[4]; +; O3-NEXT: st.param.b32 [param1], %r1; +; O3-NEXT: .param .align 4 .b8 retval0[4]; +; O3-NEXT: call.uni (retval0), test_callee, (param0, param1); +; O3-NEXT: ld.param.b32 %r3, [retval0]; +; O3-NEXT: } // callseq 2 +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %r = tail call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) ret <4 x i8> %r } define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 { -; CHECK-LABEL: test_select( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<2>; -; CHECK-NEXT: .reg .b16 %rs<3>; -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [test_select_param_2]; -; CHECK-NEXT: and.b16 %rs2, %rs1, 1; -; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; -; CHECK-NEXT: ld.param.b32 %r2, [test_select_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_select_param_0]; -; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_select( +; O0: { +; O0-NEXT: .reg .pred %p<2>; +; O0-NEXT: .reg .b16 %rs<3>; +; O0-NEXT: .reg .b32 %r<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b8 %rs1, [test_select_param_2]; +; O0-NEXT: and.b16 %rs2, %rs1, 1; +; O0-NEXT: setp.ne.b16 %p1, %rs2, 0; +; O0-NEXT: ld.param.b32 %r2, [test_select_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_select_param_0]; +; O0-NEXT: selp.b32 %r3, %r1, %r2, %p1; +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_select( +; O3: { +; O3-NEXT: .reg .pred %p<2>; +; O3-NEXT: .reg .b16 %rs<3>; +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b8 %rs1, [test_select_param_2]; +; O3-NEXT: and.b16 %rs2, %rs1, 1; +; O3-NEXT: setp.ne.b16 %p1, %rs2, 0; +; O3-NEXT: ld.param.b32 %r1, [test_select_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_select_param_1]; +; O3-NEXT: selp.b32 %r3, %r1, %r2, %p1; +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %r = select i1 %c, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r } define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> %d) #0 { -; CHECK-LABEL: test_select_cc( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<28>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r4, [test_select_cc_param_3]; -; CHECK-NEXT: ld.param.b32 %r3, [test_select_cc_param_2]; -; CHECK-NEXT: ld.param.b32 %r2, [test_select_cc_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_select_cc_param_0]; -; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r6, %r3, 0, 0x7770U; -; CHECK-NEXT: setp.ne.b32 %p1, %r6, %r5; -; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r8, %r3, 0, 0x7771U; -; CHECK-NEXT: setp.ne.b32 %p2, %r8, %r7; -; CHECK-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r10, %r3, 0, 0x7772U; -; CHECK-NEXT: setp.ne.b32 %p3, %r10, %r9; -; CHECK-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U; -; CHECK-NEXT: setp.ne.b32 %p4, %r12, %r11; -; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U; -; CHECK-NEXT: selp.b32 %r15, %r14, %r13, %p4; -; CHECK-NEXT: prmt.b32 %r16, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U; -; CHECK-NEXT: selp.b32 %r18, %r17, %r16, %p3; -; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U; -; CHECK-NEXT: prmt.b32 %r20, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U; -; CHECK-NEXT: selp.b32 %r22, %r21, %r20, %p2; -; CHECK-NEXT: prmt.b32 %r23, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U; -; CHECK-NEXT: selp.b32 %r25, %r24, %r23, %p1; -; CHECK-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U; -; CHECK-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r27; -; CHECK-NEXT: ret; +; O0-LABEL: test_select_cc( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<28>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r4, [test_select_cc_param_3]; +; O0-NEXT: ld.param.b32 %r3, [test_select_cc_param_2]; +; O0-NEXT: ld.param.b32 %r2, [test_select_cc_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_select_cc_param_0]; +; O0-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r6, %r3, 0, 0x7770U; +; O0-NEXT: setp.ne.b32 %p1, %r6, %r5; +; O0-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r8, %r3, 0, 0x7771U; +; O0-NEXT: setp.ne.b32 %p2, %r8, %r7; +; O0-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r10, %r3, 0, 0x7772U; +; O0-NEXT: setp.ne.b32 %p3, %r10, %r9; +; O0-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U; +; O0-NEXT: setp.ne.b32 %p4, %r12, %r11; +; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U; +; O0-NEXT: selp.b32 %r15, %r14, %r13, %p4; +; O0-NEXT: prmt.b32 %r16, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U; +; O0-NEXT: selp.b32 %r18, %r17, %r16, %p3; +; O0-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U; +; O0-NEXT: prmt.b32 %r20, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U; +; O0-NEXT: selp.b32 %r22, %r21, %r20, %p2; +; O0-NEXT: prmt.b32 %r23, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U; +; O0-NEXT: selp.b32 %r25, %r24, %r23, %p1; +; O0-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U; +; O0-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r27; +; O0-NEXT: ret; +; +; O3-LABEL: test_select_cc( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<28>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_select_cc_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_select_cc_param_3]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: ld.param.b32 %r4, [test_select_cc_param_2]; +; O3-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U; +; O3-NEXT: setp.ne.b32 %p1, %r5, %r3; +; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U; +; O3-NEXT: setp.ne.b32 %p2, %r7, %r6; +; O3-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U; +; O3-NEXT: setp.ne.b32 %p3, %r9, %r8; +; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U; +; O3-NEXT: setp.ne.b32 %p4, %r11, %r10; +; O3-NEXT: ld.param.b32 %r12, [test_select_cc_param_1]; +; O3-NEXT: prmt.b32 %r13, %r12, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U; +; O3-NEXT: selp.b32 %r15, %r14, %r13, %p4; +; O3-NEXT: prmt.b32 %r16, %r12, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U; +; O3-NEXT: selp.b32 %r18, %r17, %r16, %p3; +; O3-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U; +; O3-NEXT: prmt.b32 %r20, %r12, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U; +; O3-NEXT: selp.b32 %r22, %r21, %r20, %p2; +; O3-NEXT: prmt.b32 %r23, %r12, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U; +; O3-NEXT: selp.b32 %r25, %r24, %r23, %p1; +; O3-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U; +; O3-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r27; +; O3-NEXT: ret; %cc = icmp ne <4 x i8> %c, %d %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r } define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b, -; CHECK-LABEL: test_select_cc_i32_i8( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<23>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1]; -; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0]; -; CHECK-NEXT: ld.param.b32 %r10, [test_select_cc_i32_i8_param_3]; -; CHECK-NEXT: ld.param.b32 %r9, [test_select_cc_i32_i8_param_2]; -; CHECK-NEXT: prmt.b32 %r11, %r10, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r12, %r9, 0, 0x7770U; -; CHECK-NEXT: setp.ne.b32 %p1, %r12, %r11; -; CHECK-NEXT: prmt.b32 %r13, %r10, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r14, %r9, 0, 0x7771U; -; CHECK-NEXT: setp.ne.b32 %p2, %r14, %r13; -; CHECK-NEXT: prmt.b32 %r15, %r10, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r16, %r9, 0, 0x7772U; -; CHECK-NEXT: setp.ne.b32 %p3, %r16, %r15; -; CHECK-NEXT: prmt.b32 %r17, %r10, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r18, %r9, 0, 0x7773U; -; CHECK-NEXT: setp.ne.b32 %p4, %r18, %r17; -; CHECK-NEXT: selp.b32 %r19, %r4, %r8, %p4; -; CHECK-NEXT: selp.b32 %r20, %r3, %r7, %p3; -; CHECK-NEXT: selp.b32 %r21, %r2, %r6, %p2; -; CHECK-NEXT: selp.b32 %r22, %r1, %r5, %p1; -; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19}; -; CHECK-NEXT: ret; +; O0-LABEL: test_select_cc_i32_i8( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<23>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1]; +; O0-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0]; +; O0-NEXT: ld.param.b32 %r10, [test_select_cc_i32_i8_param_3]; +; O0-NEXT: ld.param.b32 %r9, [test_select_cc_i32_i8_param_2]; +; O0-NEXT: prmt.b32 %r11, %r10, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r12, %r9, 0, 0x7770U; +; O0-NEXT: setp.ne.b32 %p1, %r12, %r11; +; O0-NEXT: prmt.b32 %r13, %r10, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r14, %r9, 0, 0x7771U; +; O0-NEXT: setp.ne.b32 %p2, %r14, %r13; +; O0-NEXT: prmt.b32 %r15, %r10, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r16, %r9, 0, 0x7772U; +; O0-NEXT: setp.ne.b32 %p3, %r16, %r15; +; O0-NEXT: prmt.b32 %r17, %r10, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r18, %r9, 0, 0x7773U; +; O0-NEXT: setp.ne.b32 %p4, %r18, %r17; +; O0-NEXT: selp.b32 %r19, %r4, %r8, %p4; +; O0-NEXT: selp.b32 %r20, %r3, %r7, %p3; +; O0-NEXT: selp.b32 %r21, %r2, %r6, %p2; +; O0-NEXT: selp.b32 %r22, %r1, %r5, %p1; +; O0-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19}; +; O0-NEXT: ret; +; +; O3-LABEL: test_select_cc_i32_i8( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<23>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0]; +; O3-NEXT: ld.param.b32 %r5, [test_select_cc_i32_i8_param_3]; +; O3-NEXT: prmt.b32 %r6, %r5, 0, 0x7770U; +; O3-NEXT: ld.param.b32 %r7, [test_select_cc_i32_i8_param_2]; +; O3-NEXT: prmt.b32 %r8, %r7, 0, 0x7770U; +; O3-NEXT: setp.ne.b32 %p1, %r8, %r6; +; O3-NEXT: prmt.b32 %r9, %r5, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r10, %r7, 0, 0x7771U; +; O3-NEXT: setp.ne.b32 %p2, %r10, %r9; +; O3-NEXT: prmt.b32 %r11, %r5, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r12, %r7, 0, 0x7772U; +; O3-NEXT: setp.ne.b32 %p3, %r12, %r11; +; O3-NEXT: prmt.b32 %r13, %r5, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r14, %r7, 0, 0x7773U; +; O3-NEXT: setp.ne.b32 %p4, %r14, %r13; +; O3-NEXT: ld.param.v4.b32 {%r15, %r16, %r17, %r18}, [test_select_cc_i32_i8_param_1]; +; O3-NEXT: selp.b32 %r19, %r4, %r18, %p4; +; O3-NEXT: selp.b32 %r20, %r3, %r17, %p3; +; O3-NEXT: selp.b32 %r21, %r2, %r16, %p2; +; O3-NEXT: selp.b32 %r22, %r1, %r15, %p1; +; O3-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19}; +; O3-NEXT: ret; <4 x i8> %c, <4 x i8> %d) #0 { %cc = icmp ne <4 x i8> %c, %d %r = select <4 x i1> %cc, <4 x i32> %a, <4 x i32> %b @@ -972,37 +1575,69 @@ define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b, } define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, -; CHECK-LABEL: test_select_cc_i8_i32( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<26>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.v4.b32 {%r7, %r8, %r9, %r10}, [test_select_cc_i8_i32_param_3]; -; CHECK-NEXT: ld.param.v4.b32 {%r3, %r4, %r5, %r6}, [test_select_cc_i8_i32_param_2]; -; CHECK-NEXT: ld.param.b32 %r2, [test_select_cc_i8_i32_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0]; -; CHECK-NEXT: setp.ne.b32 %p1, %r3, %r7; -; CHECK-NEXT: setp.ne.b32 %p2, %r4, %r8; -; CHECK-NEXT: setp.ne.b32 %p3, %r5, %r9; -; CHECK-NEXT: setp.ne.b32 %p4, %r6, %r10; -; CHECK-NEXT: prmt.b32 %r11, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U; -; CHECK-NEXT: selp.b32 %r13, %r12, %r11, %p4; -; CHECK-NEXT: prmt.b32 %r14, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U; -; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U; -; CHECK-NEXT: prmt.b32 %r18, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U; -; CHECK-NEXT: selp.b32 %r20, %r19, %r18, %p2; -; CHECK-NEXT: prmt.b32 %r21, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U; -; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p1; -; CHECK-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U; -; CHECK-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r25; -; CHECK-NEXT: ret; +; O0-LABEL: test_select_cc_i8_i32( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<26>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.v4.b32 {%r7, %r8, %r9, %r10}, [test_select_cc_i8_i32_param_3]; +; O0-NEXT: ld.param.v4.b32 {%r3, %r4, %r5, %r6}, [test_select_cc_i8_i32_param_2]; +; O0-NEXT: ld.param.b32 %r2, [test_select_cc_i8_i32_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0]; +; O0-NEXT: setp.ne.b32 %p1, %r3, %r7; +; O0-NEXT: setp.ne.b32 %p2, %r4, %r8; +; O0-NEXT: setp.ne.b32 %p3, %r5, %r9; +; O0-NEXT: setp.ne.b32 %p4, %r6, %r10; +; O0-NEXT: prmt.b32 %r11, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U; +; O0-NEXT: selp.b32 %r13, %r12, %r11, %p4; +; O0-NEXT: prmt.b32 %r14, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U; +; O0-NEXT: selp.b32 %r16, %r15, %r14, %p3; +; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U; +; O0-NEXT: prmt.b32 %r18, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U; +; O0-NEXT: selp.b32 %r20, %r19, %r18, %p2; +; O0-NEXT: prmt.b32 %r21, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U; +; O0-NEXT: selp.b32 %r23, %r22, %r21, %p1; +; O0-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U; +; O0-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r25; +; O0-NEXT: ret; +; +; O3-LABEL: test_select_cc_i8_i32( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<26>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0]; +; O3-NEXT: ld.param.v4.b32 {%r2, %r3, %r4, %r5}, [test_select_cc_i8_i32_param_2]; +; O3-NEXT: ld.param.v4.b32 {%r6, %r7, %r8, %r9}, [test_select_cc_i8_i32_param_3]; +; O3-NEXT: setp.ne.b32 %p1, %r2, %r6; +; O3-NEXT: setp.ne.b32 %p2, %r3, %r7; +; O3-NEXT: setp.ne.b32 %p3, %r4, %r8; +; O3-NEXT: setp.ne.b32 %p4, %r5, %r9; +; O3-NEXT: ld.param.b32 %r10, [test_select_cc_i8_i32_param_1]; +; O3-NEXT: prmt.b32 %r11, %r10, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U; +; O3-NEXT: selp.b32 %r13, %r12, %r11, %p4; +; O3-NEXT: prmt.b32 %r14, %r10, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U; +; O3-NEXT: selp.b32 %r16, %r15, %r14, %p3; +; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U; +; O3-NEXT: prmt.b32 %r18, %r10, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U; +; O3-NEXT: selp.b32 %r20, %r19, %r18, %p2; +; O3-NEXT: prmt.b32 %r21, %r10, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U; +; O3-NEXT: selp.b32 %r23, %r22, %r21, %p1; +; O3-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U; +; O3-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r25; +; O3-NEXT: ret; <4 x i32> %c, <4 x i32> %d) #0 { %cc = icmp ne <4 x i32> %c, %d %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b @@ -1027,23 +1662,41 @@ define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 { } define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 { -; CHECK-LABEL: test_trunc_2xi64( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<8>; -; CHECK-NEXT: .reg .b64 %rd<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16]; -; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0]; -; CHECK-NEXT: cvt.u32.u64 %r1, %rd4; -; CHECK-NEXT: cvt.u32.u64 %r2, %rd3; -; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; -; CHECK-NEXT: cvt.u32.u64 %r4, %rd2; -; CHECK-NEXT: cvt.u32.u64 %r5, %rd1; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U; -; CHECK-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r7; -; CHECK-NEXT: ret; +; O0-LABEL: test_trunc_2xi64( +; O0: { +; O0-NEXT: .reg .b32 %r<8>; +; O0-NEXT: .reg .b64 %rd<5>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16]; +; O0-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0]; +; O0-NEXT: cvt.u32.u64 %r1, %rd4; +; O0-NEXT: cvt.u32.u64 %r2, %rd3; +; O0-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; O0-NEXT: cvt.u32.u64 %r4, %rd2; +; O0-NEXT: cvt.u32.u64 %r5, %rd1; +; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U; +; O0-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r7; +; O0-NEXT: ret; +; +; O3-LABEL: test_trunc_2xi64( +; O3: { +; O3-NEXT: .reg .b32 %r<8>; +; O3-NEXT: .reg .b64 %rd<5>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0]; +; O3-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16]; +; O3-NEXT: cvt.u32.u64 %r1, %rd4; +; O3-NEXT: cvt.u32.u64 %r2, %rd3; +; O3-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U; +; O3-NEXT: cvt.u32.u64 %r4, %rd2; +; O3-NEXT: cvt.u32.u64 %r5, %rd1; +; O3-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U; +; O3-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r7; +; O3-NEXT: ret; %r = trunc <4 x i64> %a to <4 x i8> ret <4 x i8> %r } @@ -1066,24 +1719,43 @@ define <4 x i32> @test_zext_2xi32(<4 x i8> %a) #0 { } define <4 x i64> @test_zext_2xi64(<4 x i8> %a) #0 { -; CHECK-LABEL: test_zext_2xi64( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<6>; -; CHECK-NEXT: .reg .b64 %rd<5>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0]; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U; -; CHECK-NEXT: cvt.u64.u32 %rd1, %r2; -; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0x7772U; -; CHECK-NEXT: cvt.u64.u32 %rd2, %r3; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7771U; -; CHECK-NEXT: cvt.u64.u32 %rd3, %r4; -; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U; -; CHECK-NEXT: cvt.u64.u32 %rd4, %r5; -; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd3}; -; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd2, %rd1}; -; CHECK-NEXT: ret; +; O0-LABEL: test_zext_2xi64( +; O0: { +; O0-NEXT: .reg .b32 %r<6>; +; O0-NEXT: .reg .b64 %rd<5>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0]; +; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U; +; O0-NEXT: cvt.u64.u32 %rd1, %r2; +; O0-NEXT: prmt.b32 %r3, %r1, 0, 0x7772U; +; O0-NEXT: cvt.u64.u32 %rd2, %r3; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7771U; +; O0-NEXT: cvt.u64.u32 %rd3, %r4; +; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U; +; O0-NEXT: cvt.u64.u32 %rd4, %r5; +; O0-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd3}; +; O0-NEXT: st.param.v2.b64 [func_retval0+16], {%rd2, %rd1}; +; O0-NEXT: ret; +; +; O3-LABEL: test_zext_2xi64( +; O3: { +; O3-NEXT: .reg .b32 %r<6>; +; O3-NEXT: .reg .b64 %rd<5>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0]; +; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x7771U; +; O3-NEXT: cvt.u64.u32 %rd1, %r2; +; O3-NEXT: prmt.b32 %r3, %r1, 0, 0x7770U; +; O3-NEXT: cvt.u64.u32 %rd2, %r3; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U; +; O3-NEXT: cvt.u64.u32 %rd3, %r4; +; O3-NEXT: prmt.b32 %r5, %r1, 0, 0x7772U; +; O3-NEXT: cvt.u64.u32 %rd4, %r5; +; O3-NEXT: st.param.v2.b64 [func_retval0+16], {%rd4, %rd3}; +; O3-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1}; +; O3-NEXT: ret; %r = zext <4 x i8> %a to <4 x i64> ret <4 x i64> %r } @@ -1142,20 +1814,31 @@ define float @test_bitcast_4xi8_to_float(<4 x i8> %a) #0 { define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 { -; CHECK-LABEL: test_bitcast_4xi8_to_2xhalf( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<6>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0]; -; CHECK-NEXT: mov.b32 %r1, 6; -; CHECK-NEXT: prmt.b32 %r2, %r1, 7, 0x3340U; -; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; -; CHECK-NEXT: prmt.b32 %r4, %r3, 5, 0x3340U; -; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r5; -; CHECK-NEXT: ret; +; O0-LABEL: test_bitcast_4xi8_to_2xhalf( +; O0: { +; O0-NEXT: .reg .b16 %rs<2>; +; O0-NEXT: .reg .b32 %r<6>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0]; +; O0-NEXT: mov.b32 %r1, 6; +; O0-NEXT: prmt.b32 %r2, %r1, 7, 0x3340U; +; O0-NEXT: cvt.u32.u16 %r3, %rs1; +; O0-NEXT: prmt.b32 %r4, %r3, 5, 0x3340U; +; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r5; +; O0-NEXT: ret; +; +; O3-LABEL: test_bitcast_4xi8_to_2xhalf( +; O3: { +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b8 %r1, [test_bitcast_4xi8_to_2xhalf_param_0]; +; O3-NEXT: prmt.b32 %r2, %r1, 5, 0x3340U; +; O3-NEXT: prmt.b32 %r3, %r2, 1798, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %ins.0 = insertelement <4 x i8> undef, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 %ins.2 = insertelement <4 x i8> %ins.1, i8 6, i32 2 @@ -1166,153 +1849,277 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 { define <4 x i8> @test_shufflevector(<4 x i8> %a) #0 { -; CHECK-LABEL: test_shufflevector( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0]; -; CHECK-NEXT: // implicit-def: %r3 -; CHECK-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; -; CHECK-NEXT: ret; +; O0-LABEL: test_shufflevector( +; O0: { +; O0-NEXT: .reg .b32 %r<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0]; +; O0-NEXT: // implicit-def: %r3 +; O0-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U; +; O0-NEXT: st.param.b32 [func_retval0], %r2; +; O0-NEXT: ret; +; +; O3-LABEL: test_shufflevector( +; O3: { +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0]; +; O3-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U; +; O3-NEXT: st.param.b32 [func_retval0], %r2; +; O3-NEXT: ret; %s = shufflevector <4 x i8> %a, <4 x i8> undef, <4 x i32> <i32 3, i32 2, i32 1, i32 0> ret <4 x i8> %s } define <4 x i8> @test_shufflevector_2(<4 x i8> %a, <4 x i8> %b) #0 { -; CHECK-LABEL: test_shufflevector_2( -; CHECK: { -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0]; -; CHECK-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_shufflevector_2( +; O0: { +; O0-NEXT: .reg .b32 %r<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0]; +; O0-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U; +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_shufflevector_2( +; O3: { +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1]; +; O3-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U; +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %s = shufflevector <4 x i8> %a, <4 x i8> %b, <4 x i32> <i32 7, i32 3, i32 5, i32 2> ret <4 x i8> %s } define <4 x i8> @test_insertelement(<4 x i8> %a, i8 %x) #0 { -; CHECK-LABEL: test_insertelement( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.b8 %rs1, [test_insertelement_param_1]; -; CHECK-NEXT: ld.param.b32 %r1, [test_insertelement_param_0]; -; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; -; CHECK-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; -; CHECK-NEXT: ret; +; O0-LABEL: test_insertelement( +; O0: { +; O0-NEXT: .reg .b16 %rs<2>; +; O0-NEXT: .reg .b32 %r<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.b8 %rs1, [test_insertelement_param_1]; +; O0-NEXT: ld.param.b32 %r1, [test_insertelement_param_0]; +; O0-NEXT: cvt.u32.u16 %r2, %rs1; +; O0-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; +; O0-NEXT: st.param.b32 [func_retval0], %r3; +; O0-NEXT: ret; +; +; O3-LABEL: test_insertelement( +; O3: { +; O3-NEXT: .reg .b32 %r<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.b32 %r1, [test_insertelement_param_0]; +; O3-NEXT: ld.param.b8 %r2, [test_insertelement_param_1]; +; O3-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; +; O3-NEXT: st.param.b32 [func_retval0], %r3; +; O3-NEXT: ret; %i = insertelement <4 x i8> %a, i8 %x, i64 1 ret <4 x i8> %i } define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 { -; CHECK-LABEL: test_fptosi_4xhalf_to_4xi8( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<12>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0]; -; CHECK-NEXT: cvt.rzi.s16.f16 %rs5, %rs4; -; CHECK-NEXT: cvt.rzi.s16.f16 %rs6, %rs3; -; CHECK-NEXT: mov.b32 %r3, {%rs6, %rs5}; -; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3; -; CHECK-NEXT: cvt.u32.u16 %r4, %rs8; -; CHECK-NEXT: cvt.u32.u16 %r5, %rs7; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U; -; CHECK-NEXT: cvt.rzi.s16.f16 %rs9, %rs2; -; CHECK-NEXT: cvt.rzi.s16.f16 %rs10, %rs1; -; CHECK-NEXT: mov.b32 %r7, {%rs10, %rs9}; -; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r7; -; CHECK-NEXT: cvt.u32.u16 %r8, %rs12; -; CHECK-NEXT: cvt.u32.u16 %r9, %rs11; -; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U; -; CHECK-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r11; -; CHECK-NEXT: ret; +; O0-LABEL: test_fptosi_4xhalf_to_4xi8( +; O0: { +; O0-NEXT: .reg .b16 %rs<13>; +; O0-NEXT: .reg .b32 %r<12>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0]; +; O0-NEXT: cvt.rzi.s16.f16 %rs5, %rs4; +; O0-NEXT: cvt.rzi.s16.f16 %rs6, %rs3; +; O0-NEXT: mov.b32 %r3, {%rs6, %rs5}; +; O0-NEXT: mov.b32 {%rs7, %rs8}, %r3; +; O0-NEXT: cvt.u32.u16 %r4, %rs8; +; O0-NEXT: cvt.u32.u16 %r5, %rs7; +; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U; +; O0-NEXT: cvt.rzi.s16.f16 %rs9, %rs2; +; O0-NEXT: cvt.rzi.s16.f16 %rs10, %rs1; +; O0-NEXT: mov.b32 %r7, {%rs10, %rs9}; +; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7; +; O0-NEXT: cvt.u32.u16 %r8, %rs12; +; O0-NEXT: cvt.u32.u16 %r9, %rs11; +; O0-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U; +; O0-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r11; +; O0-NEXT: ret; +; +; O3-LABEL: test_fptosi_4xhalf_to_4xi8( +; O3: { +; O3-NEXT: .reg .b16 %rs<13>; +; O3-NEXT: .reg .b32 %r<10>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0]; +; O3-NEXT: cvt.rzi.s16.f16 %rs5, %rs4; +; O3-NEXT: cvt.rzi.s16.f16 %rs6, %rs3; +; O3-NEXT: mov.b32 %r1, {%rs6, %rs5}; +; O3-NEXT: mov.b32 {%rs7, %rs8}, %r1; +; O3-NEXT: cvt.u32.u16 %r2, %rs8; +; O3-NEXT: cvt.u32.u16 %r3, %rs7; +; O3-NEXT: prmt.b32 %r4, %r3, %r2, 0x3340U; +; O3-NEXT: cvt.rzi.s16.f16 %rs9, %rs2; +; O3-NEXT: cvt.rzi.s16.f16 %rs10, %rs1; +; O3-NEXT: mov.b32 %r5, {%rs10, %rs9}; +; O3-NEXT: mov.b32 {%rs11, %rs12}, %r5; +; O3-NEXT: cvt.u32.u16 %r6, %rs12; +; O3-NEXT: cvt.u32.u16 %r7, %rs11; +; O3-NEXT: prmt.b32 %r8, %r7, %r6, 0x3340U; +; O3-NEXT: prmt.b32 %r9, %r8, %r4, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r9; +; O3-NEXT: ret; %r = fptosi <4 x half> %a to <4 x i8> ret <4 x i8> %r } define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 { -; CHECK-LABEL: test_fptoui_4xhalf_to_4xi8( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<12>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: -; CHECK-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0]; -; CHECK-NEXT: cvt.rzi.u16.f16 %rs5, %rs4; -; CHECK-NEXT: cvt.rzi.u16.f16 %rs6, %rs3; -; CHECK-NEXT: mov.b32 %r3, {%rs6, %rs5}; -; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3; -; CHECK-NEXT: cvt.u32.u16 %r4, %rs8; -; CHECK-NEXT: cvt.u32.u16 %r5, %rs7; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U; -; CHECK-NEXT: cvt.rzi.u16.f16 %rs9, %rs2; -; CHECK-NEXT: cvt.rzi.u16.f16 %rs10, %rs1; -; CHECK-NEXT: mov.b32 %r7, {%rs10, %rs9}; -; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r7; -; CHECK-NEXT: cvt.u32.u16 %r8, %rs12; -; CHECK-NEXT: cvt.u32.u16 %r9, %rs11; -; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U; -; CHECK-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U; -; CHECK-NEXT: st.param.b32 [func_retval0], %r11; -; CHECK-NEXT: ret; +; O0-LABEL: test_fptoui_4xhalf_to_4xi8( +; O0: { +; O0-NEXT: .reg .b16 %rs<13>; +; O0-NEXT: .reg .b32 %r<12>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: +; O0-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0]; +; O0-NEXT: cvt.rzi.u16.f16 %rs5, %rs4; +; O0-NEXT: cvt.rzi.u16.f16 %rs6, %rs3; +; O0-NEXT: mov.b32 %r3, {%rs6, %rs5}; +; O0-NEXT: mov.b32 {%rs7, %rs8}, %r3; +; O0-NEXT: cvt.u32.u16 %r4, %rs8; +; O0-NEXT: cvt.u32.u16 %r5, %rs7; +; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U; +; O0-NEXT: cvt.rzi.u16.f16 %rs9, %rs2; +; O0-NEXT: cvt.rzi.u16.f16 %rs10, %rs1; +; O0-NEXT: mov.b32 %r7, {%rs10, %rs9}; +; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7; +; O0-NEXT: cvt.u32.u16 %r8, %rs12; +; O0-NEXT: cvt.u32.u16 %r9, %rs11; +; O0-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U; +; O0-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U; +; O0-NEXT: st.param.b32 [func_retval0], %r11; +; O0-NEXT: ret; +; +; O3-LABEL: test_fptoui_4xhalf_to_4xi8( +; O3: { +; O3-NEXT: .reg .b16 %rs<13>; +; O3-NEXT: .reg .b32 %r<10>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: +; O3-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0]; +; O3-NEXT: cvt.rzi.u16.f16 %rs5, %rs4; +; O3-NEXT: cvt.rzi.u16.f16 %rs6, %rs3; +; O3-NEXT: mov.b32 %r1, {%rs6, %rs5}; +; O3-NEXT: mov.b32 {%rs7, %rs8}, %r1; +; O3-NEXT: cvt.u32.u16 %r2, %rs8; +; O3-NEXT: cvt.u32.u16 %r3, %rs7; +; O3-NEXT: prmt.b32 %r4, %r3, %r2, 0x3340U; +; O3-NEXT: cvt.rzi.u16.f16 %rs9, %rs2; +; O3-NEXT: cvt.rzi.u16.f16 %rs10, %rs1; +; O3-NEXT: mov.b32 %r5, {%rs10, %rs9}; +; O3-NEXT: mov.b32 {%rs11, %rs12}, %r5; +; O3-NEXT: cvt.u32.u16 %r6, %rs12; +; O3-NEXT: cvt.u32.u16 %r7, %rs11; +; O3-NEXT: prmt.b32 %r8, %r7, %r6, 0x3340U; +; O3-NEXT: prmt.b32 %r9, %r8, %r4, 0x5410U; +; O3-NEXT: st.param.b32 [func_retval0], %r9; +; O3-NEXT: ret; %r = fptoui <4 x half> %a to <4 x i8> ret <4 x i8> %r } define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) { -; CHECK-LABEL: test_srem_v4i8( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-NEXT: .reg .b64 %rd<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: // %entry -; CHECK-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2]; -; CHECK-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0]; -; CHECK-NEXT: ld.b32 %r1, [%rd1]; -; CHECK-NEXT: ld.b32 %r2, [%rd2]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U; -; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U; -; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; -; CHECK-NEXT: rem.s16 %rs3, %rs2, %rs1; -; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U; -; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U; -; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; -; CHECK-NEXT: rem.s16 %rs6, %rs5, %rs4; -; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; -; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x9991U; -; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x9991U; -; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; -; CHECK-NEXT: rem.s16 %rs9, %rs8, %rs7; -; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x8880U; -; CHECK-NEXT: cvt.u16.u32 %rs10, %r13; -; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x8880U; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r14; -; CHECK-NEXT: rem.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; -; CHECK-NEXT: st.b32 [%rd3], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_srem_v4i8( +; O0: { +; O0-NEXT: .reg .b16 %rs<13>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-NEXT: .reg .b64 %rd<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: // %entry +; O0-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2]; +; O0-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0]; +; O0-NEXT: ld.b32 %r1, [%rd1]; +; O0-NEXT: ld.b32 %r2, [%rd2]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U; +; O0-NEXT: cvt.u16.u32 %rs1, %r3; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U; +; O0-NEXT: cvt.u16.u32 %rs2, %r4; +; O0-NEXT: rem.s16 %rs3, %rs2, %rs1; +; O0-NEXT: cvt.u32.u16 %r5, %rs3; +; O0-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U; +; O0-NEXT: cvt.u16.u32 %rs4, %r6; +; O0-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U; +; O0-NEXT: cvt.u16.u32 %rs5, %r7; +; O0-NEXT: rem.s16 %rs6, %rs5, %rs4; +; O0-NEXT: cvt.u32.u16 %r8, %rs6; +; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x9991U; +; O0-NEXT: cvt.u16.u32 %rs7, %r10; +; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x9991U; +; O0-NEXT: cvt.u16.u32 %rs8, %r11; +; O0-NEXT: rem.s16 %rs9, %rs8, %rs7; +; O0-NEXT: cvt.u32.u16 %r12, %rs9; +; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x8880U; +; O0-NEXT: cvt.u16.u32 %rs10, %r13; +; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x8880U; +; O0-NEXT: cvt.u16.u32 %rs11, %r14; +; O0-NEXT: rem.s16 %rs12, %rs11, %rs10; +; O0-NEXT: cvt.u32.u16 %r15, %rs12; +; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O0-NEXT: st.b32 [%rd3], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_srem_v4i8( +; O3: { +; O3-NEXT: .reg .b16 %rs<13>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-NEXT: .reg .b64 %rd<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: // %entry +; O3-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0]; +; O3-NEXT: ld.b32 %r1, [%rd1]; +; O3-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1]; +; O3-NEXT: ld.b32 %r2, [%rd2]; +; O3-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U; +; O3-NEXT: cvt.u16.u32 %rs1, %r3; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U; +; O3-NEXT: cvt.u16.u32 %rs2, %r4; +; O3-NEXT: rem.s16 %rs3, %rs2, %rs1; +; O3-NEXT: cvt.u32.u16 %r5, %rs3; +; O3-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U; +; O3-NEXT: cvt.u16.u32 %rs4, %r6; +; O3-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U; +; O3-NEXT: cvt.u16.u32 %rs5, %r7; +; O3-NEXT: rem.s16 %rs6, %rs5, %rs4; +; O3-NEXT: cvt.u32.u16 %r8, %rs6; +; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x9991U; +; O3-NEXT: cvt.u16.u32 %rs7, %r10; +; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x9991U; +; O3-NEXT: cvt.u16.u32 %rs8, %r11; +; O3-NEXT: rem.s16 %rs9, %rs8, %rs7; +; O3-NEXT: cvt.u32.u16 %r12, %rs9; +; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x8880U; +; O3-NEXT: cvt.u16.u32 %rs10, %r13; +; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x8880U; +; O3-NEXT: cvt.u16.u32 %rs11, %r14; +; O3-NEXT: rem.s16 %rs12, %rs11, %rs10; +; O3-NEXT: cvt.u32.u16 %r15, %rs12; +; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U; +; O3-NEXT: st.b32 [%rd3], %r17; +; O3-NEXT: ret; entry: %t57 = load <4 x i8>, ptr %a, align 4 %t59 = load <4 x i8>, ptr %b, align 4 @@ -1328,52 +2135,97 @@ entry: ;; Ideally we want to split it into element-wise ops, but legalizer can't handle ;; odd-sized vectors. TL;DR; don't use odd-sized vectors of v8. define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) { -; CHECK-LABEL: test_srem_v3i8( -; CHECK: { -; CHECK-NEXT: .reg .b16 %rs<20>; -; CHECK-NEXT: .reg .b32 %r<14>; -; CHECK-NEXT: .reg .b64 %rd<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: // %entry -; CHECK-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2]; -; CHECK-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0]; -; CHECK-NEXT: ld.b8 %rs1, [%rd1]; -; CHECK-NEXT: ld.b8 %rs2, [%rd1+1]; -; CHECK-NEXT: shl.b16 %rs3, %rs2, 8; -; CHECK-NEXT: or.b16 %rs4, %rs3, %rs1; -; CHECK-NEXT: cvt.u32.u16 %r1, %rs4; -; CHECK-NEXT: ld.s8 %rs5, [%rd1+2]; -; CHECK-NEXT: ld.b8 %rs6, [%rd2]; -; CHECK-NEXT: ld.b8 %rs7, [%rd2+1]; -; CHECK-NEXT: shl.b16 %rs8, %rs7, 8; -; CHECK-NEXT: or.b16 %rs9, %rs8, %rs6; -; CHECK-NEXT: cvt.u32.u16 %r2, %rs9; -; CHECK-NEXT: ld.s8 %rs10, [%rd2+2]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x9991U; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r3; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x9991U; -; CHECK-NEXT: cvt.u16.u32 %rs12, %r4; -; CHECK-NEXT: rem.s16 %rs13, %rs12, %rs11; -; CHECK-NEXT: cvt.u32.u16 %r5, %rs13; -; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x8880U; -; CHECK-NEXT: cvt.u16.u32 %rs14, %r6; -; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x8880U; -; CHECK-NEXT: cvt.u16.u32 %rs15, %r7; -; CHECK-NEXT: rem.s16 %rs16, %rs15, %rs14; -; CHECK-NEXT: cvt.u32.u16 %r8, %rs16; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; -; CHECK-NEXT: // implicit-def: %r11 -; CHECK-NEXT: // implicit-def: %r12 -; CHECK-NEXT: prmt.b32 %r10, %r11, %r12, 0x3340U; -; CHECK-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U; -; CHECK-NEXT: rem.s16 %rs17, %rs5, %rs10; -; CHECK-NEXT: mov.b32 {%rs18, _}, %r13; -; CHECK-NEXT: st.b8 [%rd3], %rs18; -; CHECK-NEXT: shr.u16 %rs19, %rs18, 8; -; CHECK-NEXT: st.b8 [%rd3+1], %rs19; -; CHECK-NEXT: st.b8 [%rd3+2], %rs17; -; CHECK-NEXT: ret; +; O0-LABEL: test_srem_v3i8( +; O0: { +; O0-NEXT: .reg .b16 %rs<20>; +; O0-NEXT: .reg .b32 %r<14>; +; O0-NEXT: .reg .b64 %rd<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: // %entry +; O0-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2]; +; O0-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0]; +; O0-NEXT: ld.b8 %rs1, [%rd1]; +; O0-NEXT: ld.b8 %rs2, [%rd1+1]; +; O0-NEXT: shl.b16 %rs3, %rs2, 8; +; O0-NEXT: or.b16 %rs4, %rs3, %rs1; +; O0-NEXT: cvt.u32.u16 %r1, %rs4; +; O0-NEXT: ld.s8 %rs5, [%rd1+2]; +; O0-NEXT: ld.b8 %rs6, [%rd2]; +; O0-NEXT: ld.b8 %rs7, [%rd2+1]; +; O0-NEXT: shl.b16 %rs8, %rs7, 8; +; O0-NEXT: or.b16 %rs9, %rs8, %rs6; +; O0-NEXT: cvt.u32.u16 %r2, %rs9; +; O0-NEXT: ld.s8 %rs10, [%rd2+2]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x9991U; +; O0-NEXT: cvt.u16.u32 %rs11, %r3; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x9991U; +; O0-NEXT: cvt.u16.u32 %rs12, %r4; +; O0-NEXT: rem.s16 %rs13, %rs12, %rs11; +; O0-NEXT: cvt.u32.u16 %r5, %rs13; +; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x8880U; +; O0-NEXT: cvt.u16.u32 %rs14, %r6; +; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x8880U; +; O0-NEXT: cvt.u16.u32 %rs15, %r7; +; O0-NEXT: rem.s16 %rs16, %rs15, %rs14; +; O0-NEXT: cvt.u32.u16 %r8, %rs16; +; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O0-NEXT: // implicit-def: %r11 +; O0-NEXT: // implicit-def: %r12 +; O0-NEXT: prmt.b32 %r10, %r11, %r12, 0x3340U; +; O0-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U; +; O0-NEXT: rem.s16 %rs17, %rs5, %rs10; +; O0-NEXT: cvt.u16.u32 %rs18, %r13; +; O0-NEXT: st.b8 [%rd3], %rs18; +; O0-NEXT: shr.u16 %rs19, %rs18, 8; +; O0-NEXT: st.b8 [%rd3+1], %rs19; +; O0-NEXT: st.b8 [%rd3+2], %rs17; +; O0-NEXT: ret; +; +; O3-LABEL: test_srem_v3i8( +; O3: { +; O3-NEXT: .reg .b16 %rs<20>; +; O3-NEXT: .reg .b32 %r<14>; +; O3-NEXT: .reg .b64 %rd<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: // %entry +; O3-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0]; +; O3-NEXT: ld.b8 %rs1, [%rd1]; +; O3-NEXT: ld.b8 %rs2, [%rd1+1]; +; O3-NEXT: shl.b16 %rs3, %rs2, 8; +; O3-NEXT: or.b16 %rs4, %rs3, %rs1; +; O3-NEXT: cvt.u32.u16 %r1, %rs4; +; O3-NEXT: ld.s8 %rs5, [%rd1+2]; +; O3-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1]; +; O3-NEXT: ld.b8 %rs6, [%rd2]; +; O3-NEXT: ld.b8 %rs7, [%rd2+1]; +; O3-NEXT: shl.b16 %rs8, %rs7, 8; +; O3-NEXT: or.b16 %rs9, %rs8, %rs6; +; O3-NEXT: cvt.u32.u16 %r2, %rs9; +; O3-NEXT: ld.s8 %rs10, [%rd2+2]; +; O3-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x9991U; +; O3-NEXT: cvt.u16.u32 %rs11, %r3; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x9991U; +; O3-NEXT: cvt.u16.u32 %rs12, %r4; +; O3-NEXT: rem.s16 %rs13, %rs12, %rs11; +; O3-NEXT: cvt.u32.u16 %r5, %rs13; +; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x8880U; +; O3-NEXT: cvt.u16.u32 %rs14, %r6; +; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x8880U; +; O3-NEXT: cvt.u16.u32 %rs15, %r7; +; O3-NEXT: rem.s16 %rs16, %rs15, %rs14; +; O3-NEXT: cvt.u32.u16 %r8, %rs16; +; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U; +; O3-NEXT: prmt.b32 %r10, %r11, %r12, 0x3340U; +; O3-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U; +; O3-NEXT: rem.s16 %rs17, %rs5, %rs10; +; O3-NEXT: st.b8 [%rd3+2], %rs17; +; O3-NEXT: cvt.u16.u32 %rs18, %r13; +; O3-NEXT: st.b8 [%rd3], %rs18; +; O3-NEXT: shr.u16 %rs19, %rs18, 8; +; O3-NEXT: st.b8 [%rd3+1], %rs19; +; O3-NEXT: ret; entry: %t57 = load <3 x i8>, ptr %a, align 1 %t59 = load <3 x i8>, ptr %b, align 1 @@ -1383,39 +2235,73 @@ entry: } define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) { -; CHECK-LABEL: test_sext_v4i1_to_v4i8( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<18>; -; CHECK-NEXT: .reg .b64 %rd<4>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: // %entry -; CHECK-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2]; -; CHECK-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1]; -; CHECK-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0]; -; CHECK-NEXT: ld.b32 %r1, [%rd1]; -; CHECK-NEXT: ld.b32 %r2, [%rd2]; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; -; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; -; CHECK-NEXT: setp.gt.u32 %p1, %r4, %r3; -; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; -; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; -; CHECK-NEXT: setp.gt.u32 %p2, %r6, %r5; -; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; -; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; -; CHECK-NEXT: setp.gt.u32 %p3, %r8, %r7; -; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; -; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; -; CHECK-NEXT: setp.gt.u32 %p4, %r10, %r9; -; CHECK-NEXT: selp.b32 %r11, -1, 0, %p4; -; CHECK-NEXT: selp.b32 %r12, -1, 0, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; -; CHECK-NEXT: selp.b32 %r14, -1, 0, %p2; -; CHECK-NEXT: selp.b32 %r15, -1, 0, %p1; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; -; CHECK-NEXT: st.b32 [%rd3], %r17; -; CHECK-NEXT: ret; +; O0-LABEL: test_sext_v4i1_to_v4i8( +; O0: { +; O0-NEXT: .reg .pred %p<5>; +; O0-NEXT: .reg .b32 %r<18>; +; O0-NEXT: .reg .b64 %rd<4>; +; O0-EMPTY: +; O0-NEXT: // %bb.0: // %entry +; O0-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2]; +; O0-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1]; +; O0-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0]; +; O0-NEXT: ld.b32 %r1, [%rd1]; +; O0-NEXT: ld.b32 %r2, [%rd2]; +; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O0-NEXT: setp.gt.u32 %p1, %r4, %r3; +; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O0-NEXT: setp.gt.u32 %p2, %r6, %r5; +; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O0-NEXT: setp.gt.u32 %p3, %r8, %r7; +; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O0-NEXT: setp.gt.u32 %p4, %r10, %r9; +; O0-NEXT: selp.b32 %r11, -1, 0, %p4; +; O0-NEXT: selp.b32 %r12, -1, 0, %p3; +; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O0-NEXT: selp.b32 %r14, -1, 0, %p2; +; O0-NEXT: selp.b32 %r15, -1, 0, %p1; +; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O0-NEXT: st.b32 [%rd3], %r17; +; O0-NEXT: ret; +; +; O3-LABEL: test_sext_v4i1_to_v4i8( +; O3: { +; O3-NEXT: .reg .pred %p<5>; +; O3-NEXT: .reg .b32 %r<18>; +; O3-NEXT: .reg .b64 %rd<4>; +; O3-EMPTY: +; O3-NEXT: // %bb.0: // %entry +; O3-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0]; +; O3-NEXT: ld.b32 %r1, [%rd1]; +; O3-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1]; +; O3-NEXT: ld.b32 %r2, [%rd2]; +; O3-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2]; +; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U; +; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U; +; O3-NEXT: setp.gt.u32 %p1, %r4, %r3; +; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U; +; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U; +; O3-NEXT: setp.gt.u32 %p2, %r6, %r5; +; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U; +; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U; +; O3-NEXT: setp.gt.u32 %p3, %r8, %r7; +; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U; +; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U; +; O3-NEXT: setp.gt.u32 %p4, %r10, %r9; +; O3-NEXT: selp.b32 %r11, -1, 0, %p4; +; O3-NEXT: selp.b32 %r12, -1, 0, %p3; +; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U; +; O3-NEXT: selp.b32 %r14, -1, 0, %p2; +; O3-NEXT: selp.b32 %r15, -1, 0, %p1; +; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U; +; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U; +; O3-NEXT: st.b32 [%rd3], %r17; +; O3-NEXT: ret; entry: %t1 = load <4 x i8>, ptr %a, align 4 %t2 = load <4 x i8>, ptr %b, align 4 diff --git a/llvm/test/CodeGen/NVPTX/prmt-const-folding.ll b/llvm/test/CodeGen/NVPTX/prmt-const-folding.ll new file mode 100644 index 0000000..7afead1 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/prmt-const-folding.ll @@ -0,0 +1,171 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 +; RUN: llc < %s -verify-machineinstrs -mcpu=sm_50 | FileCheck %s +; RUN: %if ptxas %{ llc < %s -verify-machineinstrs -mcpu=sm_50 | %ptxas-verify %} + +target triple = "nvptx64-nvidia-cuda" + +@g = global i32 0 + +define void @test_prmt_f4e() { +; CHECK-LABEL: test_prmt_f4e( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: st.volatile.global.b32 [g], 50462976; +; CHECK-NEXT: st.volatile.global.b32 [g], 67305985; +; CHECK-NEXT: st.volatile.global.b32 [g], 84148994; +; CHECK-NEXT: st.volatile.global.b32 [g], 100992003; +; CHECK-NEXT: st.volatile.global.b32 [g], 50462976; +; CHECK-NEXT: ret; + %v1 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x0) + store volatile i32 %v1, ptr @g + %v2 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x1) + store volatile i32 %v2, ptr @g + %v3 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x2) + store volatile i32 %v3, ptr @g + %v4 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x3) + store volatile i32 %v4, ptr @g + %v5 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x4) + store volatile i32 %v5, ptr @g + ret void +} + +define void @test_prmt_b4e() { +; CHECK-LABEL: test_prmt_b4e( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: st.volatile.global.b32 [g], 84281088; +; CHECK-NEXT: st.volatile.global.b32 [g], 101122049; +; CHECK-NEXT: st.volatile.global.b32 [g], 117440770; +; CHECK-NEXT: st.volatile.global.b32 [g], 66051; +; CHECK-NEXT: st.volatile.global.b32 [g], 84281088; +; CHECK-NEXT: ret; + %v1 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x0) + store volatile i32 %v1, ptr @g + %v2 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x1) + store volatile i32 %v2, ptr @g + %v3 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x2) + store volatile i32 %v3, ptr @g + %v4 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x3) + store volatile i32 %v4, ptr @g + %v5 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x4) + store volatile i32 %v5, ptr @g + ret void +} + +define void @test_prmt_ecl() { +; CHECK-LABEL: test_prmt_ecl( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: st.volatile.global.b32 [g], 50462976; +; CHECK-NEXT: st.volatile.global.b32 [g], 50462977; +; CHECK-NEXT: st.volatile.global.b32 [g], 50463234; +; CHECK-NEXT: st.volatile.global.b32 [g], 50529027; +; CHECK-NEXT: ret; + %v1 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x0) + store volatile i32 %v1, ptr @g + %v2 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x1) + store volatile i32 %v2, ptr @g + %v3 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x2) + store volatile i32 %v3, ptr @g + %v4 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x3) + store volatile i32 %v4, ptr @g + ret void +} + +define void @test_prmt_ecr() { +; CHECK-LABEL: test_prmt_ecr( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: st.volatile.global.b32 [g], 0; +; CHECK-NEXT: st.volatile.global.b32 [g], 16843008; +; CHECK-NEXT: st.volatile.global.b32 [g], 33685760; +; CHECK-NEXT: st.volatile.global.b32 [g], 50462976; +; CHECK-NEXT: ret; + %v1 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x0) + store volatile i32 %v1, ptr @g + %v2 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x1) + store volatile i32 %v2, ptr @g + %v3 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x2) + store volatile i32 %v3, ptr @g + %v4 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x3) + store volatile i32 %v4, ptr @g + ret void +} + +define void @test_prmt_rc8() { +; CHECK-LABEL: test_prmt_rc8( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: st.volatile.global.b32 [g], 0; +; CHECK-NEXT: st.volatile.global.b32 [g], 16843009; +; CHECK-NEXT: st.volatile.global.b32 [g], 33686018; +; CHECK-NEXT: st.volatile.global.b32 [g], 50529027; +; CHECK-NEXT: st.volatile.global.b32 [g], 0; +; CHECK-NEXT: ret; + %v1 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x0) + store volatile i32 %v1, ptr @g + %v2 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x1) + store volatile i32 %v2, ptr @g + %v3 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x2) + store volatile i32 %v3, ptr @g + %v4 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x3) + store volatile i32 %v4, ptr @g + %v5 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x4) + store volatile i32 %v5, ptr @g + ret void +} + +define void @test_prmt_rc16() { +; CHECK-LABEL: test_prmt_rc16( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: st.volatile.global.b32 [g], 16777472; +; CHECK-NEXT: st.volatile.global.b32 [g], 50463490; +; CHECK-NEXT: st.volatile.global.b32 [g], 16777472; +; CHECK-NEXT: ret; + %v1 = call i32 @llvm.nvvm.prmt.rc16(i32 u0x03020100, i32 u0x0) + store volatile i32 %v1, ptr @g + %v2 = call i32 @llvm.nvvm.prmt.rc16(i32 u0x03020100, i32 u0x1) + store volatile i32 %v2, ptr @g + %v3 = call i32 @llvm.nvvm.prmt.rc16(i32 u0x03020100, i32 u0x2) + store volatile i32 %v3, ptr @g + ret void +} + + +define void @test_prmt_basic() { +; CHECK-LABEL: test_prmt_basic( +; CHECK: { +; CHECK-EMPTY: +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: st.volatile.global.b32 [g], 66051; +; CHECK-NEXT: st.volatile.global.b32 [g], 117507841; +; CHECK-NEXT: st.volatile.global.b32 [g], 1146447479; +; CHECK-NEXT: st.volatile.global.b32 [g], 0; +; CHECK-NEXT: st.volatile.global.b32 [g], -16711936; +; CHECK-NEXT: ret; + %v1 = call i32 @llvm.nvvm.prmt(i32 u0x03020100, i32 u0x07060504, i32 u0x0123) + store volatile i32 %v1, ptr @g + %v2 = call i32 @llvm.nvvm.prmt(i32 u0x03020100, i32 u0x07060504, i32 u0x7171) + store volatile i32 %v2, ptr @g + %v3 = call i32 @llvm.nvvm.prmt(i32 u0x33221100, i32 u0x77665544, i32 u0x4567) + store volatile i32 %v3, ptr @g + %v4 = call i32 @llvm.nvvm.prmt(i32 u0x33221100, i32 u0x77665544, i32 u0xBA98) + store volatile i32 %v4, ptr @g + %v5 = call i32 @llvm.nvvm.prmt(i32 u0xF322F100, i32 u0x77665544, i32 u0xBA98) + store volatile i32 %v5, ptr @g + ret void +} |