VCVTBIASPH2BF8

VCVTBIASPH2BF8 converts half-precision values to E5M2 by adding a byte to each before cutting it short, so that the bytes choose the rounding. The halves come from the third operand, which may be memory or a broadcast half; the bias of half i is the low byte of word i of the second operand, an unsigned number. A zmm register of halves and one of biases give 32 bytes in a ymm register, and the bits above are cleared. It belongs to AVX10.2 and to the AVX10_V1_AUX subset that ACE lists, opcode 74h without a prefix in EVEX map 0F38, with a write mask of one bit per result byte. VCVTPH2BF8 is the same conversion rounded to nearest even.

A byte below the last place#

E5M2 is the upper byte of a half, so the instruction adds the bias to the half’s bits and keeps the upper byte: the bias lies just under the result’s last place. A bias of 0 truncates toward zero, 255 rounds every inexact value away from zero, and 127 and 128 round to nearest with ties toward and away from zero. With uniformly random biases, a half rounds away from zero with a probability equal to the fraction its dropped byte stands for, d/256 for a dropped byte d: stochastic rounding, which the specification names as the purpose. MXCSR is neither read nor written, and nothing is flagged.

The carry can reach infinity: FP16’s largest value, 65,504 (7BFFh), gives 57,344 with a bias of 0 and infinity with any other. Infinities stay infinite, and NaNs come out quiet. VCVTBIASPH2BF8S saturates at 57,344 instead.

Denormal halves#

For a half below 2^-14 the documents give three readings. AVX10.2’s helper function and its description of the instruction add the bias to the half’s bits as for any other value; ACE’s helper first normalizes the denormal, so the bias lands lower and carries less often; and AVX10.2’s introduction to FP8 (section 3.2.2) rounds a denormal input to nearest even, without the bias. The Operation block computes them and leaves the result UNDEFINED wherever they differ.

Recorded and built#

Zen 5, whose CPUID reports no AVX10, raised #UD for all seven recorded forms. GCC 14 rejects -mavx10.2. Clang 23 compiles _mm512_cvtbiasph_bf8(b, a), the biases first, to vcvtbiasph2bf8 ymm0, zmm0, zmm1.

Forms

Data flow

Where the elements of the sources go, element 0 on the right; results that depend on a control value are drawn for an example. How to read.

Operation

What the forms do, in pseudocode that our tests run against results recorded on real processors. How to read.

if not cpuid_leaf(7, 1)[3] >> 19 & 1:
    fault(UD)                    # CPUID: no AVX10
ver = cpuid_leaf(0x24, 0)[1] & 0xFF
aux = cpuid_leaf(0x24, 1)[2] >> 2 & 1
if ver < 2 and not aux:
    fault(UD)                    # neither AVX10.2 nor AVX10_V1_AUX
for i in range(VL // 16):
    if k1[i]:
        x = src3.f16[i]          # FP16 element i
        b = zmm2.b[2 * i]        # its bias: word i's low byte
        a = x & 0x7FFF           # |x|, as bits
        r = (x + b) >> 8         # E5M2 is FP16's high byte: the
                                 # bias below its last place, the
                                 # sum cut there; past the
                                 # largest: infinity
        if a > 0x7C00:
            r = x >> 8 | 2       # a NaN, quieted
        elif 0 < a < 0x400:      # a denormal; r is AVX10.2's
                                 # reading (5.1, 9.3.2)
            n = to_bf8(x)        # AVX10.2's 3.2.2: to nearest
                                 # even, no bias
            k = lzcnt(a, 10) + 1
            u = ((a << k) + b) >> (8 + k)
                                 # ACE (16.2): a normalized
                                 # first, the bias k bits lower
            if r != n or (r & 0x7F) != u:
                r = UNDEFINED    # the readings differ
        ymm1.bf8[i] = r
    elif zeroing:
        ymm1.bf8[i] = 0
    # else: ymm1.bf8[i] keeps its value
for i in range(VL // 16, 16):
    ymm1.b[i] = 0
# m16bcst: src3.f16[i] is the m16 for every i
# no flags; MXCSR unused, denormals kept
# documents: AVX10.2 rev. 7, ACE 1.16.2
# ymm1 is zeroed from bit 128 (xmm) or 256 (ymm) up

Operands

Flags

None: XED records no flag that the forms test or change.

Exceptions

The exception class is XED's name for the set of exception conditions that a form follows. How to read.

All forms

AVX512-E4NF

Attributes

What else XED records about the forms. How to read.

All forms

write masking; always rounds to nearest even

Memory forms

embedded broadcast (tuple Full)

Intrinsics

The C intrinsics of Clang 23.1.2 that compile to VCVTBIASPH2BF8, each with the form it compiles to and the Clang options it needs. The forms come from compiling every intrinsic and decoding the result with Intel XED; all intrinsics A–Z explains the method.

CPUID and processors

Each form belongs to an XED ISA set. A form can be used when the processor reports every CPUID bit of one of its ISA set's alternatives.

More AVX10.2 instructions

• VADDBF16
• VCMPBF16
• VCOMISBF16
• VCOMXSD
• VCOMXSH
• VCOMXSS
• VCVT2PH2BF8
• VCVT2PH2BF8S
• VCVT2PH2HF8
• VCVT2PH2HF8S
• VCVT2PS2PHX
• VCVTBF162IBS
• VCVTBF162IUBS
• VCVTBIASPH2BF8S
• VCVTBIASPH2HF8
• VCVTBIASPH2HF8S
• VCVTHF82PH
• VCVTPH2BF8
• VCVTPH2BF8S
• VCVTPH2HF8
• VCVTPH2HF8S
• VCVTPH2IBS
• VCVTPH2IUBS
• VCVTPS2IBS
• VCVTPS2IUBS
• VCVTTBF162IBS
• VCVTTBF162IUBS
• VCVTTPD2DQS
• VCVTTPD2QQS
• VCVTTPD2UDQS
• VCVTTPD2UQQS
• VCVTTPH2IBS
• VCVTTPH2IUBS
• VCVTTPS2DQS
• VCVTTPS2IBS
• VCVTTPS2IUBS
• VCVTTPS2QQS
• VCVTTPS2UDQS
• VCVTTPS2UQQS
• VCVTTSD2SIS

All 96 AVX10.2 instructions