VCVTBIASPS2BF8

VCVTBIASPS2BF8 converts singles to E5M2 as VCVTPS2BF8 does, into the low end of an xmm register under a write mask of one bit per byte, but it rounds with a bias that the program supplies instead of to nearest even. The singles come from the third operand, a register or memory, with a broadcast single allowed; the second operand, always a register of the same width, holds one bias doubleword for each single. It belongs to AVX10_V2_AUX of the ACE specification, opcode 39h with no prefix in EVEX map 5, W0. VCVTBIASPS2BF8S saturates, and VCVTBIASPH2BF8 is AVX10.2’s form for halves.

How the bias rounds#

E5M2 keeps 2 of the single’s 23 fraction bits, so 21 are cut away. The instruction first adds the low 21 bits of the bias doubleword to the single’s bits below the sign, so that a carry can reach the kept bits and even the exponent, and then cuts. Bits 31:21 of the bias are ignored. A bias of 0 truncates toward zero, 1FFFFFh rounds every inexact magnitude away from zero, and random bits round up with a probability equal to the part cut away: stochastic rounding, the use that ACE names. The single nearest 1.1 (3F8CCCCDh) lies 0.4 of the way from 1.0 to 1.25, so it gives 3Ch (1.0) with a bias of 0, 3Dh (1.25) with 1FFFFFh, and 3Dh for 40% of random biases.

The rest of the range#

Overflow gives infinity, as an infinite single does: 60,000 gives 57,344 (7Bh) with a bias of 0 but infinity with 1FFFFFh. A NaN gives a quiet NaN, a denormal single a zero. Results below 2^-14 are E5M2 denormals, cut at 2^-16, and ACE’s pseudocode still adds the bias below the last place of a normal result, so it covers only the lower part of the bits cut there and rounds up less often than their share. MXCSR is neither read nor written, and nothing is flagged.

In our builds#

GCC 14 and Clang 23 have no option for AVX10_V2_AUX: Clang 23’s assembler and GNU as 2.44 reject the mnemonic, and ACE’s intrinsic _mm512_cvtbiasps_bf8(a, b), with the singles in a and the biases in b, is missing from Clang 23’s headers. Zen 5, whose CPUID reports no AVX10, raised #UD in all seven recorded cases.

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
if not cpuid_leaf(0x24, 1)[2] >> 3 & 1:
    fault(UD)                    # no AVX10_V2_AUX
for i in range(VL // 32):
    if k1[i]:
        x = src3.f32[i]          # single i
        c = zmm2.d[i] & 0x1FFFFF # its bias: doubleword i's
                                 # bits 20:0
        s = x >> 24 & 0x80       # the sign, in place
        a = x & 0x7FFFFFFF       # |x|, as bits
        t = a + c                # the bias below E5M2's last
                                 # place at a's exponent
        e = t >> 23              # the sum's exponent
        if a > 0x7F800000:       # a NaN: quieted, the top bit
                                 # of its payload kept
            r = s | 0x7E | x >> 21 & 1
        elif a < 0x800000:       # zero or denormal: a zero
            r = s
        elif e > 112:            # 2^-14 and up: a normal
            r = (t >> 21) - 0x1C0
                                 # cut, the exponent rebiased
                                 # (127 to 15)
            r = s | min(r, 0x7C) # infinity, and past the
                                 # largest: infinity
        else:                    # below: an E5M2 denormal
            m = t & 0x7FFFFF | 0x800000
            r = s | m >> (134 - e)  # cut at 2^-16
        xmm1.bf8[i] = r
    elif zeroing:
        xmm1.bf8[i] = 0
    # else: xmm1.bf8[i] keeps its value
for i in range(VL // 32, 16):
    xmm1.b[i] = 0
# m32bcst: src3.f32[i] is the m32 for every i
# no flags; MXCSR unused, a denormal single a zero
# document: ACE 1.16.2
# xmm1 is zeroed from bit 128 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-E4

Attributes

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

All forms

write masking

Memory forms

embedded broadcast (tuple Full); memory fault suppression

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_V2_AUX instructions

• VCVTBF42HF8
• VCVTBF62HF8
• VCVTBF82BF4S
• VCVTBF82BF6S
• VCVTBF82PS
• VCVTBIASPS2BF8S
• VCVTBIASPS2HF8
• VCVTBIASPS2HF8S
• VCVTHF62HF8
• VCVTHF82BF4S
• VCVTHF82HF6S
• VCVTHF82PS
• VCVTPS2BF8
• VCVTPS2BF8S
• VCVTPS2HF8
• VCVTPS2HF8S
• VCVTROPS2HF8
• VCVTROPS2HF8S
• VPMOVSSDB
• VUNPACKB