[llvm] [SelectionDAG] Add expansion for llvm.convert.to.arbitrary.fp (PR #193595)
Dmitry Sidorov via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 22 13:58:59 PDT 2026
https://github.com/MrSidims created https://github.com/llvm/llvm-project/pull/193595
The expansion converts a native IEEE float to an arbitrary-precision FP format, returning the result as an integer, following this algorithm:
1. Bitcast the source float to an integer and extract sign, exponent, and mantissa bit fields via masks and shifts.
2. Classify the input (zero/denormal/normal/Inf/NaN).
3. Normalize source denormals by finding the MSB position of the mantissa and adjusting the effective exponent.
4. Normal path: adjust the exponent bias from source to destination format and truncate the mantissa with rounding (supports NearestTiesToEven, TowardZero, TowardPositive, TowardNegative, NearestTiesToAway).
5. Denormal destination path: when the biased destination exponent is <= 0, shift the mantissa right to produce a denormalized result with rounding.
6. Handle mantissa overflow from rounding and exponent overflow. Produce Inf or saturate to max finite, depending on format and saturation flag.
7. Build special-value results (canonical qNaN, signed Inf, signed zero) adapted to the destination format's non-finite behavior (IEEE754, NanOnly, FiniteOnly).
8. Final selection in priority order: NaN > Inf > Zero > Overflow >
Normal/Denorm.
Currently only conversions to OCP floats are covered, in LLVM terms these are: Float8E5M2, Float8E4M3FN, Float6E3M2FN, Float6E2M3FN, Float4E2M1FN.
OCP spec:
https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf
E2E testing on X86 done with an assistance of Claude Code Opus 4.6.
>From 894b5d07d96409cd4b2887e76b2d235a7f95492c Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <Dmitry.Sidorov at amd.com>
Date: Thu, 19 Feb 2026 10:22:21 +0100
Subject: [PATCH] [SelectionDAG] Add expansion for llvm.convert.to.arbitrary.fp
The expansion converts a native IEEE float to an arbitrary-precision
FP format, returning the result as an integer, following this algorithm:
1. Bitcast the source float to an integer and extract sign, exponent,
and mantissa bit fields via masks and shifts.
2. Classify the input (zero/denormal/normal/Inf/NaN).
3. Normalize source denormals by finding the MSB position of the
mantissa and adjusting the effective exponent.
4. Normal path: adjust the exponent bias from source to destination
format and truncate the mantissa with rounding (supports
NearestTiesToEven, TowardZero, TowardPositive, TowardNegative,
NearestTiesToAway).
5. Denormal destination path: when the biased destination exponent is
<= 0, shift the mantissa right to produce a denormalized result
with rounding.
6. Handle mantissa overflow from rounding and exponent overflow. Produce Inf
or saturate to max finite, depending on format and saturation flag.
7. Build special-value results (canonical qNaN, signed Inf, signed
zero) adapted to the destination format's non-finite behavior
(IEEE754, NanOnly, FiniteOnly).
8. Final selection in priority order: NaN > Inf > Zero > Overflow >
Normal/Denorm.
Currently only conversions to OCP floats are covered, in LLVM terms
these are: Float8E5M2, Float8E4M3FN, Float6E3M2FN, Float6E2M3FN,
Float4E2M1FN.
OCP spec:
https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf
E2E testing on X86 done with an assistance of Claude Code Opus 4.6.
---
llvm/include/llvm/CodeGen/ISDOpcodes.h | 9 +
llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp | 473 ++++++
.../SelectionDAG/LegalizeFloatTypes.cpp | 20 +
.../SelectionDAG/LegalizeIntegerTypes.cpp | 13 +
llvm/lib/CodeGen/SelectionDAG/LegalizeTypes.h | 3 +
.../SelectionDAG/LegalizeVectorOps.cpp | 1 +
.../SelectionDAG/LegalizeVectorTypes.cpp | 96 +-
.../SelectionDAG/SelectionDAGBuilder.cpp | 43 +
.../SelectionDAG/SelectionDAGDumper.cpp | 1 +
llvm/lib/CodeGen/TargetLoweringBase.cpp | 3 +-
.../CodeGen/AMDGPU/float-to-arbitrary-fp.ll | 905 +++++++++++
.../CodeGen/NVPTX/float-to-arbitrary-fp.ll | 1028 +++++++++++++
.../X86/float-to-arbitrary-fp-error.ll | 76 +
.../test/CodeGen/X86/float-to-arbitrary-fp.ll | 1323 +++++++++++++++++
14 files changed, 3970 insertions(+), 24 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/float-to-arbitrary-fp.ll
create mode 100644 llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll
create mode 100644 llvm/test/CodeGen/X86/float-to-arbitrary-fp-error.ll
create mode 100644 llvm/test/CodeGen/X86/float-to-arbitrary-fp.ll
diff --git a/llvm/include/llvm/CodeGen/ISDOpcodes.h b/llvm/include/llvm/CodeGen/ISDOpcodes.h
index 8a8a9ee71ca02..de9835ee43836 100644
--- a/llvm/include/llvm/CodeGen/ISDOpcodes.h
+++ b/llvm/include/llvm/CodeGen/ISDOpcodes.h
@@ -1020,6 +1020,15 @@ enum NodeType {
/// The second operand is a constant indicating the source FP semantics.
CONVERT_FROM_ARBITRARY_FP,
+ /// CONVERT_TO_ARBITRARY_FP - Converts a native FP value to an arbitrary
+ /// floating-point format, returning the result as an integer.
+ /// The first operand is the source value.
+ /// The second operand is a constant indicating the destination FP semantics.
+ /// The third operand is a constant indication the rounding mode.
+ /// The last operand is a boolean consant indicating whether the result has
+ /// to be saturated.
+ CONVERT_TO_ARBITRARY_FP,
+
/// Perform various unary floating-point operations inspired by libm. For
/// FPOWI, the result is undefined if the integer operand doesn't fit into
/// sizeof(int).
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
index 54d86dfbfa303..7dcc3a1f1c753 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeDAG.cpp
@@ -3782,6 +3782,479 @@ bool SelectionDAGLegalize::ExpandNode(SDNode *Node) {
Results.push_back(Result);
break;
}
+ case ISD::CONVERT_TO_ARBITRARY_FP: {
+ // Expand conversion from a native IEEE float type to an arbitrary FP
+ // format, returning the result as an integer using bit manipulation.
+ //
+ // TODO: currently only conversions to FP4, FP6 and FP8 formats from OCP
+ // specification are expanded. Remaining arbitrary FP types: Float8E4M3,
+ // Float8E3M4, Float8E5M2FNUZ, Float8E4M3FNUZ, Float8E4M3B11FNUZ,
+ // Float8E8M0FNU.
+ EVT ResVT = Node->getValueType(0);
+
+ SDValue FloatVal = Node->getOperand(0);
+ const uint64_t SemEnum = Node->getConstantOperandVal(1);
+ const auto Sem = static_cast<APFloatBase::Semantics>(SemEnum);
+ const auto RoundMode =
+ static_cast<RoundingMode>(Node->getConstantOperandVal(2));
+ const bool Saturate = Node->getConstantOperandVal(3) != 0;
+
+ // Supported destination 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_TO_ARBITRARY_FP: not implemented "
+ "destination format (semantics enum " +
+ Twine(SemEnum) + ")");
+ Results.push_back(DAG.getPOISON(ResVT));
+ break;
+ }
+ if (!Results.empty())
+ break;
+
+ // Destination format parameters.
+ const fltSemantics &DstSem = APFloatBase::EnumToSemantics(Sem);
+ const unsigned DstBits = APFloat::getSizeInBits(DstSem);
+ const unsigned DstPrecision = APFloat::semanticsPrecision(DstSem);
+ const unsigned DstMant = DstPrecision - 1;
+ const unsigned DstExpBits = DstBits - DstMant - 1;
+ const int DstBias = 1 - APFloat::semanticsMinExponent(DstSem);
+ const unsigned DstExpMax = (1U << DstExpBits) - 1;
+ const uint64_t DstMantMask = (DstMant > 0) ? ((1ULL << DstMant) - 1) : 0;
+ const fltNonfiniteBehavior DstNFBehavior = DstSem.nonFiniteBehavior;
+ const fltNanEncoding DstNanEnc = DstSem.nanEncoding;
+
+ // Compute the maximum normal exponent for the destination format.
+ unsigned DstExpMaxNormal;
+ if (DstNFBehavior == fltNonfiniteBehavior::IEEE754)
+ DstExpMaxNormal = DstExpMax - 1;
+ else
+ DstExpMaxNormal = DstExpMax;
+
+ // For NanOnly formats the max exponent field for finite values
+ // is DstExpMax, but the encoding with exp = DstExpMax and
+ // mant = all-ones is NaN. So DstExpMaxNormal = DstExpMax, but max
+ // mantissa at that exponent is DstMantMask - 1 (if NanEnc == AllOnes) to
+ // avoid the NaN encoding.
+ uint64_t DstMaxMantAtMaxExp = DstMantMask;
+ if (DstNFBehavior == fltNonfiniteBehavior::NanOnly &&
+ DstNanEnc == fltNanEncoding::AllOnes)
+ DstMaxMantAtMaxExp = DstMantMask - 1;
+
+ // Source format parameters.
+ EVT SrcVT = FloatVal.getValueType();
+ const fltSemantics &SrcSem = SrcVT.getFltSemantics();
+ const unsigned SrcBits = APFloat::getSizeInBits(SrcSem);
+ const unsigned SrcPrecision = APFloat::semanticsPrecision(SrcSem);
+ const unsigned SrcMant = SrcPrecision - 1;
+ const unsigned SrcExpBits = SrcBits - SrcMant - 1;
+ const int SrcBias = 1 - APFloat::semanticsMinExponent(SrcSem);
+ const uint64_t SrcMantMask = (1ULL << SrcMant) - 1;
+ const uint64_t SrcExpMask = (1ULL << SrcExpBits) - 1;
+
+ // Work in the source integer type.
+ EVT IntVT = EVT::getIntegerVT(*DAG.getContext(), SrcBits);
+ EVT SetCCVT = getSetCCResultType(IntVT);
+
+ SDValue Zero = DAG.getConstant(0, dl, IntVT);
+ SDValue One = DAG.getConstant(1, dl, IntVT);
+
+ // Bitcast source float to integer and extract bit fields.
+ SDValue Src = DAG.getNode(ISD::BITCAST, dl, IntVT, FloatVal);
+ SDValue SrcMantField = DAG.getNode(ISD::AND, dl, IntVT, Src,
+ DAG.getConstant(SrcMantMask, dl, IntVT));
+
+ SDValue SrcExpField = DAG.getNode(
+ ISD::AND, dl, IntVT,
+ DAG.getNode(ISD::SRL, dl, IntVT, Src,
+ DAG.getShiftAmountConstant(SrcMant, IntVT, dl)),
+ DAG.getConstant(SrcExpMask, dl, IntVT));
+
+ SDValue SignBit =
+ DAG.getNode(ISD::SRL, dl, IntVT, Src,
+ DAG.getShiftAmountConstant(SrcBits - 1, IntVT, dl));
+
+ // Classify the input value.
+ SDValue SrcExpAllOnes = DAG.getConstant(SrcExpMask, dl, IntVT);
+ SDValue IsExpAllOnes =
+ DAG.getSetCC(dl, SetCCVT, SrcExpField, SrcExpAllOnes, ISD::SETEQ);
+ SDValue IsExpZero =
+ DAG.getSetCC(dl, SetCCVT, SrcExpField, Zero, ISD::SETEQ);
+ SDValue IsMantZero =
+ DAG.getSetCC(dl, SetCCVT, SrcMantField, Zero, ISD::SETEQ);
+ SDValue IsMantNonZero =
+ DAG.getSetCC(dl, SetCCVT, SrcMantField, Zero, ISD::SETNE);
+
+ // If source is IEEE fp, tehn NaN = exp_all_ones && mant != 0.
+ SDValue IsNaN =
+ DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantNonZero);
+ // Inf = exp_all_ones && mant == 0.
+ SDValue IsInf =
+ DAG.getNode(ISD::AND, dl, SetCCVT, IsExpAllOnes, IsMantZero);
+ // Zero = exp == 0 && mant == 0.
+ SDValue IsZero =
+ DAG.getNode(ISD::AND, dl, SetCCVT, IsExpZero, IsMantZero);
+ // Source denorm = exp == 0 && mant != 0.
+ SDValue IsSrcDenorm =
+ DAG.getNode(ISD::AND, dl, SetCCVT, IsExpZero, IsMantNonZero);
+
+ // Source denormal normalization.
+ // For a source denormal, the true exponent is (1 - SrcBias) and the
+ // mantissa has no implicit leading 1. Normalize by finding the position
+ // of the leading 1 in the mantissa.
+ SDValue LeadingZeros =
+ DAG.getNode(ISD::CTLZ_ZERO_UNDEF, dl, IntVT, SrcMantField);
+
+ // normShift = LeadingZeros - (SrcBits - 1 - SrcMant).
+ const unsigned LZOffset = SrcBits - 1 - SrcMant;
+ SDValue NormShift = DAG.getNode(ISD::SUB, dl, IntVT, LeadingZeros,
+ DAG.getConstant(LZOffset, dl, IntVT));
+
+ // Normalized mantissa.
+ SDValue NormMant =
+ DAG.getNode(ISD::AND, dl, IntVT,
+ DAG.getNode(ISD::SHL, dl, IntVT, SrcMantField, NormShift),
+ DAG.getConstant(SrcMantMask, dl, IntVT));
+
+ // effective_exp = 1 - NormShift.
+ SDValue DenormSrcExp =
+ DAG.getNode(ISD::SUB, dl, IntVT, One, NormShift);
+
+ // Select between normal and denorm source.
+ SDValue EffSrcExp =
+ DAG.getSelect(dl, IntVT, IsSrcDenorm, DenormSrcExp, SrcExpField);
+ SDValue EffSrcMant =
+ DAG.getSelect(dl, IntVT, IsSrcDenorm, NormMant, SrcMantField);
+
+ // Compute new biased exponent for destination.
+ // new_exp = src_exp - SrcBias + DstBias
+ const int BiasAdjust = DstBias - SrcBias;
+ SDValue NewExp = DAG.getNode(
+ ISD::ADD, dl, IntVT, EffSrcExp,
+ DAG.getConstant(APInt(SrcBits, BiasAdjust, true), dl, IntVT));
+
+ // Compute rounding increment given the round bit, sticky bits, and LSB
+ // of the truncated mantissa.
+ auto ComputeRoundUp = [&](SDValue RoundBit, SDValue StickyBits,
+ SDValue LSB) -> SDValue {
+ if (RoundMode == RoundingMode::NearestTiesToEven) {
+ // Round up if round_bit && (sticky || lsb)
+ SDValue StickyOrLSB = DAG.getNode(ISD::OR, dl, IntVT, StickyBits, LSB);
+ return DAG.getNode(ISD::AND, dl, IntVT, RoundBit, StickyOrLSB);
+ }
+ if (RoundMode == RoundingMode::TowardZero)
+ return Zero;
+ if (RoundMode == RoundingMode::TowardPositive) {
+ // Round up if positive and any truncated bits are set.
+ SDValue AnyTruncBits =
+ DAG.getNode(ISD::OR, dl, IntVT, RoundBit, StickyBits);
+ SDValue HasTruncBits =
+ DAG.getSetCC(dl, SetCCVT, AnyTruncBits, Zero, ISD::SETNE);
+ SDValue IsPositive =
+ DAG.getSetCC(dl, SetCCVT, SignBit, Zero, ISD::SETEQ);
+ SDValue DoRound =
+ DAG.getNode(ISD::AND, dl, SetCCVT, HasTruncBits, IsPositive);
+ return DAG.getNode(ISD::ZERO_EXTEND, dl, IntVT, DoRound);
+ }
+ if (RoundMode == RoundingMode::TowardNegative) {
+ // Round up if negative and any truncated bits are set (to -Inf).
+ SDValue AnyTruncBits =
+ DAG.getNode(ISD::OR, dl, IntVT, RoundBit, StickyBits);
+ SDValue HasTruncBits =
+ DAG.getSetCC(dl, SetCCVT, AnyTruncBits, Zero, ISD::SETNE);
+ SDValue IsNegative =
+ DAG.getSetCC(dl, SetCCVT, SignBit, Zero, ISD::SETNE);
+ SDValue DoRound =
+ DAG.getNode(ISD::AND, dl, SetCCVT, HasTruncBits, IsNegative);
+ return DAG.getNode(ISD::ZERO_EXTEND, dl, IntVT, DoRound);
+ }
+ if (RoundMode == RoundingMode::NearestTiesToAway)
+ return RoundBit;
+ llvm_unreachable("Unsupported rounding mode");
+ };
+
+ // Round mantissa from SrcMant bits to DstMant bits.
+ SDValue TruncMant;
+ SDValue RoundUp;
+ if (SrcMant > DstMant) {
+ const unsigned Shift = SrcMant - DstMant;
+ SDValue ShiftConst = DAG.getShiftAmountConstant(Shift, IntVT, dl);
+ TruncMant =
+ DAG.getNode(ISD::SRL, dl, IntVT, EffSrcMant, ShiftConst);
+
+ // Check bit at position Shift - 1 aka the round bit.
+ SDValue RoundBit;
+ if (Shift >= 1) {
+ RoundBit = DAG.getNode(
+ ISD::AND, dl, IntVT,
+ DAG.getNode(ISD::SRL, dl, IntVT, EffSrcMant,
+ DAG.getShiftAmountConstant(Shift - 1, IntVT, dl)),
+ One);
+ } else {
+ RoundBit = Zero;
+ }
+
+ // OR of all bits below the round bit to get sticky bits.
+ SDValue StickyBits;
+ if (Shift >= 2) {
+ uint64_t StickyMask = (1ULL << (Shift - 1)) - 1;
+ StickyBits = DAG.getNode(ISD::AND, dl, IntVT, EffSrcMant,
+ DAG.getConstant(StickyMask, dl, IntVT));
+ StickyBits = DAG.getSetCC(dl, SetCCVT, StickyBits, Zero, ISD::SETNE);
+ StickyBits =
+ DAG.getNode(ISD::ZERO_EXTEND, dl, IntVT, StickyBits);
+ } else {
+ StickyBits = Zero;
+ }
+
+ // LSB of truncated mantissa.
+ SDValue LSB = DAG.getNode(ISD::AND, dl, IntVT, TruncMant, One);
+
+ RoundUp = ComputeRoundUp(RoundBit, StickyBits, LSB);
+ } else {
+ // If DstMant >= SrcMant, then no rounding needed, just shift left.
+ SDValue MantShift =
+ DAG.getShiftAmountConstant(DstMant - SrcMant, IntVT, dl);
+ TruncMant = DAG.getNode(ISD::SHL, dl, IntVT, EffSrcMant, MantShift);
+ RoundUp = Zero;
+ }
+
+ // Apply rounding.
+ SDValue RoundedMant =
+ DAG.getNode(ISD::ADD, dl, IntVT, TruncMant, RoundUp);
+
+ // Handle mantissa overflow from rounding.
+ // If rounded_mant > DstMantMask, carry into exponent.
+ SDValue MantOverflow = DAG.getSetCC(
+ dl, SetCCVT, RoundedMant,
+ DAG.getConstant(DstMantMask, dl, IntVT), ISD::SETGT);
+ // On overflow: mant = 0, exp += 1.
+ SDValue AdjMant =
+ DAG.getSelect(dl, IntVT, MantOverflow, Zero, RoundedMant);
+ SDValue AdjExp = DAG.getNode(
+ ISD::ADD, dl, IntVT, NewExp,
+ DAG.getNode(ISD::ZERO_EXTEND, dl, IntVT, MantOverflow));
+
+ // Precompute sign shifted to MSB of destination.
+ SDValue SignShifted = DAG.getNode(
+ ISD::SHL, dl, IntVT, SignBit,
+ DAG.getShiftAmountConstant(DstBits - 1, IntVT, dl));
+
+ // Destination denormal conversion (when new_exp <= 0).
+ // Shift the mantissa right by 1 - new_exp additional bits and set the
+ // exponent field to 0.
+ SDValue ExpIsNeg =
+ DAG.getSetCC(dl, SetCCVT, AdjExp,
+ DAG.getConstant(1, dl, IntVT), ISD::SETLT);
+
+ SDValue DenormResult;
+ {
+ // denorm_shift = 1 - NewExp.
+ SDValue DenormShift =
+ DAG.getNode(ISD::SUB, dl, IntVT, One, NewExp);
+
+ // full_src_mant = (1 << SrcMant) | EffSrcMant.
+ SDValue ImplicitOne = DAG.getNode(
+ ISD::SHL, dl, IntVT, One,
+ DAG.getShiftAmountConstant(SrcMant, IntVT, dl));
+ SDValue FullSrcMant =
+ DAG.getNode(ISD::OR, dl, IntVT, EffSrcMant, ImplicitOne);
+
+ // Total right shift = (SrcMant - DstMant) + DenormShift
+ SDValue TotalShift;
+ if (SrcMant >= DstMant) {
+ TotalShift =
+ DAG.getNode(ISD::ADD, dl, IntVT, DenormShift,
+ DAG.getConstant(SrcMant - DstMant, dl, IntVT));
+ } else {
+ TotalShift =
+ DAG.getNode(ISD::SUB, dl, IntVT, DenormShift,
+ DAG.getConstant(DstMant - SrcMant, dl, IntVT));
+ }
+
+ // Clamp total shift to avoid UB, then trancate denorm mantissa.
+ SDValue MaxShift = DAG.getConstant(SrcBits - 1, dl, IntVT);
+ SDValue ClampedShift = DAG.getNode(ISD::UMIN, dl, IntVT, TotalShift,
+ MaxShift);
+ SDValue DenormTruncMant =
+ DAG.getNode(ISD::SRL, dl, IntVT, FullSrcMant, ClampedShift);
+
+ // Rounding for denorm path.
+ SDValue DenormRoundUp;
+ {
+ // Round bit is at position TotalShift - 1 of FullSrcMant.
+ // Clamp to at least 1 so the subtraction doesn't underflow and create
+ // shift nodes with invalid shift amounts.
+ SDValue SafeShift =
+ DAG.getNode(ISD::UMAX, dl, IntVT, ClampedShift, One);
+ SDValue RoundBitPos =
+ DAG.getNode(ISD::SUB, dl, IntVT, SafeShift, One);
+ SDValue DenormRoundBit = DAG.getNode(
+ ISD::AND, dl, IntVT,
+ DAG.getNode(ISD::SRL, dl, IntVT, FullSrcMant, RoundBitPos), One);
+
+ // Sticky: all bits below round bit.
+ // sticky_mask = (1 << RoundBitPos) - 1
+ SDValue StickyMask = DAG.getNode(
+ ISD::SUB, dl, IntVT,
+ DAG.getNode(ISD::SHL, dl, IntVT, One, RoundBitPos), One);
+ SDValue DenormStickyBits =
+ DAG.getNode(ISD::AND, dl, IntVT, FullSrcMant, StickyMask);
+ SDValue HasSticky =
+ DAG.getNode(ISD::ZERO_EXTEND, dl, IntVT,
+ DAG.getSetCC(dl, SetCCVT, DenormStickyBits, Zero,
+ ISD::SETNE));
+
+ SDValue DenormLSB =
+ DAG.getNode(ISD::AND, dl, IntVT, DenormTruncMant, One);
+
+ DenormRoundUp =
+ ComputeRoundUp(DenormRoundBit, HasSticky, DenormLSB);
+
+ // Only apply rounding if TotalShift >= 1 (i.e., there are bits to
+ // round).
+ SDValue ShiftGEOne =
+ DAG.getSetCC(dl, SetCCVT, ClampedShift, One, ISD::SETUGE);
+ DenormRoundUp = DAG.getSelect(dl, IntVT, ShiftGEOne, DenormRoundUp,
+ Zero);
+ }
+
+ SDValue DenormRoundedMant =
+ DAG.getNode(ISD::ADD, dl, IntVT, DenormTruncMant, DenormRoundUp);
+
+ // If rounding caused overflow into the normal range, then we get the
+ // smallest normal number.
+ SDValue DenormMantOF = DAG.getSetCC(
+ dl, SetCCVT, DenormRoundedMant,
+ DAG.getConstant(DstMantMask, dl, IntVT), ISD::SETGT);
+ SDValue DenormFinalMant = DAG.getSelect(
+ dl, IntVT, DenormMantOF, Zero, DenormRoundedMant);
+ SDValue DenormFinalExp = DAG.getSelect(
+ dl, IntVT, DenormMantOF, One, Zero);
+
+ // Assemble: sign | (exp << DstMant) | mant
+ SDValue DenormExpShifted = DAG.getNode(
+ ISD::SHL, dl, IntVT, DenormFinalExp,
+ DAG.getShiftAmountConstant(DstMant, IntVT, dl));
+ DenormResult = DAG.getNode(
+ ISD::OR, dl, IntVT,
+ DAG.getNode(ISD::OR, dl, IntVT, SignShifted, DenormExpShifted),
+ DenormFinalMant);
+
+ // If the value is to small for even a denorm (all mantissa bits
+ // shifted away), handle based on rounding mode.
+ // This is covered by the DenormRoundedMant = 0 case naturally.
+ }
+
+ // Exponent overflow detection.
+ SDValue ExpOF = DAG.getSetCC(
+ dl, SetCCVT, AdjExp,
+ DAG.getConstant(DstExpMaxNormal, dl, IntVT), ISD::SETGT);
+
+ // Also check if AdjExp == DstExpMaxNormal and mantissa overflow into
+ // a value that exceeds the max allowed mantissa at that exponent.
+ SDValue ExpAtMax = DAG.getSetCC(
+ dl, SetCCVT, AdjExp,
+ DAG.getConstant(DstExpMaxNormal, dl, IntVT), ISD::SETEQ);
+ SDValue MantExceedsMax = DAG.getSetCC(
+ dl, SetCCVT, AdjMant,
+ DAG.getConstant(DstMaxMantAtMaxExp, dl, IntVT), ISD::SETGT);
+ SDValue ExpMantOF =
+ DAG.getNode(ISD::AND, dl, SetCCVT, ExpAtMax, MantExceedsMax);
+ SDValue IsOverflow =
+ DAG.getNode(ISD::OR, dl, SetCCVT, ExpOF, ExpMantOF);
+
+ // Build overflow result.
+ SDValue OverflowResult;
+
+ if (Saturate) {
+ // Clamp to max finite value:
+ // sign | (DstExpMaxNormal << DstMant) | DstMaxMantAtMaxExp
+ uint64_t MaxFinite =
+ ((uint64_t)DstExpMaxNormal << DstMant) | DstMaxMantAtMaxExp;
+ OverflowResult =
+ DAG.getNode(ISD::OR, dl, IntVT, SignShifted,
+ DAG.getConstant(MaxFinite, dl, IntVT));
+ } else if (DstNFBehavior == fltNonfiniteBehavior::IEEE754) {
+ // Produce infinity.
+ uint64_t InfBits = (uint64_t)DstExpMax << DstMant;
+ OverflowResult =
+ DAG.getNode(ISD::OR, dl, IntVT, SignShifted,
+ DAG.getConstant(InfBits, dl, IntVT));
+ } else {
+ // Emit poison if no Inf in format and not saturating.
+ OverflowResult = DAG.getPOISON(IntVT);
+ }
+
+ // Assemble normal result: sign | (AdjExp << DstMant) | AdjMant
+ SDValue NormExpShifted = DAG.getNode(
+ ISD::SHL, dl, IntVT, AdjExp,
+ DAG.getShiftAmountConstant(DstMant, IntVT, dl));
+ SDValue NormResult = DAG.getNode(
+ ISD::OR, dl, IntVT,
+ DAG.getNode(ISD::OR, dl, IntVT, SignShifted, NormExpShifted), AdjMant);
+
+ // Build special-value results.
+ SDValue NaNResult;
+ if (DstNFBehavior == fltNonfiniteBehavior::IEEE754) {
+ // Produce canonical NaN.
+ const uint64_t QNaNBit = (DstMant > 0) ? (1ULL << (DstMant - 1)) : 0;
+ NaNResult =
+ DAG.getConstant(((uint64_t)DstExpMax << DstMant) | QNaNBit, dl,
+ IntVT);
+ } else if (DstNFBehavior == fltNonfiniteBehavior::NanOnly &&
+ DstNanEnc == fltNanEncoding::AllOnes) {
+ // E4M3FN-style: NaN is exp=all-ones, mant=all-ones.
+ NaNResult =
+ DAG.getConstant(((uint64_t)DstExpMax << DstMant) | DstMantMask, dl,
+ IntVT);
+ } else {
+ // NaN -> poison for finite only values.
+ NaNResult = DAG.getPOISON(IntVT);
+ }
+
+ // Inf handling.
+ SDValue InfResult;
+ if (DstNFBehavior == fltNonfiniteBehavior::IEEE754) {
+ // Produce signed infinity.
+ uint64_t InfBits = (uint64_t)DstExpMax << DstMant;
+ InfResult = DAG.getNode(ISD::OR, dl, IntVT, SignShifted,
+ DAG.getConstant(InfBits, dl, IntVT));
+ } else if (Saturate) {
+ // Inf saturates to max finite.
+ uint64_t MaxFinite =
+ ((uint64_t)DstExpMaxNormal << DstMant) | DstMaxMantAtMaxExp;
+ InfResult = DAG.getNode(ISD::OR, dl, IntVT, SignShifted,
+ DAG.getConstant(MaxFinite, dl, IntVT));
+ } else {
+ // No Inf and not saturating -> poison.
+ InfResult = DAG.getPOISON(IntVT);
+ }
+
+ SDValue ZeroResult = SignShifted;
+
+ // Final selection in an order: NaN takes priority, then Inf, then Zero.
+ SDValue FiniteResult = DAG.getSelect(dl, IntVT, ExpIsNeg, DenormResult,
+ NormResult);
+ FiniteResult =
+ DAG.getSelect(dl, IntVT, IsOverflow, OverflowResult, FiniteResult);
+
+ 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);
+
+ // Truncate to destination integer type.
+ Result = DAG.getZExtOrTrunc(Result, dl, ResVT);
+
+ Results.push_back(Result);
+ break;
+ }
case ISD::FCANONICALIZE: {
SDValue Mul = TLI.expandFCANONICALIZE(Node, DAG);
Results.push_back(Mul);
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
index 25f4f75eaedea..65a031027b4b1 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeFloatTypes.cpp
@@ -3194,6 +3194,8 @@ bool DAGTypeLegalizer::SoftPromoteHalfOperand(SDNode *N, unsigned OpNo) {
case ISD::FP_TO_SINT_SAT:
case ISD::FP_TO_UINT_SAT:
Res = SoftPromoteHalfOp_FP_TO_XINT_SAT(N); break;
+ case ISD::CONVERT_TO_ARBITRARY_FP:
+ Res = SoftPromoteHalfOp_CONVERT_TO_ARBITRARY_FP(N); break;
case ISD::STRICT_FP_EXTEND:
case ISD::FP_EXTEND: Res = SoftPromoteHalfOp_FP_EXTEND(N); break;
case ISD::SELECT_CC: Res = SoftPromoteHalfOp_SELECT_CC(N, OpNo); break;
@@ -3309,6 +3311,24 @@ SDValue DAGTypeLegalizer::SoftPromoteHalfOp_FP_TO_XINT_SAT(SDNode *N) {
N->getOperand(1));
}
+// TODO: CONVERT_TO_ARBITRARY_FP also needs SoftenFloatOperand and
+// ExpandFloatOperand handlers for targets with software float or ppcf128
+// source types. Same gap exists for CONVERT_FROM_ARBITRARY_FP.
+SDValue DAGTypeLegalizer::SoftPromoteHalfOp_CONVERT_TO_ARBITRARY_FP(
+ SDNode *N) {
+ EVT RVT = N->getValueType(0);
+ SDValue Op = N->getOperand(0);
+ EVT SVT = Op.getValueType();
+ SDLoc dl(N);
+
+ EVT NVT = TLI.getTypeToTransformTo(*DAG.getContext(), Op.getValueType());
+ Op = GetSoftPromotedHalf(Op);
+ SDValue Res = DAG.getNode(GetPromotionOpcode(SVT, RVT), dl, NVT, Op);
+
+ return DAG.getNode(ISD::CONVERT_TO_ARBITRARY_FP, dl, N->getValueType(0), Res,
+ N->getOperand(1), N->getOperand(2), N->getOperand(3));
+}
+
SDValue DAGTypeLegalizer::SoftPromoteHalfOp_SELECT_CC(SDNode *N,
unsigned OpNo) {
assert(OpNo == 0 && "Can only soften the comparison values");
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
index 238eba021c124..a87471c766911 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeIntegerTypes.cpp
@@ -196,6 +196,9 @@ void DAGTypeLegalizer::PromoteIntegerResult(SDNode *N, unsigned ResNo) {
case ISD::FP_TO_FP16:
Res = PromoteIntRes_FP_TO_FP16_BF16(N);
break;
+ case ISD::CONVERT_TO_ARBITRARY_FP:
+ Res = PromoteIntRes_CONVERT_TO_ARBITRARY_FP(N);
+ break;
case ISD::STRICT_FP_TO_BF16:
case ISD::STRICT_FP_TO_FP16:
Res = PromoteIntRes_STRICT_FP_TO_FP16_BF16(N);
@@ -952,6 +955,16 @@ SDValue DAGTypeLegalizer::PromoteIntRes_FP_TO_FP16_BF16(SDNode *N) {
return DAG.getNode(N->getOpcode(), dl, NVT, N->getOperand(0));
}
+// TODO: CONVERT_TO_ARBITRARY_FP also needs an ExpandIntegerResult handler for
+// wider arbitrary FP formats whose integer result requires expansion.
+SDValue DAGTypeLegalizer::PromoteIntRes_CONVERT_TO_ARBITRARY_FP(SDNode *N) {
+ EVT NVT = TLI.getTypeToTransformTo(*DAG.getContext(), N->getValueType(0));
+ SDLoc dl(N);
+
+ return DAG.getNode(ISD::CONVERT_TO_ARBITRARY_FP, dl, NVT, N->getOperand(0),
+ N->getOperand(1), N->getOperand(2), N->getOperand(3));
+}
+
SDValue DAGTypeLegalizer::PromoteIntRes_STRICT_FP_TO_FP16_BF16(SDNode *N) {
EVT NVT = TLI.getTypeToTransformTo(*DAG.getContext(), N->getValueType(0));
SDLoc dl(N);
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeTypes.h b/llvm/lib/CodeGen/SelectionDAG/LegalizeTypes.h
index 84c91a80ade79..52c6a2bc1ac43 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeTypes.h
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeTypes.h
@@ -313,6 +313,7 @@ class LLVM_LIBRARY_VISIBILITY DAGTypeLegalizer {
SDValue PromoteIntRes_FP_TO_XINT(SDNode *N);
SDValue PromoteIntRes_FP_TO_XINT_SAT(SDNode *N);
SDValue PromoteIntRes_FP_TO_FP16_BF16(SDNode *N);
+ SDValue PromoteIntRes_CONVERT_TO_ARBITRARY_FP(SDNode *N);
SDValue PromoteIntRes_STRICT_FP_TO_FP16_BF16(SDNode *N);
SDValue PromoteIntRes_XRINT(SDNode *N);
SDValue PromoteIntRes_FREEZE(SDNode *N);
@@ -804,6 +805,7 @@ class LLVM_LIBRARY_VISIBILITY DAGTypeLegalizer {
SDValue SoftPromoteHalfOp_FP_EXTEND(SDNode *N);
SDValue SoftPromoteHalfOp_Op0WithStrict(SDNode *N);
SDValue SoftPromoteHalfOp_FP_TO_XINT_SAT(SDNode *N);
+ SDValue SoftPromoteHalfOp_CONVERT_TO_ARBITRARY_FP(SDNode *N);
SDValue SoftPromoteHalfOp_SETCC(SDNode *N);
SDValue SoftPromoteHalfOp_SELECT_CC(SDNode *N, unsigned OpNo);
SDValue SoftPromoteHalfOp_STORE(SDNode *N, unsigned OpNo);
@@ -846,6 +848,7 @@ class LLVM_LIBRARY_VISIBILITY DAGTypeLegalizer {
SDValue ScalarizeVecRes_EXTRACT_SUBVECTOR(SDNode *N);
SDValue ScalarizeVecRes_FP_ROUND(SDNode *N);
SDValue ScalarizeVecRes_CONVERT_FROM_ARBITRARY_FP(SDNode *N);
+ SDValue ScalarizeVecRes_CONVERT_TO_ARBITRARY_FP(SDNode *N);
SDValue ScalarizeVecRes_UnaryOpWithExtraInput(SDNode *N);
SDValue ScalarizeVecRes_INSERT_VECTOR_ELT(SDNode *N);
SDValue ScalarizeVecRes_LOAD(LoadSDNode *N);
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp
index ccad9354fc820..49165abfc3223 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorOps.cpp
@@ -464,6 +464,7 @@ SDValue VectorLegalizer::LegalizeOp(SDValue Op) {
case ISD::SMULO:
case ISD::UMULO:
case ISD::CONVERT_FROM_ARBITRARY_FP:
+ case ISD::CONVERT_TO_ARBITRARY_FP:
case ISD::FCANONICALIZE:
case ISD::FFREXP:
case ISD::FMODF:
diff --git a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
index f3ddb9e4b811e..7a14d855e275c 100644
--- a/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/LegalizeVectorTypes.cpp
@@ -65,6 +65,9 @@ void DAGTypeLegalizer::ScalarizeVectorResult(SDNode *N, unsigned ResNo) {
case ISD::CONVERT_FROM_ARBITRARY_FP:
R = ScalarizeVecRes_CONVERT_FROM_ARBITRARY_FP(N);
break;
+ case ISD::CONVERT_TO_ARBITRARY_FP:
+ R = ScalarizeVecRes_CONVERT_TO_ARBITRARY_FP(N);
+ break;
case ISD::AssertZext:
case ISD::AssertSext:
case ISD::FPOWI:
@@ -529,6 +532,22 @@ SDValue DAGTypeLegalizer::ScalarizeVecRes_CONVERT_FROM_ARBITRARY_FP(SDNode *N) {
N->getOperand(1));
}
+SDValue DAGTypeLegalizer::ScalarizeVecRes_CONVERT_TO_ARBITRARY_FP(SDNode *N) {
+ SDLoc DL(N);
+ SDValue Op = N->getOperand(0);
+ EVT OpVT = Op.getValueType();
+ // The result needs scalarizing, but it's not a given that the source does.
+ if (getTypeAction(OpVT) == TargetLowering::TypeScalarizeVector) {
+ Op = GetScalarizedVector(Op);
+ } else {
+ EVT VT = OpVT.getVectorElementType();
+ Op = DAG.getExtractVectorElt(DL, VT, Op, 0);
+ }
+ return DAG.getNode(ISD::CONVERT_TO_ARBITRARY_FP, DL,
+ N->getValueType(0).getVectorElementType(), Op,
+ N->getOperand(1), N->getOperand(2), N->getOperand(3));
+}
+
SDValue DAGTypeLegalizer::ScalarizeVecRes_UnaryOpWithExtraInput(SDNode *N) {
SDValue Op = GetScalarizedVector(N->getOperand(0));
return DAG.getNode(N->getOpcode(), SDLoc(N), Op.getValueType(), Op,
@@ -872,6 +891,17 @@ bool DAGTypeLegalizer::ScalarizeVectorOperand(SDNode *N, unsigned OpNo) {
case ISD::CONVERT_FROM_ARBITRARY_FP:
Res = ScalarizeVecOp_UnaryOpWithExtraInput(N);
break;
+ case ISD::CONVERT_TO_ARBITRARY_FP: {
+ assert(N->getValueType(0).getVectorNumElements() == 1 &&
+ "Unexpected vector type!");
+ SDValue Elt = GetScalarizedVector(N->getOperand(0));
+ SDValue Op = DAG.getNode(N->getOpcode(), SDLoc(N),
+ N->getValueType(0).getScalarType(), Elt,
+ N->getOperand(1), N->getOperand(2),
+ N->getOperand(3));
+ Res = DAG.getNode(ISD::SCALAR_TO_VECTOR, SDLoc(N), N->getValueType(0), Op);
+ break;
+ }
case ISD::STRICT_SINT_TO_FP:
case ISD::STRICT_UINT_TO_FP:
case ISD::STRICT_FP_TO_SINT:
@@ -1475,6 +1505,7 @@ void DAGTypeLegalizer::SplitVectorResult(SDNode *N, unsigned ResNo) {
case ISD::FCANONICALIZE:
case ISD::AssertNoFPClass:
case ISD::CONVERT_FROM_ARBITRARY_FP:
+ case ISD::CONVERT_TO_ARBITRARY_FP:
SplitVecRes_UnaryOp(N, Lo, Hi);
break;
case ISD::ADDRSPACECAST:
@@ -2898,6 +2929,13 @@ void DAGTypeLegalizer::SplitVecRes_UnaryOp(SDNode *N, SDValue &Lo,
const SDNodeFlags Flags = N->getFlags();
unsigned Opcode = N->getOpcode();
+ if (Opcode == ISD::CONVERT_TO_ARBITRARY_FP) {
+ Lo = DAG.getNode(Opcode, dl, LoVT, Lo, N->getOperand(1),
+ N->getOperand(2), N->getOperand(3), Flags);
+ Hi = DAG.getNode(Opcode, dl, HiVT, Hi, N->getOperand(1),
+ N->getOperand(2), N->getOperand(3), Flags);
+ return;
+ }
if (N->getNumOperands() <= 2) {
if (Opcode == ISD::FP_ROUND || Opcode == ISD::AssertNoFPClass ||
Opcode == ISD::CONVERT_FROM_ARBITRARY_FP) {
@@ -3715,6 +3753,7 @@ bool DAGTypeLegalizer::SplitVectorOperand(SDNode *N, unsigned OpNo) {
case ISD::VP_FP_ROUND:
case ISD::FP_ROUND:
case ISD::CONVERT_FROM_ARBITRARY_FP:
+ case ISD::CONVERT_TO_ARBITRARY_FP:
Res = SplitVecOp_FP_ROUND(N);
break;
case ISD::FCOPYSIGN: Res = SplitVecOp_FPOpDifferentTypes(N); break;
@@ -4855,6 +4894,11 @@ SDValue DAGTypeLegalizer::SplitVecOp_FP_ROUND(SDNode *N) {
DAG.SplitEVL(N->getOperand(2), N->getValueType(0), DL);
Lo = DAG.getNode(ISD::VP_FP_ROUND, DL, OutVT, Lo, MaskLo, EVLLo);
Hi = DAG.getNode(ISD::VP_FP_ROUND, DL, OutVT, Hi, MaskHi, EVLHi);
+ } else if (N->getOpcode() == ISD::CONVERT_TO_ARBITRARY_FP) {
+ Lo = DAG.getNode(N->getOpcode(), DL, OutVT, Lo, N->getOperand(1),
+ N->getOperand(2), N->getOperand(3));
+ Hi = DAG.getNode(N->getOpcode(), DL, OutVT, Hi, N->getOperand(1),
+ N->getOperand(2), N->getOperand(3));
} else {
Lo = DAG.getNode(N->getOpcode(), DL, OutVT, Lo, N->getOperand(1));
Hi = DAG.getNode(N->getOpcode(), DL, OutVT, Hi, N->getOperand(1));
@@ -5294,6 +5338,7 @@ void DAGTypeLegalizer::WidenVectorResult(SDNode *N, unsigned ResNo) {
case ISD::ZERO_EXTEND:
case ISD::VP_ZERO_EXTEND:
case ISD::CONVERT_FROM_ARBITRARY_FP:
+ case ISD::CONVERT_TO_ARBITRARY_FP:
Res = WidenVecRes_Convert(N);
break;
@@ -5863,20 +5908,27 @@ SDValue DAGTypeLegalizer::WidenVecRes_Convert(SDNode *N) {
EVT InWidenVT = EVT::getVectorVT(Ctx, InEltVT, WidenEC);
ElementCount InVTEC = InVT.getVectorElementCount();
+ // Helper to build node with all scalar trailing operands.
+ auto MakeConvertNode = [&](EVT VT, SDValue Op) -> SDValue {
+ if (N->getNumOperands() == 1)
+ return DAG.getNode(Opcode, DL, VT, Op, Flags);
+ if (Opcode == ISD::CONVERT_TO_ARBITRARY_FP)
+ return DAG.getNode(Opcode, DL, VT, Op, N->getOperand(1),
+ N->getOperand(2), N->getOperand(3), Flags);
+ return DAG.getNode(Opcode, DL, VT, Op, N->getOperand(1), Flags);
+ };
+
if (getTypeAction(InVT) == TargetLowering::TypeWidenVector) {
InOp = GetWidenedVector(N->getOperand(0));
InVT = InOp.getValueType();
InVTEC = InVT.getVectorElementCount();
if (InVTEC == WidenEC) {
- if (N->getNumOperands() == 1)
- return DAG.getNode(Opcode, DL, WidenVT, InOp, Flags);
- if (N->getNumOperands() == 3) {
- assert(N->isVPOpcode() && "Expected VP opcode");
+ if (N->getNumOperands() == 3 && N->isVPOpcode()) {
SDValue Mask =
GetWidenedMask(N->getOperand(1), WidenVT.getVectorElementCount());
return DAG.getNode(Opcode, DL, WidenVT, InOp, Mask, N->getOperand(2));
}
- return DAG.getNode(Opcode, DL, WidenVT, InOp, N->getOperand(1), Flags);
+ return MakeConvertNode(WidenVT, InOp);
}
if (WidenVT.getSizeInBits() == InVT.getSizeInBits()) {
// If both input and result vector types are of same width, extend
@@ -5919,17 +5971,13 @@ SDValue DAGTypeLegalizer::WidenVecRes_Convert(SDNode *N) {
SmallVector<SDValue, 16> Ops(NumConcat, DAG.getPOISON(InVT));
Ops[0] = InOp;
SDValue InVec = DAG.getNode(ISD::CONCAT_VECTORS, DL, InWidenVT, Ops);
- if (N->getNumOperands() == 1)
- return DAG.getNode(Opcode, DL, WidenVT, InVec, Flags);
- return DAG.getNode(Opcode, DL, WidenVT, InVec, N->getOperand(1), Flags);
+ return MakeConvertNode(WidenVT, InVec);
}
if (InVTEC.isKnownMultipleOf(WidenEC.getKnownMinValue())) {
SDValue InVal = DAG.getExtractSubvector(DL, InWidenVT, InOp, 0);
// Extract the input and convert the shorten input vector.
- if (N->getNumOperands() == 1)
- return DAG.getNode(Opcode, DL, WidenVT, InVal, Flags);
- return DAG.getNode(Opcode, DL, WidenVT, InVal, N->getOperand(1), Flags);
+ return MakeConvertNode(WidenVT, InVal);
}
}
@@ -5941,10 +5989,7 @@ SDValue DAGTypeLegalizer::WidenVecRes_Convert(SDNode *N) {
unsigned MinElts = N->getValueType(0).getVectorNumElements();
for (unsigned i=0; i < MinElts; ++i) {
SDValue Val = DAG.getExtractVectorElt(DL, InEltVT, InOp, i);
- if (N->getNumOperands() == 1)
- Ops[i] = DAG.getNode(Opcode, DL, EltVT, Val, Flags);
- else
- Ops[i] = DAG.getNode(Opcode, DL, EltVT, Val, N->getOperand(1), Flags);
+ Ops[i] = MakeConvertNode(EltVT, Val);
}
return DAG.getBuildVector(WidenVT, DL, Ops);
@@ -7444,6 +7489,7 @@ bool DAGTypeLegalizer::WidenVectorOperand(SDNode *N, unsigned OpNo) {
case ISD::STRICT_UINT_TO_FP:
case ISD::TRUNCATE:
case ISD::CONVERT_FROM_ARBITRARY_FP:
+ case ISD::CONVERT_TO_ARBITRARY_FP:
Res = WidenVecOp_Convert(N);
break;
@@ -7648,6 +7694,16 @@ SDValue DAGTypeLegalizer::WidenVecOp_Convert(SDNode *N) {
EVT InVT = InOp.getValueType();
unsigned Opcode = N->getOpcode();
+ // Helper to build a convert node with all scalar trailing operands.
+ auto MakeConvertNode = [&](EVT VT, SDValue Op) -> SDValue {
+ if (Opcode == ISD::CONVERT_TO_ARBITRARY_FP)
+ return DAG.getNode(Opcode, dl, VT, Op, N->getOperand(1),
+ N->getOperand(2), N->getOperand(3));
+ if (Opcode == ISD::FP_ROUND || Opcode == ISD::CONVERT_FROM_ARBITRARY_FP)
+ return DAG.getNode(Opcode, dl, VT, Op, N->getOperand(1));
+ return DAG.getNode(Opcode, dl, VT, Op);
+ };
+
// See if a widened result type would be legal, if so widen the node.
// FIXME: This isn't safe for StrictFP. Other optimization here is needed.
EVT WideVT = EVT::getVectorVT(*DAG.getContext(), EltVT,
@@ -7665,10 +7721,7 @@ SDValue DAGTypeLegalizer::WidenVecOp_Convert(SDNode *N) {
// use the new one.
ReplaceValueWith(SDValue(N, 1), Res.getValue(1));
} else {
- if (Opcode == ISD::FP_ROUND || Opcode == ISD::CONVERT_FROM_ARBITRARY_FP)
- Res = DAG.getNode(Opcode, dl, WideVT, InOp, N->getOperand(1));
- else
- Res = DAG.getNode(Opcode, dl, WideVT, InOp);
+ Res = MakeConvertNode(WideVT, InOp);
}
return DAG.getExtractSubvector(dl, VT, Res, 0);
}
@@ -7691,10 +7744,7 @@ SDValue DAGTypeLegalizer::WidenVecOp_Convert(SDNode *N) {
} else {
for (unsigned i = 0; i < NumElts; ++i) {
SDValue Elt = DAG.getExtractVectorElt(dl, InEltVT, InOp, i);
- if (Opcode == ISD::FP_ROUND || Opcode == ISD::CONVERT_FROM_ARBITRARY_FP)
- Ops[i] = DAG.getNode(Opcode, dl, EltVT, Elt, N->getOperand(1));
- else
- Ops[i] = DAG.getNode(Opcode, dl, EltVT, Elt);
+ Ops[i] = MakeConvertNode(EltVT, Elt);
}
}
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
index 1f3b099c9c577..cc15e0c85148d 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
@@ -7194,6 +7194,49 @@ void SelectionDAGBuilder::visitIntrinsicCall(const CallInst &I,
SemConst));
return;
}
+ case Intrinsic::convert_to_arbitrary_fp: {
+ // Extract format metadata and convert to semantics enum.
+ EVT DstVT = TLI.getValueType(DAG.getDataLayout(), I.getType());
+ Metadata *MD = cast<MetadataAsValue>(I.getArgOperand(1))->getMetadata();
+ StringRef FormatStr = cast<MDString>(MD)->getString();
+ const fltSemantics *DstSem =
+ APFloatBase::getArbitraryFPSemantics(FormatStr);
+ if (!DstSem) {
+ DAG.getContext()->emitError(
+ "convert_to_arbitrary_fp: not implemented format '" + FormatStr +
+ "'");
+ setValue(&I, DAG.getPOISON(DstVT));
+ return;
+ }
+ APFloatBase::Semantics SemEnum = APFloatBase::SemanticsToEnum(*DstSem);
+
+ Metadata *RoundMD =
+ cast<MetadataAsValue>(I.getArgOperand(2))->getMetadata();
+ StringRef RoundStr = cast<MDString>(RoundMD)->getString();
+ std::optional<RoundingMode> RoundMode =
+ convertStrToRoundingMode(RoundStr);
+ if (!RoundMode || *RoundMode == RoundingMode::Dynamic) {
+ DAG.getContext()->emitError(
+ "convert_to_arbitrary_fp: unsupported rounding mode '" + RoundStr +
+ "'");
+ setValue(&I, DAG.getPOISON(DstVT));
+ return;
+ }
+
+ uint64_t Saturate =
+ cast<ConstantInt>(I.getArgOperand(3))->getZExtValue() ? 1 : 0;
+
+ SDValue FloatVal = getValue(I.getArgOperand(0));
+
+ SDValue SemConst =
+ DAG.getTargetConstant(static_cast<int>(SemEnum), sdl, MVT::i32);
+ SDValue RoundConst =
+ DAG.getTargetConstant(static_cast<int>(*RoundMode), sdl, MVT::i32);
+ SDValue SatConst = DAG.getTargetConstant(Saturate, sdl, MVT::i32);
+ setValue(&I, DAG.getNode(ISD::CONVERT_TO_ARBITRARY_FP, sdl, DstVT,
+ FloatVal, SemConst, RoundConst, SatConst));
+ return;
+ }
case Intrinsic::set_rounding:
Res = DAG.getNode(ISD::SET_ROUNDING, sdl, MVT::Other,
{getRoot(), getValue(I.getArgOperand(0))});
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp
index ce78072d21114..ebfa895d2ca91 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGDumper.cpp
@@ -436,6 +436,7 @@ std::string SDNode::getOperationName(const SelectionDAG *G) const {
case ISD::FP_TO_BF16: return "fp_to_bf16";
case ISD::STRICT_FP_TO_BF16: return "strict_fp_to_bf16";
case ISD::CONVERT_FROM_ARBITRARY_FP: return "convert_from_arbitrary_fp";
+ case ISD::CONVERT_TO_ARBITRARY_FP: return "convert_to_arbitrary_fp";
case ISD::LROUND: return "lround";
case ISD::STRICT_LROUND: return "strict_lround";
case ISD::LLROUND: return "llround";
diff --git a/llvm/lib/CodeGen/TargetLoweringBase.cpp b/llvm/lib/CodeGen/TargetLoweringBase.cpp
index 2f1e3f2f3ff7a..fbc821a42d899 100644
--- a/llvm/lib/CodeGen/TargetLoweringBase.cpp
+++ b/llvm/lib/CodeGen/TargetLoweringBase.cpp
@@ -1149,7 +1149,8 @@ void TargetLoweringBase::initActions() {
ISD::FASIN, ISD::FATAN,
ISD::FCOSH, ISD::FSINH,
ISD::FTANH, ISD::FATAN2,
- ISD::FMULADD, ISD::CONVERT_FROM_ARBITRARY_FP},
+ ISD::FMULADD, ISD::CONVERT_FROM_ARBITRARY_FP,
+ ISD::CONVERT_TO_ARBITRARY_FP},
VT, Expand);
// Overflow operations default to expand
diff --git a/llvm/test/CodeGen/AMDGPU/float-to-arbitrary-fp.ll b/llvm/test/CodeGen/AMDGPU/float-to-arbitrary-fp.ll
new file mode 100644
index 0000000000000..45639b9d50ca9
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/float-to-arbitrary-fp.ll
@@ -0,0 +1,905 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 | FileCheck %s
+
+; Test llvm.convert.to.arbitrary.fp intrinsic expansion.
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+declare i6 @llvm.convert.to.arbitrary.fp.i6.f32(float, metadata, metadata, i1)
+declare i4 @llvm.convert.to.arbitrary.fp.i4.f32(float, metadata, metadata, i1)
+declare <4 x i4> @llvm.convert.to.arbitrary.fp.v4i4.v4f32(<4 x float>, metadata, metadata, i1)
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f16(half, metadata, metadata, i1)
+declare i8 @llvm.convert.to.arbitrary.fp.i8.bf16(bfloat, metadata, metadata, i1)
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f64(double, metadata, metadata, i1)
+
+; Float8E5M2
+; Layout: sign(1) exp(5) mant(2), bias=15
+; Supports: Inf, NaN, signed zero, denormals
+
+; Float8E5M2 normal: f32 1.0 -> 0_01111_00 = 0x3C
+define i8 @to_f8e5m2_normal() {
+; CHECK-LABEL: to_f8e5m2_normal:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 60
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 zero: f32 0.0 -> 0_00000_00 = 0x00
+define i8 @to_f8e5m2_zero() {
+; CHECK-LABEL: to_f8e5m2_zero:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 negative zero: f32 -0.0 -> 1_00000_00 = 0x80
+define i8 @to_f8e5m2_neg_zero() {
+; CHECK-LABEL: to_f8e5m2_neg_zero:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x80
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float -0.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 +Inf: f32 Inf -> 0_11111_00 = 0x7C
+define i8 @to_f8e5m2_inf() {
+; CHECK-LABEL: to_f8e5m2_inf:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7c
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF0000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 NaN: f32 NaN -> qNaN
+define i8 @to_f8e5m2_nan() {
+; CHECK-LABEL: to_f8e5m2_nan:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7e
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF8000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 max: 57344.0 -> 0_11110_11 = 0x7B
+define i8 @to_f8e5m2_max() {
+; CHECK-LABEL: to_f8e5m2_max:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7b
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 57344.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 overflow (no saturation): 100000.0 -> Inf = 0x7C
+define i8 @to_f8e5m2_overflow() {
+; CHECK-LABEL: to_f8e5m2_overflow:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7c
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 100000.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 overflow with saturation: 100000.0 -> max = 0x7B
+define i8 @to_f8e5m2_overflow_sat() {
+; CHECK-LABEL: to_f8e5m2_overflow_sat:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7b
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 100000.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 true)
+ ret i8 %r
+}
+
+; Float8E5M2 denorm: very small value -> dst denorm
+define i8 @to_f8e5m2_denorm() {
+ ; 2^(-16) = 0x37800000 -> smallest E5M2 denorm 0_00000_01 = 0x01
+; CHECK-LABEL: to_f8e5m2_denorm:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 1
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x3EF0000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 runtime arg test
+define i8 @to_f8e5m2_dynamic(float %x) {
+; CHECK-LABEL: to_f8e5m2_dynamic:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v1, 0x7fffff, v0
+; CHECK-NEXT: v_ffbh_u32_e32 v3, v1
+; CHECK-NEXT: v_bfe_u32 v2, v0, 23, 8
+; CHECK-NEXT: v_add_u32_e32 v4, -8, v3
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v1
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, v4, v1
+; CHECK-NEXT: v_and_b32_e32 v4, 0x7fffff, v4
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v1, v4, s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v6, 0xfffff, v4
+; CHECK-NEXT: v_cmp_ne_u32_e64 s[6:7], 0, v6
+; CHECK-NEXT: v_lshrrev_b32_e32 v5, 21, v4
+; CHECK-NEXT: v_cndmask_b32_e64 v6, 0, 1, s[6:7]
+; CHECK-NEXT: v_and_or_b32 v6, v5, 1, v6
+; CHECK-NEXT: v_lshrrev_b32_e32 v7, 20, v4
+; CHECK-NEXT: v_and_b32_e32 v6, v7, v6
+; CHECK-NEXT: v_add_u32_e32 v5, v5, v6
+; CHECK-NEXT: v_sub_u32_e32 v3, 9, v3
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[6:7], 3, v5
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v2, v3, s[4:5]
+; CHECK-NEXT: v_mov_b32_e32 v6, 0xffffff90
+; CHECK-NEXT: v_addc_co_u32_e64 v6, s[4:5], v3, v6, s[6:7]
+; CHECK-NEXT: v_sub_u32_e32 v3, 0x86, v3
+; CHECK-NEXT: v_min_u32_e32 v3, 31, v3
+; CHECK-NEXT: v_or_b32_e32 v4, 0x800000, v4
+; CHECK-NEXT: v_sub_u32_e64 v9, v3, 1 clamp
+; CHECK-NEXT: s_movk_i32 s4, 0x80
+; CHECK-NEXT: v_bfe_u32 v10, v4, 0, v9
+; CHECK-NEXT: v_and_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:DWORD
+; CHECK-NEXT: v_cmp_ne_u32_e64 s[4:5], 0, v10
+; CHECK-NEXT: v_lshrrev_b32_e32 v8, v3, v4
+; CHECK-NEXT: v_cndmask_b32_e64 v10, 0, 1, s[4:5]
+; CHECK-NEXT: v_and_or_b32 v10, v8, 1, v10
+; CHECK-NEXT: v_lshrrev_b32_e32 v4, v9, v4
+; CHECK-NEXT: v_and_b32_e32 v4, v4, v10
+; CHECK-NEXT: v_cmp_ne_u32_e64 s[4:5], 0, v3
+; CHECK-NEXT: v_cndmask_b32_e64 v3, 0, v4, s[4:5]
+; CHECK-NEXT: v_add_u32_e32 v3, v8, v3
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 3, v3
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, 0, s[6:7]
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, 2, v6
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, 0, s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v4, 0, 4, s[4:5]
+; CHECK-NEXT: v_or3_b32 v7, v0, v7, v5
+; CHECK-NEXT: v_or3_b32 v3, v0, v4, v3
+; CHECK-NEXT: v_cmp_gt_i32_e64 s[4:5], 1, v6
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v7, v3, s[4:5]
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 3, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 30, v6
+; CHECK-NEXT: s_and_b64 s[6:7], s[6:7], s[4:5]
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 30, v6
+; CHECK-NEXT: v_or_b32_e32 v4, 0x7c, v0
+; CHECK-NEXT: s_or_b64 s[4:5], s[4:5], s[6:7]
+; CHECK-NEXT: v_or_b32_e32 v5, v2, v1
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, v4, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5
+; CHECK-NEXT: s_movk_i32 s6, 0xff
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v3, v0, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], s6, v2
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v0, v4, s[4:5]
+; CHECK-NEXT: v_mov_b32_e32 v1, 0x7e
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v0, v1, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN (NanOnly, NanEncoding=AllOnes)
+; Layout: sign(1) exp(4) mant(3), maxExp=8, minExp=-6, bias=7
+; Only 0_1111_111 and 1_1111_111 are NaN; all other exp=15 values are finite.
+
+; Float8E4M3FN normal: f32 1.0 -> 0_0111_000 = 0x38
+define i8 @to_f8e4m3fn_normal() {
+; CHECK-LABEL: to_f8e4m3fn_normal:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 56
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN max: 448.0 -> 0_1111_110 = 0x7E
+define i8 @to_f8e4m3fn_max() {
+; CHECK-LABEL: to_f8e4m3fn_max:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7e
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 448.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN NaN: f32 NaN -> 0_1111_111 = 0x7F
+define i8 @to_f8e4m3fn_nan() {
+; CHECK-LABEL: to_f8e4m3fn_nan:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7f
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF8000000000000, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN overflow with saturation: 1000.0 -> max = 0x7E
+define i8 @to_f8e4m3fn_overflow_sat() {
+; CHECK-LABEL: to_f8e4m3fn_overflow_sat:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7e
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1000.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 true)
+ ret i8 %r
+}
+
+; Float6E3M2FN (FiniteOnly)
+; Layout: sign(1) exp(3) mant(2), bias=3, maxExp=4
+; No Inf, no NaN. All bit patterns are finite.
+
+; Float6E3M2FN normal: f32 1.0 -> 0_011_00 = 0x0C
+define i6 @to_f6e3m2fn_normal() {
+; CHECK-LABEL: to_f6e3m2fn_normal:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 12
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 1.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN max: 28.0 -> 0_111_11 = 0x1F
+define i6 @to_f6e3m2fn_max() {
+; CHECK-LABEL: to_f6e3m2fn_max:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 31
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 28.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN zero: f32 0.0 -> 0_000_00 = 0x00
+define i6 @to_f6e3m2fn_zero() {
+; CHECK-LABEL: to_f6e3m2fn_zero:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 0.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN negative: f32 -1.0 -> 1_011_00
+define i6 @to_f6e3m2fn_negative() {
+; CHECK-LABEL: to_f6e3m2fn_negative:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 44
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float -1.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E2M3FN (FiniteOnly)
+; Layout: sign(1) exp(2) mant(3), bias=1, maxExp=2
+; No Inf, no NaN. All bit patterns are finite.
+
+; Float6E2M3FN normal: f32 1.0 -> 0_01_000 = 0x08
+define i6 @to_f6e2m3fn_normal() {
+; CHECK-LABEL: to_f6e2m3fn_normal:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 8
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 1.0, metadata !"Float6E2M3FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E2M3FN max: 7.5 -> 0_11_111 = 0x1F
+define i6 @to_f6e2m3fn_max() {
+; CHECK-LABEL: to_f6e2m3fn_max:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 31
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 7.5, metadata !"Float6E2M3FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float4E2M1FN (FiniteOnly)
+; Layout: sign(1) exp(2) mant(1), bias=1, maxExp=2
+; No Inf, no NaN.
+
+; Float4E2M1FN normal: f32 1.0 -> 0_01_0 = 0x2
+define i4 @to_f4e2m1fn_normal() {
+; CHECK-LABEL: to_f4e2m1fn_normal:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 2
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 1.0, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Float4E2M1FN max: 6.0 -> 0_11_1 = 0x7
+define i4 @to_f4e2m1fn_max() {
+; CHECK-LABEL: to_f4e2m1fn_max:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 7
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 6.0, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Float4E2M1FN denorm: 0.5 -> 0_00_1 = 0x1
+define i4 @to_f4e2m1fn_denorm() {
+; CHECK-LABEL: to_f4e2m1fn_denorm:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 1
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 0.5, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Rounding mode tests
+
+; Round to nearest
+define i8 @to_f8e5m2_round_nearest() {
+; CHECK-LABEL: to_f8e5m2_round_nearest:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, 62
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.5, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Round toward zero with runtime arg
+define i8 @to_f8e5m2_round_towardzero(float %x) {
+; CHECK-LABEL: to_f8e5m2_round_towardzero:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v1, 0x7fffff, v0
+; CHECK-NEXT: v_bfe_u32 v2, v0, 23, 8
+; CHECK-NEXT: v_ffbh_u32_e32 v3, v1
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v1
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_add_u32_e32 v4, -8, v3
+; CHECK-NEXT: v_lshlrev_b32_e32 v4, v4, v1
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_sub_u32_e32 v3, 9, v3
+; CHECK-NEXT: v_and_b32_e32 v4, 0x7fffff, v4
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v2, v3, s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v1, v4, s[4:5]
+; CHECK-NEXT: v_sub_u32_e32 v6, 0x86, v3
+; CHECK-NEXT: v_or_b32_e32 v5, 0x800000, v4
+; CHECK-NEXT: v_min_u32_e32 v6, 31, v6
+; CHECK-NEXT: v_lshrrev_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 3, v5
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, 0, s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v6, 0, 4, s[4:5]
+; CHECK-NEXT: s_movk_i32 s4, 0x80
+; CHECK-NEXT: v_and_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:DWORD
+; CHECK-NEXT: v_lshrrev_b32_e32 v4, 21, v4
+; CHECK-NEXT: v_or3_b32 v5, v0, v6, v5
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 3, v4
+; CHECK-NEXT: v_mov_b32_e32 v6, 0xffffff90
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v4, 0, s[4:5]
+; CHECK-NEXT: v_addc_co_u32_e64 v3, s[4:5], v3, v6, s[4:5]
+; CHECK-NEXT: v_lshlrev_b32_e32 v6, 2, v3
+; CHECK-NEXT: v_or3_b32 v6, v0, v6, v4
+; CHECK-NEXT: v_cmp_gt_i32_e64 s[4:5], 1, v3
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v6, v5, s[4:5]
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 3, v4
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], 30, v3
+; CHECK-NEXT: s_and_b64 s[6:7], s[6:7], s[4:5]
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 30, v3
+; CHECK-NEXT: v_or_b32_e32 v3, 0x7c, v0
+; CHECK-NEXT: s_or_b64 s[4:5], s[4:5], s[6:7]
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v5, v3, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v5, v2, v1
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v5
+; CHECK-NEXT: s_movk_i32 s6, 0xff
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v4, v0, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[6:7], s6, v2
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v0, v3, s[4:5]
+; CHECK-NEXT: v_mov_b32_e32 v1, 0x7e
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v0, v1, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float %x, metadata !"Float8E5M2", metadata !"round.towardzero", i1 false)
+ ret i8 %r
+}
+
+; Vector test: f32 <4 x float> -> <4 x i4> Float4E2M1FN
+define <4 x i4> @to_f4e2m1fn_v4f32(<4 x float> %x) {
+; CHECK-LABEL: to_f4e2m1fn_v4f32:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v5, 0x7fffff, v0
+; CHECK-NEXT: v_ffbh_u32_e32 v4, v5
+; CHECK-NEXT: v_bfe_u32 v6, v0, 23, 8
+; CHECK-NEXT: v_add_u32_e32 v7, -8, v4
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v7, v7, v5
+; CHECK-NEXT: v_and_b32_e32 v7, 0x7fffff, v7
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v7, v5, v7, vcc
+; CHECK-NEXT: v_and_b32_e32 v9, 0x1fffff, v7
+; CHECK-NEXT: v_cmp_ne_u32_e64 s[4:5], 0, v9
+; CHECK-NEXT: v_lshrrev_b32_e32 v8, 22, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v9, 0, 1, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v9, v9, v8
+; CHECK-NEXT: v_lshrrev_b32_e32 v10, 21, v7
+; CHECK-NEXT: v_and_b32_e32 v9, v10, v9
+; CHECK-NEXT: v_add_u32_e32 v8, v8, v9
+; CHECK-NEXT: v_sub_u32_e32 v4, 9, v4
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 1, v8
+; CHECK-NEXT: v_cndmask_b32_e32 v9, v6, v4, vcc
+; CHECK-NEXT: v_mov_b32_e32 v4, 0xffffff82
+; CHECK-NEXT: v_addc_co_u32_e64 v10, vcc, v9, v4, s[4:5]
+; CHECK-NEXT: v_sub_u32_e32 v9, 0x95, v9
+; CHECK-NEXT: v_min_u32_e32 v9, 31, v9
+; CHECK-NEXT: v_or_b32_e32 v7, 0x800000, v7
+; CHECK-NEXT: v_sub_u32_e64 v12, v9, 1 clamp
+; CHECK-NEXT: v_lshrrev_b32_e32 v0, 28, v0
+; CHECK-NEXT: v_bfe_u32 v13, v7, 0, v12
+; CHECK-NEXT: v_cndmask_b32_e64 v8, v8, 0, s[4:5]
+; CHECK-NEXT: v_lshlrev_b32_e32 v11, 1, v10
+; CHECK-NEXT: v_and_b32_e32 v0, 8, v0
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v13
+; CHECK-NEXT: v_or3_b32 v8, v0, v11, v8
+; CHECK-NEXT: v_lshrrev_b32_e32 v11, v9, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v13, 0, 1, vcc
+; CHECK-NEXT: v_and_or_b32 v13, v11, 1, v13
+; CHECK-NEXT: v_lshrrev_b32_e32 v7, v12, v7
+; CHECK-NEXT: v_and_b32_e32 v7, v7, v13
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v9
+; CHECK-NEXT: v_cndmask_b32_e32 v7, 0, v7, vcc
+; CHECK-NEXT: v_add_u32_e32 v7, v11, v7
+; CHECK-NEXT: v_cmp_lt_i32_e32 vcc, 1, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v7, v7, 0, vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v9, 0, 2, vcc
+; CHECK-NEXT: v_or3_b32 v7, v0, v9, v7
+; CHECK-NEXT: v_cmp_gt_i32_e32 vcc, 1, v10
+; CHECK-NEXT: v_or_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_cndmask_b32_e32 v7, v8, v7, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_and_b32_e32 v5, 0x7fffff, v1
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v7, v0, vcc
+; CHECK-NEXT: v_ffbh_u32_e32 v7, v5
+; CHECK-NEXT: v_bfe_u32 v6, v1, 23, 8
+; CHECK-NEXT: v_add_u32_e32 v8, -8, v7
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, v8, v5
+; CHECK-NEXT: v_and_b32_e32 v8, 0x7fffff, v8
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v8, v5, v8, vcc
+; CHECK-NEXT: v_and_b32_e32 v10, 0x1fffff, v8
+; CHECK-NEXT: v_cmp_ne_u32_e64 s[4:5], 0, v10
+; CHECK-NEXT: v_lshrrev_b32_e32 v9, 22, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v10, 0, 1, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v10, v10, v9
+; CHECK-NEXT: v_lshrrev_b32_e32 v11, 21, v8
+; CHECK-NEXT: v_and_b32_e32 v10, v11, v10
+; CHECK-NEXT: v_add_u32_e32 v9, v9, v10
+; CHECK-NEXT: v_sub_u32_e32 v7, 9, v7
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 1, v9
+; CHECK-NEXT: v_cndmask_b32_e32 v7, v6, v7, vcc
+; CHECK-NEXT: v_addc_co_u32_e64 v10, vcc, v7, v4, s[4:5]
+; CHECK-NEXT: v_sub_u32_e32 v7, 0x95, v7
+; CHECK-NEXT: v_min_u32_e32 v7, 31, v7
+; CHECK-NEXT: v_or_b32_e32 v8, 0x800000, v8
+; CHECK-NEXT: v_sub_u32_e64 v12, v7, 1 clamp
+; CHECK-NEXT: v_lshrrev_b32_e32 v1, 28, v1
+; CHECK-NEXT: v_bfe_u32 v13, v8, 0, v12
+; CHECK-NEXT: v_cndmask_b32_e64 v9, v9, 0, s[4:5]
+; CHECK-NEXT: v_lshlrev_b32_e32 v11, 1, v10
+; CHECK-NEXT: v_and_b32_e32 v1, 8, v1
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v13
+; CHECK-NEXT: v_or3_b32 v9, v1, v11, v9
+; CHECK-NEXT: v_lshrrev_b32_e32 v11, v7, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v13, 0, 1, vcc
+; CHECK-NEXT: v_and_or_b32 v13, v11, 1, v13
+; CHECK-NEXT: v_lshrrev_b32_e32 v8, v12, v8
+; CHECK-NEXT: v_and_b32_e32 v8, v8, v13
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v7
+; CHECK-NEXT: v_cndmask_b32_e32 v7, 0, v8, vcc
+; CHECK-NEXT: v_add_u32_e32 v7, v11, v7
+; CHECK-NEXT: v_cmp_lt_i32_e32 vcc, 1, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v7, v7, 0, vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v8, 0, 2, vcc
+; CHECK-NEXT: v_or3_b32 v7, v1, v8, v7
+; CHECK-NEXT: v_cmp_gt_i32_e32 vcc, 1, v10
+; CHECK-NEXT: v_or_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_cndmask_b32_e32 v7, v9, v7, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_and_b32_e32 v5, 0x7fffff, v2
+; CHECK-NEXT: v_cndmask_b32_e32 v1, v7, v1, vcc
+; CHECK-NEXT: v_ffbh_u32_e32 v7, v5
+; CHECK-NEXT: v_bfe_u32 v6, v2, 23, 8
+; CHECK-NEXT: v_add_u32_e32 v8, -8, v7
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, v8, v5
+; CHECK-NEXT: v_and_b32_e32 v8, 0x7fffff, v8
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v8, v5, v8, vcc
+; CHECK-NEXT: v_and_b32_e32 v10, 0x1fffff, v8
+; CHECK-NEXT: v_cmp_ne_u32_e64 s[4:5], 0, v10
+; CHECK-NEXT: v_lshrrev_b32_e32 v9, 22, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v10, 0, 1, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v10, v10, v9
+; CHECK-NEXT: v_lshrrev_b32_e32 v11, 21, v8
+; CHECK-NEXT: v_and_b32_e32 v10, v11, v10
+; CHECK-NEXT: v_add_u32_e32 v9, v9, v10
+; CHECK-NEXT: v_sub_u32_e32 v7, 9, v7
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 1, v9
+; CHECK-NEXT: v_cndmask_b32_e32 v7, v6, v7, vcc
+; CHECK-NEXT: v_addc_co_u32_e64 v10, vcc, v7, v4, s[4:5]
+; CHECK-NEXT: v_sub_u32_e32 v7, 0x95, v7
+; CHECK-NEXT: v_min_u32_e32 v7, 31, v7
+; CHECK-NEXT: v_or_b32_e32 v8, 0x800000, v8
+; CHECK-NEXT: v_sub_u32_e64 v12, v7, 1 clamp
+; CHECK-NEXT: v_lshrrev_b32_e32 v2, 28, v2
+; CHECK-NEXT: v_bfe_u32 v13, v8, 0, v12
+; CHECK-NEXT: v_cndmask_b32_e64 v9, v9, 0, s[4:5]
+; CHECK-NEXT: v_lshlrev_b32_e32 v11, 1, v10
+; CHECK-NEXT: v_and_b32_e32 v2, 8, v2
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v13
+; CHECK-NEXT: v_or3_b32 v9, v2, v11, v9
+; CHECK-NEXT: v_lshrrev_b32_e32 v11, v7, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v13, 0, 1, vcc
+; CHECK-NEXT: v_and_or_b32 v13, v11, 1, v13
+; CHECK-NEXT: v_lshrrev_b32_e32 v8, v12, v8
+; CHECK-NEXT: v_and_b32_e32 v8, v8, v13
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v7
+; CHECK-NEXT: v_cndmask_b32_e32 v7, 0, v8, vcc
+; CHECK-NEXT: v_add_u32_e32 v7, v11, v7
+; CHECK-NEXT: v_cmp_lt_i32_e32 vcc, 1, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v7, v7, 0, vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v8, 0, 2, vcc
+; CHECK-NEXT: v_or3_b32 v7, v2, v8, v7
+; CHECK-NEXT: v_cmp_gt_i32_e32 vcc, 1, v10
+; CHECK-NEXT: v_or_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_cndmask_b32_e32 v7, v9, v7, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_and_b32_e32 v5, 0x7fffff, v3
+; CHECK-NEXT: v_cndmask_b32_e32 v2, v7, v2, vcc
+; CHECK-NEXT: v_ffbh_u32_e32 v7, v5
+; CHECK-NEXT: v_bfe_u32 v6, v3, 23, 8
+; CHECK-NEXT: v_add_u32_e32 v8, -8, v7
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cmp_eq_u32_e64 s[4:5], 0, v6
+; CHECK-NEXT: v_lshlrev_b32_e32 v8, v8, v5
+; CHECK-NEXT: v_and_b32_e32 v8, 0x7fffff, v8
+; CHECK-NEXT: s_and_b64 vcc, s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v8, v5, v8, vcc
+; CHECK-NEXT: v_and_b32_e32 v10, 0x1fffff, v8
+; CHECK-NEXT: v_cmp_ne_u32_e64 s[4:5], 0, v10
+; CHECK-NEXT: v_lshrrev_b32_e32 v9, 22, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v10, 0, 1, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v10, v10, v9
+; CHECK-NEXT: v_lshrrev_b32_e32 v11, 21, v8
+; CHECK-NEXT: v_and_b32_e32 v10, v11, v10
+; CHECK-NEXT: v_add_u32_e32 v9, v9, v10
+; CHECK-NEXT: v_sub_u32_e32 v7, 9, v7
+; CHECK-NEXT: v_cmp_lt_i32_e64 s[4:5], 1, v9
+; CHECK-NEXT: v_cndmask_b32_e32 v7, v6, v7, vcc
+; CHECK-NEXT: v_addc_co_u32_e64 v4, vcc, v7, v4, s[4:5]
+; CHECK-NEXT: v_sub_u32_e32 v7, 0x95, v7
+; CHECK-NEXT: v_min_u32_e32 v7, 31, v7
+; CHECK-NEXT: v_or_b32_e32 v8, 0x800000, v8
+; CHECK-NEXT: v_sub_u32_e64 v11, v7, 1 clamp
+; CHECK-NEXT: v_lshrrev_b32_e32 v3, 28, v3
+; CHECK-NEXT: v_bfe_u32 v12, v8, 0, v11
+; CHECK-NEXT: v_cndmask_b32_e64 v9, v9, 0, s[4:5]
+; CHECK-NEXT: v_lshlrev_b32_e32 v10, 1, v4
+; CHECK-NEXT: v_and_b32_e32 v3, 8, v3
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v12
+; CHECK-NEXT: v_or3_b32 v9, v3, v10, v9
+; CHECK-NEXT: v_lshrrev_b32_e32 v10, v7, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v12, 0, 1, vcc
+; CHECK-NEXT: v_and_or_b32 v12, v10, 1, v12
+; CHECK-NEXT: v_lshrrev_b32_e32 v8, v11, v8
+; CHECK-NEXT: v_and_b32_e32 v8, v8, v12
+; CHECK-NEXT: v_cmp_ne_u32_e32 vcc, 0, v7
+; CHECK-NEXT: v_cndmask_b32_e32 v7, 0, v8, vcc
+; CHECK-NEXT: v_add_u32_e32 v7, v10, v7
+; CHECK-NEXT: v_cmp_lt_i32_e32 vcc, 1, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v7, v7, 0, vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v8, 0, 2, vcc
+; CHECK-NEXT: v_or3_b32 v7, v3, v8, v7
+; CHECK-NEXT: v_cmp_gt_i32_e32 vcc, 1, v4
+; CHECK-NEXT: v_or_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_cndmask_b32_e32 v4, v9, v7, vcc
+; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 0, v5
+; CHECK-NEXT: v_cndmask_b32_e32 v3, v4, v3, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call <4 x i4> @llvm.convert.to.arbitrary.fp.v4i4.v4f32(<4 x float> %x, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret <4 x i4> %r
+}
+
+; Float8E5M2 from f16: half -> i8
+define i8 @to_f8e5m2_from_f16(half %x) {
+; CHECK-LABEL: to_f8e5m2_from_f16:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v1, 0x3ff, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v3, 16, v1
+; CHECK-NEXT: v_ffbh_u32_e32 v3, v3
+; CHECK-NEXT: v_bfe_u32 v2, v0, 10, 5
+; CHECK-NEXT: v_add_u16_e32 v4, -5, v3
+; CHECK-NEXT: v_cmp_ne_u16_e32 vcc, 0, v1
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_lshlrev_b16_e32 v4, v4, v1
+; CHECK-NEXT: v_and_b32_e32 v4, 0x3ff, v4
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v1, v4, s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v6, 0x7f, v4
+; CHECK-NEXT: v_mov_b32_e32 v5, 1
+; CHECK-NEXT: v_cmp_ne_u16_e64 s[6:7], 0, v6
+; CHECK-NEXT: v_and_b32_sdwa v5, v4, v5 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:DWORD
+; CHECK-NEXT: v_cndmask_b32_e64 v6, 0, 1, s[6:7]
+; CHECK-NEXT: v_or_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_lshrrev_b16_e32 v6, 7, v4
+; CHECK-NEXT: v_and_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_add_u16_sdwa v5, v4, v5 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:DWORD
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[6:7], 3, v5
+; CHECK-NEXT: v_sub_u16_e32 v3, 6, v3
+; CHECK-NEXT: v_cndmask_b32_e64 v6, 0, 1, s[6:7]
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v2, v3, s[4:5]
+; CHECK-NEXT: v_add_u16_e32 v6, v3, v6
+; CHECK-NEXT: v_sub_u16_e32 v3, 9, v3
+; CHECK-NEXT: v_min_u16_e32 v3, 15, v3
+; CHECK-NEXT: v_sub_u16_e64 v10, v3, 1 clamp
+; CHECK-NEXT: v_lshlrev_b16_e64 v11, v10, 1
+; CHECK-NEXT: v_or_b32_e32 v4, 0x400, v4
+; CHECK-NEXT: v_add_u16_e32 v11, -1, v11
+; CHECK-NEXT: s_movk_i32 s4, 0x80
+; CHECK-NEXT: v_and_b32_e32 v11, v4, v11
+; CHECK-NEXT: v_and_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:DWORD
+; CHECK-NEXT: v_lshrrev_b16_e32 v8, v3, v4
+; CHECK-NEXT: v_cmp_ne_u16_e64 s[4:5], 0, v11
+; CHECK-NEXT: v_and_b32_e32 v9, 1, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v11, 0, 1, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v9, v11, v9
+; CHECK-NEXT: v_lshrrev_b16_e32 v4, v10, v4
+; CHECK-NEXT: v_and_b32_e32 v4, v4, v9
+; CHECK-NEXT: v_cmp_ne_u16_e64 s[4:5], 0, v3
+; CHECK-NEXT: v_cndmask_b32_e64 v3, 0, v4, s[4:5]
+; CHECK-NEXT: v_add_u16_e32 v3, v8, v3
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[4:5], 3, v3
+; CHECK-NEXT: v_lshlrev_b16_e32 v7, 2, v6
+; CHECK-NEXT: v_cndmask_b32_e64 v4, 0, 4, s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, 0, s[6:7]
+; CHECK-NEXT: v_or_b32_e32 v7, v0, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, 0, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v4, v0, v4
+; CHECK-NEXT: v_or_b32_e32 v7, v7, v5
+; CHECK-NEXT: v_or_b32_e32 v3, v4, v3
+; CHECK-NEXT: v_cmp_gt_i16_e64 s[4:5], 1, v6
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v7, v3, s[4:5]
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[4:5], 3, v5
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[6:7], 30, v6
+; CHECK-NEXT: s_and_b64 s[6:7], s[6:7], s[4:5]
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[4:5], 30, v6
+; CHECK-NEXT: v_or_b32_e32 v4, 0x7c, v0
+; CHECK-NEXT: s_or_b64 s[4:5], s[4:5], s[6:7]
+; CHECK-NEXT: v_or_b32_e32 v5, v2, v1
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, v4, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v5
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v3, v0, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[6:7], 31, v2
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v0, v4, s[4:5]
+; CHECK-NEXT: v_mov_b32_e32 v1, 0x7e
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v0, v1, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f16(half %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 from bf16: bfloat -> i8
+define i8 @to_f8e5m2_from_bf16(bfloat %x) {
+; CHECK-LABEL: to_f8e5m2_from_bf16:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_and_b32_e32 v1, 0x7f, v0
+; CHECK-NEXT: v_lshlrev_b32_e32 v3, 16, v1
+; CHECK-NEXT: v_ffbh_u32_e32 v3, v3
+; CHECK-NEXT: v_bfe_u32 v2, v0, 7, 8
+; CHECK-NEXT: v_add_u16_e32 v4, -8, v3
+; CHECK-NEXT: v_cmp_ne_u16_e32 vcc, 0, v1
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v2
+; CHECK-NEXT: v_lshlrev_b16_e32 v4, v4, v1
+; CHECK-NEXT: v_and_b32_e32 v4, 0x7f, v4
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v4, v1, v4, s[4:5]
+; CHECK-NEXT: v_and_b32_e32 v7, 15, v4
+; CHECK-NEXT: v_lshrrev_b16_e32 v5, 5, v4
+; CHECK-NEXT: v_cmp_ne_u16_e64 s[6:7], 0, v7
+; CHECK-NEXT: v_and_b32_e32 v6, 1, v5
+; CHECK-NEXT: v_cndmask_b32_e64 v7, 0, 1, s[6:7]
+; CHECK-NEXT: v_or_b32_e32 v6, v7, v6
+; CHECK-NEXT: v_lshrrev_b16_e32 v7, 4, v4
+; CHECK-NEXT: v_and_b32_e32 v6, v7, v6
+; CHECK-NEXT: v_add_u16_e32 v5, v5, v6
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[6:7], 3, v5
+; CHECK-NEXT: v_sub_u16_e32 v3, 9, v3
+; CHECK-NEXT: v_cndmask_b32_e64 v6, 0, 1, s[6:7]
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v2, v3, s[4:5]
+; CHECK-NEXT: v_add_u16_e32 v6, v3, v6
+; CHECK-NEXT: v_sub_u16_e32 v3, 0x76, v3
+; CHECK-NEXT: v_min_u16_e32 v3, 15, v3
+; CHECK-NEXT: v_sub_u16_e64 v10, v3, 1 clamp
+; CHECK-NEXT: v_lshlrev_b16_e64 v11, v10, 1
+; CHECK-NEXT: v_or_b32_e32 v4, 0x80, v4
+; CHECK-NEXT: v_add_u16_e32 v11, -1, v11
+; CHECK-NEXT: s_movk_i32 s4, 0x80
+; CHECK-NEXT: v_and_b32_e32 v11, v4, v11
+; CHECK-NEXT: v_and_b32_sdwa v0, v0, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:DWORD
+; CHECK-NEXT: v_lshrrev_b16_e32 v8, v3, v4
+; CHECK-NEXT: v_cmp_ne_u16_e64 s[4:5], 0, v11
+; CHECK-NEXT: v_and_b32_e32 v9, 1, v8
+; CHECK-NEXT: v_cndmask_b32_e64 v11, 0, 1, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v9, v11, v9
+; CHECK-NEXT: v_lshrrev_b16_e32 v4, v10, v4
+; CHECK-NEXT: v_and_b32_e32 v4, v4, v9
+; CHECK-NEXT: v_cmp_ne_u16_e64 s[4:5], 0, v3
+; CHECK-NEXT: v_cndmask_b32_e64 v3, 0, v4, s[4:5]
+; CHECK-NEXT: v_add_u16_e32 v3, v8, v3
+; CHECK-NEXT: v_add_u16_e32 v6, 0xff90, v6
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[4:5], 3, v3
+; CHECK-NEXT: v_lshlrev_b16_e32 v7, 2, v6
+; CHECK-NEXT: v_cndmask_b32_e64 v4, 0, 4, s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, 0, s[6:7]
+; CHECK-NEXT: v_or_b32_e32 v7, v0, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, 0, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v4, v0, v4
+; CHECK-NEXT: v_or_b32_e32 v7, v7, v5
+; CHECK-NEXT: v_or_b32_e32 v3, v4, v3
+; CHECK-NEXT: v_cmp_gt_i16_e64 s[4:5], 1, v6
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v7, v3, s[4:5]
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[4:5], 3, v5
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[6:7], 30, v6
+; CHECK-NEXT: s_and_b64 s[6:7], s[6:7], s[4:5]
+; CHECK-NEXT: v_cmp_lt_i16_e64 s[4:5], 30, v6
+; CHECK-NEXT: v_or_b32_e32 v4, 0x7c, v0
+; CHECK-NEXT: s_or_b64 s[4:5], s[4:5], s[6:7]
+; CHECK-NEXT: v_or_b32_e32 v5, v2, v1
+; CHECK-NEXT: v_cndmask_b32_e64 v3, v3, v4, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v5
+; CHECK-NEXT: s_movk_i32 s6, 0xff
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v3, v0, s[4:5]
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[4:5], 0, v1
+; CHECK-NEXT: v_cmp_eq_u16_e64 s[6:7], s6, v2
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v0, v4, s[4:5]
+; CHECK-NEXT: v_mov_b32_e32 v1, 0x7e
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v0, v1, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.bf16(bfloat %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 from f64: double -> i8
+define i8 @to_f8e5m2_from_f64(double %x) {
+; CHECK-LABEL: to_f8e5m2_from_f64:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v2, v1
+; CHECK-NEXT: v_and_b32_e32 v1, 0xfffff, v2
+; CHECK-NEXT: v_ffbh_u32_e32 v5, v0
+; CHECK-NEXT: v_bfe_u32 v3, v2, 20, 11
+; CHECK-NEXT: v_mov_b32_e32 v4, 0
+; CHECK-NEXT: v_add_u32_e32 v5, 32, v5
+; CHECK-NEXT: v_ffbh_u32_e32 v6, v1
+; CHECK-NEXT: v_cmp_ne_u64_e32 vcc, 0, v[0:1]
+; CHECK-NEXT: v_cmp_eq_u64_e64 s[4:5], 0, v[3:4]
+; CHECK-NEXT: v_min_u32_e32 v9, v5, v6
+; CHECK-NEXT: v_sub_co_u32_e64 v5, s[6:7], 12, v9
+; CHECK-NEXT: v_subb_co_u32_e64 v6, s[6:7], 0, 0, s[6:7]
+; CHECK-NEXT: s_and_b64 s[4:5], s[4:5], vcc
+; CHECK-NEXT: v_cndmask_b32_e64 v14, v3, v5, s[4:5]
+; CHECK-NEXT: s_movk_i32 s6, 0x423
+; CHECK-NEXT: v_cndmask_b32_e64 v13, 0, v6, s[4:5]
+; CHECK-NEXT: v_sub_co_u32_e64 v5, s[6:7], s6, v14
+; CHECK-NEXT: v_subb_co_u32_e64 v6, s[6:7], 0, v13, s[6:7]
+; CHECK-NEXT: v_cmp_gt_u64_e64 s[6:7], 63, v[5:6]
+; CHECK-NEXT: v_cndmask_b32_e64 v6, 0, v6, s[6:7]
+; CHECK-NEXT: v_cndmask_b32_e64 v5, 63, v5, s[6:7]
+; CHECK-NEXT: v_cmp_lt_u64_e64 s[6:7], 1, v[5:6]
+; CHECK-NEXT: v_cndmask_b32_e64 v7, 1, v5, s[6:7]
+; CHECK-NEXT: v_add_u32_e32 v15, -1, v7
+; CHECK-NEXT: v_lshlrev_b64 v[7:8], v15, 1
+; CHECK-NEXT: v_add_co_u32_e64 v9, s[6:7], -11, v9
+; CHECK-NEXT: v_lshlrev_b64 v[9:10], v9, v[0:1]
+; CHECK-NEXT: v_add_co_u32_e64 v11, s[6:7], -1, v7
+; CHECK-NEXT: v_addc_co_u32_e64 v12, s[6:7], -1, v8, s[6:7]
+; CHECK-NEXT: v_and_b32_e32 v8, 0xfffff, v10
+; CHECK-NEXT: v_cndmask_b32_e64 v16, v1, v8, s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v7, v0, v9, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v8, 0x100000, v16
+; CHECK-NEXT: v_and_b32_e32 v10, v8, v12
+; CHECK-NEXT: v_and_b32_e32 v9, v7, v11
+; CHECK-NEXT: v_lshrrev_b64 v[11:12], v5, v[7:8]
+; CHECK-NEXT: v_cmp_ne_u64_e64 s[4:5], 0, v[9:10]
+; CHECK-NEXT: v_and_b32_e32 v17, 1, v11
+; CHECK-NEXT: v_cndmask_b32_e64 v10, 0, 1, s[4:5]
+; CHECK-NEXT: v_lshrrev_b64 v[8:9], v15, v[7:8]
+; CHECK-NEXT: v_or_b32_e32 v9, v10, v17
+; CHECK-NEXT: v_cmp_ne_u64_e64 s[4:5], 0, v[5:6]
+; CHECK-NEXT: v_and_b32_e32 v8, v8, v9
+; CHECK-NEXT: v_cndmask_b32_e64 v5, 0, v8, s[4:5]
+; CHECK-NEXT: v_add_co_u32_e64 v5, s[4:5], v11, v5
+; CHECK-NEXT: v_addc_co_u32_e64 v6, s[4:5], 0, v12, s[4:5]
+; CHECK-NEXT: v_cmp_lt_i64_e64 s[4:5], 3, v[5:6]
+; CHECK-NEXT: v_and_b32_e32 v8, 0x1ffff, v16
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, 0, s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v6, 0, 4, s[4:5]
+; CHECK-NEXT: s_movk_i32 s4, 0x80
+; CHECK-NEXT: v_and_b32_sdwa v2, v2, s4 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:DWORD
+; CHECK-NEXT: v_cmp_ne_u64_e64 s[4:5], 0, v[7:8]
+; CHECK-NEXT: v_or3_b32 v9, v2, v6, v5
+; CHECK-NEXT: v_cndmask_b32_e64 v5, 0, 1, s[4:5]
+; CHECK-NEXT: v_bfe_u32 v6, v16, 18, 1
+; CHECK-NEXT: v_or_b32_e32 v5, v5, v6
+; CHECK-NEXT: v_lshrrev_b32_e32 v6, 17, v16
+; CHECK-NEXT: v_and_b32_e32 v5, v6, v5
+; CHECK-NEXT: v_lshrrev_b32_e32 v6, 18, v16
+; CHECK-NEXT: v_add_co_u32_e64 v5, s[4:5], v6, v5
+; CHECK-NEXT: v_addc_co_u32_e64 v6, s[4:5], 0, 0, s[4:5]
+; CHECK-NEXT: v_cmp_lt_i64_e64 s[4:5], 3, v[5:6]
+; CHECK-NEXT: v_cndmask_b32_e64 v7, 0, 1, s[4:5]
+; CHECK-NEXT: v_add_co_u32_e64 v7, s[6:7], v14, v7
+; CHECK-NEXT: v_addc_co_u32_e64 v8, s[6:7], 0, v13, s[6:7]
+; CHECK-NEXT: s_movk_i32 s6, 0xfc10
+; CHECK-NEXT: v_add_co_u32_e64 v7, s[6:7], s6, v7
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v5, 0, s[4:5]
+; CHECK-NEXT: v_addc_co_u32_e64 v8, s[6:7], -1, v8, s[6:7]
+; CHECK-NEXT: v_cndmask_b32_e64 v6, v6, 0, s[4:5]
+; CHECK-NEXT: v_cmp_lt_i64_e64 s[4:5], 3, v[5:6]
+; CHECK-NEXT: v_cmp_eq_u64_e64 s[6:7], 30, v[7:8]
+; CHECK-NEXT: v_lshlrev_b32_e32 v10, 2, v7
+; CHECK-NEXT: v_cmp_gt_i64_e64 s[8:9], 1, v[7:8]
+; CHECK-NEXT: s_and_b64 s[6:7], s[6:7], s[4:5]
+; CHECK-NEXT: v_cmp_lt_i64_e64 s[4:5], 30, v[7:8]
+; CHECK-NEXT: v_or3_b32 v10, v10, v2, v5
+; CHECK-NEXT: v_cndmask_b32_e64 v5, v10, v9, s[8:9]
+; CHECK-NEXT: v_or_b32_e32 v7, 0x7c, v2
+; CHECK-NEXT: s_or_b64 s[4:5], s[4:5], s[6:7]
+; CHECK-NEXT: v_cndmask_b32_e64 v8, v5, v7, s[4:5]
+; CHECK-NEXT: v_or_b32_e32 v5, v3, v0
+; CHECK-NEXT: v_mov_b32_e32 v6, v1
+; CHECK-NEXT: s_mov_b64 s[6:7], 0x7ff
+; CHECK-NEXT: v_cmp_eq_u64_e64 s[4:5], 0, v[0:1]
+; CHECK-NEXT: v_cmp_eq_u64_e64 s[6:7], s[6:7], v[3:4]
+; CHECK-NEXT: v_cmp_eq_u64_e64 s[8:9], 0, v[5:6]
+; CHECK-NEXT: s_and_b64 s[4:5], s[6:7], s[4:5]
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v8, v2, s[8:9]
+; CHECK-NEXT: v_cndmask_b32_e64 v0, v0, v7, s[4:5]
+; CHECK-NEXT: v_mov_b32_e32 v1, 0x7e
+; CHECK-NEXT: s_and_b64 vcc, s[6:7], vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v0, v0, v1, vcc
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f64(double %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
diff --git a/llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll b/llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll
new file mode 100644
index 0000000000000..0c98d9ca21713
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/float-to-arbitrary-fp.ll
@@ -0,0 +1,1028 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64-unknown-unknown | FileCheck %s
+
+; Test llvm.convert.to.arbitrary.fp intrinsic expansion.
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+declare i6 @llvm.convert.to.arbitrary.fp.i6.f32(float, metadata, metadata, i1)
+declare i4 @llvm.convert.to.arbitrary.fp.i4.f32(float, metadata, metadata, i1)
+declare <4 x i4> @llvm.convert.to.arbitrary.fp.v4i4.v4f32(<4 x float>, metadata, metadata, i1)
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f16(half, metadata, metadata, i1)
+declare i8 @llvm.convert.to.arbitrary.fp.i8.bf16(bfloat, metadata, metadata, i1)
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f64(double, metadata, metadata, i1)
+
+; Float8E5M2
+; Layout: sign(1) exp(5) mant(2), bias=15
+; Supports: Inf, NaN, signed zero, denormals
+
+; Float8E5M2 normal: f32 1.0 -> 0_01111_00 = 0x3C
+define i8 @to_f8e5m2_normal() {
+; CHECK-LABEL: to_f8e5m2_normal(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 60;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 zero: f32 0.0 -> 0_00000_00 = 0x00
+define i8 @to_f8e5m2_zero() {
+; CHECK-LABEL: to_f8e5m2_zero(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 0;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 negative zero: f32 -0.0 -> 1_00000_00 = 0x80
+define i8 @to_f8e5m2_neg_zero() {
+; CHECK-LABEL: to_f8e5m2_neg_zero(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 128;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float -0.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 +Inf: f32 Inf -> 0_11111_00 = 0x7C
+define i8 @to_f8e5m2_inf() {
+; CHECK-LABEL: to_f8e5m2_inf(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 124;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF0000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 NaN: f32 NaN -> qNaN
+define i8 @to_f8e5m2_nan() {
+; CHECK-LABEL: to_f8e5m2_nan(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 126;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF8000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 max: 57344.0 -> 0_11110_11 = 0x7B
+define i8 @to_f8e5m2_max() {
+; CHECK-LABEL: to_f8e5m2_max(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 123;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 57344.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 overflow (no saturation): 100000.0 -> Inf = 0x7C
+define i8 @to_f8e5m2_overflow() {
+; CHECK-LABEL: to_f8e5m2_overflow(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 124;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 100000.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 overflow with saturation: 100000.0 -> max = 0x7B
+define i8 @to_f8e5m2_overflow_sat() {
+; CHECK-LABEL: to_f8e5m2_overflow_sat(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 123;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 100000.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 true)
+ ret i8 %r
+}
+
+; Float8E5M2 denorm: very small value -> dst denorm
+define i8 @to_f8e5m2_denorm() {
+ ; 2^(-16) = 0x37800000 -> smallest E5M2 denorm 0_00000_01 = 0x01
+; CHECK-LABEL: to_f8e5m2_denorm(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 1;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x3EF0000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 runtime arg test
+define i8 @to_f8e5m2_dynamic(float %x) {
+; CHECK-LABEL: to_f8e5m2_dynamic(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<16>;
+; CHECK-NEXT: .reg .b32 %r<60>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [to_f8e5m2_dynamic_param_0];
+; CHECK-NEXT: and.b32 %r2, %r1, 8388607;
+; CHECK-NEXT: setp.ne.b32 %p1, %r2, 0;
+; CHECK-NEXT: bfe.u32 %r3, %r1, 23, 8;
+; CHECK-NEXT: setp.eq.b32 %p2, %r3, 0;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: clz.b32 %r4, %r2;
+; CHECK-NEXT: add.s32 %r5, %r4, -8;
+; CHECK-NEXT: shl.b32 %r6, %r2, %r5;
+; CHECK-NEXT: and.b32 %r7, %r6, 8388607;
+; CHECK-NEXT: selp.b32 %r8, %r7, %r2, %p3;
+; CHECK-NEXT: bfe.u32 %r9, %r8, 21, 1;
+; CHECK-NEXT: and.b32 %r10, %r8, 1048575;
+; CHECK-NEXT: setp.ne.b32 %p4, %r10, 0;
+; CHECK-NEXT: selp.b32 %r11, 1, 0, %p4;
+; CHECK-NEXT: or.b32 %r12, %r11, %r9;
+; CHECK-NEXT: shr.u32 %r13, %r8, 20;
+; CHECK-NEXT: and.b32 %r14, %r13, %r12;
+; CHECK-NEXT: shr.u32 %r15, %r8, 21;
+; CHECK-NEXT: add.s32 %r16, %r15, %r14;
+; CHECK-NEXT: setp.gt.s32 %p5, %r16, 3;
+; CHECK-NEXT: selp.b32 %r17, 0, %r16, %p5;
+; CHECK-NEXT: selp.b32 %r18, 1, 0, %p5;
+; CHECK-NEXT: sub.s32 %r19, 9, %r4;
+; CHECK-NEXT: selp.b32 %r20, %r19, %r3, %p3;
+; CHECK-NEXT: add.s32 %r21, %r20, %r18;
+; CHECK-NEXT: add.s32 %r22, %r21, -112;
+; CHECK-NEXT: shl.b32 %r23, %r22, 2;
+; CHECK-NEXT: shr.u32 %r24, %r1, 24;
+; CHECK-NEXT: and.b32 %r25, %r24, 128;
+; CHECK-NEXT: or.b32 %r26, %r25, %r23;
+; CHECK-NEXT: or.b32 %r27, %r26, %r17;
+; CHECK-NEXT: or.b32 %r28, %r8, 8388608;
+; CHECK-NEXT: sub.s32 %r29, 134, %r20;
+; CHECK-NEXT: min.u32 %r30, %r29, 31;
+; CHECK-NEXT: shr.u32 %r31, %r28, %r30;
+; CHECK-NEXT: and.b32 %r32, %r31, 1;
+; CHECK-NEXT: max.u32 %r33, %r30, 1;
+; CHECK-NEXT: add.s32 %r34, %r33, -1;
+; CHECK-NEXT: mov.b32 %r35, 1;
+; CHECK-NEXT: shl.b32 %r36, %r35, %r34;
+; CHECK-NEXT: add.s32 %r37, %r36, -1;
+; CHECK-NEXT: and.b32 %r38, %r28, %r37;
+; CHECK-NEXT: setp.ne.b32 %p6, %r38, 0;
+; CHECK-NEXT: selp.b32 %r39, 1, 0, %p6;
+; CHECK-NEXT: or.b32 %r40, %r39, %r32;
+; CHECK-NEXT: shr.u32 %r41, %r28, %r34;
+; CHECK-NEXT: and.b32 %r42, %r41, %r40;
+; CHECK-NEXT: setp.ne.b32 %p7, %r30, 0;
+; CHECK-NEXT: selp.b32 %r43, %r42, 0, %p7;
+; CHECK-NEXT: add.s32 %r44, %r31, %r43;
+; CHECK-NEXT: setp.gt.s32 %p8, %r44, 3;
+; CHECK-NEXT: selp.b32 %r45, 0, %r44, %p8;
+; CHECK-NEXT: selp.b32 %r46, 4, 0, %p8;
+; CHECK-NEXT: or.b32 %r47, %r25, %r46;
+; CHECK-NEXT: or.b32 %r48, %r47, %r45;
+; CHECK-NEXT: setp.lt.s32 %p9, %r22, 1;
+; CHECK-NEXT: selp.b32 %r49, %r48, %r27, %p9;
+; CHECK-NEXT: setp.gt.s32 %p10, %r17, 3;
+; CHECK-NEXT: or.b32 %r50, %r25, 124;
+; CHECK-NEXT: selp.b32 %r51, %r50, %r49, %p10;
+; CHECK-NEXT: setp.eq.b32 %p11, %r22, 30;
+; CHECK-NEXT: selp.b32 %r52, %r51, %r49, %p11;
+; CHECK-NEXT: setp.gt.s32 %p12, %r22, 30;
+; CHECK-NEXT: selp.b32 %r53, %r50, %r52, %p12;
+; CHECK-NEXT: or.b32 %r54, %r3, %r2;
+; CHECK-NEXT: setp.eq.b32 %p13, %r54, 0;
+; CHECK-NEXT: selp.b32 %r55, %r25, %r53, %p13;
+; CHECK-NEXT: setp.eq.b32 %p14, %r2, 0;
+; CHECK-NEXT: selp.b32 %r56, %r50, %r55, %p14;
+; CHECK-NEXT: setp.eq.b32 %p15, %r3, 255;
+; CHECK-NEXT: selp.b32 %r57, %r56, %r55, %p15;
+; CHECK-NEXT: selp.b32 %r58, 126, %r57, %p1;
+; CHECK-NEXT: selp.b32 %r59, %r58, %r57, %p15;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r59;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN (NanOnly, NanEncoding=AllOnes)
+; Layout: sign(1) exp(4) mant(3), maxExp=8, minExp=-6, bias=7
+; Only 0_1111_111 and 1_1111_111 are NaN; all other exp=15 values are finite.
+
+; Float8E4M3FN normal: f32 1.0 -> 0_0111_000 = 0x38
+define i8 @to_f8e4m3fn_normal() {
+; CHECK-LABEL: to_f8e4m3fn_normal(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 56;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN max: 448.0 -> 0_1111_110 = 0x7E
+define i8 @to_f8e4m3fn_max() {
+; CHECK-LABEL: to_f8e4m3fn_max(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 126;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 448.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN NaN: f32 NaN -> 0_1111_111 = 0x7F
+define i8 @to_f8e4m3fn_nan() {
+; CHECK-LABEL: to_f8e4m3fn_nan(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 127;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF8000000000000, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN overflow with saturation: 1000.0 -> max = 0x7E
+define i8 @to_f8e4m3fn_overflow_sat() {
+; CHECK-LABEL: to_f8e4m3fn_overflow_sat(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 126;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1000.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 true)
+ ret i8 %r
+}
+
+; Float6E3M2FN (FiniteOnly)
+; Layout: sign(1) exp(3) mant(2), bias=3, maxExp=4
+; No Inf, no NaN. All bit patterns are finite.
+
+; Float6E3M2FN normal: f32 1.0 -> 0_011_00 = 0x0C
+define i6 @to_f6e3m2fn_normal() {
+; CHECK-LABEL: to_f6e3m2fn_normal(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 12;
+; CHECK-NEXT: ret;
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 1.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN max: 28.0 -> 0_111_11 = 0x1F
+define i6 @to_f6e3m2fn_max() {
+; CHECK-LABEL: to_f6e3m2fn_max(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 31;
+; CHECK-NEXT: ret;
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 28.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN zero: f32 0.0 -> 0_000_00 = 0x00
+define i6 @to_f6e3m2fn_zero() {
+; CHECK-LABEL: to_f6e3m2fn_zero(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 0;
+; CHECK-NEXT: ret;
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 0.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN negative: f32 -1.0 -> 1_011_00
+define i6 @to_f6e3m2fn_negative() {
+; CHECK-LABEL: to_f6e3m2fn_negative(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 44;
+; CHECK-NEXT: ret;
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float -1.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E2M3FN (FiniteOnly)
+; Layout: sign(1) exp(2) mant(3), bias=1, maxExp=2
+; No Inf, no NaN. All bit patterns are finite.
+
+; Float6E2M3FN normal: f32 1.0 -> 0_01_000 = 0x08
+define i6 @to_f6e2m3fn_normal() {
+; CHECK-LABEL: to_f6e2m3fn_normal(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 8;
+; CHECK-NEXT: ret;
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 1.0, metadata !"Float6E2M3FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E2M3FN max: 7.5 -> 0_11_111 = 0x1F
+define i6 @to_f6e2m3fn_max() {
+; CHECK-LABEL: to_f6e2m3fn_max(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 31;
+; CHECK-NEXT: ret;
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 7.5, metadata !"Float6E2M3FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float4E2M1FN (FiniteOnly)
+; Layout: sign(1) exp(2) mant(1), bias=1, maxExp=2
+; No Inf, no NaN.
+
+; Float4E2M1FN normal: f32 1.0 -> 0_01_0 = 0x2
+define i4 @to_f4e2m1fn_normal() {
+; CHECK-LABEL: to_f4e2m1fn_normal(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 2;
+; CHECK-NEXT: ret;
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 1.0, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Float4E2M1FN max: 6.0 -> 0_11_1 = 0x7
+define i4 @to_f4e2m1fn_max() {
+; CHECK-LABEL: to_f4e2m1fn_max(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 7;
+; CHECK-NEXT: ret;
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 6.0, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Float4E2M1FN denorm: 0.5 -> 0_00_1 = 0x1
+define i4 @to_f4e2m1fn_denorm() {
+; CHECK-LABEL: to_f4e2m1fn_denorm(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 1;
+; CHECK-NEXT: ret;
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 0.5, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Rounding mode tests
+
+; Round to nearest
+define i8 @to_f8e5m2_round_nearest() {
+; CHECK-LABEL: to_f8e5m2_round_nearest(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.param.b32 [func_retval0], 62;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.5, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Round toward zero with runtime arg
+define i8 @to_f8e5m2_round_towardzero(float %x) {
+; CHECK-LABEL: to_f8e5m2_round_towardzero(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<13>;
+; CHECK-NEXT: .reg .b32 %r<40>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [to_f8e5m2_round_towardzero_param_0];
+; CHECK-NEXT: and.b32 %r2, %r1, 8388607;
+; CHECK-NEXT: setp.ne.b32 %p1, %r2, 0;
+; CHECK-NEXT: bfe.u32 %r3, %r1, 23, 8;
+; CHECK-NEXT: setp.eq.b32 %p2, %r3, 0;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: clz.b32 %r4, %r2;
+; CHECK-NEXT: add.s32 %r5, %r4, -8;
+; CHECK-NEXT: shl.b32 %r6, %r2, %r5;
+; CHECK-NEXT: and.b32 %r7, %r6, 8388607;
+; CHECK-NEXT: selp.b32 %r8, %r7, %r2, %p3;
+; CHECK-NEXT: or.b32 %r9, %r8, 8388608;
+; CHECK-NEXT: sub.s32 %r10, 9, %r4;
+; CHECK-NEXT: selp.b32 %r11, %r10, %r3, %p3;
+; CHECK-NEXT: sub.s32 %r12, 134, %r11;
+; CHECK-NEXT: min.u32 %r13, %r12, 31;
+; CHECK-NEXT: shr.u32 %r14, %r9, %r13;
+; CHECK-NEXT: setp.gt.s32 %p4, %r14, 3;
+; CHECK-NEXT: selp.b32 %r15, 0, %r14, %p4;
+; CHECK-NEXT: selp.b32 %r16, 4, 0, %p4;
+; CHECK-NEXT: shr.u32 %r17, %r1, 24;
+; CHECK-NEXT: and.b32 %r18, %r17, 128;
+; CHECK-NEXT: or.b32 %r19, %r18, %r16;
+; CHECK-NEXT: or.b32 %r20, %r19, %r15;
+; CHECK-NEXT: shr.u32 %r21, %r8, 21;
+; CHECK-NEXT: setp.gt.s32 %p5, %r21, 3;
+; CHECK-NEXT: selp.b32 %r22, 0, %r21, %p5;
+; CHECK-NEXT: selp.b32 %r23, 1, 0, %p5;
+; CHECK-NEXT: add.s32 %r24, %r11, %r23;
+; CHECK-NEXT: add.s32 %r25, %r24, -112;
+; CHECK-NEXT: shl.b32 %r26, %r25, 2;
+; CHECK-NEXT: or.b32 %r27, %r18, %r26;
+; CHECK-NEXT: or.b32 %r28, %r27, %r22;
+; CHECK-NEXT: setp.lt.s32 %p6, %r25, 1;
+; CHECK-NEXT: selp.b32 %r29, %r20, %r28, %p6;
+; CHECK-NEXT: setp.gt.s32 %p7, %r22, 3;
+; CHECK-NEXT: or.b32 %r30, %r18, 124;
+; CHECK-NEXT: selp.b32 %r31, %r30, %r29, %p7;
+; CHECK-NEXT: setp.eq.b32 %p8, %r25, 30;
+; CHECK-NEXT: selp.b32 %r32, %r31, %r29, %p8;
+; CHECK-NEXT: setp.gt.s32 %p9, %r25, 30;
+; CHECK-NEXT: selp.b32 %r33, %r30, %r32, %p9;
+; CHECK-NEXT: or.b32 %r34, %r3, %r2;
+; CHECK-NEXT: setp.eq.b32 %p10, %r34, 0;
+; CHECK-NEXT: selp.b32 %r35, %r18, %r33, %p10;
+; CHECK-NEXT: setp.eq.b32 %p11, %r2, 0;
+; CHECK-NEXT: selp.b32 %r36, %r30, %r35, %p11;
+; CHECK-NEXT: setp.eq.b32 %p12, %r3, 255;
+; CHECK-NEXT: selp.b32 %r37, %r36, %r35, %p12;
+; CHECK-NEXT: selp.b32 %r38, 126, %r37, %p1;
+; CHECK-NEXT: selp.b32 %r39, %r38, %r37, %p12;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r39;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float %x, metadata !"Float8E5M2", metadata !"round.towardzero", i1 false)
+ ret i8 %r
+}
+
+; Vector test: f32 <4 x float> -> <4 x i4> Float4E2M1FN
+define <4 x i4> @to_f4e2m1fn_v4f32(<4 x float> %x) {
+; CHECK-LABEL: to_f4e2m1fn_v4f32(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<41>;
+; CHECK-NEXT: .reg .b32 %r<202>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [to_f4e2m1fn_v4f32_param_0];
+; CHECK-NEXT: and.b32 %r5, %r4, 8388607;
+; CHECK-NEXT: setp.ne.b32 %p1, %r5, 0;
+; CHECK-NEXT: bfe.u32 %r6, %r4, 23, 8;
+; CHECK-NEXT: setp.eq.b32 %p2, %r6, 0;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: clz.b32 %r7, %r5;
+; CHECK-NEXT: add.s32 %r8, %r7, -8;
+; CHECK-NEXT: shl.b32 %r9, %r5, %r8;
+; CHECK-NEXT: and.b32 %r10, %r9, 8388607;
+; CHECK-NEXT: selp.b32 %r11, %r10, %r5, %p3;
+; CHECK-NEXT: shr.u32 %r12, %r11, 22;
+; CHECK-NEXT: and.b32 %r13, %r11, 2097151;
+; CHECK-NEXT: setp.ne.b32 %p4, %r13, 0;
+; CHECK-NEXT: selp.b32 %r14, 1, 0, %p4;
+; CHECK-NEXT: or.b32 %r15, %r14, %r12;
+; CHECK-NEXT: shr.u32 %r16, %r11, 21;
+; CHECK-NEXT: and.b32 %r17, %r16, %r15;
+; CHECK-NEXT: add.s32 %r18, %r12, %r17;
+; CHECK-NEXT: setp.gt.s32 %p5, %r18, 1;
+; CHECK-NEXT: selp.b32 %r19, 0, %r18, %p5;
+; CHECK-NEXT: selp.b32 %r20, 1, 0, %p5;
+; CHECK-NEXT: sub.s32 %r21, 9, %r7;
+; CHECK-NEXT: selp.b32 %r22, %r21, %r6, %p3;
+; CHECK-NEXT: add.s32 %r23, %r22, %r20;
+; CHECK-NEXT: add.s32 %r24, %r23, -126;
+; CHECK-NEXT: shl.b32 %r25, %r24, 1;
+; CHECK-NEXT: shr.u32 %r26, %r4, 28;
+; CHECK-NEXT: and.b32 %r27, %r26, 8;
+; CHECK-NEXT: or.b32 %r28, %r27, %r25;
+; CHECK-NEXT: or.b32 %r29, %r28, %r19;
+; CHECK-NEXT: or.b32 %r30, %r11, 8388608;
+; CHECK-NEXT: sub.s32 %r31, 149, %r22;
+; CHECK-NEXT: min.u32 %r32, %r31, 31;
+; CHECK-NEXT: shr.u32 %r33, %r30, %r32;
+; CHECK-NEXT: and.b32 %r34, %r33, 1;
+; CHECK-NEXT: max.u32 %r35, %r32, 1;
+; CHECK-NEXT: add.s32 %r36, %r35, -1;
+; CHECK-NEXT: mov.b32 %r37, 1;
+; CHECK-NEXT: shl.b32 %r38, %r37, %r36;
+; CHECK-NEXT: add.s32 %r39, %r38, -1;
+; CHECK-NEXT: and.b32 %r40, %r30, %r39;
+; CHECK-NEXT: setp.ne.b32 %p6, %r40, 0;
+; CHECK-NEXT: selp.b32 %r41, 1, 0, %p6;
+; CHECK-NEXT: or.b32 %r42, %r41, %r34;
+; CHECK-NEXT: shr.u32 %r43, %r30, %r36;
+; CHECK-NEXT: and.b32 %r44, %r43, %r42;
+; CHECK-NEXT: setp.ne.b32 %p7, %r32, 0;
+; CHECK-NEXT: selp.b32 %r45, %r44, 0, %p7;
+; CHECK-NEXT: add.s32 %r46, %r33, %r45;
+; CHECK-NEXT: setp.gt.s32 %p8, %r46, 1;
+; CHECK-NEXT: selp.b32 %r47, 0, %r46, %p8;
+; CHECK-NEXT: selp.b32 %r48, 2, 0, %p8;
+; CHECK-NEXT: or.b32 %r49, %r27, %r48;
+; CHECK-NEXT: or.b32 %r50, %r49, %r47;
+; CHECK-NEXT: setp.lt.s32 %p9, %r24, 1;
+; CHECK-NEXT: selp.b32 %r51, %r50, %r29, %p9;
+; CHECK-NEXT: or.b32 %r52, %r6, %r5;
+; CHECK-NEXT: setp.eq.b32 %p10, %r52, 0;
+; CHECK-NEXT: selp.b32 %r53, %r27, %r51, %p10;
+; CHECK-NEXT: and.b32 %r54, %r3, 8388607;
+; CHECK-NEXT: setp.ne.b32 %p11, %r54, 0;
+; CHECK-NEXT: bfe.u32 %r55, %r3, 23, 8;
+; CHECK-NEXT: setp.eq.b32 %p12, %r55, 0;
+; CHECK-NEXT: and.pred %p13, %p12, %p11;
+; CHECK-NEXT: clz.b32 %r56, %r54;
+; CHECK-NEXT: add.s32 %r57, %r56, -8;
+; CHECK-NEXT: shl.b32 %r58, %r54, %r57;
+; CHECK-NEXT: and.b32 %r59, %r58, 8388607;
+; CHECK-NEXT: selp.b32 %r60, %r59, %r54, %p13;
+; CHECK-NEXT: shr.u32 %r61, %r60, 22;
+; CHECK-NEXT: and.b32 %r62, %r60, 2097151;
+; CHECK-NEXT: setp.ne.b32 %p14, %r62, 0;
+; CHECK-NEXT: selp.b32 %r63, 1, 0, %p14;
+; CHECK-NEXT: or.b32 %r64, %r63, %r61;
+; CHECK-NEXT: shr.u32 %r65, %r60, 21;
+; CHECK-NEXT: and.b32 %r66, %r65, %r64;
+; CHECK-NEXT: add.s32 %r67, %r61, %r66;
+; CHECK-NEXT: setp.gt.s32 %p15, %r67, 1;
+; CHECK-NEXT: selp.b32 %r68, 0, %r67, %p15;
+; CHECK-NEXT: selp.b32 %r69, 1, 0, %p15;
+; CHECK-NEXT: sub.s32 %r70, 9, %r56;
+; CHECK-NEXT: selp.b32 %r71, %r70, %r55, %p13;
+; CHECK-NEXT: add.s32 %r72, %r71, %r69;
+; CHECK-NEXT: add.s32 %r73, %r72, -126;
+; CHECK-NEXT: shl.b32 %r74, %r73, 1;
+; CHECK-NEXT: shr.u32 %r75, %r3, 28;
+; CHECK-NEXT: and.b32 %r76, %r75, 8;
+; CHECK-NEXT: or.b32 %r77, %r76, %r74;
+; CHECK-NEXT: or.b32 %r78, %r77, %r68;
+; CHECK-NEXT: or.b32 %r79, %r60, 8388608;
+; CHECK-NEXT: sub.s32 %r80, 149, %r71;
+; CHECK-NEXT: min.u32 %r81, %r80, 31;
+; CHECK-NEXT: shr.u32 %r82, %r79, %r81;
+; CHECK-NEXT: and.b32 %r83, %r82, 1;
+; CHECK-NEXT: max.u32 %r84, %r81, 1;
+; CHECK-NEXT: add.s32 %r85, %r84, -1;
+; CHECK-NEXT: shl.b32 %r86, %r37, %r85;
+; CHECK-NEXT: add.s32 %r87, %r86, -1;
+; CHECK-NEXT: and.b32 %r88, %r79, %r87;
+; CHECK-NEXT: setp.ne.b32 %p16, %r88, 0;
+; CHECK-NEXT: selp.b32 %r89, 1, 0, %p16;
+; CHECK-NEXT: or.b32 %r90, %r89, %r83;
+; CHECK-NEXT: shr.u32 %r91, %r79, %r85;
+; CHECK-NEXT: and.b32 %r92, %r91, %r90;
+; CHECK-NEXT: setp.ne.b32 %p17, %r81, 0;
+; CHECK-NEXT: selp.b32 %r93, %r92, 0, %p17;
+; CHECK-NEXT: add.s32 %r94, %r82, %r93;
+; CHECK-NEXT: setp.gt.s32 %p18, %r94, 1;
+; CHECK-NEXT: selp.b32 %r95, 0, %r94, %p18;
+; CHECK-NEXT: selp.b32 %r96, 2, 0, %p18;
+; CHECK-NEXT: or.b32 %r97, %r76, %r96;
+; CHECK-NEXT: or.b32 %r98, %r97, %r95;
+; CHECK-NEXT: setp.lt.s32 %p19, %r73, 1;
+; CHECK-NEXT: selp.b32 %r99, %r98, %r78, %p19;
+; CHECK-NEXT: or.b32 %r100, %r55, %r54;
+; CHECK-NEXT: setp.eq.b32 %p20, %r100, 0;
+; CHECK-NEXT: selp.b32 %r101, %r76, %r99, %p20;
+; CHECK-NEXT: prmt.b32 %r102, %r101, %r53, 0x3340U;
+; CHECK-NEXT: and.b32 %r103, %r2, 8388607;
+; CHECK-NEXT: setp.ne.b32 %p21, %r103, 0;
+; CHECK-NEXT: bfe.u32 %r104, %r2, 23, 8;
+; CHECK-NEXT: setp.eq.b32 %p22, %r104, 0;
+; CHECK-NEXT: and.pred %p23, %p22, %p21;
+; CHECK-NEXT: clz.b32 %r105, %r103;
+; CHECK-NEXT: add.s32 %r106, %r105, -8;
+; CHECK-NEXT: shl.b32 %r107, %r103, %r106;
+; CHECK-NEXT: and.b32 %r108, %r107, 8388607;
+; CHECK-NEXT: selp.b32 %r109, %r108, %r103, %p23;
+; CHECK-NEXT: shr.u32 %r110, %r109, 22;
+; CHECK-NEXT: and.b32 %r111, %r109, 2097151;
+; CHECK-NEXT: setp.ne.b32 %p24, %r111, 0;
+; CHECK-NEXT: selp.b32 %r112, 1, 0, %p24;
+; CHECK-NEXT: or.b32 %r113, %r112, %r110;
+; CHECK-NEXT: shr.u32 %r114, %r109, 21;
+; CHECK-NEXT: and.b32 %r115, %r114, %r113;
+; CHECK-NEXT: add.s32 %r116, %r110, %r115;
+; CHECK-NEXT: setp.gt.s32 %p25, %r116, 1;
+; CHECK-NEXT: selp.b32 %r117, 0, %r116, %p25;
+; CHECK-NEXT: selp.b32 %r118, 1, 0, %p25;
+; CHECK-NEXT: sub.s32 %r119, 9, %r105;
+; CHECK-NEXT: selp.b32 %r120, %r119, %r104, %p23;
+; CHECK-NEXT: add.s32 %r121, %r120, %r118;
+; CHECK-NEXT: add.s32 %r122, %r121, -126;
+; CHECK-NEXT: shl.b32 %r123, %r122, 1;
+; CHECK-NEXT: shr.u32 %r124, %r2, 28;
+; CHECK-NEXT: and.b32 %r125, %r124, 8;
+; CHECK-NEXT: or.b32 %r126, %r125, %r123;
+; CHECK-NEXT: or.b32 %r127, %r126, %r117;
+; CHECK-NEXT: or.b32 %r128, %r109, 8388608;
+; CHECK-NEXT: sub.s32 %r129, 149, %r120;
+; CHECK-NEXT: min.u32 %r130, %r129, 31;
+; CHECK-NEXT: shr.u32 %r131, %r128, %r130;
+; CHECK-NEXT: and.b32 %r132, %r131, 1;
+; CHECK-NEXT: max.u32 %r133, %r130, 1;
+; CHECK-NEXT: add.s32 %r134, %r133, -1;
+; CHECK-NEXT: shl.b32 %r135, %r37, %r134;
+; CHECK-NEXT: add.s32 %r136, %r135, -1;
+; CHECK-NEXT: and.b32 %r137, %r128, %r136;
+; CHECK-NEXT: setp.ne.b32 %p26, %r137, 0;
+; CHECK-NEXT: selp.b32 %r138, 1, 0, %p26;
+; CHECK-NEXT: or.b32 %r139, %r138, %r132;
+; CHECK-NEXT: shr.u32 %r140, %r128, %r134;
+; CHECK-NEXT: and.b32 %r141, %r140, %r139;
+; CHECK-NEXT: setp.ne.b32 %p27, %r130, 0;
+; CHECK-NEXT: selp.b32 %r142, %r141, 0, %p27;
+; CHECK-NEXT: add.s32 %r143, %r131, %r142;
+; CHECK-NEXT: setp.gt.s32 %p28, %r143, 1;
+; CHECK-NEXT: selp.b32 %r144, 0, %r143, %p28;
+; CHECK-NEXT: selp.b32 %r145, 2, 0, %p28;
+; CHECK-NEXT: or.b32 %r146, %r125, %r145;
+; CHECK-NEXT: or.b32 %r147, %r146, %r144;
+; CHECK-NEXT: setp.lt.s32 %p29, %r122, 1;
+; CHECK-NEXT: selp.b32 %r148, %r147, %r127, %p29;
+; CHECK-NEXT: or.b32 %r149, %r104, %r103;
+; CHECK-NEXT: setp.eq.b32 %p30, %r149, 0;
+; CHECK-NEXT: selp.b32 %r150, %r125, %r148, %p30;
+; CHECK-NEXT: and.b32 %r151, %r1, 8388607;
+; CHECK-NEXT: setp.ne.b32 %p31, %r151, 0;
+; CHECK-NEXT: bfe.u32 %r152, %r1, 23, 8;
+; CHECK-NEXT: setp.eq.b32 %p32, %r152, 0;
+; CHECK-NEXT: and.pred %p33, %p32, %p31;
+; CHECK-NEXT: clz.b32 %r153, %r151;
+; CHECK-NEXT: add.s32 %r154, %r153, -8;
+; CHECK-NEXT: shl.b32 %r155, %r151, %r154;
+; CHECK-NEXT: and.b32 %r156, %r155, 8388607;
+; CHECK-NEXT: selp.b32 %r157, %r156, %r151, %p33;
+; CHECK-NEXT: shr.u32 %r158, %r157, 22;
+; CHECK-NEXT: and.b32 %r159, %r157, 2097151;
+; CHECK-NEXT: setp.ne.b32 %p34, %r159, 0;
+; CHECK-NEXT: selp.b32 %r160, 1, 0, %p34;
+; CHECK-NEXT: or.b32 %r161, %r160, %r158;
+; CHECK-NEXT: shr.u32 %r162, %r157, 21;
+; CHECK-NEXT: and.b32 %r163, %r162, %r161;
+; CHECK-NEXT: add.s32 %r164, %r158, %r163;
+; CHECK-NEXT: setp.gt.s32 %p35, %r164, 1;
+; CHECK-NEXT: selp.b32 %r165, 0, %r164, %p35;
+; CHECK-NEXT: selp.b32 %r166, 1, 0, %p35;
+; CHECK-NEXT: sub.s32 %r167, 9, %r153;
+; CHECK-NEXT: selp.b32 %r168, %r167, %r152, %p33;
+; CHECK-NEXT: add.s32 %r169, %r168, %r166;
+; CHECK-NEXT: add.s32 %r170, %r169, -126;
+; CHECK-NEXT: shl.b32 %r171, %r170, 1;
+; CHECK-NEXT: shr.u32 %r172, %r1, 28;
+; CHECK-NEXT: and.b32 %r173, %r172, 8;
+; CHECK-NEXT: or.b32 %r174, %r173, %r171;
+; CHECK-NEXT: or.b32 %r175, %r174, %r165;
+; CHECK-NEXT: or.b32 %r176, %r157, 8388608;
+; CHECK-NEXT: sub.s32 %r177, 149, %r168;
+; CHECK-NEXT: min.u32 %r178, %r177, 31;
+; CHECK-NEXT: shr.u32 %r179, %r176, %r178;
+; CHECK-NEXT: and.b32 %r180, %r179, 1;
+; CHECK-NEXT: max.u32 %r181, %r178, 1;
+; CHECK-NEXT: add.s32 %r182, %r181, -1;
+; CHECK-NEXT: shl.b32 %r183, %r37, %r182;
+; CHECK-NEXT: add.s32 %r184, %r183, -1;
+; CHECK-NEXT: and.b32 %r185, %r176, %r184;
+; CHECK-NEXT: setp.ne.b32 %p36, %r185, 0;
+; CHECK-NEXT: selp.b32 %r186, 1, 0, %p36;
+; CHECK-NEXT: or.b32 %r187, %r186, %r180;
+; CHECK-NEXT: shr.u32 %r188, %r176, %r182;
+; CHECK-NEXT: and.b32 %r189, %r188, %r187;
+; CHECK-NEXT: setp.ne.b32 %p37, %r178, 0;
+; CHECK-NEXT: selp.b32 %r190, %r189, 0, %p37;
+; CHECK-NEXT: add.s32 %r191, %r179, %r190;
+; CHECK-NEXT: setp.gt.s32 %p38, %r191, 1;
+; CHECK-NEXT: selp.b32 %r192, 0, %r191, %p38;
+; CHECK-NEXT: selp.b32 %r193, 2, 0, %p38;
+; CHECK-NEXT: or.b32 %r194, %r173, %r193;
+; CHECK-NEXT: or.b32 %r195, %r194, %r192;
+; CHECK-NEXT: setp.lt.s32 %p39, %r170, 1;
+; CHECK-NEXT: selp.b32 %r196, %r195, %r175, %p39;
+; CHECK-NEXT: or.b32 %r197, %r152, %r151;
+; CHECK-NEXT: setp.eq.b32 %p40, %r197, 0;
+; CHECK-NEXT: selp.b32 %r198, %r173, %r196, %p40;
+; CHECK-NEXT: prmt.b32 %r199, %r198, %r150, 0x3340U;
+; CHECK-NEXT: prmt.b32 %r200, %r199, %r102, 0x5410U;
+; CHECK-NEXT: st.param.b16 [func_retval0], %r200;
+; CHECK-NEXT: shr.u32 %r201, %r200, 16;
+; CHECK-NEXT: st.param.b16 [func_retval0+2], %r201;
+; CHECK-NEXT: ret;
+ %r = call <4 x i4> @llvm.convert.to.arbitrary.fp.v4i4.v4f32(<4 x float> %x, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret <4 x i4> %r
+}
+
+; Float8E5M2 from f16: half -> i8
+define i8 @to_f8e5m2_from_f16(half %x) {
+; CHECK-LABEL: to_f8e5m2_from_f16(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<16>;
+; CHECK-NEXT: .reg .b16 %rs<60>;
+; CHECK-NEXT: .reg .b32 %r<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [to_f8e5m2_from_f16_param_0];
+; CHECK-NEXT: and.b16 %rs2, %rs1, 1023;
+; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT: shr.u16 %rs3, %rs1, 10;
+; CHECK-NEXT: and.b16 %rs4, %rs3, 31;
+; CHECK-NEXT: setp.eq.b16 %p2, %rs4, 0;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
+; CHECK-NEXT: shl.b32 %r2, %r1, 16;
+; CHECK-NEXT: clz.b32 %r3, %r2;
+; CHECK-NEXT: cvt.u16.u32 %rs5, %r3;
+; CHECK-NEXT: add.s16 %rs6, %rs5, -5;
+; CHECK-NEXT: cvt.u32.u16 %r4, %rs6;
+; CHECK-NEXT: shl.b16 %rs7, %rs2, %r4;
+; CHECK-NEXT: and.b16 %rs8, %rs7, 1023;
+; CHECK-NEXT: selp.b16 %rs9, %rs8, %rs2, %p3;
+; CHECK-NEXT: shr.u16 %rs10, %rs9, 8;
+; CHECK-NEXT: and.b16 %rs11, %rs10, 1;
+; CHECK-NEXT: and.b16 %rs12, %rs9, 127;
+; CHECK-NEXT: setp.ne.b16 %p4, %rs12, 0;
+; CHECK-NEXT: selp.b16 %rs13, 1, 0, %p4;
+; CHECK-NEXT: or.b16 %rs14, %rs13, %rs11;
+; CHECK-NEXT: shr.u16 %rs15, %rs9, 7;
+; CHECK-NEXT: and.b16 %rs16, %rs15, %rs14;
+; CHECK-NEXT: add.s16 %rs17, %rs10, %rs16;
+; CHECK-NEXT: setp.gt.s16 %p5, %rs17, 3;
+; CHECK-NEXT: selp.b16 %rs18, 0, %rs17, %p5;
+; CHECK-NEXT: selp.b16 %rs19, 1, 0, %p5;
+; CHECK-NEXT: sub.s16 %rs20, 6, %rs5;
+; CHECK-NEXT: selp.b16 %rs21, %rs20, %rs4, %p3;
+; CHECK-NEXT: add.s16 %rs22, %rs21, %rs19;
+; CHECK-NEXT: shl.b16 %rs23, %rs22, 2;
+; CHECK-NEXT: shr.u16 %rs24, %rs1, 8;
+; CHECK-NEXT: and.b16 %rs25, %rs24, 128;
+; CHECK-NEXT: or.b16 %rs26, %rs25, %rs23;
+; CHECK-NEXT: or.b16 %rs27, %rs26, %rs18;
+; CHECK-NEXT: or.b16 %rs28, %rs9, 1024;
+; CHECK-NEXT: sub.s16 %rs29, 9, %rs21;
+; CHECK-NEXT: min.u16 %rs30, %rs29, 15;
+; CHECK-NEXT: cvt.u32.u16 %r5, %rs30;
+; CHECK-NEXT: shr.u16 %rs31, %rs28, %r5;
+; CHECK-NEXT: and.b16 %rs32, %rs31, 1;
+; CHECK-NEXT: max.u16 %rs33, %rs30, 1;
+; CHECK-NEXT: add.s16 %rs34, %rs33, -1;
+; CHECK-NEXT: cvt.u32.u16 %r6, %rs34;
+; CHECK-NEXT: mov.b16 %rs35, 1;
+; CHECK-NEXT: shl.b16 %rs36, %rs35, %r6;
+; CHECK-NEXT: add.s16 %rs37, %rs36, -1;
+; CHECK-NEXT: and.b16 %rs38, %rs28, %rs37;
+; CHECK-NEXT: setp.ne.b16 %p6, %rs38, 0;
+; CHECK-NEXT: selp.b16 %rs39, 1, 0, %p6;
+; CHECK-NEXT: or.b16 %rs40, %rs39, %rs32;
+; CHECK-NEXT: shr.u16 %rs41, %rs28, %r6;
+; CHECK-NEXT: and.b16 %rs42, %rs41, %rs40;
+; CHECK-NEXT: setp.ne.b16 %p7, %rs30, 0;
+; CHECK-NEXT: selp.b16 %rs43, %rs42, 0, %p7;
+; CHECK-NEXT: add.s16 %rs44, %rs31, %rs43;
+; CHECK-NEXT: setp.gt.s16 %p8, %rs44, 3;
+; CHECK-NEXT: selp.b16 %rs45, 0, %rs44, %p8;
+; CHECK-NEXT: selp.b16 %rs46, 4, 0, %p8;
+; CHECK-NEXT: or.b16 %rs47, %rs25, %rs46;
+; CHECK-NEXT: or.b16 %rs48, %rs47, %rs45;
+; CHECK-NEXT: setp.lt.s16 %p9, %rs22, 1;
+; CHECK-NEXT: selp.b16 %rs49, %rs48, %rs27, %p9;
+; CHECK-NEXT: setp.gt.s16 %p10, %rs18, 3;
+; CHECK-NEXT: or.b16 %rs50, %rs25, 124;
+; CHECK-NEXT: selp.b16 %rs51, %rs50, %rs49, %p10;
+; CHECK-NEXT: setp.eq.b16 %p11, %rs22, 30;
+; CHECK-NEXT: selp.b16 %rs52, %rs51, %rs49, %p11;
+; CHECK-NEXT: setp.gt.s16 %p12, %rs22, 30;
+; CHECK-NEXT: selp.b16 %rs53, %rs50, %rs52, %p12;
+; CHECK-NEXT: or.b16 %rs54, %rs4, %rs2;
+; CHECK-NEXT: setp.eq.b16 %p13, %rs54, 0;
+; CHECK-NEXT: selp.b16 %rs55, %rs25, %rs53, %p13;
+; CHECK-NEXT: setp.eq.b16 %p14, %rs2, 0;
+; CHECK-NEXT: selp.b16 %rs56, %rs50, %rs55, %p14;
+; CHECK-NEXT: setp.eq.b16 %p15, %rs4, 31;
+; CHECK-NEXT: selp.b16 %rs57, %rs56, %rs55, %p15;
+; CHECK-NEXT: selp.b16 %rs58, 126, %rs57, %p1;
+; CHECK-NEXT: selp.b16 %rs59, %rs58, %rs57, %p15;
+; CHECK-NEXT: cvt.u32.u16 %r7, %rs59;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f16(half %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 from bf16: bfloat -> i8
+define i8 @to_f8e5m2_from_bf16(bfloat %x) {
+; CHECK-LABEL: to_f8e5m2_from_bf16(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<16>;
+; CHECK-NEXT: .reg .b16 %rs<61>;
+; CHECK-NEXT: .reg .b32 %r<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [to_f8e5m2_from_bf16_param_0];
+; CHECK-NEXT: and.b16 %rs2, %rs1, 127;
+; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
+; CHECK-NEXT: shr.u16 %rs3, %rs1, 7;
+; CHECK-NEXT: and.b16 %rs4, %rs3, 255;
+; CHECK-NEXT: setp.eq.b16 %p2, %rs4, 0;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
+; CHECK-NEXT: shl.b32 %r2, %r1, 16;
+; CHECK-NEXT: clz.b32 %r3, %r2;
+; CHECK-NEXT: cvt.u16.u32 %rs5, %r3;
+; CHECK-NEXT: add.s16 %rs6, %rs5, -8;
+; CHECK-NEXT: cvt.u32.u16 %r4, %rs6;
+; CHECK-NEXT: shl.b16 %rs7, %rs2, %r4;
+; CHECK-NEXT: and.b16 %rs8, %rs7, 127;
+; CHECK-NEXT: selp.b16 %rs9, %rs8, %rs2, %p3;
+; CHECK-NEXT: shr.u16 %rs10, %rs9, 5;
+; CHECK-NEXT: and.b16 %rs11, %rs10, 1;
+; CHECK-NEXT: and.b16 %rs12, %rs9, 15;
+; CHECK-NEXT: setp.ne.b16 %p4, %rs12, 0;
+; CHECK-NEXT: selp.b16 %rs13, 1, 0, %p4;
+; CHECK-NEXT: or.b16 %rs14, %rs13, %rs11;
+; CHECK-NEXT: shr.u16 %rs15, %rs9, 4;
+; CHECK-NEXT: and.b16 %rs16, %rs15, %rs14;
+; CHECK-NEXT: add.s16 %rs17, %rs10, %rs16;
+; CHECK-NEXT: setp.gt.s16 %p5, %rs17, 3;
+; CHECK-NEXT: selp.b16 %rs18, 0, %rs17, %p5;
+; CHECK-NEXT: selp.b16 %rs19, 1, 0, %p5;
+; CHECK-NEXT: sub.s16 %rs20, 9, %rs5;
+; CHECK-NEXT: selp.b16 %rs21, %rs20, %rs4, %p3;
+; CHECK-NEXT: add.s16 %rs22, %rs21, %rs19;
+; CHECK-NEXT: add.s16 %rs23, %rs22, -112;
+; CHECK-NEXT: shl.b16 %rs24, %rs23, 2;
+; CHECK-NEXT: shr.u16 %rs25, %rs1, 8;
+; CHECK-NEXT: and.b16 %rs26, %rs25, 128;
+; CHECK-NEXT: or.b16 %rs27, %rs26, %rs24;
+; CHECK-NEXT: or.b16 %rs28, %rs27, %rs18;
+; CHECK-NEXT: or.b16 %rs29, %rs9, 128;
+; CHECK-NEXT: sub.s16 %rs30, 118, %rs21;
+; CHECK-NEXT: min.u16 %rs31, %rs30, 15;
+; CHECK-NEXT: cvt.u32.u16 %r5, %rs31;
+; CHECK-NEXT: shr.u16 %rs32, %rs29, %r5;
+; CHECK-NEXT: and.b16 %rs33, %rs32, 1;
+; CHECK-NEXT: max.u16 %rs34, %rs31, 1;
+; CHECK-NEXT: add.s16 %rs35, %rs34, -1;
+; CHECK-NEXT: cvt.u32.u16 %r6, %rs35;
+; CHECK-NEXT: mov.b16 %rs36, 1;
+; CHECK-NEXT: shl.b16 %rs37, %rs36, %r6;
+; CHECK-NEXT: add.s16 %rs38, %rs37, -1;
+; CHECK-NEXT: and.b16 %rs39, %rs29, %rs38;
+; CHECK-NEXT: setp.ne.b16 %p6, %rs39, 0;
+; CHECK-NEXT: selp.b16 %rs40, 1, 0, %p6;
+; CHECK-NEXT: or.b16 %rs41, %rs40, %rs33;
+; CHECK-NEXT: shr.u16 %rs42, %rs29, %r6;
+; CHECK-NEXT: and.b16 %rs43, %rs42, %rs41;
+; CHECK-NEXT: setp.ne.b16 %p7, %rs31, 0;
+; CHECK-NEXT: selp.b16 %rs44, %rs43, 0, %p7;
+; CHECK-NEXT: add.s16 %rs45, %rs32, %rs44;
+; CHECK-NEXT: setp.gt.s16 %p8, %rs45, 3;
+; CHECK-NEXT: selp.b16 %rs46, 0, %rs45, %p8;
+; CHECK-NEXT: selp.b16 %rs47, 4, 0, %p8;
+; CHECK-NEXT: or.b16 %rs48, %rs26, %rs47;
+; CHECK-NEXT: or.b16 %rs49, %rs48, %rs46;
+; CHECK-NEXT: setp.lt.s16 %p9, %rs23, 1;
+; CHECK-NEXT: selp.b16 %rs50, %rs49, %rs28, %p9;
+; CHECK-NEXT: setp.gt.s16 %p10, %rs18, 3;
+; CHECK-NEXT: or.b16 %rs51, %rs26, 124;
+; CHECK-NEXT: selp.b16 %rs52, %rs51, %rs50, %p10;
+; CHECK-NEXT: setp.eq.b16 %p11, %rs23, 30;
+; CHECK-NEXT: selp.b16 %rs53, %rs52, %rs50, %p11;
+; CHECK-NEXT: setp.gt.s16 %p12, %rs23, 30;
+; CHECK-NEXT: selp.b16 %rs54, %rs51, %rs53, %p12;
+; CHECK-NEXT: or.b16 %rs55, %rs4, %rs2;
+; CHECK-NEXT: setp.eq.b16 %p13, %rs55, 0;
+; CHECK-NEXT: selp.b16 %rs56, %rs26, %rs54, %p13;
+; CHECK-NEXT: setp.eq.b16 %p14, %rs2, 0;
+; CHECK-NEXT: selp.b16 %rs57, %rs51, %rs56, %p14;
+; CHECK-NEXT: setp.eq.b16 %p15, %rs4, 255;
+; CHECK-NEXT: selp.b16 %rs58, %rs57, %rs56, %p15;
+; CHECK-NEXT: selp.b16 %rs59, 126, %rs58, %p1;
+; CHECK-NEXT: selp.b16 %rs60, %rs59, %rs58, %p15;
+; CHECK-NEXT: cvt.u32.u16 %r7, %rs60;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.bf16(bfloat %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 from f64: double -> i8
+define i8 @to_f8e5m2_from_f64(double %x) {
+; CHECK-LABEL: to_f8e5m2_from_f64(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<16>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-NEXT: .reg .b64 %rd<61>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [to_f8e5m2_from_f64_param_0];
+; CHECK-NEXT: and.b64 %rd2, %rd1, 4503599627370495;
+; CHECK-NEXT: setp.ne.b64 %p1, %rd2, 0;
+; CHECK-NEXT: shr.u64 %rd3, %rd1, 52;
+; CHECK-NEXT: and.b64 %rd4, %rd3, 2047;
+; CHECK-NEXT: setp.eq.b64 %p2, %rd4, 0;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: clz.b64 %r1, %rd2;
+; CHECK-NEXT: cvt.u64.u32 %rd5, %r1;
+; CHECK-NEXT: add.s64 %rd6, %rd5, -11;
+; CHECK-NEXT: cvt.u32.u64 %r2, %rd6;
+; CHECK-NEXT: shl.b64 %rd7, %rd2, %r2;
+; CHECK-NEXT: and.b64 %rd8, %rd7, 4503599627370495;
+; CHECK-NEXT: selp.b64 %rd9, %rd8, %rd2, %p3;
+; CHECK-NEXT: shr.u64 %rd10, %rd9, 50;
+; CHECK-NEXT: and.b64 %rd11, %rd10, 1;
+; CHECK-NEXT: and.b64 %rd12, %rd9, 562949953421311;
+; CHECK-NEXT: setp.ne.b64 %p4, %rd12, 0;
+; CHECK-NEXT: selp.b64 %rd13, 1, 0, %p4;
+; CHECK-NEXT: or.b64 %rd14, %rd13, %rd11;
+; CHECK-NEXT: shr.u64 %rd15, %rd9, 49;
+; CHECK-NEXT: and.b64 %rd16, %rd15, %rd14;
+; CHECK-NEXT: add.s64 %rd17, %rd10, %rd16;
+; CHECK-NEXT: setp.gt.s64 %p5, %rd17, 3;
+; CHECK-NEXT: selp.b64 %rd18, 0, %rd17, %p5;
+; CHECK-NEXT: selp.b64 %rd19, 1, 0, %p5;
+; CHECK-NEXT: sub.s64 %rd20, 12, %rd5;
+; CHECK-NEXT: selp.b64 %rd21, %rd20, %rd4, %p3;
+; CHECK-NEXT: add.s64 %rd22, %rd21, %rd19;
+; CHECK-NEXT: add.s64 %rd23, %rd22, -1008;
+; CHECK-NEXT: shl.b64 %rd24, %rd23, 2;
+; CHECK-NEXT: shr.u64 %rd25, %rd1, 56;
+; CHECK-NEXT: and.b64 %rd26, %rd25, 128;
+; CHECK-NEXT: or.b64 %rd27, %rd26, %rd24;
+; CHECK-NEXT: or.b64 %rd28, %rd27, %rd18;
+; CHECK-NEXT: or.b64 %rd29, %rd9, 4503599627370496;
+; CHECK-NEXT: sub.s64 %rd30, 1059, %rd21;
+; CHECK-NEXT: min.u64 %rd31, %rd30, 63;
+; CHECK-NEXT: cvt.u32.u64 %r3, %rd31;
+; CHECK-NEXT: shr.u64 %rd32, %rd29, %r3;
+; CHECK-NEXT: and.b64 %rd33, %rd32, 1;
+; CHECK-NEXT: max.u64 %rd34, %rd31, 1;
+; CHECK-NEXT: add.s64 %rd35, %rd34, -1;
+; CHECK-NEXT: cvt.u32.u64 %r4, %rd35;
+; CHECK-NEXT: mov.b64 %rd36, 1;
+; CHECK-NEXT: shl.b64 %rd37, %rd36, %r4;
+; CHECK-NEXT: add.s64 %rd38, %rd37, -1;
+; CHECK-NEXT: and.b64 %rd39, %rd29, %rd38;
+; CHECK-NEXT: setp.ne.b64 %p6, %rd39, 0;
+; CHECK-NEXT: selp.b64 %rd40, 1, 0, %p6;
+; CHECK-NEXT: or.b64 %rd41, %rd40, %rd33;
+; CHECK-NEXT: shr.u64 %rd42, %rd29, %r4;
+; CHECK-NEXT: and.b64 %rd43, %rd42, %rd41;
+; CHECK-NEXT: setp.ne.b64 %p7, %rd31, 0;
+; CHECK-NEXT: selp.b64 %rd44, %rd43, 0, %p7;
+; CHECK-NEXT: add.s64 %rd45, %rd32, %rd44;
+; CHECK-NEXT: setp.gt.s64 %p8, %rd45, 3;
+; CHECK-NEXT: selp.b64 %rd46, 0, %rd45, %p8;
+; CHECK-NEXT: selp.b64 %rd47, 4, 0, %p8;
+; CHECK-NEXT: or.b64 %rd48, %rd26, %rd47;
+; CHECK-NEXT: or.b64 %rd49, %rd48, %rd46;
+; CHECK-NEXT: setp.lt.s64 %p9, %rd23, 1;
+; CHECK-NEXT: selp.b64 %rd50, %rd49, %rd28, %p9;
+; CHECK-NEXT: setp.gt.s64 %p10, %rd18, 3;
+; CHECK-NEXT: or.b64 %rd51, %rd26, 124;
+; CHECK-NEXT: selp.b64 %rd52, %rd51, %rd50, %p10;
+; CHECK-NEXT: setp.eq.b64 %p11, %rd23, 30;
+; CHECK-NEXT: selp.b64 %rd53, %rd52, %rd50, %p11;
+; CHECK-NEXT: setp.gt.s64 %p12, %rd23, 30;
+; CHECK-NEXT: selp.b64 %rd54, %rd51, %rd53, %p12;
+; CHECK-NEXT: or.b64 %rd55, %rd4, %rd2;
+; CHECK-NEXT: setp.eq.b64 %p13, %rd55, 0;
+; CHECK-NEXT: selp.b64 %rd56, %rd26, %rd54, %p13;
+; CHECK-NEXT: setp.eq.b64 %p14, %rd2, 0;
+; CHECK-NEXT: selp.b64 %rd57, %rd51, %rd56, %p14;
+; CHECK-NEXT: setp.eq.b64 %p15, %rd4, 2047;
+; CHECK-NEXT: selp.b64 %rd58, %rd57, %rd56, %p15;
+; CHECK-NEXT: selp.b64 %rd59, 126, %rd58, %p1;
+; CHECK-NEXT: selp.b64 %rd60, %rd59, %rd58, %p15;
+; CHECK-NEXT: st.param.b32 [func_retval0], %rd60;
+; CHECK-NEXT: ret;
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f64(double %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
diff --git a/llvm/test/CodeGen/X86/float-to-arbitrary-fp-error.ll b/llvm/test/CodeGen/X86/float-to-arbitrary-fp-error.ll
new file mode 100644
index 0000000000000..b1c5139ebede7
--- /dev/null
+++ b/llvm/test/CodeGen/X86/float-to-arbitrary-fp-error.ll
@@ -0,0 +1,76 @@
+; RUN: split-file %s %t
+; RUN: not llc < %t/float8e4m3.ll -mtriple=x86_64-unknown-unknown 2>&1 | FileCheck %s --check-prefix=E4M3
+; RUN: not llc < %t/float8e3m4.ll -mtriple=x86_64-unknown-unknown 2>&1 | FileCheck %s --check-prefix=E3M4
+; RUN: not llc < %t/float8e5m2fnuz.ll -mtriple=x86_64-unknown-unknown 2>&1 | FileCheck %s --check-prefix=E5M2FNUZ
+; RUN: not llc < %t/float8e4m3fnuz.ll -mtriple=x86_64-unknown-unknown 2>&1 | FileCheck %s --check-prefix=E4M3FNUZ
+; RUN: not llc < %t/float8e4m3b11fnuz.ll -mtriple=x86_64-unknown-unknown 2>&1 | FileCheck %s --check-prefix=E4M3B11FNUZ
+; RUN: not llc < %t/float8e8m0fnu.ll -mtriple=x86_64-unknown-unknown 2>&1 | FileCheck %s --check-prefix=E8M0FNU
+
+; Test that llvm.convert.to.arbitrary.fp emits an error for formats that pass
+; verifier validation but are not yet implemented in SelectionDAGBuilder.
+
+;--- float8e4m3.ll
+; E4M3: error: convert_to_arbitrary_fp: not implemented format 'Float8E4M3'
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+
+define i8 @to_f8e4m3(float %v) {
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(
+ float %v, metadata !"Float8E4M3", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+;--- float8e3m4.ll
+; E3M4: error: convert_to_arbitrary_fp: not implemented format 'Float8E3M4'
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+
+define i8 @to_f8e3m4(float %v) {
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(
+ float %v, metadata !"Float8E3M4", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+;--- float8e5m2fnuz.ll
+; E5M2FNUZ: error: convert_to_arbitrary_fp: not implemented format 'Float8E5M2FNUZ'
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+
+define i8 @to_f8e5m2fnuz(float %v) {
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(
+ float %v, metadata !"Float8E5M2FNUZ", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+;--- float8e4m3fnuz.ll
+; E4M3FNUZ: error: convert_to_arbitrary_fp: not implemented format 'Float8E4M3FNUZ'
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+
+define i8 @to_f8e4m3fnuz(float %v) {
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(
+ float %v, metadata !"Float8E4M3FNUZ", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+;--- float8e4m3b11fnuz.ll
+; E4M3B11FNUZ: error: convert_to_arbitrary_fp: not implemented format 'Float8E4M3B11FNUZ'
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+
+define i8 @to_f8e4m3b11fnuz(float %v) {
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(
+ float %v, metadata !"Float8E4M3B11FNUZ", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+;--- float8e8m0fnu.ll
+; E8M0FNU: error: convert_to_arbitrary_fp: not implemented format 'Float8E8M0FNU'
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+
+define i8 @to_f8e8m0fnu(float %v) {
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(
+ float %v, metadata !"Float8E8M0FNU", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
diff --git a/llvm/test/CodeGen/X86/float-to-arbitrary-fp.ll b/llvm/test/CodeGen/X86/float-to-arbitrary-fp.ll
new file mode 100644
index 0000000000000..619e42ab03718
--- /dev/null
+++ b/llvm/test/CodeGen/X86/float-to-arbitrary-fp.ll
@@ -0,0 +1,1323 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=x86_64-unknown-unknown | FileCheck %s
+
+; Test llvm.convert.to.arbitrary.fp intrinsic expansion.
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f32(float, metadata, metadata, i1)
+declare i6 @llvm.convert.to.arbitrary.fp.i6.f32(float, metadata, metadata, i1)
+declare i4 @llvm.convert.to.arbitrary.fp.i4.f32(float, metadata, metadata, i1)
+declare <4 x i4> @llvm.convert.to.arbitrary.fp.v4i4.v4f32(<4 x float>, metadata, metadata, i1)
+
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f16(half, metadata, metadata, i1)
+declare i8 @llvm.convert.to.arbitrary.fp.i8.bf16(bfloat, metadata, metadata, i1)
+declare i8 @llvm.convert.to.arbitrary.fp.i8.f64(double, metadata, metadata, i1)
+
+; Float8E5M2
+; Layout: sign(1) exp(5) mant(2), bias=15
+; Supports: Inf, NaN, signed zero, denormals
+
+; Float8E5M2 normal: f32 1.0 -> 0_01111_00 = 0x3C
+define i8 @to_f8e5m2_normal() {
+; CHECK-LABEL: to_f8e5m2_normal:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $60, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 zero: f32 0.0 -> 0_00000_00 = 0x00
+define i8 @to_f8e5m2_zero() {
+; CHECK-LABEL: to_f8e5m2_zero:
+; CHECK: # %bb.0:
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 negative zero: f32 -0.0 -> 1_00000_00 = 0x80
+define i8 @to_f8e5m2_neg_zero() {
+; CHECK-LABEL: to_f8e5m2_neg_zero:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $-128, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float -0.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 +Inf: f32 Inf -> 0_11111_00 = 0x7C
+define i8 @to_f8e5m2_inf() {
+; CHECK-LABEL: to_f8e5m2_inf:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $124, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF0000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 NaN: f32 NaN -> 0_11111_10 = 0x7E
+define i8 @to_f8e5m2_nan() {
+; CHECK-LABEL: to_f8e5m2_nan:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $126, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF8000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 max: 57344.0 -> 0_11110_11 = 0x7B
+define i8 @to_f8e5m2_max() {
+; CHECK-LABEL: to_f8e5m2_max:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $123, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 57344.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 overflow (no saturation): 100000.0 -> Inf = 0x7C
+define i8 @to_f8e5m2_overflow() {
+; CHECK-LABEL: to_f8e5m2_overflow:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $124, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 100000.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 overflow with saturation: 100000.0 -> max = 0x7B
+define i8 @to_f8e5m2_overflow_sat() {
+; CHECK-LABEL: to_f8e5m2_overflow_sat:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $123, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 100000.0, metadata !"Float8E5M2", metadata !"round.tonearest", i1 true)
+ ret i8 %r
+}
+
+; Float8E5M2 denorm: very small value -> dst denorm
+define i8 @to_f8e5m2_denorm() {
+ ; 2^(-16) = 0x00010000 in f32 = 1.52587891e-5
+; CHECK-LABEL: to_f8e5m2_denorm:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $1, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x3EF0000000000000, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E5M2 dynamic input
+define i8 @to_f8e5m2_dynamic(float %x) {
+; CHECK-LABEL: to_f8e5m2_dynamic:
+; CHECK: # %bb.0:
+; CHECK-NEXT: pushq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: pushq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: pushq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: pushq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: .cfi_offset %rbx, -40
+; CHECK-NEXT: .cfi_offset %r14, -32
+; CHECK-NEXT: .cfi_offset %r15, -24
+; CHECK-NEXT: .cfi_offset %rbp, -16
+; CHECK-NEXT: movd %xmm0, %edx
+; CHECK-NEXT: movl %edx, %eax
+; CHECK-NEXT: andl $8388607, %eax # imm = 0x7FFFFF
+; CHECK-NEXT: bsrl %eax, %ecx
+; CHECK-NEXT: xorl $31, %ecx
+; CHECK-NEXT: movl $9, %r11d
+; CHECK-NEXT: subl %ecx, %r11d
+; CHECK-NEXT: addl $-8, %ecx
+; CHECK-NEXT: movl %eax, %r10d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r10d
+; CHECK-NEXT: andl $8388607, %r10d # imm = 0x7FFFFF
+; CHECK-NEXT: testl %eax, %eax
+; CHECK-NEXT: sete %dil
+; CHECK-NEXT: setne %sil
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: shrl $23, %ecx
+; CHECK-NEXT: movzbl %cl, %r8d
+; CHECK-NEXT: testl %r8d, %r8d
+; CHECK-NEXT: sete %r9b
+; CHECK-NEXT: testb %sil, %r9b
+; CHECK-NEXT: cmovel %eax, %r10d
+; CHECK-NEXT: cmovel %r8d, %r11d
+; CHECK-NEXT: movl $134, %ecx
+; CHECK-NEXT: subl %r11d, %ecx
+; CHECK-NEXT: cmpl $31, %ecx
+; CHECK-NEXT: movl $31, %eax
+; CHECK-NEXT: cmovbl %ecx, %eax
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: adcl $-1, %ecx
+; CHECK-NEXT: leal 8388608(%r10), %ebx
+; CHECK-NEXT: movl %ebx, %ebp
+; CHECK-NEXT: shrl %cl, %ebp
+; CHECK-NEXT: movl $1, %r14d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r14d
+; CHECK-NEXT: decl %r14d
+; CHECK-NEXT: xorl %r15d, %r15d
+; CHECK-NEXT: testl %r14d, %ebx
+; CHECK-NEXT: setne %r15b
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: shrl %cl, %ebx
+; CHECK-NEXT: movl %ebx, %r14d
+; CHECK-NEXT: andl $1, %r14d
+; CHECK-NEXT: orl %r15d, %r14d
+; CHECK-NEXT: andl %ebp, %r14d
+; CHECK-NEXT: xorl %ecx, %ecx
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: cmovbl %ecx, %r14d
+; CHECK-NEXT: addl %ebx, %r14d
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: cmpl $4, %r14d
+; CHECK-NEXT: cmovgel %ecx, %r14d
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: shrl $24, %edx
+; CHECK-NEXT: andl $-128, %edx
+; CHECK-NEXT: leal (%rdx,%rax,4), %eax
+; CHECK-NEXT: orl %r14d, %eax
+; CHECK-NEXT: movl %r10d, %ebx
+; CHECK-NEXT: shrl $21, %ebx
+; CHECK-NEXT: movl %ebx, %ebp
+; CHECK-NEXT: andl $1, %ebp
+; CHECK-NEXT: xorl %r14d, %r14d
+; CHECK-NEXT: testl $1048575, %r10d # imm = 0xFFFFF
+; CHECK-NEXT: setne %r14b
+; CHECK-NEXT: orl %ebp, %r14d
+; CHECK-NEXT: shrl $20, %r10d
+; CHECK-NEXT: andl %r14d, %r10d
+; CHECK-NEXT: addl %ebx, %r10d
+; CHECK-NEXT: xorl %ebx, %ebx
+; CHECK-NEXT: cmpl $4, %r10d
+; CHECK-NEXT: cmovgel %ecx, %r10d
+; CHECK-NEXT: setge %bl
+; CHECK-NEXT: leal -112(%r11,%rbx), %ebp
+; CHECK-NEXT: addl %ebx, %r11d
+; CHECK-NEXT: leal -448(,%r11,4), %r11d
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: orl %r10d, %ecx
+; CHECK-NEXT: orl %r11d, %ecx
+; CHECK-NEXT: testl %ebp, %ebp
+; CHECK-NEXT: cmovlel %eax, %ecx
+; CHECK-NEXT: cmpl $4, %r10d
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: cmpl $30, %ebp
+; CHECK-NEXT: sete %r10b
+; CHECK-NEXT: andb %al, %r10b
+; CHECK-NEXT: cmpl $31, %ebp
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: orb %r10b, %al
+; CHECK-NEXT: leal 124(%rdx), %r10d
+; CHECK-NEXT: testb %al, %al
+; CHECK-NEXT: cmovnel %r10d, %ecx
+; CHECK-NEXT: testb %dil, %r9b
+; CHECK-NEXT: cmovnel %edx, %ecx
+; CHECK-NEXT: cmpl $255, %r8d
+; CHECK-NEXT: sete %al
+; CHECK-NEXT: testb %dil, %al
+; CHECK-NEXT: cmovnel %r10d, %ecx
+; CHECK-NEXT: testb %sil, %al
+; CHECK-NEXT: movl $126, %eax
+; CHECK-NEXT: cmovel %ecx, %eax
+; CHECK-NEXT: # kill: def $al killed $al killed $eax
+; CHECK-NEXT: popq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: popq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: popq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: popq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 8
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN
+; Layout: sign(1) exp(4) mant(3), bias=7
+; Supports: NaN (special encoding: 0_1111_111 only), no Inf, signed zero
+
+; Float8E4M3FN normal: f32 1.0 -> 0_0111_000 = 0x38
+define i8 @to_f8e4m3fn_normal() {
+; CHECK-LABEL: to_f8e4m3fn_normal:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $56, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN max: 448.0 -> 0_1111_110 = 0x7E
+define i8 @to_f8e4m3fn_max() {
+; CHECK-LABEL: to_f8e4m3fn_max:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $126, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 448.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN NaN: f32 NaN -> 0_1111_111 = 0x7F
+define i8 @to_f8e4m3fn_nan() {
+; CHECK-LABEL: to_f8e4m3fn_nan:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $127, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 0x7FF8000000000000, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Float8E4M3FN overflow with saturation: large value -> max = 0x7E
+define i8 @to_f8e4m3fn_overflow_sat() {
+; CHECK-LABEL: to_f8e4m3fn_overflow_sat:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $126, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1000.0, metadata !"Float8E4M3FN", metadata !"round.tonearest", i1 true)
+ ret i8 %r
+}
+
+; Float6E3M2FN
+; Layout: sign(1) exp(3) mant(2), bias=3
+; Supports: no Inf, no NaN, signed zero (FiniteOnly)
+
+; Float6E3M2FN normal: f32 1.0 -> 0_011_00 = 0x0C
+define i6 @to_f6e3m2fn_normal() {
+; CHECK-LABEL: to_f6e3m2fn_normal:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $12, %al
+; CHECK-NEXT: retq
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 1.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN max: 28.0 -> 0_111_11 = 0x1F
+define i6 @to_f6e3m2fn_max() {
+; CHECK-LABEL: to_f6e3m2fn_max:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $31, %al
+; CHECK-NEXT: retq
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 28.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN zero: f32 0.0 -> 0_000_00 = 0x00
+define i6 @to_f6e3m2fn_zero() {
+; CHECK-LABEL: to_f6e3m2fn_zero:
+; CHECK: # %bb.0:
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: retq
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 0.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E3M2FN negative: f32 -2.0 -> 1_100_00 = 0x30
+define i6 @to_f6e3m2fn_negative() {
+; CHECK-LABEL: to_f6e3m2fn_negative:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $48, %al
+; CHECK-NEXT: retq
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float -2.0, metadata !"Float6E3M2FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E2M3FN
+; Layout: sign(1) exp(2) mant(3), bias=1
+; Supports: no Inf, no NaN, signed zero (FiniteOnly)
+
+; Float6E2M3FN normal: f32 1.0 -> 0_01_000 = 0x08
+define i6 @to_f6e2m3fn_normal() {
+; CHECK-LABEL: to_f6e2m3fn_normal:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $8, %al
+; CHECK-NEXT: retq
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 1.0, metadata !"Float6E2M3FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float6E2M3FN max: 7.5 -> 0_11_111 = 0x1F
+define i6 @to_f6e2m3fn_max() {
+; CHECK-LABEL: to_f6e2m3fn_max:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $31, %al
+; CHECK-NEXT: retq
+ %r = call i6 @llvm.convert.to.arbitrary.fp.i6.f32(float 7.5, metadata !"Float6E2M3FN", metadata !"round.tonearest", i1 false)
+ ret i6 %r
+}
+
+; Float4E2M1FN
+; Layout: sign(1) exp(2) mant(1), bias=1
+; Supports: no Inf, no NaN, signed zero (FiniteOnly)
+
+; Float4E2M1FN normal: f32 1.0 -> 0_01_0 = 0x2
+define i4 @to_f4e2m1fn_normal() {
+; CHECK-LABEL: to_f4e2m1fn_normal:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $2, %al
+; CHECK-NEXT: retq
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 1.0, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Float4E2M1FN max: 6.0 -> 0_11_1 = 0x7
+define i4 @to_f4e2m1fn_max() {
+; CHECK-LABEL: to_f4e2m1fn_max:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $7, %al
+; CHECK-NEXT: retq
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 6.0, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Float4E2M1FN denorm: 0.5 -> 0_00_1 = 0x1
+define i4 @to_f4e2m1fn_denorm() {
+; CHECK-LABEL: to_f4e2m1fn_denorm:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $1, %al
+; CHECK-NEXT: retq
+ %r = call i4 @llvm.convert.to.arbitrary.fp.i4.f32(float 0.5, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret i4 %r
+}
+
+; Rounding tests
+
+; Round to nearest even: 1.5 in f32 with E5M2 should round to 1.5 (0_01111_10 = 0x3E)
+define i8 @to_f8e5m2_round_nearest() {
+; CHECK-LABEL: to_f8e5m2_round_nearest:
+; CHECK: # %bb.0:
+; CHECK-NEXT: movb $62, %al
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float 1.5, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; Round toward zero
+define i8 @to_f8e5m2_round_towardzero(float %x) {
+; CHECK-LABEL: to_f8e5m2_round_towardzero:
+; CHECK: # %bb.0:
+; CHECK-NEXT: pushq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: pushq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: pushq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: pushq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: .cfi_offset %rbx, -40
+; CHECK-NEXT: .cfi_offset %r14, -32
+; CHECK-NEXT: .cfi_offset %r15, -24
+; CHECK-NEXT: .cfi_offset %rbp, -16
+; CHECK-NEXT: movd %xmm0, %eax
+; CHECK-NEXT: movl %eax, %r11d
+; CHECK-NEXT: andl $8388607, %r11d # imm = 0x7FFFFF
+; CHECK-NEXT: bsrl %r11d, %edx
+; CHECK-NEXT: xorl $31, %edx
+; CHECK-NEXT: leal -8(%rdx), %ecx
+; CHECK-NEXT: movl %r11d, %r9d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r9d
+; CHECK-NEXT: andl $8388607, %r9d # imm = 0x7FFFFF
+; CHECK-NEXT: movl $9, %r10d
+; CHECK-NEXT: subl %edx, %r10d
+; CHECK-NEXT: testl %r11d, %r11d
+; CHECK-NEXT: sete %sil
+; CHECK-NEXT: setne %dl
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: shrl $23, %ecx
+; CHECK-NEXT: movzbl %cl, %edi
+; CHECK-NEXT: testl %edi, %edi
+; CHECK-NEXT: sete %r8b
+; CHECK-NEXT: testb %dl, %r8b
+; CHECK-NEXT: cmovel %edi, %r10d
+; CHECK-NEXT: cmovel %r11d, %r9d
+; CHECK-NEXT: movl %r9d, %r11d
+; CHECK-NEXT: shrl $21, %r11d
+; CHECK-NEXT: xorl %ebp, %ebp
+; CHECK-NEXT: xorl %r14d, %r14d
+; CHECK-NEXT: cmpl $4, %r11d
+; CHECK-NEXT: cmovgel %ebp, %r11d
+; CHECK-NEXT: setge %r14b
+; CHECK-NEXT: leal (%r10,%r14), %ecx
+; CHECK-NEXT: leal -448(,%rcx,4), %ebx
+; CHECK-NEXT: shrl $24, %eax
+; CHECK-NEXT: andl $-128, %eax
+; CHECK-NEXT: orl %eax, %ebx
+; CHECK-NEXT: orl %r11d, %ebx
+; CHECK-NEXT: movl $134, %r15d
+; CHECK-NEXT: subl %r10d, %r15d
+; CHECK-NEXT: cmpl $31, %r15d
+; CHECK-NEXT: movl $31, %ecx
+; CHECK-NEXT: cmovbl %r15d, %ecx
+; CHECK-NEXT: orl $8388608, %r9d # imm = 0x800000
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shrl %cl, %r9d
+; CHECK-NEXT: xorl %ecx, %ecx
+; CHECK-NEXT: cmpl $4, %r9d
+; CHECK-NEXT: cmovgel %ebp, %r9d
+; CHECK-NEXT: leal -112(%r10,%r14), %r10d
+; CHECK-NEXT: setge %cl
+; CHECK-NEXT: leal (%rax,%rcx,4), %ecx
+; CHECK-NEXT: orl %r9d, %ecx
+; CHECK-NEXT: testl %r10d, %r10d
+; CHECK-NEXT: cmovgl %ebx, %ecx
+; CHECK-NEXT: cmpl $4, %r11d
+; CHECK-NEXT: setge %r9b
+; CHECK-NEXT: cmpl $30, %r10d
+; CHECK-NEXT: sete %r11b
+; CHECK-NEXT: andb %r9b, %r11b
+; CHECK-NEXT: cmpl $31, %r10d
+; CHECK-NEXT: setge %r9b
+; CHECK-NEXT: orb %r11b, %r9b
+; CHECK-NEXT: leal 124(%rax), %r10d
+; CHECK-NEXT: testb %r9b, %r9b
+; CHECK-NEXT: cmovnel %r10d, %ecx
+; CHECK-NEXT: testb %sil, %r8b
+; CHECK-NEXT: cmovnel %eax, %ecx
+; CHECK-NEXT: cmpl $255, %edi
+; CHECK-NEXT: sete %al
+; CHECK-NEXT: testb %sil, %al
+; CHECK-NEXT: cmovnel %r10d, %ecx
+; CHECK-NEXT: testb %dl, %al
+; CHECK-NEXT: movl $126, %eax
+; CHECK-NEXT: cmovel %ecx, %eax
+; CHECK-NEXT: # kill: def $al killed $al killed $eax
+; CHECK-NEXT: popq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: popq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: popq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: popq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 8
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f32(float %x, metadata !"Float8E5M2", metadata !"round.towardzero", i1 false)
+ ret i8 %r
+}
+
+; Vector test
+
+; <4 x float> to <4 x i4> (Float4E2M1FN)
+define <4 x i4> @to_f4e2m1fn_v4f32(<4 x float> %x) {
+; CHECK-LABEL: to_f4e2m1fn_v4f32:
+; CHECK: # %bb.0:
+; CHECK-NEXT: pushq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: pushq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: pushq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: pushq %r13
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: pushq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 48
+; CHECK-NEXT: pushq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 56
+; CHECK-NEXT: .cfi_offset %rbx, -56
+; CHECK-NEXT: .cfi_offset %r12, -48
+; CHECK-NEXT: .cfi_offset %r13, -40
+; CHECK-NEXT: .cfi_offset %r14, -32
+; CHECK-NEXT: .cfi_offset %r15, -24
+; CHECK-NEXT: .cfi_offset %rbp, -16
+; CHECK-NEXT: movaps %xmm0, %xmm1
+; CHECK-NEXT: shufps {{.*#+}} xmm1 = xmm1[3,3],xmm0[3,3]
+; CHECK-NEXT: movd %xmm1, %ecx
+; CHECK-NEXT: movl %ecx, %eax
+; CHECK-NEXT: movq %rcx, %rdx
+; CHECK-NEXT: movq %rcx, {{[-0-9]+}}(%r{{[sb]}}p) # 8-byte Spill
+; CHECK-NEXT: andl $8388607, %eax # imm = 0x7FFFFF
+; CHECK-NEXT: bsrl %eax, %ecx
+; CHECK-NEXT: xorl $31, %ecx
+; CHECK-NEXT: movl $9, %esi
+; CHECK-NEXT: subl %ecx, %esi
+; CHECK-NEXT: addl $-8, %ecx
+; CHECK-NEXT: movl %eax, %r10d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r10d
+; CHECK-NEXT: andl $8388607, %r10d # imm = 0x7FFFFF
+; CHECK-NEXT: testl %eax, %eax
+; CHECK-NEXT: sete {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Spill
+; CHECK-NEXT: setne %cl
+; CHECK-NEXT: # kill: def $edx killed $edx killed $rdx
+; CHECK-NEXT: shrl $23, %edx
+; CHECK-NEXT: movzbl %dl, %edx
+; CHECK-NEXT: testl %edx, %edx
+; CHECK-NEXT: sete %dil
+; CHECK-NEXT: movb %dil, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Spill
+; CHECK-NEXT: testb %cl, %dil
+; CHECK-NEXT: cmovel %eax, %r10d
+; CHECK-NEXT: cmovel %edx, %esi
+; CHECK-NEXT: movq %rsi, {{[-0-9]+}}(%r{{[sb]}}p) # 8-byte Spill
+; CHECK-NEXT: movl $149, %eax
+; CHECK-NEXT: subl %esi, %eax
+; CHECK-NEXT: cmpl $31, %eax
+; CHECK-NEXT: movl $31, %r13d
+; CHECK-NEXT: cmovael %r13d, %eax
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: adcl $-1, %ecx
+; CHECK-NEXT: leal 8388608(%r10), %ebp
+; CHECK-NEXT: movl %ebp, %edx
+; CHECK-NEXT: shrl %cl, %edx
+; CHECK-NEXT: movl $1, %esi
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %esi
+; CHECK-NEXT: decl %esi
+; CHECK-NEXT: xorl %edi, %edi
+; CHECK-NEXT: testl %esi, %ebp
+; CHECK-NEXT: setne %dil
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: shrl %cl, %ebp
+; CHECK-NEXT: movl %ebp, {{[-0-9]+}}(%r{{[sb]}}p) # 4-byte Spill
+; CHECK-NEXT: andl $1, %ebp
+; CHECK-NEXT: orl %edi, %ebp
+; CHECK-NEXT: andl %edx, %ebp
+; CHECK-NEXT: movaps %xmm0, %xmm1
+; CHECK-NEXT: unpckhpd {{.*#+}} xmm1 = xmm1[1],xmm0[1]
+; CHECK-NEXT: movd %xmm1, %ecx
+; CHECK-NEXT: movl %ecx, %esi
+; CHECK-NEXT: movq %rcx, %rbx
+; CHECK-NEXT: movq %rcx, {{[-0-9]+}}(%r{{[sb]}}p) # 8-byte Spill
+; CHECK-NEXT: andl $8388607, %esi # imm = 0x7FFFFF
+; CHECK-NEXT: sete {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Spill
+; CHECK-NEXT: setne %r8b
+; CHECK-NEXT: movd %xmm0, %ecx
+; CHECK-NEXT: movl %ecx, %edx
+; CHECK-NEXT: movq %rcx, %r14
+; CHECK-NEXT: movq %rcx, {{[-0-9]+}}(%r{{[sb]}}p) # 8-byte Spill
+; CHECK-NEXT: andl $8388607, %edx # imm = 0x7FFFFF
+; CHECK-NEXT: sete {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Spill
+; CHECK-NEXT: setne %dil
+; CHECK-NEXT: xorl %r9d, %r9d
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: cmovbl %r9d, %ebp
+; CHECK-NEXT: bsrl %esi, %ecx
+; CHECK-NEXT: xorl $31, %ecx
+; CHECK-NEXT: movl $9, %r11d
+; CHECK-NEXT: subl %ecx, %r11d
+; CHECK-NEXT: addl $-8, %ecx
+; CHECK-NEXT: movl %esi, %r12d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r12d
+; CHECK-NEXT: andl $8388607, %r12d # imm = 0x7FFFFF
+; CHECK-NEXT: movl %ebx, %eax
+; CHECK-NEXT: shrl $23, %eax
+; CHECK-NEXT: movzbl %al, %eax
+; CHECK-NEXT: testl %eax, %eax
+; CHECK-NEXT: sete %cl
+; CHECK-NEXT: movb %cl, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Spill
+; CHECK-NEXT: testb %r8b, %cl
+; CHECK-NEXT: cmovel %esi, %r12d
+; CHECK-NEXT: cmovel %eax, %r11d
+; CHECK-NEXT: movq %r11, {{[-0-9]+}}(%r{{[sb]}}p) # 8-byte Spill
+; CHECK-NEXT: movl $149, %eax
+; CHECK-NEXT: subl %r11d, %eax
+; CHECK-NEXT: cmpl $31, %eax
+; CHECK-NEXT: cmovael %r13d, %eax
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: adcl $-1, %ecx
+; CHECK-NEXT: leal 8388608(%r12), %r11d
+; CHECK-NEXT: movl %r11d, %esi
+; CHECK-NEXT: shrl %cl, %esi
+; CHECK-NEXT: movl $1, %r8d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r8d
+; CHECK-NEXT: decl %r8d
+; CHECK-NEXT: xorl %ebx, %ebx
+; CHECK-NEXT: testl %r8d, %r11d
+; CHECK-NEXT: setne %bl
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: shrl %cl, %r11d
+; CHECK-NEXT: movl %r11d, {{[-0-9]+}}(%r{{[sb]}}p) # 4-byte Spill
+; CHECK-NEXT: andl $1, %r11d
+; CHECK-NEXT: orl %ebx, %r11d
+; CHECK-NEXT: andl %esi, %r11d
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: cmovbl %r9d, %r11d
+; CHECK-NEXT: bsrl %edx, %ecx
+; CHECK-NEXT: xorl $31, %ecx
+; CHECK-NEXT: movl $9, %esi
+; CHECK-NEXT: subl %ecx, %esi
+; CHECK-NEXT: addl $-8, %ecx
+; CHECK-NEXT: movl %edx, %r8d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r8d
+; CHECK-NEXT: andl $8388607, %r8d # imm = 0x7FFFFF
+; CHECK-NEXT: movl %r14d, %eax
+; CHECK-NEXT: shrl $23, %eax
+; CHECK-NEXT: movzbl %al, %eax
+; CHECK-NEXT: testl %eax, %eax
+; CHECK-NEXT: sete %cl
+; CHECK-NEXT: movb %cl, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Spill
+; CHECK-NEXT: testb %dil, %cl
+; CHECK-NEXT: cmovel %edx, %r8d
+; CHECK-NEXT: cmovel %eax, %esi
+; CHECK-NEXT: movq %rsi, {{[-0-9]+}}(%r{{[sb]}}p) # 8-byte Spill
+; CHECK-NEXT: movl $149, %eax
+; CHECK-NEXT: subl %esi, %eax
+; CHECK-NEXT: cmpl $31, %eax
+; CHECK-NEXT: cmovael %r13d, %eax
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: adcl $-1, %ecx
+; CHECK-NEXT: leal 8388608(%r8), %r14d
+; CHECK-NEXT: movl %r14d, %edx
+; CHECK-NEXT: shrl %cl, %edx
+; CHECK-NEXT: movl $1, %edi
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %edi
+; CHECK-NEXT: decl %edi
+; CHECK-NEXT: xorl %ebx, %ebx
+; CHECK-NEXT: testl %edi, %r14d
+; CHECK-NEXT: setne %bl
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: shrl %cl, %r14d
+; CHECK-NEXT: movl %r14d, {{[-0-9]+}}(%r{{[sb]}}p) # 4-byte Spill
+; CHECK-NEXT: andl $1, %r14d
+; CHECK-NEXT: orl %ebx, %r14d
+; CHECK-NEXT: andl %edx, %r14d
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: cmovbl %r9d, %r14d
+; CHECK-NEXT: shufps {{.*#+}} xmm0 = xmm0[1,1,1,1]
+; CHECK-NEXT: movd %xmm0, %ecx
+; CHECK-NEXT: movl %ecx, %eax
+; CHECK-NEXT: movq %rcx, %rdx
+; CHECK-NEXT: movq %rcx, {{[-0-9]+}}(%r{{[sb]}}p) # 8-byte Spill
+; CHECK-NEXT: andl $8388607, %eax # imm = 0x7FFFFF
+; CHECK-NEXT: bsrl %eax, %ecx
+; CHECK-NEXT: xorl $31, %ecx
+; CHECK-NEXT: movl $9, %r15d
+; CHECK-NEXT: subl %ecx, %r15d
+; CHECK-NEXT: addl $-8, %ecx
+; CHECK-NEXT: movl %eax, %edi
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %edi
+; CHECK-NEXT: andl $8388607, %edi # imm = 0x7FFFFF
+; CHECK-NEXT: testl %eax, %eax
+; CHECK-NEXT: sete {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Spill
+; CHECK-NEXT: setne %cl
+; CHECK-NEXT: movl %edx, %ebx
+; CHECK-NEXT: shrl $23, %ebx
+; CHECK-NEXT: movzbl %bl, %ebx
+; CHECK-NEXT: testl %ebx, %ebx
+; CHECK-NEXT: sete %dl
+; CHECK-NEXT: movb %dl, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Spill
+; CHECK-NEXT: testb %cl, %dl
+; CHECK-NEXT: cmovel %eax, %edi
+; CHECK-NEXT: cmovel %ebx, %r15d
+; CHECK-NEXT: movl $149, %edx
+; CHECK-NEXT: subl %r15d, %edx
+; CHECK-NEXT: cmpl $31, %edx
+; CHECK-NEXT: cmovael %r13d, %edx
+; CHECK-NEXT: cmpl $1, %edx
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: adcl $-1, %ecx
+; CHECK-NEXT: leal 8388608(%rdi), %r13d
+; CHECK-NEXT: movl %r13d, %ebx
+; CHECK-NEXT: shrl %cl, %ebx
+; CHECK-NEXT: movl $1, %eax
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %eax
+; CHECK-NEXT: decl %eax
+; CHECK-NEXT: xorl %esi, %esi
+; CHECK-NEXT: testl %eax, %r13d
+; CHECK-NEXT: setne %sil
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: shrl %cl, %r13d
+; CHECK-NEXT: movl %r13d, %ecx
+; CHECK-NEXT: andl $1, %ecx
+; CHECK-NEXT: orl %esi, %ecx
+; CHECK-NEXT: andl %ebx, %ecx
+; CHECK-NEXT: cmpl $1, %edx
+; CHECK-NEXT: movl $0, %edx
+; CHECK-NEXT: cmovbl %edx, %ecx
+; CHECK-NEXT: addl {{[-0-9]+}}(%r{{[sb]}}p), %ebp # 4-byte Folded Reload
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: cmpl $2, %ebp
+; CHECK-NEXT: cmovgel %edx, %ebp
+; CHECK-NEXT: movl $0, %ebx
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: addl %eax, %eax
+; CHECK-NEXT: movq {{[-0-9]+}}(%r{{[sb]}}p), %rdx # 8-byte Reload
+; CHECK-NEXT: shrl $31, %edx
+; CHECK-NEXT: leal (%rax,%rdx,8), %r9d
+; CHECK-NEXT: movq %rdx, %rax
+; CHECK-NEXT: orl %ebp, %r9d
+; CHECK-NEXT: movl %r10d, %edx
+; CHECK-NEXT: shrl $22, %edx
+; CHECK-NEXT: xorl %esi, %esi
+; CHECK-NEXT: testl $2097151, %r10d # imm = 0x1FFFFF
+; CHECK-NEXT: setne %sil
+; CHECK-NEXT: orl %edx, %esi
+; CHECK-NEXT: shrl $21, %r10d
+; CHECK-NEXT: andl %esi, %r10d
+; CHECK-NEXT: addl %edx, %r10d
+; CHECK-NEXT: xorl %edx, %edx
+; CHECK-NEXT: cmpl $2, %r10d
+; CHECK-NEXT: cmovgel %ebx, %r10d
+; CHECK-NEXT: setge %dl
+; CHECK-NEXT: movq {{[-0-9]+}}(%r{{[sb]}}p), %rbp # 8-byte Reload
+; CHECK-NEXT: leal (%rbp,%rdx), %esi
+; CHECK-NEXT: leal -252(%rsi,%rsi), %esi
+; CHECK-NEXT: shll $3, %eax
+; CHECK-NEXT: orl %eax, %r10d
+; CHECK-NEXT: orl %esi, %r10d
+; CHECK-NEXT: leal -126(%rbp,%rdx), %edx
+; CHECK-NEXT: testl %edx, %edx
+; CHECK-NEXT: cmovlel %r9d, %r10d
+; CHECK-NEXT: movzbl {{[-0-9]+}}(%r{{[sb]}}p), %edx # 1-byte Folded Reload
+; CHECK-NEXT: testb %dl, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Reload
+; CHECK-NEXT: cmovnel %eax, %r10d
+; CHECK-NEXT: movd %r10d, %xmm1
+; CHECK-NEXT: addl {{[-0-9]+}}(%r{{[sb]}}p), %r11d # 4-byte Folded Reload
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: cmpl $2, %r11d
+; CHECK-NEXT: cmovgel %ebx, %r11d
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: addl %eax, %eax
+; CHECK-NEXT: movq {{[-0-9]+}}(%r{{[sb]}}p), %r9 # 8-byte Reload
+; CHECK-NEXT: shrl $31, %r9d
+; CHECK-NEXT: leal (%rax,%r9,8), %eax
+; CHECK-NEXT: orl %r11d, %eax
+; CHECK-NEXT: movl %r12d, %edx
+; CHECK-NEXT: shrl $22, %edx
+; CHECK-NEXT: xorl %esi, %esi
+; CHECK-NEXT: testl $2097151, %r12d # imm = 0x1FFFFF
+; CHECK-NEXT: setne %sil
+; CHECK-NEXT: orl %edx, %esi
+; CHECK-NEXT: shrl $21, %r12d
+; CHECK-NEXT: andl %esi, %r12d
+; CHECK-NEXT: addl %edx, %r12d
+; CHECK-NEXT: xorl %edx, %edx
+; CHECK-NEXT: cmpl $2, %r12d
+; CHECK-NEXT: cmovgel %ebx, %r12d
+; CHECK-NEXT: setge %dl
+; CHECK-NEXT: movq {{[-0-9]+}}(%r{{[sb]}}p), %r10 # 8-byte Reload
+; CHECK-NEXT: leal (%r10,%rdx), %esi
+; CHECK-NEXT: leal -252(%rsi,%rsi), %esi
+; CHECK-NEXT: shll $3, %r9d
+; CHECK-NEXT: orl %r9d, %r12d
+; CHECK-NEXT: orl %esi, %r12d
+; CHECK-NEXT: leal -126(%r10,%rdx), %edx
+; CHECK-NEXT: testl %edx, %edx
+; CHECK-NEXT: cmovlel %eax, %r12d
+; CHECK-NEXT: movzbl {{[-0-9]+}}(%r{{[sb]}}p), %eax # 1-byte Folded Reload
+; CHECK-NEXT: testb %al, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Reload
+; CHECK-NEXT: cmovnel %r9d, %r12d
+; CHECK-NEXT: movd %r12d, %xmm2
+; CHECK-NEXT: addl {{[-0-9]+}}(%r{{[sb]}}p), %r14d # 4-byte Folded Reload
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: cmpl $2, %r14d
+; CHECK-NEXT: cmovgel %ebx, %r14d
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: addl %eax, %eax
+; CHECK-NEXT: movq {{[-0-9]+}}(%r{{[sb]}}p), %r9 # 8-byte Reload
+; CHECK-NEXT: shrl $31, %r9d
+; CHECK-NEXT: leal (%rax,%r9,8), %eax
+; CHECK-NEXT: orl %r14d, %eax
+; CHECK-NEXT: movl %r8d, %edx
+; CHECK-NEXT: shrl $22, %edx
+; CHECK-NEXT: xorl %esi, %esi
+; CHECK-NEXT: testl $2097151, %r8d # imm = 0x1FFFFF
+; CHECK-NEXT: setne %sil
+; CHECK-NEXT: orl %edx, %esi
+; CHECK-NEXT: shrl $21, %r8d
+; CHECK-NEXT: andl %esi, %r8d
+; CHECK-NEXT: addl %edx, %r8d
+; CHECK-NEXT: xorl %edx, %edx
+; CHECK-NEXT: cmpl $2, %r8d
+; CHECK-NEXT: cmovgel %ebx, %r8d
+; CHECK-NEXT: setge %dl
+; CHECK-NEXT: movq {{[-0-9]+}}(%r{{[sb]}}p), %r10 # 8-byte Reload
+; CHECK-NEXT: leal (%r10,%rdx), %esi
+; CHECK-NEXT: leal -252(%rsi,%rsi), %esi
+; CHECK-NEXT: shll $3, %r9d
+; CHECK-NEXT: orl %r9d, %r8d
+; CHECK-NEXT: orl %esi, %r8d
+; CHECK-NEXT: leal -126(%r10,%rdx), %edx
+; CHECK-NEXT: testl %edx, %edx
+; CHECK-NEXT: cmovlel %eax, %r8d
+; CHECK-NEXT: movzbl {{[-0-9]+}}(%r{{[sb]}}p), %eax # 1-byte Folded Reload
+; CHECK-NEXT: testb %al, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Reload
+; CHECK-NEXT: cmovnel %r9d, %r8d
+; CHECK-NEXT: movd %r8d, %xmm0
+; CHECK-NEXT: addl %r13d, %ecx
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: cmpl $2, %ecx
+; CHECK-NEXT: cmovgel %ebx, %ecx
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: addl %eax, %eax
+; CHECK-NEXT: movq {{[-0-9]+}}(%r{{[sb]}}p), %rsi # 8-byte Reload
+; CHECK-NEXT: shrl $31, %esi
+; CHECK-NEXT: leal (%rax,%rsi,8), %eax
+; CHECK-NEXT: orl %ecx, %eax
+; CHECK-NEXT: movl %edi, %ecx
+; CHECK-NEXT: shrl $22, %ecx
+; CHECK-NEXT: xorl %edx, %edx
+; CHECK-NEXT: testl $2097151, %edi # imm = 0x1FFFFF
+; CHECK-NEXT: setne %dl
+; CHECK-NEXT: orl %ecx, %edx
+; CHECK-NEXT: shrl $21, %edi
+; CHECK-NEXT: andl %edx, %edi
+; CHECK-NEXT: addl %ecx, %edi
+; CHECK-NEXT: xorl %ecx, %ecx
+; CHECK-NEXT: cmpl $2, %edi
+; CHECK-NEXT: cmovgel %ebx, %edi
+; CHECK-NEXT: setge %cl
+; CHECK-NEXT: leal -126(%r15,%rcx), %edx
+; CHECK-NEXT: addl %ecx, %r15d
+; CHECK-NEXT: leal -252(%r15,%r15), %ecx
+; CHECK-NEXT: shll $3, %esi
+; CHECK-NEXT: orl %esi, %edi
+; CHECK-NEXT: orl %ecx, %edi
+; CHECK-NEXT: testl %edx, %edx
+; CHECK-NEXT: cmovlel %eax, %edi
+; CHECK-NEXT: movzbl {{[-0-9]+}}(%r{{[sb]}}p), %eax # 1-byte Folded Reload
+; CHECK-NEXT: testb %al, {{[-0-9]+}}(%r{{[sb]}}p) # 1-byte Folded Reload
+; CHECK-NEXT: cmovnel %esi, %edi
+; CHECK-NEXT: movd %edi, %xmm3
+; CHECK-NEXT: punpckldq {{.*#+}} xmm2 = xmm2[0],xmm1[0],xmm2[1],xmm1[1]
+; CHECK-NEXT: punpckldq {{.*#+}} xmm0 = xmm0[0],xmm3[0],xmm0[1],xmm3[1]
+; CHECK-NEXT: punpcklqdq {{.*#+}} xmm0 = xmm0[0],xmm2[0]
+; CHECK-NEXT: popq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 48
+; CHECK-NEXT: popq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: popq %r13
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: popq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: popq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: popq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 8
+; CHECK-NEXT: retq
+ %r = call <4 x i4> @llvm.convert.to.arbitrary.fp.v4i4.v4f32(<4 x float> %x, metadata !"Float4E2M1FN", metadata !"round.tonearest", i1 false)
+ ret <4 x i4> %r
+}
+
+; Different source types
+
+; f16 -> i8 (Float8E5M2)
+define i8 @to_f8e5m2_from_f16(half %x) {
+; CHECK-LABEL: to_f8e5m2_from_f16:
+; CHECK: # %bb.0:
+; CHECK-NEXT: pushq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: pushq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: pushq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: pushq %r13
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: pushq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 48
+; CHECK-NEXT: pushq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 56
+; CHECK-NEXT: .cfi_offset %rbx, -56
+; CHECK-NEXT: .cfi_offset %r12, -48
+; CHECK-NEXT: .cfi_offset %r13, -40
+; CHECK-NEXT: .cfi_offset %r14, -32
+; CHECK-NEXT: .cfi_offset %r15, -24
+; CHECK-NEXT: .cfi_offset %rbp, -16
+; CHECK-NEXT: pextrw $0, %xmm0, %eax
+; CHECK-NEXT: movl %eax, %r11d
+; CHECK-NEXT: andl $1023, %r11d # imm = 0x3FF
+; CHECK-NEXT: bsrw %r11w, %si
+; CHECK-NEXT: xorl $15, %esi
+; CHECK-NEXT: leal -5(%rsi), %ecx
+; CHECK-NEXT: movl %r11d, %edx
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %edx
+; CHECK-NEXT: andl $1023, %edx # imm = 0x3FF
+; CHECK-NEXT: movl $6, %r10d
+; CHECK-NEXT: subl %esi, %r10d
+; CHECK-NEXT: testl %r11d, %r11d
+; CHECK-NEXT: sete %dil
+; CHECK-NEXT: setne %r13b
+; CHECK-NEXT: movl %eax, %r8d
+; CHECK-NEXT: shrl $10, %r8d
+; CHECK-NEXT: andl $31, %r8d
+; CHECK-NEXT: sete %r9b
+; CHECK-NEXT: testb %r13b, %r9b
+; CHECK-NEXT: cmovel %r11d, %edx
+; CHECK-NEXT: cmovel %r8d, %r10d
+; CHECK-NEXT: movl $9, %r11d
+; CHECK-NEXT: subl %r10d, %r11d
+; CHECK-NEXT: cmpw $15, %r11w
+; CHECK-NEXT: movl $15, %ecx
+; CHECK-NEXT: cmovbl %r11d, %ecx
+; CHECK-NEXT: movl %ecx, %ebx
+; CHECK-NEXT: leal 1024(%rdx), %r11d
+; CHECK-NEXT: movl %r11d, %ebp
+; CHECK-NEXT: shrl %cl, %ebp
+; CHECK-NEXT: movl %ebp, %r14d
+; CHECK-NEXT: andl $1, %r14d
+; CHECK-NEXT: cmpw $1, %cx
+; CHECK-NEXT: adcl $-1, %ecx
+; CHECK-NEXT: movl $1, %r15d
+; CHECK-NEXT: shll %cl, %r15d
+; CHECK-NEXT: decl %r15d
+; CHECK-NEXT: xorl %r12d, %r12d
+; CHECK-NEXT: testw %r15w, %r11w
+; CHECK-NEXT: setne %r12b
+; CHECK-NEXT: orl %r14d, %r12d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shrl %cl, %r11d
+; CHECK-NEXT: andl %r12d, %r11d
+; CHECK-NEXT: xorl %ecx, %ecx
+; CHECK-NEXT: cmpw $1, %bx
+; CHECK-NEXT: cmovbl %ecx, %r11d
+; CHECK-NEXT: addl %ebp, %r11d
+; CHECK-NEXT: xorl %ebx, %ebx
+; CHECK-NEXT: cmpw $4, %r11w
+; CHECK-NEXT: cmovgel %ecx, %r11d
+; CHECK-NEXT: setge %bl
+; CHECK-NEXT: shrl $8, %eax
+; CHECK-NEXT: andl $128, %eax
+; CHECK-NEXT: leal (%rax,%rbx,4), %ebx
+; CHECK-NEXT: orl %r11d, %ebx
+; CHECK-NEXT: movzbl %dh, %esi
+; CHECK-NEXT: movl %esi, %ebp
+; CHECK-NEXT: andl $1, %ebp
+; CHECK-NEXT: xorl %r11d, %r11d
+; CHECK-NEXT: testb $127, %dl
+; CHECK-NEXT: setne %r11b
+; CHECK-NEXT: orl %ebp, %r11d
+; CHECK-NEXT: shrl $7, %edx
+; CHECK-NEXT: andl %r11d, %edx
+; CHECK-NEXT: addl %esi, %edx
+; CHECK-NEXT: xorl %esi, %esi
+; CHECK-NEXT: cmpw $4, %dx
+; CHECK-NEXT: cmovgel %ecx, %edx
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: setge %sil
+; CHECK-NEXT: addl %r10d, %esi
+; CHECK-NEXT: movl %esi, %r10d
+; CHECK-NEXT: shll $2, %r10d
+; CHECK-NEXT: orl %eax, %edx
+; CHECK-NEXT: orl %r10d, %edx
+; CHECK-NEXT: testw %si, %si
+; CHECK-NEXT: cmovlel %ebx, %edx
+; CHECK-NEXT: cmpw $4, %cx
+; CHECK-NEXT: setge %cl
+; CHECK-NEXT: cmpw $30, %si
+; CHECK-NEXT: sete %r10b
+; CHECK-NEXT: andb %cl, %r10b
+; CHECK-NEXT: cmpw $31, %si
+; CHECK-NEXT: setge %cl
+; CHECK-NEXT: orb %r10b, %cl
+; CHECK-NEXT: leal 124(%rax), %esi
+; CHECK-NEXT: testb %cl, %cl
+; CHECK-NEXT: cmovnel %esi, %edx
+; CHECK-NEXT: testb %dil, %r9b
+; CHECK-NEXT: cmovnel %eax, %edx
+; CHECK-NEXT: cmpw $31, %r8w
+; CHECK-NEXT: sete %al
+; CHECK-NEXT: testb %dil, %al
+; CHECK-NEXT: cmovnel %esi, %edx
+; CHECK-NEXT: testb %r13b, %al
+; CHECK-NEXT: movl $126, %eax
+; CHECK-NEXT: cmovel %edx, %eax
+; CHECK-NEXT: # kill: def $al killed $al killed $eax
+; CHECK-NEXT: popq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 48
+; CHECK-NEXT: popq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: popq %r13
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: popq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: popq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: popq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 8
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f16(half %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; bf16 -> i8 (Float8E5M2)
+define i8 @to_f8e5m2_from_bf16(bfloat %x) {
+; CHECK-LABEL: to_f8e5m2_from_bf16:
+; CHECK: # %bb.0:
+; CHECK-NEXT: pushq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: pushq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: pushq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: pushq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: pushq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 48
+; CHECK-NEXT: .cfi_offset %rbx, -48
+; CHECK-NEXT: .cfi_offset %r12, -40
+; CHECK-NEXT: .cfi_offset %r14, -32
+; CHECK-NEXT: .cfi_offset %r15, -24
+; CHECK-NEXT: .cfi_offset %rbp, -16
+; CHECK-NEXT: pextrw $0, %xmm0, %edx
+; CHECK-NEXT: movl %edx, %eax
+; CHECK-NEXT: andl $127, %eax
+; CHECK-NEXT: shll $16, %eax
+; CHECK-NEXT: bsrl %eax, %esi
+; CHECK-NEXT: xorl $31, %esi
+; CHECK-NEXT: leal -8(%rsi), %ecx
+; CHECK-NEXT: movl %eax, %r9d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shll %cl, %r9d
+; CHECK-NEXT: andl $8323072, %r9d # imm = 0x7F0000
+; CHECK-NEXT: movl $9, %r11d
+; CHECK-NEXT: subl %esi, %r11d
+; CHECK-NEXT: testl %eax, %eax
+; CHECK-NEXT: sete %dil
+; CHECK-NEXT: setne %sil
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: shrl $7, %ecx
+; CHECK-NEXT: movzbl %cl, %r8d
+; CHECK-NEXT: testl %r8d, %r8d
+; CHECK-NEXT: sete %r10b
+; CHECK-NEXT: testb %sil, %r10b
+; CHECK-NEXT: cmovel %eax, %r9d
+; CHECK-NEXT: cmovel %r8d, %r11d
+; CHECK-NEXT: leal 8388608(%r9), %ebx
+; CHECK-NEXT: movl $134, %ecx
+; CHECK-NEXT: subl %r11d, %ecx
+; CHECK-NEXT: cmpl $31, %ecx
+; CHECK-NEXT: movl $31, %eax
+; CHECK-NEXT: cmovbl %ecx, %eax
+; CHECK-NEXT: movl %ebx, %ebp
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: shrl %cl, %ebp
+; CHECK-NEXT: movl %ebp, %r14d
+; CHECK-NEXT: andl $1, %r14d
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: adcl $-1, %ecx
+; CHECK-NEXT: movl $1, %r15d
+; CHECK-NEXT: shll %cl, %r15d
+; CHECK-NEXT: decl %r15d
+; CHECK-NEXT: xorl %r12d, %r12d
+; CHECK-NEXT: testl %r15d, %ebx
+; CHECK-NEXT: setne %r12b
+; CHECK-NEXT: orl %r14d, %r12d
+; CHECK-NEXT: # kill: def $cl killed $cl killed $ecx
+; CHECK-NEXT: shrl %cl, %ebx
+; CHECK-NEXT: andl %r12d, %ebx
+; CHECK-NEXT: xorl %ecx, %ecx
+; CHECK-NEXT: cmpl $1, %eax
+; CHECK-NEXT: cmovbl %ecx, %ebx
+; CHECK-NEXT: addl %ebp, %ebx
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: cmpl $4, %ebx
+; CHECK-NEXT: cmovgel %ecx, %ebx
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: shrl $8, %edx
+; CHECK-NEXT: andl $128, %edx
+; CHECK-NEXT: leal (%rdx,%rax,4), %eax
+; CHECK-NEXT: orl %ebx, %eax
+; CHECK-NEXT: movl %r9d, %ebx
+; CHECK-NEXT: shrl $21, %ebx
+; CHECK-NEXT: movl %ebx, %ebp
+; CHECK-NEXT: andl $1, %ebp
+; CHECK-NEXT: xorl %r14d, %r14d
+; CHECK-NEXT: testl $983040, %r9d # imm = 0xF0000
+; CHECK-NEXT: setne %r14b
+; CHECK-NEXT: orl %ebp, %r14d
+; CHECK-NEXT: shrl $20, %r9d
+; CHECK-NEXT: andl %r14d, %r9d
+; CHECK-NEXT: addl %ebx, %r9d
+; CHECK-NEXT: xorl %ebx, %ebx
+; CHECK-NEXT: cmpl $4, %r9d
+; CHECK-NEXT: cmovgel %ecx, %r9d
+; CHECK-NEXT: setge %bl
+; CHECK-NEXT: leal -112(%r11,%rbx), %ebp
+; CHECK-NEXT: addl %ebx, %r11d
+; CHECK-NEXT: leal -448(,%r11,4), %r11d
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: orl %r9d, %ecx
+; CHECK-NEXT: orl %r11d, %ecx
+; CHECK-NEXT: testl %ebp, %ebp
+; CHECK-NEXT: cmovlel %eax, %ecx
+; CHECK-NEXT: cmpl $4, %r9d
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: cmpl $30, %ebp
+; CHECK-NEXT: sete %r9b
+; CHECK-NEXT: andb %al, %r9b
+; CHECK-NEXT: cmpl $31, %ebp
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: orb %r9b, %al
+; CHECK-NEXT: leal 124(%rdx), %r9d
+; CHECK-NEXT: testb %al, %al
+; CHECK-NEXT: cmovnel %r9d, %ecx
+; CHECK-NEXT: testb %dil, %r10b
+; CHECK-NEXT: cmovnel %edx, %ecx
+; CHECK-NEXT: cmpl $255, %r8d
+; CHECK-NEXT: sete %al
+; CHECK-NEXT: testb %dil, %al
+; CHECK-NEXT: cmovnel %r9d, %ecx
+; CHECK-NEXT: testb %sil, %al
+; CHECK-NEXT: movl $126, %eax
+; CHECK-NEXT: cmovel %ecx, %eax
+; CHECK-NEXT: # kill: def $al killed $al killed $eax
+; CHECK-NEXT: popq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: popq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: popq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: popq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: popq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 8
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.bf16(bfloat %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
+
+; f64 -> i8 (Float8E5M2)
+define i8 @to_f8e5m2_from_f64(double %x) {
+; CHECK-LABEL: to_f8e5m2_from_f64:
+; CHECK: # %bb.0:
+; CHECK-NEXT: pushq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: pushq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: pushq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: pushq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: pushq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 48
+; CHECK-NEXT: .cfi_offset %rbx, -48
+; CHECK-NEXT: .cfi_offset %r12, -40
+; CHECK-NEXT: .cfi_offset %r14, -32
+; CHECK-NEXT: .cfi_offset %r15, -24
+; CHECK-NEXT: .cfi_offset %rbp, -16
+; CHECK-NEXT: movq %xmm0, %rsi
+; CHECK-NEXT: movabsq $4503599627370495, %rax # imm = 0xFFFFFFFFFFFFF
+; CHECK-NEXT: movq %rsi, %rdx
+; CHECK-NEXT: andq %rax, %rdx
+; CHECK-NEXT: bsrq %rdx, %rdi
+; CHECK-NEXT: xorq $63, %rdi
+; CHECK-NEXT: leaq -11(%rdi), %rcx
+; CHECK-NEXT: movq %rdx, %r10
+; CHECK-NEXT: # kill: def $cl killed $cl killed $rcx
+; CHECK-NEXT: shlq %cl, %r10
+; CHECK-NEXT: andq %rax, %r10
+; CHECK-NEXT: movl $12, %ebx
+; CHECK-NEXT: subq %rdi, %rbx
+; CHECK-NEXT: testq %rdx, %rdx
+; CHECK-NEXT: sete %r8b
+; CHECK-NEXT: setne %dil
+; CHECK-NEXT: movq %rsi, %r9
+; CHECK-NEXT: shrq $52, %r9
+; CHECK-NEXT: andl $2047, %r9d # imm = 0x7FF
+; CHECK-NEXT: sete %r11b
+; CHECK-NEXT: testb %dil, %r11b
+; CHECK-NEXT: cmoveq %rdx, %r10
+; CHECK-NEXT: cmoveq %r9, %rbx
+; CHECK-NEXT: movl $1059, %ecx # imm = 0x423
+; CHECK-NEXT: subq %rbx, %rcx
+; CHECK-NEXT: cmpq $63, %rcx
+; CHECK-NEXT: movl $63, %eax
+; CHECK-NEXT: cmovbq %rcx, %rax
+; CHECK-NEXT: cmpq $1, %rax
+; CHECK-NEXT: movq %rax, %rdx
+; CHECK-NEXT: adcq $0, %rdx
+; CHECK-NEXT: decl %edx
+; CHECK-NEXT: movl $1, %r15d
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: shlq %cl, %r15
+; CHECK-NEXT: decq %r15
+; CHECK-NEXT: movabsq $4503599627370496, %r14 # imm = 0x10000000000000
+; CHECK-NEXT: orq %r10, %r14
+; CHECK-NEXT: xorl %ebp, %ebp
+; CHECK-NEXT: testq %r15, %r14
+; CHECK-NEXT: setne %bpl
+; CHECK-NEXT: movq %r14, %r15
+; CHECK-NEXT: movl %eax, %ecx
+; CHECK-NEXT: shrq %cl, %r15
+; CHECK-NEXT: movl %r15d, %r12d
+; CHECK-NEXT: andl $1, %r12d
+; CHECK-NEXT: orl %ebp, %r12d
+; CHECK-NEXT: movl %edx, %ecx
+; CHECK-NEXT: shrq %cl, %r14
+; CHECK-NEXT: andl %r12d, %r14d
+; CHECK-NEXT: xorl %ecx, %ecx
+; CHECK-NEXT: cmpq $1, %rax
+; CHECK-NEXT: cmovbq %rcx, %r14
+; CHECK-NEXT: addq %r15, %r14
+; CHECK-NEXT: xorl %eax, %eax
+; CHECK-NEXT: cmpq $4, %r14
+; CHECK-NEXT: cmovgeq %rcx, %r14
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: shrq $63, %rsi
+; CHECK-NEXT: shll $7, %esi
+; CHECK-NEXT: leal (%rsi,%rax,4), %eax
+; CHECK-NEXT: orq %r14, %rax
+; CHECK-NEXT: movq %r10, %rdx
+; CHECK-NEXT: shrq $50, %rdx
+; CHECK-NEXT: movl %edx, %ebp
+; CHECK-NEXT: andl $1, %ebp
+; CHECK-NEXT: xorl %r14d, %r14d
+; CHECK-NEXT: movq %r10, %r15
+; CHECK-NEXT: shlq $15, %r15
+; CHECK-NEXT: setne %r14b
+; CHECK-NEXT: orl %ebp, %r14d
+; CHECK-NEXT: shrq $49, %r10
+; CHECK-NEXT: andl %r14d, %r10d
+; CHECK-NEXT: addq %rdx, %r10
+; CHECK-NEXT: xorl %edx, %edx
+; CHECK-NEXT: cmpq $4, %r10
+; CHECK-NEXT: cmovgeq %rcx, %r10
+; CHECK-NEXT: setge %dl
+; CHECK-NEXT: leaq -1008(%rbx,%rdx), %r14
+; CHECK-NEXT: addq %rdx, %rbx
+; CHECK-NEXT: leaq -4032(,%rbx,4), %rdx
+; CHECK-NEXT: movq %rsi, %rcx
+; CHECK-NEXT: orq %r10, %rcx
+; CHECK-NEXT: orq %rdx, %rcx
+; CHECK-NEXT: testq %r14, %r14
+; CHECK-NEXT: cmovleq %rax, %rcx
+; CHECK-NEXT: cmpq $4, %r10
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: cmpq $30, %r14
+; CHECK-NEXT: sete %dl
+; CHECK-NEXT: andb %al, %dl
+; CHECK-NEXT: cmpq $31, %r14
+; CHECK-NEXT: setge %al
+; CHECK-NEXT: orb %dl, %al
+; CHECK-NEXT: leaq 124(%rsi), %rdx
+; CHECK-NEXT: testb %al, %al
+; CHECK-NEXT: cmovneq %rdx, %rcx
+; CHECK-NEXT: testb %r8b, %r11b
+; CHECK-NEXT: cmovneq %rsi, %rcx
+; CHECK-NEXT: cmpq $2047, %r9 # imm = 0x7FF
+; CHECK-NEXT: sete %al
+; CHECK-NEXT: testb %r8b, %al
+; CHECK-NEXT: cmovneq %rdx, %rcx
+; CHECK-NEXT: testb %dil, %al
+; CHECK-NEXT: movl $126, %eax
+; CHECK-NEXT: cmoveq %rcx, %rax
+; CHECK-NEXT: # kill: def $al killed $al killed $rax
+; CHECK-NEXT: popq %rbx
+; CHECK-NEXT: .cfi_def_cfa_offset 40
+; CHECK-NEXT: popq %r12
+; CHECK-NEXT: .cfi_def_cfa_offset 32
+; CHECK-NEXT: popq %r14
+; CHECK-NEXT: .cfi_def_cfa_offset 24
+; CHECK-NEXT: popq %r15
+; CHECK-NEXT: .cfi_def_cfa_offset 16
+; CHECK-NEXT: popq %rbp
+; CHECK-NEXT: .cfi_def_cfa_offset 8
+; CHECK-NEXT: retq
+ %r = call i8 @llvm.convert.to.arbitrary.fp.i8.f64(double %x, metadata !"Float8E5M2", metadata !"round.tonearest", i1 false)
+ ret i8 %r
+}
More information about the llvm-commits
mailing list