VCVTBIASPS2HF8

VCVTBIASPS2HF8 converts singles to E4M3 with a rounding bias, as VCVTBIASPS2BF8 does for E5M2. The singles come from the third operand, a register, memory or a broadcast single, and the biases from the second, one doubleword per single; the bytes go to the low end of an xmm register under a write mask of one bit per byte. E4M3 keeps 3 of the single’s 23 fraction bits, so the low 20 bits of each bias are added below its last place before the cut: 0 truncates, FFFFFh rounds every inexact magnitude up, and random bits give stochastic rounding. The single nearest 1.1 gives 38h (1.0) with a bias of 0 and 39h (1.125) with FFFFFh. It belongs to AVX10_V2_AUX of the ACE specification, opcode 38h with no prefix in EVEX map 5, W0. VCVTBIASPS2HF8S saturates.

Large and small results#

E4M3 has no infinity, so a sum that reaches 480, the value 7Fh would otherwise hold, gives the NaN; 448 and up can get there with a large enough bias. Infinities and NaNs give the NaN as well.

Below 2^-6, where E4M3 results would be denormals, ACE contradicts itself. Its description of these conversions and their table say that results smaller than the smallest denormal become zeros and other results are kept, while its pseudocode helper returns a zero for every result below 2^-6. The Operation section gives UNDEFINED there rather than choose. A denormal single counts as a zero either way. 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_hf8(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] & 0xFFFFF  # its bias: doubleword i's
                                 # bits 19:0
        s = x >> 24 & 0x80       # the sign, in place
        a = x & 0x7FFFFFFF       # |x|, as bits
        t = a + c                # the bias below E4M3's last
                                 # place at a's exponent
        e = t >> 23              # the sum's exponent
        if a >= 0x7F800000:      # infinity or NaN: the NaN
            r = s | 0x7F
        elif a < 0x800000:       # zero or denormal: a zero
            r = s
        elif e > 120:            # 2^-6 and up: a normal
            r = (t >> 20) - 0x3C0
                                 # cut, the exponent rebiased
                                 # (127 to 7)
            r = s | min(r, 0x7F) # 480 and up: the NaN
        else:                    # below 2^-6: ACE's helper
                                 # (16.1) gives a zero, its text
                                 # (9.2.1 and its table) keeps
                                 # denormal results
            r = UNDEFINED
        xmm1.hf8[i] = r
    elif zeroing:
        xmm1.hf8[i] = 0
    # else: xmm1.hf8[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
• VCVTBIASPS2BF8
• VCVTBIASPS2BF8S
• VCVTBIASPS2HF8S
• VCVTHF62HF8
• VCVTHF82BF4S
• VCVTHF82HF6S
• VCVTHF82PS
• VCVTPS2BF8
• VCVTPS2BF8S
• VCVTPS2HF8
• VCVTPS2HF8S
• VCVTROPS2HF8
• VCVTROPS2HF8S
• VPMOVSSDB
• VUNPACKB