[llvm] 535814f - [SelectionDAG] Support vector types in llvm.convert.from.arbitrary.fp expansion (#196185)
via llvm-commits
llvm-commits at lists.llvm.org
Thu May 7 02:14:58 PDT 2026
Author: Dmitry Sidorov
Date: 2026-05-07T11:14:52+02:00
New Revision: 535814f14f1d23f7ae1aa692c53d943bace471c4
URL: https://github.com/llvm/llvm-project/commit/535814f14f1d23f7ae1aa692c53d943bace471c4
DIFF: https://github.com/llvm/llvm-project/commit/535814f14f1d23f7ae1aa692c53d943bace471c4.diff
LOG: [SelectionDAG] Support vector types in llvm.convert.from.arbitrary.fp expansion (#196185)
Move the bit-twiddling expansion of ISD::CONVERT_FROM_ARBITRARY_FP from
LegalizeDAG to TargetLowering and make it working on vector destination
types.
Added:
Modified:
llvm/include/llvm/CodeGen/TargetLowering.h
llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp
llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
llvm/test/CodeGen/AMDGPU/arbitrary-fp-to-float.ll
llvm/test/CodeGen/NVPTX/arbitrary-fp-to-float.ll
llvm/test/CodeGen/X86/arbitrary-fp-to-float.ll
Removed:
################################################################################
diff --git a/llvm/include/llvm/CodeGen/TargetLowering.h b/llvm/include/llvm/CodeGen/TargetLowering.h
index 9b0bfa111f2d5..450c4bab5b734 100644
--- a/llvm/include/llvm/CodeGen/TargetLowering.h
+++ b/llvm/include/llvm/CodeGen/TargetLowering.h
@@ -5648,6 +5648,12 @@ class LLVM_ABI TargetLowering : public TargetLoweringBase {
/// \returns The expansion result
SDValue expandFCANONICALIZE(SDNode *Node, SelectionDAG &DAG) const;
+ /// Expand CONVERT_FROM_ARBITRARY_FP using bit manipulation.
+ /// \param Node Node to expand.
+ /// \returns The expansion result, or SDValue() if fails.
+ SDValue expandCONVERT_FROM_ARBITRARY_FP(SDNode *Node,
+ SelectionDAG &DAG) const;
+
/// Expand CTPOP nodes. Expands vector/scalar CTPOP nodes,
/// vector nodes can only succeed if all operations are legal/custom.
/// \param N Node to expand
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
index 9e8e0a684f2f6..4550d9e67d8f3 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
@@ -3555,231 +3555,10 @@ bool SelectionDAGLegalize::ExpandNode(SDNode *Node) {
// Float8E3M4, Float8E5M2FNUZ, Float8E4M3FNUZ, Float8E4M3B11FNUZ,
// Float8E8M0FNU.
EVT DstVT = Node->getValueType(0);
-
- SDValue IntVal = Node->getOperand(0);
- const uint64_t SemEnum = Node->getConstantOperandVal(1);
- const auto Sem = static_cast<APFloatBase::Semantics>(SemEnum);
-
- // Supported source formats.
- switch (Sem) {
- case APFloatBase::S_Float8E5M2:
- case APFloatBase::S_Float8E4M3FN:
- case APFloatBase::S_Float6E3M2FN:
- case APFloatBase::S_Float6E2M3FN:
- case APFloatBase::S_Float4E2M1FN:
- break;
- default:
- DAG.getContext()->emitError("CONVERT_FROM_ARBITRARY_FP: not implemented "
- "source format (semantics enum " +
- Twine(SemEnum) + ")");
+ if (SDValue Expanded = TLI.expandCONVERT_FROM_ARBITRARY_FP(Node, DAG))
+ Results.push_back(Expanded);
+ else
Results.push_back(DAG.getPOISON(DstVT));
- break;
- }
- if (!Results.empty())
- break;
-
- const fltSemantics &SrcSem = APFloatBase::EnumToSemantics(Sem);
-
- const unsigned SrcBits = APFloat::getSizeInBits(SrcSem);
- const unsigned SrcPrecision = APFloat::semanticsPrecision(SrcSem);
- const unsigned SrcMant = SrcPrecision - 1;
- const unsigned SrcExp = SrcBits - SrcMant - 1;
- const int SrcBias = 1 - APFloat::semanticsMinExponent(SrcSem);
-
- const fltNonfiniteBehavior NFBehavior = SrcSem.nonFiniteBehavior;
-
- // Destination format parameters.
- const fltSemantics &DstSem = DstVT.getFltSemantics();
-
- const unsigned DstBits = APFloat::getSizeInBits(DstSem);
- const unsigned DstMant = APFloat::semanticsPrecision(DstSem) - 1;
- const unsigned DstExpBits = DstBits - DstMant - 1;
- const int DstMinExp = APFloat::semanticsMinExponent(DstSem);
- const int DstBias = 1 - DstMinExp;
- const uint64_t DstExpAllOnes = (1ULL << DstExpBits) - 1;
-
- // Work in an integer type matching the destination float width.
- // Use zero-extend to preserve the raw bit-pattern.
- EVT IntVT = EVT::getIntegerVT(*DAG.getContext(), DstBits);
- SDValue Src = DAG.getZExtOrTrunc(IntVal, dl, IntVT);
-
- EVT SetCCVT = getSetCCResultType(IntVT);
-
- SDValue Zero = DAG.getConstant(0, dl, IntVT);
- SDValue One = DAG.getConstant(1, dl, IntVT);
-
- // Extract bit fields.
- const uint64_t MantMask = (SrcMant > 0) ? ((1ULL << SrcMant) - 1) : 0;
- const uint64_t ExpMask = (1ULL << SrcExp) - 1;
-
- SDValue MantField = DAG.getNode(ISD::AND, dl, IntVT, Src,
- DAG.getConstant(MantMask, dl, IntVT));
-
- SDValue ExpField =
- DAG.getNode(ISD::AND, dl, IntVT,
- DAG.getNode(ISD::SRL, dl, IntVT, Src,
- DAG.getShiftAmountConstant(SrcMant, IntVT, dl)),
- DAG.getConstant(ExpMask, dl, IntVT));
-
- SDValue SignBit =
- DAG.getNode(ISD::SRL, dl, IntVT, Src,
- DAG.getShiftAmountConstant(SrcBits - 1, IntVT, dl));
-
- // Precompute sign shifted to MSB of destination.
- SDValue SignShifted =
- DAG.getNode(ISD::SHL, dl, IntVT, SignBit,
- DAG.getShiftAmountConstant(DstBits - 1, IntVT, dl));
-
- // Classify the input value based on compile-time format properties.
- SDValue ExpAllOnes = DAG.getConstant(ExpMask, dl, IntVT);
- SDValue IsExpAllOnes =
- DAG.getSetCC(dl, SetCCVT, ExpField, ExpAllOnes, ISD::SETEQ);
- SDValue IsExpZero = DAG.getSetCC(dl, SetCCVT, ExpField, Zero, ISD::SETEQ);
- SDValue IsMantZero = DAG.getSetCC(dl, SetCCVT, MantField, Zero, ISD::SETEQ);
- SDValue IsMantNonZero =
- DAG.getSetCC(dl, SetCCVT, MantField, Zero, ISD::SETNE);
-
- // NaN detection.
- SDValue IsNaN;
- if (NFBehavior == fltNonfiniteBehavior::FiniteOnly) {
- // FiniteOnly formats (E2M1FN, E3M2FN, E2M3FN) never produce NaN.
- IsNaN = DAG.getBoolConstant(false, dl, SetCCVT, IntVT);
- } else if (NFBehavior == fltNonfiniteBehavior::IEEE754) {
- // E5M2 produces NaN when exp == all-ones AND mantissa != 0.
- IsNaN = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantNonZero);
- } else {
- // NanOnly + AllOnes (E4M3FN): NaN when all exp and mantissa bits are 1.
- assert(SrcSem.nanEncoding == fltNanEncoding::AllOnes);
- SDValue MantAllOnes = DAG.getConstant(MantMask, dl, IntVT);
- SDValue IsMantAllOnes =
- DAG.getSetCC(dl, SetCCVT, MantField, MantAllOnes, ISD::SETEQ);
- IsNaN = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantAllOnes);
- }
-
- // Inf detection.
- SDValue IsInf;
- if (NFBehavior == fltNonfiniteBehavior::IEEE754) {
- // E5M2: Inf when exp == all-ones AND mantissa == 0.
- IsInf = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantZero);
- } else {
- // NanOnly and FiniteOnly formats have no Inf representation.
- IsInf = DAG.getBoolConstant(false, dl, SetCCVT, IntVT);
- }
-
- // Zero detection.
- SDValue IsZero = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpZero, IsMantZero);
-
- // Denorm detection: exp == 0 AND mant != 0.
- SDValue IsDenorm =
- DAG.getNode(ISD::AND, dl, SetCCVT, IsExpZero, IsMantNonZero);
-
- // Normal value conversion.
- // dst_exp = exp_field + (DstBias - SrcBias)
- // dst_mant = mant << (DstMant - SrcMant)
- const int BiasAdjust = DstBias - SrcBias;
- SDValue NormDstExp = DAG.getNode(
- ISD::ADD, dl, IntVT, ExpField,
- DAG.getConstant(APInt(DstBits, BiasAdjust, true), dl, IntVT));
-
- SDValue NormDstMant;
- if (DstMant > SrcMant) {
- SDValue NormDstMantShift =
- DAG.getShiftAmountConstant(DstMant - SrcMant, IntVT, dl);
- NormDstMant =
- DAG.getNode(ISD::SHL, dl, IntVT, MantField, NormDstMantShift);
- } else {
- NormDstMant = MantField;
- }
-
- // Assemble normal result.
- SDValue DstMantShift = DAG.getShiftAmountConstant(DstMant, IntVT, dl);
- SDValue NormExpShifted =
- DAG.getNode(ISD::SHL, dl, IntVT, NormDstExp, DstMantShift);
- SDValue NormResult = DAG.getNode(
- ISD::OR, dl, IntVT,
- DAG.getNode(ISD::OR, dl, IntVT, SignShifted, NormExpShifted),
- NormDstMant);
-
- // Denormal value conversion.
- // For a denormal source (exp_field == 0, mant != 0), normalize by finding
- // the MSB position of mant using CTLZ, then compute the correct
- // exponent and mantissa for the destination format.
- SDValue DenormResult;
- {
- const unsigned IntVTBits = DstBits;
- SDValue LeadingZeros =
- DAG.getNode(ISD::CTLZ_ZERO_UNDEF, dl, IntVT, MantField);
-
- // dst_exp_denorm = (IntVTBits + DstBias - SrcBias - SrcMant) -
- // LeadingZeros
- const int DenormExpConst =
- (int)IntVTBits + DstBias - SrcBias - (int)SrcMant;
- SDValue DenormDstExp = DAG.getNode(
- ISD::SUB, dl, IntVT,
- DAG.getConstant(APInt(DstBits, DenormExpConst, true), dl, IntVT),
- LeadingZeros);
-
- // MSB position of the mantissa (0-indexed from LSB).
- SDValue MantMSB =
- DAG.getNode(ISD::SUB, dl, IntVT,
- DAG.getConstant(IntVTBits - 1, dl, IntVT), LeadingZeros);
-
- // leading_one = 1 << MantMSB
- SDValue LeadingOne = DAG.getNode(ISD::SHL, dl, IntVT, One, MantMSB);
-
- // frac = mant XOR leading_one (strip the implicit 1)
- SDValue Frac = DAG.getNode(ISD::XOR, dl, IntVT, MantField, LeadingOne);
-
- // shift_amount = DstMant - MantMSB
- // = DstMant - (IntVTBits - 1 - LeadingZeros)
- // = LeadingZeros - (IntVTBits - 1 - DstMant)
- const unsigned ShiftSub = IntVTBits - 1 - DstMant; // always >= 0
- SDValue ShiftAmount = DAG.getNode(ISD::SUB, dl, IntVT, LeadingZeros,
- DAG.getConstant(ShiftSub, dl, IntVT));
-
- SDValue DenormDstMant =
- DAG.getNode(ISD::SHL, dl, IntVT, Frac, ShiftAmount);
-
- // Assemble denorm as sign | (denorm_dst_exp << DstMant) | denorm_dst_mant
- SDValue DenormExpShifted =
- DAG.getNode(ISD::SHL, dl, IntVT, DenormDstExp, DstMantShift);
- DenormResult = DAG.getNode(
- ISD::OR, dl, IntVT,
- DAG.getNode(ISD::OR, dl, IntVT, SignShifted, DenormExpShifted),
- DenormDstMant);
- }
-
- // Select between normal and denorm paths.
- SDValue FiniteResult =
- DAG.getSelect(dl, IntVT, IsDenorm, DenormResult, NormResult);
-
- // Build special-value results.
- // NaN -> canonical quiet NaN: sign=0, exp=all-ones, qNaN bit set.
- // Encoding: (DstExpAllOnes << DstMant) | (1 << (DstMant - 1))
- const uint64_t QNaNBit = (DstMant > 0) ? (1ULL << (DstMant - 1)) : 0;
- SDValue NaNResult =
- DAG.getConstant((DstExpAllOnes << DstMant) | QNaNBit, dl, IntVT);
-
- // Inf -> destination Inf.
- // sign | (DstExpAllOnes << DstMant)
- SDValue InfResult =
- DAG.getNode(ISD::OR, dl, IntVT, SignShifted,
- DAG.getConstant(DstExpAllOnes << DstMant, dl, IntVT));
-
- // Zero -> signed zero.
- // Sign bit only.
- SDValue ZeroResult = SignShifted;
-
- // Final selection goes in order: NaN takes priority, then Inf, then Zero.
- SDValue Result = FiniteResult;
- Result = DAG.getSelect(dl, IntVT, IsZero, ZeroResult, Result);
- Result = DAG.getSelect(dl, IntVT, IsInf, InfResult, Result);
- Result = DAG.getSelect(dl, IntVT, IsNaN, NaNResult, Result);
-
- // Bitcast integer result to destination float type.
- Result = DAG.getNode(ISD::BITCAST, dl, DstVT, Result);
-
- Results.push_back(Result);
break;
}
case ISD::FCANONICALIZE: {
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp
index 1aa21fc636a9d..f054ab3fdc911 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp
@@ -1411,6 +1411,12 @@ void VectorLegalizer::Expand(SDNode *Node, SmallVectorImpl<SDValue> &Results) {
return;
}
break;
+ case ISD::CONVERT_FROM_ARBITRARY_FP:
+ if (SDValue Expanded = TLI.expandCONVERT_FROM_ARBITRARY_FP(Node, DAG))
+ Results.push_back(Expanded);
+ else
+ Results.push_back(DAG.getPOISON(Node->getValueType(0)));
+ return;
case ISD::MASKED_UDIV:
case ISD::MASKED_SDIV:
case ISD::MASKED_UREM:
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index 867fd87d3c137..9b05582e471ff 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -9024,6 +9024,196 @@ SDValue TargetLowering::expandFCANONICALIZE(SDNode *Node,
return Mul;
}
+SDValue
+TargetLowering::expandCONVERT_FROM_ARBITRARY_FP(SDNode *Node,
+ SelectionDAG &DAG) const {
+ SDLoc dl(Node);
+ EVT DstVT = Node->getValueType(0);
+ EVT DstScalarVT = DstVT.getScalarType();
+
+ SDValue IntVal = Node->getOperand(0);
+ const uint64_t SemEnum = Node->getConstantOperandVal(1);
+ const auto Sem = static_cast<APFloatBase::Semantics>(SemEnum);
+
+ // Supported source formats.
+ switch (Sem) {
+ case APFloatBase::S_Float8E5M2:
+ case APFloatBase::S_Float8E4M3FN:
+ case APFloatBase::S_Float6E3M2FN:
+ case APFloatBase::S_Float6E2M3FN:
+ case APFloatBase::S_Float4E2M1FN:
+ break;
+ default:
+ DAG.getContext()->emitError("CONVERT_FROM_ARBITRARY_FP: not implemented "
+ "source format (semantics enum " +
+ Twine(SemEnum) + ")");
+ return SDValue();
+ }
+
+ const fltSemantics &SrcSem = APFloatBase::EnumToSemantics(Sem);
+ const unsigned SrcBits = APFloat::getSizeInBits(SrcSem);
+ const unsigned SrcPrecision = APFloat::semanticsPrecision(SrcSem);
+ const unsigned SrcMant = SrcPrecision - 1;
+ const unsigned SrcExp = SrcBits - SrcMant - 1;
+ const int SrcBias = 1 - APFloat::semanticsMinExponent(SrcSem);
+ const fltNonfiniteBehavior NFBehavior = SrcSem.nonFiniteBehavior;
+
+ // Destination format parameters.
+ const fltSemantics &DstSem = DstScalarVT.getFltSemantics();
+ const unsigned DstBits = APFloat::getSizeInBits(DstSem);
+ const unsigned DstMant = APFloat::semanticsPrecision(DstSem) - 1;
+ const unsigned DstExpBits = DstBits - DstMant - 1;
+ const int DstMinExp = APFloat::semanticsMinExponent(DstSem);
+ const int DstBias = 1 - DstMinExp;
+ const uint64_t DstExpAllOnes = (1ULL << DstExpBits) - 1;
+
+ // Work in an integer type matching the destination float width.
+ EVT IntScalarVT = EVT::getIntegerVT(*DAG.getContext(), DstBits);
+ EVT IntVT = DstVT.isVector()
+ ? EVT::getVectorVT(*DAG.getContext(), IntScalarVT,
+ DstVT.getVectorElementCount())
+ : IntScalarVT;
+
+ SDValue Src = DAG.getZExtOrTrunc(IntVal, dl, IntVT);
+
+ EVT SetCCVT =
+ getSetCCResultType(DAG.getDataLayout(), *DAG.getContext(), IntVT);
+
+ SDValue Zero = DAG.getConstant(0, dl, IntVT);
+ SDValue One = DAG.getConstant(1, dl, IntVT);
+
+ // Extract bit fields.
+ const uint64_t MantMask = (SrcMant > 0) ? ((1ULL << SrcMant) - 1) : 0;
+ const uint64_t ExpMask = (1ULL << SrcExp) - 1;
+
+ SDValue MantField = DAG.getNode(ISD::AND, dl, IntVT, Src,
+ DAG.getConstant(MantMask, dl, IntVT));
+
+ SDValue ExpField =
+ DAG.getNode(ISD::AND, dl, IntVT,
+ DAG.getNode(ISD::SRL, dl, IntVT, Src,
+ DAG.getShiftAmountConstant(SrcMant, IntVT, dl)),
+ DAG.getConstant(ExpMask, dl, IntVT));
+
+ SDValue SignBit =
+ DAG.getNode(ISD::SRL, dl, IntVT, Src,
+ DAG.getShiftAmountConstant(SrcBits - 1, IntVT, dl));
+
+ SDValue SignShifted =
+ DAG.getNode(ISD::SHL, dl, IntVT, SignBit,
+ DAG.getShiftAmountConstant(DstBits - 1, IntVT, dl));
+
+ // Classify the input.
+ SDValue ExpAllOnes = DAG.getConstant(ExpMask, dl, IntVT);
+ SDValue IsExpAllOnes =
+ DAG.getSetCC(dl, SetCCVT, ExpField, ExpAllOnes, ISD::SETEQ);
+ SDValue IsExpZero = DAG.getSetCC(dl, SetCCVT, ExpField, Zero, ISD::SETEQ);
+ SDValue IsMantZero = DAG.getSetCC(dl, SetCCVT, MantField, Zero, ISD::SETEQ);
+ SDValue IsMantNonZero =
+ DAG.getSetCC(dl, SetCCVT, MantField, Zero, ISD::SETNE);
+
+ SDValue IsNaN;
+ if (NFBehavior == fltNonfiniteBehavior::FiniteOnly) {
+ IsNaN = DAG.getBoolConstant(false, dl, SetCCVT, IntVT);
+ } else if (NFBehavior == fltNonfiniteBehavior::IEEE754) {
+ IsNaN = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantNonZero);
+ } else {
+ assert(SrcSem.nanEncoding == fltNanEncoding::AllOnes);
+ SDValue MantAllOnes = DAG.getConstant(MantMask, dl, IntVT);
+ SDValue IsMantAllOnes =
+ DAG.getSetCC(dl, SetCCVT, MantField, MantAllOnes, ISD::SETEQ);
+ IsNaN = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantAllOnes);
+ }
+
+ SDValue IsInf;
+ if (NFBehavior == fltNonfiniteBehavior::IEEE754)
+ IsInf = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantZero);
+ else
+ IsInf = DAG.getBoolConstant(false, dl, SetCCVT, IntVT);
+
+ SDValue IsZero = DAG.getNode(ISD::AND, dl, SetCCVT, IsExpZero, IsMantZero);
+ SDValue IsDenorm =
+ DAG.getNode(ISD::AND, dl, SetCCVT, IsExpZero, IsMantNonZero);
+
+ // Normal value conversion.
+ const int BiasAdjust = DstBias - SrcBias;
+ SDValue NormDstExp =
+ DAG.getNode(ISD::ADD, dl, IntVT, ExpField,
+ DAG.getConstant(APInt(DstBits, BiasAdjust, true), dl, IntVT));
+
+ SDValue NormDstMant;
+ if (DstMant > SrcMant) {
+ SDValue NormDstMantShift =
+ DAG.getShiftAmountConstant(DstMant - SrcMant, IntVT, dl);
+ NormDstMant = DAG.getNode(ISD::SHL, dl, IntVT, MantField, NormDstMantShift);
+ } else {
+ NormDstMant = MantField;
+ }
+
+ SDValue DstMantShift = DAG.getShiftAmountConstant(DstMant, IntVT, dl);
+ SDValue NormExpShifted =
+ DAG.getNode(ISD::SHL, dl, IntVT, NormDstExp, DstMantShift);
+ SDValue NormResult =
+ DAG.getNode(ISD::OR, dl, IntVT,
+ DAG.getNode(ISD::OR, dl, IntVT, SignShifted, NormExpShifted),
+ NormDstMant);
+
+ // Denormal value conversion.
+ SDValue DenormResult;
+ {
+ const unsigned IntVTBits = DstBits;
+ SDValue LeadingZeros =
+ DAG.getNode(ISD::CTLZ_ZERO_UNDEF, dl, IntVT, MantField);
+
+ const int DenormExpConst =
+ (int)IntVTBits + DstBias - SrcBias - (int)SrcMant;
+ SDValue DenormDstExp = DAG.getNode(
+ ISD::SUB, dl, IntVT,
+ DAG.getConstant(APInt(DstBits, DenormExpConst, true), dl, IntVT),
+ LeadingZeros);
+
+ SDValue MantMSB =
+ DAG.getNode(ISD::SUB, dl, IntVT,
+ DAG.getConstant(IntVTBits - 1, dl, IntVT), LeadingZeros);
+
+ SDValue LeadingOne = DAG.getNode(ISD::SHL, dl, IntVT, One, MantMSB);
+ SDValue Frac = DAG.getNode(ISD::XOR, dl, IntVT, MantField, LeadingOne);
+
+ const unsigned ShiftSub = IntVTBits - 1 - DstMant;
+ SDValue ShiftAmount = DAG.getNode(ISD::SUB, dl, IntVT, LeadingZeros,
+ DAG.getConstant(ShiftSub, dl, IntVT));
+
+ SDValue DenormDstMant = DAG.getNode(ISD::SHL, dl, IntVT, Frac, ShiftAmount);
+
+ SDValue DenormExpShifted =
+ DAG.getNode(ISD::SHL, dl, IntVT, DenormDstExp, DstMantShift);
+ DenormResult = DAG.getNode(
+ ISD::OR, dl, IntVT,
+ DAG.getNode(ISD::OR, dl, IntVT, SignShifted, DenormExpShifted),
+ DenormDstMant);
+ }
+
+ SDValue FiniteResult =
+ DAG.getSelect(dl, IntVT, IsDenorm, DenormResult, NormResult);
+
+ const uint64_t QNaNBit = (DstMant > 0) ? (1ULL << (DstMant - 1)) : 0;
+ SDValue NaNResult =
+ DAG.getConstant((DstExpAllOnes << DstMant) | QNaNBit, dl, IntVT);
+
+ SDValue InfResult =
+ DAG.getNode(ISD::OR, dl, IntVT, SignShifted,
+ DAG.getConstant(DstExpAllOnes << DstMant, dl, IntVT));
+
+ SDValue ZeroResult = SignShifted;
+
+ SDValue Result = FiniteResult;
+ Result = DAG.getSelect(dl, IntVT, IsZero, ZeroResult, Result);
+ Result = DAG.getSelect(dl, IntVT, IsInf, InfResult, Result);
+ Result = DAG.getSelect(dl, IntVT, IsNaN, NaNResult, Result);
+
+ return DAG.getNode(ISD::BITCAST, dl, DstVT, Result);
+}
+
bool TargetLowering::expandFP_TO_SINT(SDNode *Node, SDValue &Result,
SelectionDAG &DAG) const {
unsigned OpNo = Node->isStrictFPOpcode() ? 1 : 0;
diff --git a/llvm/test/CodeGen/AMDGPU/arbitrary-fp-to-float.ll b/llvm/test/CodeGen/AMDGPU/arbitrary-fp-to-float.ll
index e7bb8825fec05..35100626e0d14 100644
--- a/llvm/test/CodeGen/AMDGPU/arbitrary-fp-to-float.ll
+++ b/llvm/test/CodeGen/AMDGPU/arbitrary-fp-to-float.ll
@@ -7,8 +7,12 @@ declare float @llvm.convert.from.arbitrary.fp.f32.i8(i8, metadata)
declare float @llvm.convert.from.arbitrary.fp.f32.i6(i6, metadata)
declare float @llvm.convert.from.arbitrary.fp.f32.i4(i4, metadata)
declare <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4>, metadata)
+declare <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8>, metadata)
+declare <3 x float> @llvm.convert.from.arbitrary.fp.v3f32.v3i8(<3 x i8>, metadata)
+declare <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8>, metadata)
declare half @llvm.convert.from.arbitrary.fp.f16.i8(i8, metadata)
+declare <2 x half> @llvm.convert.from.arbitrary.fp.v2f16.v2i8(<2 x i8>, metadata)
declare double @llvm.convert.from.arbitrary.fp.f64.i8(i8, metadata)
; Float8E5M2
@@ -557,8 +561,8 @@ define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 1, v0
; CHECK-NEXT: v_add_u32_e32 v6, 0.5, v6
; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
-; CHECK-NEXT: s_and_b64 vcc, vcc, s[4:5]
-; CHECK-NEXT: v_or_b32_e32 v4, v4, v5
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_or_b32_e32 v4, v5, v4
; CHECK-NEXT: v_cndmask_b32_e32 v0, v6, v8, vcc
; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4
; CHECK-NEXT: v_and_b32_e32 v4, 0xffff, v1
@@ -583,8 +587,8 @@ define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 1, v1
; CHECK-NEXT: v_add_u32_e32 v6, 0.5, v6
; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
-; CHECK-NEXT: s_and_b64 vcc, vcc, s[4:5]
-; CHECK-NEXT: v_or_b32_e32 v4, v4, v5
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_or_b32_e32 v4, v5, v4
; CHECK-NEXT: v_cndmask_b32_e32 v1, v6, v8, vcc
; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4
; CHECK-NEXT: v_and_b32_e32 v4, 0xffff, v2
@@ -609,8 +613,8 @@ define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 1, v2
; CHECK-NEXT: v_add_u32_e32 v6, 0.5, v6
; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
-; CHECK-NEXT: s_and_b64 vcc, vcc, s[4:5]
-; CHECK-NEXT: v_or_b32_e32 v4, v4, v5
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_or_b32_e32 v4, v5, v4
; CHECK-NEXT: v_cndmask_b32_e32 v2, v6, v8, vcc
; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4
; CHECK-NEXT: v_and_b32_e32 v4, 0xffff, v3
@@ -635,8 +639,8 @@ define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 1, v3
; CHECK-NEXT: v_add_u32_e32 v6, 0.5, v6
; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
-; CHECK-NEXT: s_and_b64 vcc, vcc, s[4:5]
-; CHECK-NEXT: v_or_b32_e32 v4, v4, v5
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_or_b32_e32 v4, v5, v4
; CHECK-NEXT: v_cndmask_b32_e32 v3, v6, v8, vcc
; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v4
; CHECK-NEXT: v_cndmask_b32_e32 v3, v3, v7, vcc
@@ -644,3 +648,586 @@ define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
%r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4> %x, metadata !"Float4E2M1FN")
ret <4 x float> %r
}
+
+; <2 x i8> -> <2 x float> Float8E4M3FN
+define <2 x float> @from_f8e4m3fn_v2f32(<2 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v2f32:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; CHECK-NEXT: v_and_b32_e32 v2, 7, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, 24, v0
+; CHECK-NEXT: v_bfe_u32 v0, v0, 3, 4
+; CHECK-NEXT: v_lshlrev_b32_e32 v3, 20, v2
+; CHECK-NEXT: v_and_b32_e32 v4, 0x80000000, v4
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 23, v0
+; CHECK-NEXT: v_or3_b32 v3, v5, v4, v3
+; CHECK-NEXT: v_ffbh_u32_e32 v5, v2
+; CHECK-NEXT: v_sub_u32_e32 v6, 31, v5
+; CHECK-NEXT: v_lshlrev_b32_e64 v6, v6, 1
+; CHECK-NEXT: v_xor_b32_e32 v6, v2, v6
+; CHECK-NEXT: v_add_u32_e32 v7, -8, v5
+; CHECK-NEXT: v_sub_u32_e32 v5, 0x95, v5
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, v7, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 23, v5
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v2
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v0
+; CHECK-NEXT: v_add_u32_e32 v3, 0x3c000000, v3
+; CHECK-NEXT: v_or3_b32 v5, v4, v5, v6
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v3, v3, v5, vcc
+; CHECK-NEXT: v_or_b32_e32 v5, v0, v2
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cndmask_b32_e32 v3, v3, v4, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v2
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v0
+; CHECK-NEXT: v_mov_b32_e32 v2, 0x7fc00000
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v3, v2, vcc
+; CHECK-NEXT: v_and_b32_e32 v3, 7, v1
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 24, v1
+; CHECK-NEXT: v_bfe_u32 v1, v1, 3, 4
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, 20, v3
+; CHECK-NEXT: v_and_b32_e32 v5, 0x80000000, v5
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 23, v1
+; CHECK-NEXT: v_or3_b32 v4, v6, v5, v4
+; CHECK-NEXT: v_ffbh_u32_e32 v6, v3
+; CHECK-NEXT: v_sub_u32_e32 v7, 31, v6
+; CHECK-NEXT: v_lshlrev_b32_e64 v7, v7, 1
+; CHECK-NEXT: v_xor_b32_e32 v7, v3, v7
+; CHECK-NEXT: v_add_u32_e32 v8, -8, v6
+; CHECK-NEXT: v_sub_u32_e32 v6, 0x95, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, v8, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 23, v6
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v3
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_add_u32_e32 v4, 0x3c000000, v4
+; CHECK-NEXT: v_or3_b32 v6, v5, v6, v7
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v4, v4, v6, vcc
+; CHECK-NEXT: v_or_b32_e32 v6, v1, v3
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v6
+; CHECK-NEXT: v_cndmask_b32_e32 v4, v4, v5, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v3
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v1
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v1, v4, v2, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8> %x, metadata !"Float8E4M3FN")
+ ret <2 x float> %r
+}
+
+; <2 x i8> -> <2 x float> Float8E5M2
+define <2 x float> @from_f8e5m2_v2f32(<2 x i8> %x) {
+; CHECK-LABEL: from_f8e5m2_v2f32:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; CHECK-NEXT: v_and_b32_e32 v2, 3, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, 24, v0
+; CHECK-NEXT: v_bfe_u32 v0, v0, 2, 5
+; CHECK-NEXT: v_lshlrev_b32_e32 v3, 21, v2
+; CHECK-NEXT: v_and_b32_e32 v4, 0x80000000, v4
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 23, v0
+; CHECK-NEXT: v_or3_b32 v3, v5, v4, v3
+; CHECK-NEXT: v_ffbh_u32_e32 v5, v2
+; CHECK-NEXT: v_sub_u32_e32 v6, 31, v5
+; CHECK-NEXT: v_lshlrev_b32_e64 v6, v6, 1
+; CHECK-NEXT: v_xor_b32_e32 v6, v2, v6
+; CHECK-NEXT: v_add_u32_e32 v7, -8, v5
+; CHECK-NEXT: v_sub_u32_e32 v5, 0x8e, v5
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, v7, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 23, v5
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v2
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v0
+; CHECK-NEXT: v_add_u32_e32 v3, 0x38000000, v3
+; CHECK-NEXT: v_or3_b32 v5, v4, v5, v6
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, v5, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v5, v0, v2
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, v4, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 31, v0
+; CHECK-NEXT: v_or_b32_e32 v0, 0x7f800000, v4
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v3, v0, s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v3, 3, v1
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 24, v1
+; CHECK-NEXT: v_bfe_u32 v1, v1, 2, 5
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, 21, v3
+; CHECK-NEXT: v_and_b32_e32 v5, 0x80000000, v5
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 23, v1
+; CHECK-NEXT: v_or3_b32 v4, v6, v5, v4
+; CHECK-NEXT: v_ffbh_u32_e32 v6, v3
+; CHECK-NEXT: v_sub_u32_e32 v7, 31, v6
+; CHECK-NEXT: v_lshlrev_b32_e64 v7, v7, 1
+; CHECK-NEXT: v_mov_b32_e32 v2, 0x7fc00000
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_xor_b32_e32 v7, v3, v7
+; CHECK-NEXT: v_add_u32_e32 v8, -8, v6
+; CHECK-NEXT: v_sub_u32_e32 v6, 0x8e, v6
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v0, v2, vcc
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, v8, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 23, v6
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v3
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_add_u32_e32 v4, 0x38000000, v4
+; CHECK-NEXT: v_or3_b32 v6, v5, v6, v7
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v4, v6, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v6, v1, v3
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v6
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v4, v5, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v3
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 31, v1
+; CHECK-NEXT: v_or_b32_e32 v1, 0x7f800000, v5
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v1, v4, v1, s[4:5]
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v1, v1, v2, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8> %x, metadata !"Float8E5M2")
+ ret <2 x float> %r
+}
+
+; <3 x i8> -> <3 x float> Float8E4M3FN
+define <3 x float> @from_f8e4m3fn_v3f32(<3 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v3f32:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; CHECK-NEXT: v_and_b32_e32 v3, 7, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 24, v0
+; CHECK-NEXT: v_bfe_u32 v0, v0, 3, 4
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, 20, v3
+; CHECK-NEXT: v_and_b32_e32 v5, 0x80000000, v5
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 23, v0
+; CHECK-NEXT: v_or3_b32 v4, v6, v5, v4
+; CHECK-NEXT: v_ffbh_u32_e32 v6, v3
+; CHECK-NEXT: v_sub_u32_e32 v7, 31, v6
+; CHECK-NEXT: v_lshlrev_b32_e64 v7, v7, 1
+; CHECK-NEXT: v_xor_b32_e32 v7, v3, v7
+; CHECK-NEXT: v_add_u32_e32 v8, -8, v6
+; CHECK-NEXT: v_sub_u32_e32 v6, 0x95, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, v8, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 23, v6
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v3
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v0
+; CHECK-NEXT: v_add_u32_e32 v4, 0x3c000000, v4
+; CHECK-NEXT: v_or3_b32 v6, v5, v6, v7
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v4, v4, v6, vcc
+; CHECK-NEXT: v_or_b32_e32 v6, v0, v3
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v6
+; CHECK-NEXT: v_cndmask_b32_e32 v4, v4, v5, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v3
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v0
+; CHECK-NEXT: v_mov_b32_e32 v3, 0x7fc00000
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v4, v3, vcc
+; CHECK-NEXT: v_and_b32_e32 v4, 7, v1
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 24, v1
+; CHECK-NEXT: v_bfe_u32 v1, v1, 3, 4
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 20, v4
+; CHECK-NEXT: v_and_b32_e32 v6, 0x80000000, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v1
+; CHECK-NEXT: v_or3_b32 v5, v7, v6, v5
+; CHECK-NEXT: v_ffbh_u32_e32 v7, v4
+; CHECK-NEXT: v_sub_u32_e32 v8, 31, v7
+; CHECK-NEXT: v_lshlrev_b32_e64 v8, v8, 1
+; CHECK-NEXT: v_xor_b32_e32 v8, v4, v8
+; CHECK-NEXT: v_add_u32_e32 v9, -8, v7
+; CHECK-NEXT: v_sub_u32_e32 v7, 0x95, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, v9, v8
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v7
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_add_u32_e32 v5, 0x3c000000, v5
+; CHECK-NEXT: v_or3_b32 v7, v6, v7, v8
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v5, v5, v7, vcc
+; CHECK-NEXT: v_or_b32_e32 v7, v1, v4
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v7
+; CHECK-NEXT: v_cndmask_b32_e32 v5, v5, v6, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v1
+; CHECK-NEXT: v_and_b32_e32 v2, 0xffff, v2
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_and_b32_e32 v4, 7, v2
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 24, v2
+; CHECK-NEXT: v_bfe_u32 v2, v2, 3, 4
+; CHECK-NEXT: v_cndmask_b32_e32 v1, v5, v3, vcc
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 20, v4
+; CHECK-NEXT: v_and_b32_e32 v6, 0x80000000, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v2
+; CHECK-NEXT: v_or3_b32 v5, v7, v6, v5
+; CHECK-NEXT: v_ffbh_u32_e32 v7, v4
+; CHECK-NEXT: v_sub_u32_e32 v8, 31, v7
+; CHECK-NEXT: v_lshlrev_b32_e64 v8, v8, 1
+; CHECK-NEXT: v_xor_b32_e32 v8, v4, v8
+; CHECK-NEXT: v_add_u32_e32 v9, -8, v7
+; CHECK-NEXT: v_sub_u32_e32 v7, 0x95, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, v9, v8
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v7
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_add_u32_e32 v5, 0x3c000000, v5
+; CHECK-NEXT: v_or3_b32 v7, v6, v7, v8
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v5, v5, v7, vcc
+; CHECK-NEXT: v_or_b32_e32 v7, v2, v4
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v7
+; CHECK-NEXT: v_cndmask_b32_e32 v5, v5, v6, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v2
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v2, v5, v3, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call <3 x float> @llvm.convert.from.arbitrary.fp.v3f32.v3i8(<3 x i8> %x, metadata !"Float8E4M3FN")
+ ret <3 x float> %r
+}
+
+; <4 x i8> -> <4 x float> Float8E4M3FN
+define <4 x float> @from_f8e4m3fn_v4f32(<4 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v4f32:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; CHECK-NEXT: v_and_b32_e32 v4, 7, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 24, v0
+; CHECK-NEXT: v_bfe_u32 v0, v0, 3, 4
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 20, v4
+; CHECK-NEXT: v_and_b32_e32 v6, 0x80000000, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v0
+; CHECK-NEXT: v_or3_b32 v5, v7, v6, v5
+; CHECK-NEXT: v_ffbh_u32_e32 v7, v4
+; CHECK-NEXT: v_sub_u32_e32 v8, 31, v7
+; CHECK-NEXT: v_lshlrev_b32_e64 v8, v8, 1
+; CHECK-NEXT: v_xor_b32_e32 v8, v4, v8
+; CHECK-NEXT: v_add_u32_e32 v9, -8, v7
+; CHECK-NEXT: v_sub_u32_e32 v7, 0x95, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, v9, v8
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v7
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v0
+; CHECK-NEXT: v_add_u32_e32 v5, 0x3c000000, v5
+; CHECK-NEXT: v_or3_b32 v7, v6, v7, v8
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v5, v5, v7, vcc
+; CHECK-NEXT: v_or_b32_e32 v7, v0, v4
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v7
+; CHECK-NEXT: v_cndmask_b32_e32 v5, v5, v6, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v0
+; CHECK-NEXT: v_mov_b32_e32 v4, 0x7fc00000
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v5, v4, vcc
+; CHECK-NEXT: v_and_b32_e32 v5, 7, v1
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 24, v1
+; CHECK-NEXT: v_bfe_u32 v1, v1, 3, 4
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 20, v5
+; CHECK-NEXT: v_and_b32_e32 v7, 0x80000000, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v1
+; CHECK-NEXT: v_or3_b32 v6, v8, v7, v6
+; CHECK-NEXT: v_ffbh_u32_e32 v8, v5
+; CHECK-NEXT: v_sub_u32_e32 v9, 31, v8
+; CHECK-NEXT: v_lshlrev_b32_e64 v9, v9, 1
+; CHECK-NEXT: v_xor_b32_e32 v9, v5, v9
+; CHECK-NEXT: v_add_u32_e32 v10, -8, v8
+; CHECK-NEXT: v_sub_u32_e32 v8, 0x95, v8
+; CHECK-NEXT: v_lshlrev_b32_e32 v9, v10, v9
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v8
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_add_u32_e32 v6, 0x3c000000, v6
+; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v6, v6, v8, vcc
+; CHECK-NEXT: v_or_b32_e32 v8, v1, v5
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v8
+; CHECK-NEXT: v_cndmask_b32_e32 v6, v6, v7, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v1
+; CHECK-NEXT: v_and_b32_e32 v2, 0xffff, v2
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_and_b32_e32 v5, 7, v2
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 24, v2
+; CHECK-NEXT: v_bfe_u32 v2, v2, 3, 4
+; CHECK-NEXT: v_cndmask_b32_e32 v1, v6, v4, vcc
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 20, v5
+; CHECK-NEXT: v_and_b32_e32 v7, 0x80000000, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v2
+; CHECK-NEXT: v_or3_b32 v6, v8, v7, v6
+; CHECK-NEXT: v_ffbh_u32_e32 v8, v5
+; CHECK-NEXT: v_sub_u32_e32 v9, 31, v8
+; CHECK-NEXT: v_lshlrev_b32_e64 v9, v9, 1
+; CHECK-NEXT: v_xor_b32_e32 v9, v5, v9
+; CHECK-NEXT: v_add_u32_e32 v10, -8, v8
+; CHECK-NEXT: v_sub_u32_e32 v8, 0x95, v8
+; CHECK-NEXT: v_lshlrev_b32_e32 v9, v10, v9
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v8
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_add_u32_e32 v6, 0x3c000000, v6
+; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v6, v6, v8, vcc
+; CHECK-NEXT: v_or_b32_e32 v8, v2, v5
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v8
+; CHECK-NEXT: v_cndmask_b32_e32 v6, v6, v7, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v2
+; CHECK-NEXT: v_and_b32_e32 v3, 0xffff, v3
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_and_b32_e32 v5, 7, v3
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 24, v3
+; CHECK-NEXT: v_bfe_u32 v3, v3, 3, 4
+; CHECK-NEXT: v_cndmask_b32_e32 v2, v6, v4, vcc
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 20, v5
+; CHECK-NEXT: v_and_b32_e32 v7, 0x80000000, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v3
+; CHECK-NEXT: v_or3_b32 v6, v8, v7, v6
+; CHECK-NEXT: v_ffbh_u32_e32 v8, v5
+; CHECK-NEXT: v_sub_u32_e32 v9, 31, v8
+; CHECK-NEXT: v_lshlrev_b32_e64 v9, v9, 1
+; CHECK-NEXT: v_xor_b32_e32 v9, v5, v9
+; CHECK-NEXT: v_add_u32_e32 v10, -8, v8
+; CHECK-NEXT: v_sub_u32_e32 v8, 0x95, v8
+; CHECK-NEXT: v_lshlrev_b32_e32 v9, v10, v9
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v8
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v3
+; CHECK-NEXT: v_add_u32_e32 v6, 0x3c000000, v6
+; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v6, v6, v8, vcc
+; CHECK-NEXT: v_or_b32_e32 v8, v3, v5
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v8
+; CHECK-NEXT: v_cndmask_b32_e32 v6, v6, v7, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 7, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 15, v3
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v3, v6, v4, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> %x, metadata !"Float8E4M3FN")
+ ret <4 x float> %r
+}
+
+; <4 x i8> -> <4 x float> Float8E5M2
+define <4 x float> @from_f8e5m2_v4f32(<4 x i8> %x) {
+; CHECK-LABEL: from_f8e5m2_v4f32:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; CHECK-NEXT: v_and_b32_e32 v4, 3, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 24, v0
+; CHECK-NEXT: v_bfe_u32 v0, v0, 2, 5
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 21, v4
+; CHECK-NEXT: v_and_b32_e32 v6, 0x80000000, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v0
+; CHECK-NEXT: v_or3_b32 v5, v7, v6, v5
+; CHECK-NEXT: v_ffbh_u32_e32 v7, v4
+; CHECK-NEXT: v_sub_u32_e32 v8, 31, v7
+; CHECK-NEXT: v_lshlrev_b32_e64 v8, v8, 1
+; CHECK-NEXT: v_xor_b32_e32 v8, v4, v8
+; CHECK-NEXT: v_add_u32_e32 v9, -8, v7
+; CHECK-NEXT: v_sub_u32_e32 v7, 0x8e, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, v9, v8
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 23, v7
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v0
+; CHECK-NEXT: v_add_u32_e32 v5, 0x38000000, v5
+; CHECK-NEXT: v_or3_b32 v7, v6, v7, v8
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, v7, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v7, v0, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, v6, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 31, v0
+; CHECK-NEXT: v_or_b32_e32 v0, 0x7f800000, v6
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v5, v0, s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v5, 3, v1
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 24, v1
+; CHECK-NEXT: v_bfe_u32 v1, v1, 2, 5
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 21, v5
+; CHECK-NEXT: v_and_b32_e32 v7, 0x80000000, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v1
+; CHECK-NEXT: v_or3_b32 v6, v8, v7, v6
+; CHECK-NEXT: v_ffbh_u32_e32 v8, v5
+; CHECK-NEXT: v_sub_u32_e32 v9, 31, v8
+; CHECK-NEXT: v_lshlrev_b32_e64 v9, v9, 1
+; CHECK-NEXT: v_mov_b32_e32 v4, 0x7fc00000
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_xor_b32_e32 v9, v5, v9
+; CHECK-NEXT: v_add_u32_e32 v10, -8, v8
+; CHECK-NEXT: v_sub_u32_e32 v8, 0x8e, v8
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v0, v4, vcc
+; CHECK-NEXT: v_lshlrev_b32_e32 v9, v10, v9
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v8
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_add_u32_e32 v6, 0x38000000, v6
+; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v6, v6, v8, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v8, v1, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v6, v6, v7, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 31, v1
+; CHECK-NEXT: v_and_b32_e32 v2, 0xffff, v2
+; CHECK-NEXT: v_or_b32_e32 v1, 0x7f800000, v7
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v5, 3, v2
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 24, v2
+; CHECK-NEXT: v_bfe_u32 v2, v2, 2, 5
+; CHECK-NEXT: v_cndmask_b32_e64 v1, v6, v1, s[4:5]
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 21, v5
+; CHECK-NEXT: v_and_b32_e32 v7, 0x80000000, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v2
+; CHECK-NEXT: v_or3_b32 v6, v8, v7, v6
+; CHECK-NEXT: v_ffbh_u32_e32 v8, v5
+; CHECK-NEXT: v_sub_u32_e32 v9, 31, v8
+; CHECK-NEXT: v_lshlrev_b32_e64 v9, v9, 1
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_xor_b32_e32 v9, v5, v9
+; CHECK-NEXT: v_add_u32_e32 v10, -8, v8
+; CHECK-NEXT: v_sub_u32_e32 v8, 0x8e, v8
+; CHECK-NEXT: v_cndmask_b32_e32 v1, v1, v4, vcc
+; CHECK-NEXT: v_lshlrev_b32_e32 v9, v10, v9
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v8
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_add_u32_e32 v6, 0x38000000, v6
+; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v6, v6, v8, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v8, v2, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v6, v6, v7, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 31, v2
+; CHECK-NEXT: v_and_b32_e32 v3, 0xffff, v3
+; CHECK-NEXT: v_or_b32_e32 v2, 0x7f800000, v7
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v5, 3, v3
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 24, v3
+; CHECK-NEXT: v_bfe_u32 v3, v3, 2, 5
+; CHECK-NEXT: v_cndmask_b32_e64 v2, v6, v2, s[4:5]
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 21, v5
+; CHECK-NEXT: v_and_b32_e32 v7, 0x80000000, v7
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v3
+; CHECK-NEXT: v_or3_b32 v6, v8, v7, v6
+; CHECK-NEXT: v_ffbh_u32_e32 v8, v5
+; CHECK-NEXT: v_sub_u32_e32 v9, 31, v8
+; CHECK-NEXT: v_lshlrev_b32_e64 v9, v9, 1
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_xor_b32_e32 v9, v5, v9
+; CHECK-NEXT: v_add_u32_e32 v10, -8, v8
+; CHECK-NEXT: v_sub_u32_e32 v8, 0x8e, v8
+; CHECK-NEXT: v_cndmask_b32_e32 v2, v2, v4, vcc
+; CHECK-NEXT: v_lshlrev_b32_e32 v9, v10, v9
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, 23, v8
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v3
+; CHECK-NEXT: v_add_u32_e32 v6, 0x38000000, v6
+; CHECK-NEXT: v_or3_b32 v8, v7, v8, v9
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v6, v6, v8, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v8, v3, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v6, v6, v7, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 31, v3
+; CHECK-NEXT: v_or_b32_e32 v3, 0x7f800000, v7
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v6, v3, s[4:5]
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v3, v3, v4, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> %x, metadata !"Float8E5M2")
+ ret <4 x float> %r
+}
+
+; <2 x i8> -> <2 x half> Float8E4M3FN
+define <2 x half> @from_f8e4m3fn_v2f16(<2 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v2f16:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v3, 7, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, 16, v3
+; CHECK-NEXT: v_ffbh_u32_e32 v4, v4
+; CHECK-NEXT: v_sub_u16_e32 v6, 15, v4
+; CHECK-NEXT: v_lshlrev_b16_e32 v2, 8, v0
+; CHECK-NEXT: v_sub_u16_e32 v5, 21, v4
+; CHECK-NEXT: v_lshlrev_b16_e64 v6, v6, 1
+; CHECK-NEXT: v_and_b32_e32 v2, 0xffff8000, v2
+; CHECK-NEXT: v_lshlrev_b16_e32 v5, 10, v5
+; CHECK-NEXT: v_xor_b32_e32 v6, v3, v6
+; CHECK-NEXT: v_add_u16_e32 v4, -5, v4
+; CHECK-NEXT: v_or_b32_e32 v5, v2, v5
+; CHECK-NEXT: v_lshlrev_b16_e32 v4, v4, v6
+; CHECK-NEXT: v_bfe_u32 v0, v0, 3, 4
+; CHECK-NEXT: v_or_b32_e32 v4, v5, v4
+; CHECK-NEXT: v_lshlrev_b16_e32 v5, 10, v0
+; CHECK-NEXT: v_or_b32_e32 v5, v5, v2
+; CHECK-NEXT: v_lshlrev_b16_e32 v6, 7, v3
+; CHECK-NEXT: v_cmp_ne_u16_e32 vcc, 0, v3
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v0
+; CHECK-NEXT: v_or_b32_e32 v5, v5, v6
+; CHECK-NEXT: v_add_u16_e32 v5, 0x2000, v5
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v4, v5, v4, vcc
+; CHECK-NEXT: v_or_b32_e32 v5, v0, v3
+; CHECK-NEXT: v_cmp_eq_u16_e32 vcc, 0, v5
+; CHECK-NEXT: v_cndmask_b32_e32 v2, v4, v2, vcc
+; CHECK-NEXT: v_and_b32_e32 v4, 7, v1
+; CHECK-NEXT: v_lshlrev_b32_e32 v5, 16, v4
+; CHECK-NEXT: v_ffbh_u32_e32 v5, v5
+; CHECK-NEXT: v_sub_u16_e32 v7, 15, v5
+; CHECK-NEXT: v_cmp_eq_u16_e32 vcc, 7, v3
+; CHECK-NEXT: v_lshlrev_b16_e32 v3, 8, v1
+; CHECK-NEXT: v_sub_u16_e32 v6, 21, v5
+; CHECK-NEXT: v_lshlrev_b16_e64 v7, v7, 1
+; CHECK-NEXT: v_and_b32_e32 v3, 0xffff8000, v3
+; CHECK-NEXT: v_lshlrev_b16_e32 v6, 10, v6
+; CHECK-NEXT: v_xor_b32_e32 v7, v4, v7
+; CHECK-NEXT: v_add_u16_e32 v5, -5, v5
+; CHECK-NEXT: v_or_b32_e32 v6, v3, v6
+; CHECK-NEXT: v_lshlrev_b16_e32 v5, v5, v7
+; CHECK-NEXT: v_bfe_u32 v1, v1, 3, 4
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 15, v0
+; CHECK-NEXT: v_or_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_lshlrev_b16_e32 v6, 10, v1
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7e00
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_or_b32_e32 v6, v6, v3
+; CHECK-NEXT: v_lshlrev_b16_e32 v7, 7, v4
+; CHECK-NEXT: v_cndmask_b32_e32 v2, v2, v0, vcc
+; CHECK-NEXT: v_cmp_ne_u16_e32 vcc, 0, v4
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_or_b32_e32 v6, v6, v7
+; CHECK-NEXT: v_add_u16_e32 v6, 0x2000, v6
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v5, v6, v5, vcc
+; CHECK-NEXT: v_or_b32_e32 v6, v1, v4
+; CHECK-NEXT: v_cmp_eq_u16_e32 vcc, 0, v6
+; CHECK-NEXT: v_cndmask_b32_e32 v3, v5, v3, vcc
+; CHECK-NEXT: v_cmp_eq_u16_e32 vcc, 7, v4
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 15, v1
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v3, v0, vcc
+; CHECK-NEXT: s_mov_b32 s4, 0x5040100
+; CHECK-NEXT: v_perm_b32 v0, v0, v2, s4
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call <2 x half> @llvm.convert.from.arbitrary.fp.v2f16.v2i8(<2 x i8> %x, metadata !"Float8E4M3FN")
+ ret <2 x half> %r
+}
diff --git a/llvm/test/CodeGen/NVPTX/arbitrary-fp-to-float.ll b/llvm/test/CodeGen/NVPTX/arbitrary-fp-to-float.ll
index aff1bd385308d..238406ccfda2c 100644
--- a/llvm/test/CodeGen/NVPTX/arbitrary-fp-to-float.ll
+++ b/llvm/test/CodeGen/NVPTX/arbitrary-fp-to-float.ll
@@ -6,7 +6,11 @@
declare float @llvm.convert.from.arbitrary.fp.f32.i8(i8, metadata)
declare float @llvm.convert.from.arbitrary.fp.f32.i6(i6, metadata)
declare float @llvm.convert.from.arbitrary.fp.f32.i4(i4, metadata)
+declare <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8>, metadata)
+declare <3 x float> @llvm.convert.from.arbitrary.fp.v3f32.v3i8(<3 x i8>, metadata)
declare <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4>, metadata)
+declare <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8>, metadata)
+declare <2 x half> @llvm.convert.from.arbitrary.fp.v2f16.v2i8(<2 x i8>, metadata)
declare half @llvm.convert.from.arbitrary.fp.f16.i8(i8, metadata)
declare double @llvm.convert.from.arbitrary.fp.f64.i8(i8, metadata)
@@ -759,3 +763,664 @@ define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
%r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4> %x, metadata !"Float4E2M1FN")
ret <4 x float> %r
}
+
+; <2 x i8> -> <2 x float> Float8E4M3FN
+define <2 x float> @from_f8e4m3fn_v2f32(<2 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v2f32(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<11>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<54>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [from_f8e4m3fn_v2f32_param_0];
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
+; CHECK-NEXT: shl.b32 %r2, %r1, 24;
+; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
+; CHECK-NEXT: and.b32 %r4, %r1, 7;
+; CHECK-NEXT: clz.b32 %r5, %r4;
+; CHECK-NEXT: sub.s32 %r6, 149, %r5;
+; CHECK-NEXT: shl.b32 %r7, %r6, 23;
+; CHECK-NEXT: or.b32 %r8, %r3, %r7;
+; CHECK-NEXT: sub.s32 %r9, 31, %r5;
+; CHECK-NEXT: mov.b32 %r10, 1;
+; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
+; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
+; CHECK-NEXT: add.s32 %r13, %r5, -8;
+; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
+; CHECK-NEXT: or.b32 %r15, %r8, %r14;
+; CHECK-NEXT: bfe.u32 %r16, %r1, 3, 4;
+; CHECK-NEXT: shl.b32 %r17, %r16, 23;
+; CHECK-NEXT: or.b32 %r18, %r17, %r3;
+; CHECK-NEXT: shl.b32 %r19, %r4, 20;
+; CHECK-NEXT: or.b32 %r20, %r18, %r19;
+; CHECK-NEXT: add.s32 %r21, %r20, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
+; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
+; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
+; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
+; CHECK-NEXT: or.b32 %r24, %r16, %r4;
+; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
+; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
+; CHECK-NEXT: setp.eq.b32 %p4, %r4, 7;
+; CHECK-NEXT: selp.b32 %r26, 2143289344, %r25, %p4;
+; CHECK-NEXT: setp.eq.b32 %p5, %r16, 15;
+; CHECK-NEXT: selp.b32 %r27, %r26, %r25, %p5;
+; CHECK-NEXT: cvt.u32.u16 %r28, %rs1;
+; CHECK-NEXT: shl.b32 %r29, %r28, 24;
+; CHECK-NEXT: and.b32 %r30, %r29, -2147483648;
+; CHECK-NEXT: and.b32 %r31, %r28, 7;
+; CHECK-NEXT: clz.b32 %r32, %r31;
+; CHECK-NEXT: sub.s32 %r33, 149, %r32;
+; CHECK-NEXT: shl.b32 %r34, %r33, 23;
+; CHECK-NEXT: or.b32 %r35, %r30, %r34;
+; CHECK-NEXT: sub.s32 %r36, 31, %r32;
+; CHECK-NEXT: shl.b32 %r37, %r10, %r36;
+; CHECK-NEXT: xor.b32 %r38, %r31, %r37;
+; CHECK-NEXT: add.s32 %r39, %r32, -8;
+; CHECK-NEXT: shl.b32 %r40, %r38, %r39;
+; CHECK-NEXT: or.b32 %r41, %r35, %r40;
+; CHECK-NEXT: bfe.u32 %r42, %r28, 3, 4;
+; CHECK-NEXT: shl.b32 %r43, %r42, 23;
+; CHECK-NEXT: or.b32 %r44, %r43, %r30;
+; CHECK-NEXT: shl.b32 %r45, %r31, 20;
+; CHECK-NEXT: or.b32 %r46, %r44, %r45;
+; CHECK-NEXT: add.s32 %r47, %r46, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p6, %r31, 0;
+; CHECK-NEXT: selp.b32 %r48, %r41, %r47, %p6;
+; CHECK-NEXT: setp.eq.b32 %p7, %r42, 0;
+; CHECK-NEXT: selp.b32 %r49, %r48, %r47, %p7;
+; CHECK-NEXT: or.b32 %r50, %r42, %r31;
+; CHECK-NEXT: setp.eq.b32 %p8, %r50, 0;
+; CHECK-NEXT: selp.b32 %r51, %r30, %r49, %p8;
+; CHECK-NEXT: setp.eq.b32 %p9, %r31, 7;
+; CHECK-NEXT: selp.b32 %r52, 2143289344, %r51, %p9;
+; CHECK-NEXT: setp.eq.b32 %p10, %r42, 15;
+; CHECK-NEXT: selp.b32 %r53, %r52, %r51, %p10;
+; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r53, %r27};
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8> %x, metadata !"Float8E4M3FN")
+ ret <2 x float> %r
+}
+
+; <2 x i8> -> <2 x float> Float8E5M2
+define <2 x float> @from_f8e5m2_v2f32(<2 x i8> %x) {
+; CHECK-LABEL: from_f8e5m2_v2f32(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<11>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<60>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [from_f8e5m2_v2f32_param_0];
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
+; CHECK-NEXT: shl.b32 %r2, %r1, 24;
+; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
+; CHECK-NEXT: and.b32 %r4, %r1, 3;
+; CHECK-NEXT: clz.b32 %r5, %r4;
+; CHECK-NEXT: sub.s32 %r6, 142, %r5;
+; CHECK-NEXT: shl.b32 %r7, %r6, 23;
+; CHECK-NEXT: or.b32 %r8, %r3, %r7;
+; CHECK-NEXT: sub.s32 %r9, 31, %r5;
+; CHECK-NEXT: mov.b32 %r10, 1;
+; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
+; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
+; CHECK-NEXT: add.s32 %r13, %r5, -8;
+; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
+; CHECK-NEXT: or.b32 %r15, %r8, %r14;
+; CHECK-NEXT: bfe.u32 %r16, %r1, 2, 5;
+; CHECK-NEXT: shl.b32 %r17, %r16, 23;
+; CHECK-NEXT: or.b32 %r18, %r17, %r3;
+; CHECK-NEXT: shl.b32 %r19, %r4, 21;
+; CHECK-NEXT: or.b32 %r20, %r18, %r19;
+; CHECK-NEXT: add.s32 %r21, %r20, 939524096;
+; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
+; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
+; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
+; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
+; CHECK-NEXT: or.b32 %r24, %r16, %r4;
+; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
+; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
+; CHECK-NEXT: setp.eq.b32 %p4, %r4, 0;
+; CHECK-NEXT: or.b32 %r26, %r3, 2139095040;
+; CHECK-NEXT: selp.b32 %r27, %r26, %r25, %p4;
+; CHECK-NEXT: setp.eq.b32 %p5, %r16, 31;
+; CHECK-NEXT: selp.b32 %r28, %r27, %r25, %p5;
+; CHECK-NEXT: selp.b32 %r29, 2143289344, %r28, %p1;
+; CHECK-NEXT: selp.b32 %r30, %r29, %r28, %p5;
+; CHECK-NEXT: cvt.u32.u16 %r31, %rs1;
+; CHECK-NEXT: shl.b32 %r32, %r31, 24;
+; CHECK-NEXT: and.b32 %r33, %r32, -2147483648;
+; CHECK-NEXT: and.b32 %r34, %r31, 3;
+; CHECK-NEXT: clz.b32 %r35, %r34;
+; CHECK-NEXT: sub.s32 %r36, 142, %r35;
+; CHECK-NEXT: shl.b32 %r37, %r36, 23;
+; CHECK-NEXT: or.b32 %r38, %r33, %r37;
+; CHECK-NEXT: sub.s32 %r39, 31, %r35;
+; CHECK-NEXT: shl.b32 %r40, %r10, %r39;
+; CHECK-NEXT: xor.b32 %r41, %r34, %r40;
+; CHECK-NEXT: add.s32 %r42, %r35, -8;
+; CHECK-NEXT: shl.b32 %r43, %r41, %r42;
+; CHECK-NEXT: or.b32 %r44, %r38, %r43;
+; CHECK-NEXT: bfe.u32 %r45, %r31, 2, 5;
+; CHECK-NEXT: shl.b32 %r46, %r45, 23;
+; CHECK-NEXT: or.b32 %r47, %r46, %r33;
+; CHECK-NEXT: shl.b32 %r48, %r34, 21;
+; CHECK-NEXT: or.b32 %r49, %r47, %r48;
+; CHECK-NEXT: add.s32 %r50, %r49, 939524096;
+; CHECK-NEXT: setp.ne.b32 %p6, %r34, 0;
+; CHECK-NEXT: selp.b32 %r51, %r44, %r50, %p6;
+; CHECK-NEXT: setp.eq.b32 %p7, %r45, 0;
+; CHECK-NEXT: selp.b32 %r52, %r51, %r50, %p7;
+; CHECK-NEXT: or.b32 %r53, %r45, %r34;
+; CHECK-NEXT: setp.eq.b32 %p8, %r53, 0;
+; CHECK-NEXT: selp.b32 %r54, %r33, %r52, %p8;
+; CHECK-NEXT: setp.eq.b32 %p9, %r34, 0;
+; CHECK-NEXT: or.b32 %r55, %r33, 2139095040;
+; CHECK-NEXT: selp.b32 %r56, %r55, %r54, %p9;
+; CHECK-NEXT: setp.eq.b32 %p10, %r45, 31;
+; CHECK-NEXT: selp.b32 %r57, %r56, %r54, %p10;
+; CHECK-NEXT: selp.b32 %r58, 2143289344, %r57, %p6;
+; CHECK-NEXT: selp.b32 %r59, %r58, %r57, %p10;
+; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r59, %r30};
+; CHECK-NEXT: ret;
+ %r = call <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8> %x, metadata !"Float8E5M2")
+ ret <2 x float> %r
+}
+
+; <3 x i8> -> <3 x float> Float8E4M3FN
+define <3 x float> @from_f8e4m3fn_v3f32(<3 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v3f32(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<16>;
+; CHECK-NEXT: .reg .b32 %r<81>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [from_f8e4m3fn_v3f32_param_0];
+; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;
+; CHECK-NEXT: shl.b32 %r3, %r2, 24;
+; CHECK-NEXT: and.b32 %r4, %r3, -2147483648;
+; CHECK-NEXT: and.b32 %r5, %r2, 7;
+; CHECK-NEXT: clz.b32 %r6, %r5;
+; CHECK-NEXT: sub.s32 %r7, 149, %r6;
+; CHECK-NEXT: shl.b32 %r8, %r7, 23;
+; CHECK-NEXT: or.b32 %r9, %r4, %r8;
+; CHECK-NEXT: sub.s32 %r10, 31, %r6;
+; CHECK-NEXT: mov.b32 %r11, 1;
+; CHECK-NEXT: shl.b32 %r12, %r11, %r10;
+; CHECK-NEXT: xor.b32 %r13, %r5, %r12;
+; CHECK-NEXT: add.s32 %r14, %r6, -8;
+; CHECK-NEXT: shl.b32 %r15, %r13, %r14;
+; CHECK-NEXT: or.b32 %r16, %r9, %r15;
+; CHECK-NEXT: bfe.u32 %r17, %r2, 3, 4;
+; CHECK-NEXT: shl.b32 %r18, %r17, 23;
+; CHECK-NEXT: or.b32 %r19, %r18, %r4;
+; CHECK-NEXT: shl.b32 %r20, %r5, 20;
+; CHECK-NEXT: or.b32 %r21, %r19, %r20;
+; CHECK-NEXT: add.s32 %r22, %r21, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p1, %r5, 0;
+; CHECK-NEXT: selp.b32 %r23, %r16, %r22, %p1;
+; CHECK-NEXT: setp.eq.b32 %p2, %r17, 0;
+; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p2;
+; CHECK-NEXT: or.b32 %r25, %r17, %r5;
+; CHECK-NEXT: setp.eq.b32 %p3, %r25, 0;
+; CHECK-NEXT: selp.b32 %r26, %r4, %r24, %p3;
+; CHECK-NEXT: setp.eq.b32 %p4, %r5, 7;
+; CHECK-NEXT: selp.b32 %r27, 2143289344, %r26, %p4;
+; CHECK-NEXT: setp.eq.b32 %p5, %r17, 15;
+; CHECK-NEXT: selp.b32 %r28, %r27, %r26, %p5;
+; CHECK-NEXT: prmt.b32 %r29, %r1, 0, 0x7771U;
+; CHECK-NEXT: shl.b32 %r30, %r29, 24;
+; CHECK-NEXT: and.b32 %r31, %r30, -2147483648;
+; CHECK-NEXT: and.b32 %r32, %r29, 7;
+; CHECK-NEXT: clz.b32 %r33, %r32;
+; CHECK-NEXT: sub.s32 %r34, 149, %r33;
+; CHECK-NEXT: shl.b32 %r35, %r34, 23;
+; CHECK-NEXT: or.b32 %r36, %r31, %r35;
+; CHECK-NEXT: sub.s32 %r37, 31, %r33;
+; CHECK-NEXT: shl.b32 %r38, %r11, %r37;
+; CHECK-NEXT: xor.b32 %r39, %r32, %r38;
+; CHECK-NEXT: add.s32 %r40, %r33, -8;
+; CHECK-NEXT: shl.b32 %r41, %r39, %r40;
+; CHECK-NEXT: or.b32 %r42, %r36, %r41;
+; CHECK-NEXT: bfe.u32 %r43, %r29, 3, 4;
+; CHECK-NEXT: shl.b32 %r44, %r43, 23;
+; CHECK-NEXT: or.b32 %r45, %r44, %r31;
+; CHECK-NEXT: shl.b32 %r46, %r32, 20;
+; CHECK-NEXT: or.b32 %r47, %r45, %r46;
+; CHECK-NEXT: add.s32 %r48, %r47, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p6, %r32, 0;
+; CHECK-NEXT: selp.b32 %r49, %r42, %r48, %p6;
+; CHECK-NEXT: setp.eq.b32 %p7, %r43, 0;
+; CHECK-NEXT: selp.b32 %r50, %r49, %r48, %p7;
+; CHECK-NEXT: or.b32 %r51, %r43, %r32;
+; CHECK-NEXT: setp.eq.b32 %p8, %r51, 0;
+; CHECK-NEXT: selp.b32 %r52, %r31, %r50, %p8;
+; CHECK-NEXT: setp.eq.b32 %p9, %r32, 7;
+; CHECK-NEXT: selp.b32 %r53, 2143289344, %r52, %p9;
+; CHECK-NEXT: setp.eq.b32 %p10, %r43, 15;
+; CHECK-NEXT: selp.b32 %r54, %r53, %r52, %p10;
+; CHECK-NEXT: prmt.b32 %r55, %r1, 0, 0x7770U;
+; CHECK-NEXT: shl.b32 %r56, %r55, 24;
+; CHECK-NEXT: and.b32 %r57, %r56, -2147483648;
+; CHECK-NEXT: and.b32 %r58, %r55, 7;
+; CHECK-NEXT: clz.b32 %r59, %r58;
+; CHECK-NEXT: sub.s32 %r60, 149, %r59;
+; CHECK-NEXT: shl.b32 %r61, %r60, 23;
+; CHECK-NEXT: or.b32 %r62, %r57, %r61;
+; CHECK-NEXT: sub.s32 %r63, 31, %r59;
+; CHECK-NEXT: shl.b32 %r64, %r11, %r63;
+; CHECK-NEXT: xor.b32 %r65, %r58, %r64;
+; CHECK-NEXT: add.s32 %r66, %r59, -8;
+; CHECK-NEXT: shl.b32 %r67, %r65, %r66;
+; CHECK-NEXT: or.b32 %r68, %r62, %r67;
+; CHECK-NEXT: bfe.u32 %r69, %r55, 3, 4;
+; CHECK-NEXT: shl.b32 %r70, %r69, 23;
+; CHECK-NEXT: or.b32 %r71, %r70, %r57;
+; CHECK-NEXT: shl.b32 %r72, %r58, 20;
+; CHECK-NEXT: or.b32 %r73, %r71, %r72;
+; CHECK-NEXT: add.s32 %r74, %r73, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p11, %r58, 0;
+; CHECK-NEXT: selp.b32 %r75, %r68, %r74, %p11;
+; CHECK-NEXT: setp.eq.b32 %p12, %r69, 0;
+; CHECK-NEXT: selp.b32 %r76, %r75, %r74, %p12;
+; CHECK-NEXT: or.b32 %r77, %r69, %r58;
+; CHECK-NEXT: setp.eq.b32 %p13, %r77, 0;
+; CHECK-NEXT: selp.b32 %r78, %r57, %r76, %p13;
+; CHECK-NEXT: setp.eq.b32 %p14, %r58, 7;
+; CHECK-NEXT: selp.b32 %r79, 2143289344, %r78, %p14;
+; CHECK-NEXT: setp.eq.b32 %p15, %r69, 15;
+; CHECK-NEXT: selp.b32 %r80, %r79, %r78, %p15;
+; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r80, %r54};
+; CHECK-NEXT: st.param.b32 [func_retval0+8], %r28;
+; CHECK-NEXT: ret;
+ %r = call <3 x float> @llvm.convert.from.arbitrary.fp.v3f32.v3i8(<3 x i8> %x, metadata !"Float8E4M3FN")
+ ret <3 x float> %r
+}
+
+; <4 x i8> -> <4 x float> Float8E4M3FN
+define <4 x float> @from_f8e4m3fn_v4f32(<4 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v4f32(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<21>;
+; CHECK-NEXT: .reg .b32 %r<107>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [from_f8e4m3fn_v4f32_param_0];
+; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;
+; CHECK-NEXT: shl.b32 %r3, %r2, 24;
+; CHECK-NEXT: and.b32 %r4, %r3, -2147483648;
+; CHECK-NEXT: and.b32 %r5, %r2, 7;
+; CHECK-NEXT: clz.b32 %r6, %r5;
+; CHECK-NEXT: sub.s32 %r7, 149, %r6;
+; CHECK-NEXT: shl.b32 %r8, %r7, 23;
+; CHECK-NEXT: or.b32 %r9, %r4, %r8;
+; CHECK-NEXT: sub.s32 %r10, 31, %r6;
+; CHECK-NEXT: mov.b32 %r11, 1;
+; CHECK-NEXT: shl.b32 %r12, %r11, %r10;
+; CHECK-NEXT: xor.b32 %r13, %r5, %r12;
+; CHECK-NEXT: add.s32 %r14, %r6, -8;
+; CHECK-NEXT: shl.b32 %r15, %r13, %r14;
+; CHECK-NEXT: or.b32 %r16, %r9, %r15;
+; CHECK-NEXT: bfe.u32 %r17, %r2, 3, 4;
+; CHECK-NEXT: shl.b32 %r18, %r17, 23;
+; CHECK-NEXT: or.b32 %r19, %r18, %r4;
+; CHECK-NEXT: shl.b32 %r20, %r5, 20;
+; CHECK-NEXT: or.b32 %r21, %r19, %r20;
+; CHECK-NEXT: add.s32 %r22, %r21, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p1, %r5, 0;
+; CHECK-NEXT: selp.b32 %r23, %r16, %r22, %p1;
+; CHECK-NEXT: setp.eq.b32 %p2, %r17, 0;
+; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p2;
+; CHECK-NEXT: or.b32 %r25, %r17, %r5;
+; CHECK-NEXT: setp.eq.b32 %p3, %r25, 0;
+; CHECK-NEXT: selp.b32 %r26, %r4, %r24, %p3;
+; CHECK-NEXT: setp.eq.b32 %p4, %r5, 7;
+; CHECK-NEXT: selp.b32 %r27, 2143289344, %r26, %p4;
+; CHECK-NEXT: setp.eq.b32 %p5, %r17, 15;
+; CHECK-NEXT: selp.b32 %r28, %r27, %r26, %p5;
+; CHECK-NEXT: prmt.b32 %r29, %r1, 0, 0x7772U;
+; CHECK-NEXT: shl.b32 %r30, %r29, 24;
+; CHECK-NEXT: and.b32 %r31, %r30, -2147483648;
+; CHECK-NEXT: and.b32 %r32, %r29, 7;
+; CHECK-NEXT: clz.b32 %r33, %r32;
+; CHECK-NEXT: sub.s32 %r34, 149, %r33;
+; CHECK-NEXT: shl.b32 %r35, %r34, 23;
+; CHECK-NEXT: or.b32 %r36, %r31, %r35;
+; CHECK-NEXT: sub.s32 %r37, 31, %r33;
+; CHECK-NEXT: shl.b32 %r38, %r11, %r37;
+; CHECK-NEXT: xor.b32 %r39, %r32, %r38;
+; CHECK-NEXT: add.s32 %r40, %r33, -8;
+; CHECK-NEXT: shl.b32 %r41, %r39, %r40;
+; CHECK-NEXT: or.b32 %r42, %r36, %r41;
+; CHECK-NEXT: bfe.u32 %r43, %r29, 3, 4;
+; CHECK-NEXT: shl.b32 %r44, %r43, 23;
+; CHECK-NEXT: or.b32 %r45, %r44, %r31;
+; CHECK-NEXT: shl.b32 %r46, %r32, 20;
+; CHECK-NEXT: or.b32 %r47, %r45, %r46;
+; CHECK-NEXT: add.s32 %r48, %r47, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p6, %r32, 0;
+; CHECK-NEXT: selp.b32 %r49, %r42, %r48, %p6;
+; CHECK-NEXT: setp.eq.b32 %p7, %r43, 0;
+; CHECK-NEXT: selp.b32 %r50, %r49, %r48, %p7;
+; CHECK-NEXT: or.b32 %r51, %r43, %r32;
+; CHECK-NEXT: setp.eq.b32 %p8, %r51, 0;
+; CHECK-NEXT: selp.b32 %r52, %r31, %r50, %p8;
+; CHECK-NEXT: setp.eq.b32 %p9, %r32, 7;
+; CHECK-NEXT: selp.b32 %r53, 2143289344, %r52, %p9;
+; CHECK-NEXT: setp.eq.b32 %p10, %r43, 15;
+; CHECK-NEXT: selp.b32 %r54, %r53, %r52, %p10;
+; CHECK-NEXT: prmt.b32 %r55, %r1, 0, 0x7771U;
+; CHECK-NEXT: shl.b32 %r56, %r55, 24;
+; CHECK-NEXT: and.b32 %r57, %r56, -2147483648;
+; CHECK-NEXT: and.b32 %r58, %r55, 7;
+; CHECK-NEXT: clz.b32 %r59, %r58;
+; CHECK-NEXT: sub.s32 %r60, 149, %r59;
+; CHECK-NEXT: shl.b32 %r61, %r60, 23;
+; CHECK-NEXT: or.b32 %r62, %r57, %r61;
+; CHECK-NEXT: sub.s32 %r63, 31, %r59;
+; CHECK-NEXT: shl.b32 %r64, %r11, %r63;
+; CHECK-NEXT: xor.b32 %r65, %r58, %r64;
+; CHECK-NEXT: add.s32 %r66, %r59, -8;
+; CHECK-NEXT: shl.b32 %r67, %r65, %r66;
+; CHECK-NEXT: or.b32 %r68, %r62, %r67;
+; CHECK-NEXT: bfe.u32 %r69, %r55, 3, 4;
+; CHECK-NEXT: shl.b32 %r70, %r69, 23;
+; CHECK-NEXT: or.b32 %r71, %r70, %r57;
+; CHECK-NEXT: shl.b32 %r72, %r58, 20;
+; CHECK-NEXT: or.b32 %r73, %r71, %r72;
+; CHECK-NEXT: add.s32 %r74, %r73, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p11, %r58, 0;
+; CHECK-NEXT: selp.b32 %r75, %r68, %r74, %p11;
+; CHECK-NEXT: setp.eq.b32 %p12, %r69, 0;
+; CHECK-NEXT: selp.b32 %r76, %r75, %r74, %p12;
+; CHECK-NEXT: or.b32 %r77, %r69, %r58;
+; CHECK-NEXT: setp.eq.b32 %p13, %r77, 0;
+; CHECK-NEXT: selp.b32 %r78, %r57, %r76, %p13;
+; CHECK-NEXT: setp.eq.b32 %p14, %r58, 7;
+; CHECK-NEXT: selp.b32 %r79, 2143289344, %r78, %p14;
+; CHECK-NEXT: setp.eq.b32 %p15, %r69, 15;
+; CHECK-NEXT: selp.b32 %r80, %r79, %r78, %p15;
+; CHECK-NEXT: prmt.b32 %r81, %r1, 0, 0x7770U;
+; CHECK-NEXT: shl.b32 %r82, %r81, 24;
+; CHECK-NEXT: and.b32 %r83, %r82, -2147483648;
+; CHECK-NEXT: and.b32 %r84, %r81, 7;
+; CHECK-NEXT: clz.b32 %r85, %r84;
+; CHECK-NEXT: sub.s32 %r86, 149, %r85;
+; CHECK-NEXT: shl.b32 %r87, %r86, 23;
+; CHECK-NEXT: or.b32 %r88, %r83, %r87;
+; CHECK-NEXT: sub.s32 %r89, 31, %r85;
+; CHECK-NEXT: shl.b32 %r90, %r11, %r89;
+; CHECK-NEXT: xor.b32 %r91, %r84, %r90;
+; CHECK-NEXT: add.s32 %r92, %r85, -8;
+; CHECK-NEXT: shl.b32 %r93, %r91, %r92;
+; CHECK-NEXT: or.b32 %r94, %r88, %r93;
+; CHECK-NEXT: bfe.u32 %r95, %r81, 3, 4;
+; CHECK-NEXT: shl.b32 %r96, %r95, 23;
+; CHECK-NEXT: or.b32 %r97, %r96, %r83;
+; CHECK-NEXT: shl.b32 %r98, %r84, 20;
+; CHECK-NEXT: or.b32 %r99, %r97, %r98;
+; CHECK-NEXT: add.s32 %r100, %r99, 1006632960;
+; CHECK-NEXT: setp.ne.b32 %p16, %r84, 0;
+; CHECK-NEXT: selp.b32 %r101, %r94, %r100, %p16;
+; CHECK-NEXT: setp.eq.b32 %p17, %r95, 0;
+; CHECK-NEXT: selp.b32 %r102, %r101, %r100, %p17;
+; CHECK-NEXT: or.b32 %r103, %r95, %r84;
+; CHECK-NEXT: setp.eq.b32 %p18, %r103, 0;
+; CHECK-NEXT: selp.b32 %r104, %r83, %r102, %p18;
+; CHECK-NEXT: setp.eq.b32 %p19, %r84, 7;
+; CHECK-NEXT: selp.b32 %r105, 2143289344, %r104, %p19;
+; CHECK-NEXT: setp.eq.b32 %p20, %r95, 15;
+; CHECK-NEXT: selp.b32 %r106, %r105, %r104, %p20;
+; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r106, %r80, %r54, %r28};
+; CHECK-NEXT: ret;
+ %r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> %x, metadata !"Float8E4M3FN")
+ ret <4 x float> %r
+}
+
+; <4 x i8> -> <4 x float> Float8E5M2
+define <4 x float> @from_f8e5m2_v4f32(<4 x i8> %x) {
+; CHECK-LABEL: from_f8e5m2_v4f32(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<21>;
+; CHECK-NEXT: .reg .b32 %r<119>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [from_f8e5m2_v4f32_param_0];
+; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;
+; CHECK-NEXT: shl.b32 %r3, %r2, 24;
+; CHECK-NEXT: and.b32 %r4, %r3, -2147483648;
+; CHECK-NEXT: and.b32 %r5, %r2, 3;
+; CHECK-NEXT: clz.b32 %r6, %r5;
+; CHECK-NEXT: sub.s32 %r7, 142, %r6;
+; CHECK-NEXT: shl.b32 %r8, %r7, 23;
+; CHECK-NEXT: or.b32 %r9, %r4, %r8;
+; CHECK-NEXT: sub.s32 %r10, 31, %r6;
+; CHECK-NEXT: mov.b32 %r11, 1;
+; CHECK-NEXT: shl.b32 %r12, %r11, %r10;
+; CHECK-NEXT: xor.b32 %r13, %r5, %r12;
+; CHECK-NEXT: add.s32 %r14, %r6, -8;
+; CHECK-NEXT: shl.b32 %r15, %r13, %r14;
+; CHECK-NEXT: or.b32 %r16, %r9, %r15;
+; CHECK-NEXT: bfe.u32 %r17, %r2, 2, 5;
+; CHECK-NEXT: shl.b32 %r18, %r17, 23;
+; CHECK-NEXT: or.b32 %r19, %r18, %r4;
+; CHECK-NEXT: shl.b32 %r20, %r5, 21;
+; CHECK-NEXT: or.b32 %r21, %r19, %r20;
+; CHECK-NEXT: add.s32 %r22, %r21, 939524096;
+; CHECK-NEXT: setp.ne.b32 %p1, %r5, 0;
+; CHECK-NEXT: selp.b32 %r23, %r16, %r22, %p1;
+; CHECK-NEXT: setp.eq.b32 %p2, %r17, 0;
+; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p2;
+; CHECK-NEXT: or.b32 %r25, %r17, %r5;
+; CHECK-NEXT: setp.eq.b32 %p3, %r25, 0;
+; CHECK-NEXT: selp.b32 %r26, %r4, %r24, %p3;
+; CHECK-NEXT: setp.eq.b32 %p4, %r5, 0;
+; CHECK-NEXT: or.b32 %r27, %r4, 2139095040;
+; CHECK-NEXT: selp.b32 %r28, %r27, %r26, %p4;
+; CHECK-NEXT: setp.eq.b32 %p5, %r17, 31;
+; CHECK-NEXT: selp.b32 %r29, %r28, %r26, %p5;
+; CHECK-NEXT: selp.b32 %r30, 2143289344, %r29, %p1;
+; CHECK-NEXT: selp.b32 %r31, %r30, %r29, %p5;
+; CHECK-NEXT: prmt.b32 %r32, %r1, 0, 0x7772U;
+; CHECK-NEXT: shl.b32 %r33, %r32, 24;
+; CHECK-NEXT: and.b32 %r34, %r33, -2147483648;
+; CHECK-NEXT: and.b32 %r35, %r32, 3;
+; CHECK-NEXT: clz.b32 %r36, %r35;
+; CHECK-NEXT: sub.s32 %r37, 142, %r36;
+; CHECK-NEXT: shl.b32 %r38, %r37, 23;
+; CHECK-NEXT: or.b32 %r39, %r34, %r38;
+; CHECK-NEXT: sub.s32 %r40, 31, %r36;
+; CHECK-NEXT: shl.b32 %r41, %r11, %r40;
+; CHECK-NEXT: xor.b32 %r42, %r35, %r41;
+; CHECK-NEXT: add.s32 %r43, %r36, -8;
+; CHECK-NEXT: shl.b32 %r44, %r42, %r43;
+; CHECK-NEXT: or.b32 %r45, %r39, %r44;
+; CHECK-NEXT: bfe.u32 %r46, %r32, 2, 5;
+; CHECK-NEXT: shl.b32 %r47, %r46, 23;
+; CHECK-NEXT: or.b32 %r48, %r47, %r34;
+; CHECK-NEXT: shl.b32 %r49, %r35, 21;
+; CHECK-NEXT: or.b32 %r50, %r48, %r49;
+; CHECK-NEXT: add.s32 %r51, %r50, 939524096;
+; CHECK-NEXT: setp.ne.b32 %p6, %r35, 0;
+; CHECK-NEXT: selp.b32 %r52, %r45, %r51, %p6;
+; CHECK-NEXT: setp.eq.b32 %p7, %r46, 0;
+; CHECK-NEXT: selp.b32 %r53, %r52, %r51, %p7;
+; CHECK-NEXT: or.b32 %r54, %r46, %r35;
+; CHECK-NEXT: setp.eq.b32 %p8, %r54, 0;
+; CHECK-NEXT: selp.b32 %r55, %r34, %r53, %p8;
+; CHECK-NEXT: setp.eq.b32 %p9, %r35, 0;
+; CHECK-NEXT: or.b32 %r56, %r34, 2139095040;
+; CHECK-NEXT: selp.b32 %r57, %r56, %r55, %p9;
+; CHECK-NEXT: setp.eq.b32 %p10, %r46, 31;
+; CHECK-NEXT: selp.b32 %r58, %r57, %r55, %p10;
+; CHECK-NEXT: selp.b32 %r59, 2143289344, %r58, %p6;
+; CHECK-NEXT: selp.b32 %r60, %r59, %r58, %p10;
+; CHECK-NEXT: prmt.b32 %r61, %r1, 0, 0x7771U;
+; CHECK-NEXT: shl.b32 %r62, %r61, 24;
+; CHECK-NEXT: and.b32 %r63, %r62, -2147483648;
+; CHECK-NEXT: and.b32 %r64, %r61, 3;
+; CHECK-NEXT: clz.b32 %r65, %r64;
+; CHECK-NEXT: sub.s32 %r66, 142, %r65;
+; CHECK-NEXT: shl.b32 %r67, %r66, 23;
+; CHECK-NEXT: or.b32 %r68, %r63, %r67;
+; CHECK-NEXT: sub.s32 %r69, 31, %r65;
+; CHECK-NEXT: shl.b32 %r70, %r11, %r69;
+; CHECK-NEXT: xor.b32 %r71, %r64, %r70;
+; CHECK-NEXT: add.s32 %r72, %r65, -8;
+; CHECK-NEXT: shl.b32 %r73, %r71, %r72;
+; CHECK-NEXT: or.b32 %r74, %r68, %r73;
+; CHECK-NEXT: bfe.u32 %r75, %r61, 2, 5;
+; CHECK-NEXT: shl.b32 %r76, %r75, 23;
+; CHECK-NEXT: or.b32 %r77, %r76, %r63;
+; CHECK-NEXT: shl.b32 %r78, %r64, 21;
+; CHECK-NEXT: or.b32 %r79, %r77, %r78;
+; CHECK-NEXT: add.s32 %r80, %r79, 939524096;
+; CHECK-NEXT: setp.ne.b32 %p11, %r64, 0;
+; CHECK-NEXT: selp.b32 %r81, %r74, %r80, %p11;
+; CHECK-NEXT: setp.eq.b32 %p12, %r75, 0;
+; CHECK-NEXT: selp.b32 %r82, %r81, %r80, %p12;
+; CHECK-NEXT: or.b32 %r83, %r75, %r64;
+; CHECK-NEXT: setp.eq.b32 %p13, %r83, 0;
+; CHECK-NEXT: selp.b32 %r84, %r63, %r82, %p13;
+; CHECK-NEXT: setp.eq.b32 %p14, %r64, 0;
+; CHECK-NEXT: or.b32 %r85, %r63, 2139095040;
+; CHECK-NEXT: selp.b32 %r86, %r85, %r84, %p14;
+; CHECK-NEXT: setp.eq.b32 %p15, %r75, 31;
+; CHECK-NEXT: selp.b32 %r87, %r86, %r84, %p15;
+; CHECK-NEXT: selp.b32 %r88, 2143289344, %r87, %p11;
+; CHECK-NEXT: selp.b32 %r89, %r88, %r87, %p15;
+; CHECK-NEXT: prmt.b32 %r90, %r1, 0, 0x7770U;
+; CHECK-NEXT: shl.b32 %r91, %r90, 24;
+; CHECK-NEXT: and.b32 %r92, %r91, -2147483648;
+; CHECK-NEXT: and.b32 %r93, %r90, 3;
+; CHECK-NEXT: clz.b32 %r94, %r93;
+; CHECK-NEXT: sub.s32 %r95, 142, %r94;
+; CHECK-NEXT: shl.b32 %r96, %r95, 23;
+; CHECK-NEXT: or.b32 %r97, %r92, %r96;
+; CHECK-NEXT: sub.s32 %r98, 31, %r94;
+; CHECK-NEXT: shl.b32 %r99, %r11, %r98;
+; CHECK-NEXT: xor.b32 %r100, %r93, %r99;
+; CHECK-NEXT: add.s32 %r101, %r94, -8;
+; CHECK-NEXT: shl.b32 %r102, %r100, %r101;
+; CHECK-NEXT: or.b32 %r103, %r97, %r102;
+; CHECK-NEXT: bfe.u32 %r104, %r90, 2, 5;
+; CHECK-NEXT: shl.b32 %r105, %r104, 23;
+; CHECK-NEXT: or.b32 %r106, %r105, %r92;
+; CHECK-NEXT: shl.b32 %r107, %r93, 21;
+; CHECK-NEXT: or.b32 %r108, %r106, %r107;
+; CHECK-NEXT: add.s32 %r109, %r108, 939524096;
+; CHECK-NEXT: setp.ne.b32 %p16, %r93, 0;
+; CHECK-NEXT: selp.b32 %r110, %r103, %r109, %p16;
+; CHECK-NEXT: setp.eq.b32 %p17, %r104, 0;
+; CHECK-NEXT: selp.b32 %r111, %r110, %r109, %p17;
+; CHECK-NEXT: or.b32 %r112, %r104, %r93;
+; CHECK-NEXT: setp.eq.b32 %p18, %r112, 0;
+; CHECK-NEXT: selp.b32 %r113, %r92, %r111, %p18;
+; CHECK-NEXT: setp.eq.b32 %p19, %r93, 0;
+; CHECK-NEXT: or.b32 %r114, %r92, 2139095040;
+; CHECK-NEXT: selp.b32 %r115, %r114, %r113, %p19;
+; CHECK-NEXT: setp.eq.b32 %p20, %r104, 31;
+; CHECK-NEXT: selp.b32 %r116, %r115, %r113, %p20;
+; CHECK-NEXT: selp.b32 %r117, 2143289344, %r116, %p16;
+; CHECK-NEXT: selp.b32 %r118, %r117, %r116, %p20;
+; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r118, %r89, %r60, %r31};
+; CHECK-NEXT: ret;
+ %r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> %x, metadata !"Float8E5M2")
+ ret <4 x float> %r
+}
+
+; <2 x i8> -> <2 x half> Float8E4M3FN
+define <2 x half> @from_f8e4m3fn_v2f16(<2 x i8> %x) {
+; CHECK-LABEL: from_f8e4m3fn_v2f16(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<11>;
+; CHECK-NEXT: .reg .b16 %rs<52>;
+; CHECK-NEXT: .reg .b32 %r<26>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [from_f8e4m3fn_v2f16_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
+; CHECK-NEXT: and.b32 %r2, %r1, 458759;
+; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r2;
+; CHECK-NEXT: cvt.u32.u16 %r3, %rs4;
+; CHECK-NEXT: shl.b32 %r4, %r3, 16;
+; CHECK-NEXT: clz.b32 %r5, %r4;
+; CHECK-NEXT: cvt.u16.u32 %rs5, %r5;
+; CHECK-NEXT: sub.s16 %rs6, 15, %rs5;
+; CHECK-NEXT: cvt.u32.u16 %r6, %rs6;
+; CHECK-NEXT: mov.b16 %rs7, 1;
+; CHECK-NEXT: shl.b16 %rs8, %rs7, %r6;
+; CHECK-NEXT: cvt.u32.u16 %r7, %rs3;
+; CHECK-NEXT: shl.b32 %r8, %r7, 16;
+; CHECK-NEXT: clz.b32 %r9, %r8;
+; CHECK-NEXT: cvt.u16.u32 %rs9, %r9;
+; CHECK-NEXT: sub.s16 %rs10, 15, %rs9;
+; CHECK-NEXT: cvt.u32.u16 %r10, %rs10;
+; CHECK-NEXT: shl.b16 %rs11, %rs7, %r10;
+; CHECK-NEXT: mov.b32 %r11, {%rs11, %rs8};
+; CHECK-NEXT: xor.b32 %r12, %r2, %r11;
+; CHECK-NEXT: mov.b32 {%rs12, %rs13}, %r12;
+; CHECK-NEXT: add.s16 %rs14, %rs5, -5;
+; CHECK-NEXT: cvt.u32.u16 %r13, %rs14;
+; CHECK-NEXT: shl.b16 %rs15, %rs13, %r13;
+; CHECK-NEXT: add.s16 %rs16, %rs9, -5;
+; CHECK-NEXT: cvt.u32.u16 %r14, %rs16;
+; CHECK-NEXT: shl.b16 %rs17, %rs12, %r14;
+; CHECK-NEXT: mov.b32 %r15, {%rs17, %rs15};
+; CHECK-NEXT: shl.b16 %rs18, %rs2, 8;
+; CHECK-NEXT: and.b16 %rs19, %rs18, -32768;
+; CHECK-NEXT: shl.b16 %rs20, %rs1, 8;
+; CHECK-NEXT: and.b16 %rs21, %rs20, -32768;
+; CHECK-NEXT: mov.b32 %r16, {%rs21, %rs19};
+; CHECK-NEXT: sub.s16 %rs22, 21, %rs5;
+; CHECK-NEXT: shl.b16 %rs23, %rs22, 10;
+; CHECK-NEXT: sub.s16 %rs24, 21, %rs9;
+; CHECK-NEXT: shl.b16 %rs25, %rs24, 10;
+; CHECK-NEXT: mov.b32 %r17, {%rs25, %rs23};
+; CHECK-NEXT: or.b32 %r18, %r16, %r17;
+; CHECK-NEXT: or.b32 %r19, %r18, %r15;
+; CHECK-NEXT: mov.b32 {%rs26, %rs27}, %r19;
+; CHECK-NEXT: shr.u16 %rs28, %rs2, 3;
+; CHECK-NEXT: shr.u16 %rs29, %rs1, 3;
+; CHECK-NEXT: mov.b32 %r20, {%rs29, %rs28};
+; CHECK-NEXT: and.b32 %r21, %r20, 983055;
+; CHECK-NEXT: mov.b32 {%rs30, %rs31}, %r21;
+; CHECK-NEXT: shl.b16 %rs32, %rs31, 10;
+; CHECK-NEXT: add.s16 %rs33, %rs32, 8192;
+; CHECK-NEXT: shl.b16 %rs34, %rs30, 10;
+; CHECK-NEXT: add.s16 %rs35, %rs34, 8192;
+; CHECK-NEXT: mov.b32 %r22, {%rs35, %rs33};
+; CHECK-NEXT: or.b32 %r23, %r16, %r22;
+; CHECK-NEXT: shl.b16 %rs36, %rs4, 7;
+; CHECK-NEXT: shl.b16 %rs37, %rs3, 7;
+; CHECK-NEXT: mov.b32 %r24, {%rs37, %rs36};
+; CHECK-NEXT: or.b32 %r25, %r23, %r24;
+; CHECK-NEXT: mov.b32 {%rs38, %rs39}, %r25;
+; CHECK-NEXT: setp.ne.b16 %p1, %rs4, 0;
+; CHECK-NEXT: selp.b16 %rs40, %rs27, %rs39, %p1;
+; CHECK-NEXT: setp.eq.b16 %p2, %rs31, 0;
+; CHECK-NEXT: selp.b16 %rs41, %rs40, %rs39, %p2;
+; CHECK-NEXT: or.b16 %rs42, %rs31, %rs4;
+; CHECK-NEXT: setp.eq.b16 %p3, %rs42, 0;
+; CHECK-NEXT: selp.b16 %rs43, %rs19, %rs41, %p3;
+; CHECK-NEXT: setp.eq.b16 %p4, %rs4, 7;
+; CHECK-NEXT: selp.b16 %rs44, 32256, %rs43, %p4;
+; CHECK-NEXT: setp.eq.b16 %p5, %rs31, 15;
+; CHECK-NEXT: selp.b16 %rs45, %rs44, %rs43, %p5;
+; CHECK-NEXT: setp.ne.b16 %p6, %rs3, 0;
+; CHECK-NEXT: selp.b16 %rs46, %rs26, %rs38, %p6;
+; CHECK-NEXT: setp.eq.b16 %p7, %rs30, 0;
+; CHECK-NEXT: selp.b16 %rs47, %rs46, %rs38, %p7;
+; CHECK-NEXT: or.b16 %rs48, %rs30, %rs3;
+; CHECK-NEXT: setp.eq.b16 %p8, %rs48, 0;
+; CHECK-NEXT: selp.b16 %rs49, %rs21, %rs47, %p8;
+; CHECK-NEXT: setp.eq.b16 %p9, %rs3, 7;
+; CHECK-NEXT: selp.b16 %rs50, 32256, %rs49, %p9;
+; CHECK-NEXT: setp.eq.b16 %p10, %rs30, 15;
+; CHECK-NEXT: selp.b16 %rs51, %rs50, %rs49, %p10;
+; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs51, %rs45};
+; CHECK-NEXT: ret;
+ %r = call <2 x half> @llvm.convert.from.arbitrary.fp.v2f16.v2i8(<2 x i8> %x, metadata !"Float8E4M3FN")
+ ret <2 x half> %r
+}
diff --git a/llvm/test/CodeGen/X86/arbitrary-fp-to-float.ll b/llvm/test/CodeGen/X86/arbitrary-fp-to-float.ll
index e60fed64fd179..f73a7a0d33edc 100644
--- a/llvm/test/CodeGen/X86/arbitrary-fp-to-float.ll
+++ b/llvm/test/CodeGen/X86/arbitrary-fp-to-float.ll
@@ -570,157 +570,75 @@ define bfloat @from_f8e5m2_to_bf16() {
define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
; CHECK-LABEL: fp4_to_f32_vec:
; CHECK: # %bb.0:
-; CHECK-NEXT: pshufd {{.*#+}} xmm1 = xmm0[3,3,3,3]
-; CHECK-NEXT: movd %xmm1, %esi
-; CHECK-NEXT: movl %esi, %edi
-; CHECK-NEXT: andl $1, %edi
-; CHECK-NEXT: movl %edi, %eax
-; CHECK-NEXT: shll $22, %eax
-; CHECK-NEXT: movl %esi, %edx
-; CHECK-NEXT: andl $-8, %edx
-; CHECK-NEXT: shll $28, %edx
-; CHECK-NEXT: shrl %esi
-; CHECK-NEXT: andl $3, %esi
-; CHECK-NEXT: movl %esi, %ecx
-; CHECK-NEXT: shll $23, %ecx
-; CHECK-NEXT: orl %edx, %ecx
-; CHECK-NEXT: leal 1056964608(%rax,%rcx), %r8d
-; CHECK-NEXT: bsrl %edi, %r9d
-; CHECK-NEXT: movl %edi, %r10d
-; CHECK-NEXT: btcl %r9d, %r10d
-; CHECK-NEXT: xorl $31, %r9d
-; CHECK-NEXT: leal -8(%r9), %ecx
-; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
-; CHECK-NEXT: shll %cl, %r10d
-; CHECK-NEXT: movl $157, %eax
-; CHECK-NEXT: movl $157, %ecx
-; CHECK-NEXT: subl %r9d, %ecx
-; CHECK-NEXT: shll $23, %ecx
-; CHECK-NEXT: orl %edx, %ecx
-; CHECK-NEXT: orl %r10d, %ecx
-; CHECK-NEXT: testl %edi, %edi
-; CHECK-NEXT: sete %dil
-; CHECK-NEXT: setne %r9b
-; CHECK-NEXT: testl %esi, %esi
-; CHECK-NEXT: sete %sil
-; CHECK-NEXT: testb %r9b, %sil
-; CHECK-NEXT: cmovel %r8d, %ecx
-; CHECK-NEXT: testb %dil, %sil
-; CHECK-NEXT: cmovnel %edx, %ecx
-; CHECK-NEXT: movd %ecx, %xmm1
-; CHECK-NEXT: pshufd {{.*#+}} xmm2 = xmm0[2,3,2,3]
-; CHECK-NEXT: movd %xmm2, %esi
-; CHECK-NEXT: movl %esi, %edi
-; CHECK-NEXT: andl $1, %edi
-; CHECK-NEXT: movl %edi, %ecx
-; CHECK-NEXT: shll $22, %ecx
-; CHECK-NEXT: movl %esi, %edx
-; CHECK-NEXT: andl $-8, %edx
-; CHECK-NEXT: shll $28, %edx
-; CHECK-NEXT: shrl %esi
-; CHECK-NEXT: andl $3, %esi
-; CHECK-NEXT: movl %esi, %r8d
-; CHECK-NEXT: shll $23, %r8d
-; CHECK-NEXT: orl %edx, %r8d
-; CHECK-NEXT: leal 1056964608(%rcx,%r8), %r8d
-; CHECK-NEXT: bsrl %edi, %r9d
-; CHECK-NEXT: movl %edi, %r10d
-; CHECK-NEXT: btcl %r9d, %r10d
-; CHECK-NEXT: xorl $31, %r9d
-; CHECK-NEXT: leal -8(%r9), %ecx
-; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
-; CHECK-NEXT: shll %cl, %r10d
-; CHECK-NEXT: movl $157, %ecx
-; CHECK-NEXT: subl %r9d, %ecx
-; CHECK-NEXT: shll $23, %ecx
-; CHECK-NEXT: orl %edx, %ecx
-; CHECK-NEXT: orl %r10d, %ecx
-; CHECK-NEXT: testl %edi, %edi
-; CHECK-NEXT: sete %dil
-; CHECK-NEXT: setne %r9b
-; CHECK-NEXT: testl %esi, %esi
-; CHECK-NEXT: sete %sil
-; CHECK-NEXT: testb %r9b, %sil
-; CHECK-NEXT: cmovel %r8d, %ecx
-; CHECK-NEXT: testb %dil, %sil
-; CHECK-NEXT: cmovnel %edx, %ecx
-; CHECK-NEXT: movd %ecx, %xmm2
-; CHECK-NEXT: punpckldq {{.*#+}} xmm2 = xmm2[0],xmm1[0],xmm2[1],xmm1[1]
-; CHECK-NEXT: movd %xmm0, %esi
-; CHECK-NEXT: movl %esi, %edi
-; CHECK-NEXT: andl $1, %edi
-; CHECK-NEXT: movl %edi, %ecx
-; CHECK-NEXT: shll $22, %ecx
-; CHECK-NEXT: movl %esi, %edx
-; CHECK-NEXT: andl $-8, %edx
-; CHECK-NEXT: shll $28, %edx
-; CHECK-NEXT: shrl %esi
-; CHECK-NEXT: andl $3, %esi
-; CHECK-NEXT: movl %esi, %r8d
-; CHECK-NEXT: shll $23, %r8d
-; CHECK-NEXT: orl %edx, %r8d
-; CHECK-NEXT: leal 1056964608(%rcx,%r8), %r8d
-; CHECK-NEXT: bsrl %edi, %r9d
-; CHECK-NEXT: movl %edi, %r10d
-; CHECK-NEXT: btcl %r9d, %r10d
-; CHECK-NEXT: xorl $31, %r9d
-; CHECK-NEXT: leal -8(%r9), %ecx
-; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
-; CHECK-NEXT: shll %cl, %r10d
-; CHECK-NEXT: movl $157, %ecx
-; CHECK-NEXT: subl %r9d, %ecx
-; CHECK-NEXT: shll $23, %ecx
-; CHECK-NEXT: orl %edx, %ecx
-; CHECK-NEXT: orl %r10d, %ecx
-; CHECK-NEXT: testl %edi, %edi
-; CHECK-NEXT: sete %dil
-; CHECK-NEXT: setne %r9b
-; CHECK-NEXT: testl %esi, %esi
-; CHECK-NEXT: sete %sil
-; CHECK-NEXT: testb %r9b, %sil
-; CHECK-NEXT: cmovel %r8d, %ecx
-; CHECK-NEXT: testb %dil, %sil
-; CHECK-NEXT: cmovnel %edx, %ecx
-; CHECK-NEXT: movd %ecx, %xmm1
-; CHECK-NEXT: pshufd {{.*#+}} xmm0 = xmm0[1,1,1,1]
-; CHECK-NEXT: movd %xmm0, %esi
-; CHECK-NEXT: movl %esi, %edi
-; CHECK-NEXT: andl $1, %edi
-; CHECK-NEXT: movl %edi, %ecx
-; CHECK-NEXT: shll $22, %ecx
-; CHECK-NEXT: movl %esi, %edx
-; CHECK-NEXT: andl $-8, %edx
-; CHECK-NEXT: shll $28, %edx
-; CHECK-NEXT: shrl %esi
-; CHECK-NEXT: andl $3, %esi
-; CHECK-NEXT: movl %esi, %r8d
-; CHECK-NEXT: shll $23, %r8d
-; CHECK-NEXT: orl %edx, %r8d
-; CHECK-NEXT: leal 1056964608(%rcx,%r8), %r8d
-; CHECK-NEXT: bsrl %edi, %r9d
-; CHECK-NEXT: movl %edi, %r10d
-; CHECK-NEXT: btcl %r9d, %r10d
-; CHECK-NEXT: xorl $31, %r9d
-; CHECK-NEXT: leal -8(%r9), %ecx
-; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
-; CHECK-NEXT: shll %cl, %r10d
-; CHECK-NEXT: subl %r9d, %eax
-; CHECK-NEXT: shll $23, %eax
-; CHECK-NEXT: orl %edx, %eax
-; CHECK-NEXT: orl %r10d, %eax
-; CHECK-NEXT: testl %edi, %edi
-; CHECK-NEXT: sete %cl
-; CHECK-NEXT: setne %dil
-; CHECK-NEXT: testl %esi, %esi
-; CHECK-NEXT: sete %sil
-; CHECK-NEXT: testb %dil, %sil
-; CHECK-NEXT: cmovel %r8d, %eax
-; CHECK-NEXT: testb %cl, %sil
-; CHECK-NEXT: cmovnel %edx, %eax
-; CHECK-NEXT: movd %eax, %xmm0
-; CHECK-NEXT: punpckldq {{.*#+}} xmm1 = xmm1[0],xmm0[0],xmm1[1],xmm0[1]
-; CHECK-NEXT: punpcklqdq {{.*#+}} xmm1 = xmm1[0],xmm2[0]
-; CHECK-NEXT: movdqa %xmm1, %xmm0
+; CHECK-NEXT: movdqa %xmm0, %xmm1
+; CHECK-NEXT: movdqa {{.*#+}} xmm2 = [1,1,1,1]
+; CHECK-NEXT: pand %xmm2, %xmm0
+; CHECK-NEXT: pcmpeqd %xmm3, %xmm3
+; CHECK-NEXT: pxor %xmm0, %xmm3
+; CHECK-NEXT: psubb {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm3
+; CHECK-NEXT: movdqa {{.*#+}} xmm4 = [51,51,51,51,51,51,51,51,51,51,51,51,51,51,51,51]
+; CHECK-NEXT: movdqa %xmm3, %xmm5
+; CHECK-NEXT: pand %xmm4, %xmm5
+; CHECK-NEXT: psrlw $2, %xmm3
+; CHECK-NEXT: pand %xmm4, %xmm3
+; CHECK-NEXT: paddb %xmm5, %xmm3
+; CHECK-NEXT: movdqa %xmm3, %xmm5
+; CHECK-NEXT: psrlw $4, %xmm5
+; CHECK-NEXT: paddb %xmm3, %xmm5
+; CHECK-NEXT: pand {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm5
+; CHECK-NEXT: pxor %xmm3, %xmm3
+; CHECK-NEXT: movdqa %xmm5, %xmm4
+; CHECK-NEXT: punpckhdq {{.*#+}} xmm4 = xmm4[2],xmm3[2],xmm4[3],xmm3[3]
+; CHECK-NEXT: psadbw %xmm3, %xmm4
+; CHECK-NEXT: punpckldq {{.*#+}} xmm5 = xmm5[0],xmm3[0],xmm5[1],xmm3[1]
+; CHECK-NEXT: psadbw %xmm3, %xmm5
+; CHECK-NEXT: packuswb %xmm4, %xmm5
+; CHECK-NEXT: movdqa {{.*#+}} xmm4 = [31,31,31,31]
+; CHECK-NEXT: psubd %xmm5, %xmm4
+; CHECK-NEXT: pslld $23, %xmm4
+; CHECK-NEXT: movdqa {{.*#+}} xmm6 = [1065353216,1065353216,1065353216,1065353216]
+; CHECK-NEXT: paddd %xmm6, %xmm4
+; CHECK-NEXT: cvttps2dq %xmm4, %xmm7
+; CHECK-NEXT: pxor %xmm0, %xmm7
+; CHECK-NEXT: movdqa {{.*#+}} xmm4 = [157,157,157,157]
+; CHECK-NEXT: psubd %xmm5, %xmm4
+; CHECK-NEXT: psubd {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm5
+; CHECK-NEXT: pslld $23, %xmm5
+; CHECK-NEXT: paddd %xmm6, %xmm5
+; CHECK-NEXT: cvttps2dq %xmm5, %xmm5
+; CHECK-NEXT: pshufd {{.*#+}} xmm6 = xmm7[1,1,3,3]
+; CHECK-NEXT: pmuludq %xmm5, %xmm7
+; CHECK-NEXT: pshufd {{.*#+}} xmm7 = xmm7[0,2,2,3]
+; CHECK-NEXT: pshufd {{.*#+}} xmm5 = xmm5[1,1,3,3]
+; CHECK-NEXT: pmuludq %xmm6, %xmm5
+; CHECK-NEXT: pshufd {{.*#+}} xmm5 = xmm5[0,2,2,3]
+; CHECK-NEXT: punpckldq {{.*#+}} xmm7 = xmm7[0],xmm5[0],xmm7[1],xmm5[1]
+; CHECK-NEXT: pslld $23, %xmm4
+; CHECK-NEXT: movdqa %xmm1, %xmm5
+; CHECK-NEXT: psrld $3, %xmm5
+; CHECK-NEXT: pslld $31, %xmm5
+; CHECK-NEXT: por %xmm5, %xmm4
+; CHECK-NEXT: por %xmm7, %xmm4
+; CHECK-NEXT: pcmpeqd %xmm0, %xmm2
+; CHECK-NEXT: psrld $1, %xmm1
+; CHECK-NEXT: pand {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm1
+; CHECK-NEXT: movdqa %xmm1, %xmm6
+; CHECK-NEXT: pcmpeqd %xmm3, %xmm6
+; CHECK-NEXT: pand %xmm6, %xmm2
+; CHECK-NEXT: pand %xmm2, %xmm4
+; CHECK-NEXT: paddd {{\.?LCPI[0-9]+_[0-9]+}}(%rip), %xmm1 # [126,126,126,126]
+; CHECK-NEXT: pslld $23, %xmm1
+; CHECK-NEXT: movdqa %xmm0, %xmm7
+; CHECK-NEXT: pslld $22, %xmm7
+; CHECK-NEXT: por %xmm5, %xmm7
+; CHECK-NEXT: por %xmm1, %xmm7
+; CHECK-NEXT: pandn %xmm7, %xmm2
+; CHECK-NEXT: por %xmm4, %xmm2
+; CHECK-NEXT: pcmpeqd %xmm3, %xmm0
+; CHECK-NEXT: pand %xmm6, %xmm0
+; CHECK-NEXT: pand %xmm0, %xmm5
+; CHECK-NEXT: pandn %xmm2, %xmm0
+; CHECK-NEXT: por %xmm5, %xmm0
; CHECK-NEXT: retq
%r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4> %x, metadata !"Float4E2M1FN")
ret <4 x float> %r
More information about the llvm-commits
mailing list