[RFC][Clang] Elementwise builtins for converting encoded floating point values

Clang Frontend
,
1

Summary

Target-independent Clang builtins that convert integer-encoded arbitrary floating point (currently only FP8) values to native floating-point types. They lower to llvm.convert.from.arbitrary.fp. No IR change is proposed.

Each spelling names the source encoding and the destination type, and takes one argument:

__builtin_elementwise_convert_from_<source_format>_<destination_type>(bits)

The operand is a scalar integer or a fixed-length integer vector. Conversion is elementwise and keeps the element count:

float  scalar = __builtin_elementwise_convert_from_f8e5m2_f32(bits);
float4 packed = __builtin_elementwise_convert_from_f8e5m2_f32(bits4);

Three source encodings (f8e5m2, f8e4m3fn, f8e5m3fnu) and three destinations (f16 → _Float16, bf16 → __bf16, f32 → float) give nine builtins. Encoding the destination in the name keeps these ordinary CallExprs.

Motivation

C and C++ have no portable arithmetic types for these encodings, so they are stored as integers and converted to _Float16, __bf16, or float when needed.

HIP’s FP8 headers call AMDGPU builtins that match the instruction: an int holding four FP8 values, a byte index, and a second builtin for a pair.

float  one = __builtin_amdgcn_cvt_f32_fp8((int)word, 0);
float2 two = __builtin_amdgcn_cvt_pk_f32_fp8((int)word, false);

The proposed builtin takes 8-bit elements. The same spelling works for one value or two:

float  one = __builtin_elementwise_convert_from_f8e4m3fn_f32(bits);  // unsigned char
float2 two = __builtin_elementwise_convert_from_f8e4m3fn_f32(bits2); // unsigned char2

Packed-word and lane-select APIs remain target-specific. No SPIR-V or universal codegen is claimed.

Existing LLVM support

%result = call float @llvm.convert.from.arbitrary.fp.f32.i8(
    i8 %bits, metadata !"Float8E4M3FN")

The integer holds the encoding. Metadata names the interpretation. The intrinsic is overloaded on both types and has scalar and vector forms.

Source suffix Interpretation Width Proposal
f8e5m2 Float8E5M2 8 Exposed
f8e4m3fn Float8E4M3FN 8 Exposed
f8e5m3fnu Float8E5M3FNU 8 Exposed
f8e5m2fnuz Float8E5M2FNUZ 8 Deferred
f8e4m3fnuz Float8E4M3FNUZ 8 Deferred
f6e3m2fn Float6E3M2FN 6 Deferred
f6e2m3fn Float6E2M3FN 6 Deferred
f4e2m1fn Float4E2M1FN 4 Deferred

FNUZ is deferred because llvm.convert.from.arbitrary.fp does not lower it yet. The encodings are 8-bit and can be added later. Sub-byte encodings are deferred because Clang cannot spell their vector element types: there is no 6-bit element type, and <8 x _BitInt(4)> has sizeof 8 while <8 x i4> occupies 4 bytes.

The inverse intrinsic llvm.convert.to.arbitrary.fp also takes a rounding mode and a saturation flag. No builtin is proposed for it here. A narrowing family would still need one spelling per encoding, with rounding and saturation as ordinary constant arguments, so the naming question below is for the widening direction.

Scope

Widening conversions from the three exposed encodings to _Float16, __bf16, and float, for scalar and fixed-length vector operands, in C, C++, OpenCL, CUDA, and HIP.

Out of scope: IR changes, FNUZ, sub-byte encodings, f64, narrowing, rounding controls, constant evaluation, scalable vectors, and target-specific vector kinds.

Naming

The source suffix comes first so the name reads as convert-from-X-to-Y. LLVM IR mangles the destination first (llvm.convert.from.arbitrary.fp.f32.i8).

elementwise means a scalar or fixed-length vector, applied per element. Suffixes mean the same in every language mode. f16 is _Float16, not __fp16 or OpenCL half.

Putting both types in the name keeps an ordinary call. A destination type argument is not a call argument, so it needs a keyword, a parser production, and a new expression node (dependence, printing, profiling, TreeTransform, ASTImporter, serialization, AST consumers).

The suffix form costs one spelling per destination: 9 today.

Operand and result types

One argument. The source element type is an integer of exactly the width of the source suffix, not _Bool/bool or an enumeration. Wider containers are rejected even if the low bits hold the encoding. Signedness is ignored; the integer is a bit container. char, signed char, unsigned char, and _BitInt(8) are all fine where they are 8 bits.

Integer promotions and the usual arithmetic conversions do not apply, so the operand matches the intrinsic width. The downside is that C bit-manipulation expressions have type int:

unsigned char b;
__builtin_elementwise_convert_from_f8e5m2_f32(b >> 1);                  // error: 'int' is not 8 bits
__builtin_elementwise_convert_from_f8e5m2_f32((unsigned char)(b >> 1)); // ok

Accepting wider containers later would be source-compatible. Accepting a packed i32 and using the low 8 bits would bring back the instruction-shaped API.

A vector result has the same element count, the destination element type, and the same vector kind (GNU vector_size or Clang/OpenCL ext_vector_type). Scalable, sizeless, matrix, and target-specific fixed kinds are rejected.

Sema applies the usual target and language availability rules to the result type, including offload diagnostics. Dependent C++ calls are checked at instantiation.

Conversion semantics

Each element is interpreted as the named format and converted independently. The call lowers to llvm.convert.from.arbitrary.fp.

  • Every defined integer bit pattern produces a defined result. Poison and undef propagate as usual.
  • Finite values follow the intrinsic and are exact, except Float8E5M3FNU → _Float16. That format has maxExponent 16 vs 15 for _Float16, so the seven largest finite encodings (65536 through 114688) become infinity. bf16 and f32 are exact for every encoding.
  • f8e4m3fn and f8e5m3fnu have no infinity encoding. f8e5m3fnu has no sign bit.
  • A NaN encoding produces a NaN. LangRef says quiet/signaling is preserved and the payload is only truncated or extended. The builtin does not provide that: hardware conversions often canonicalize NaNs, so sign, quiet vs signaling, and payload are unspecified.
  • No dynamic rounding mode, no floating-environment side effects. The call may be speculated.

The vector form does not require a particular instruction or packing.

Builtin support queries

__has_builtin means Clang knows the spelling. That is how an encoding is queried; encodings have no other preprocessor name. It does not mean the destination type exists, and it does not mean the backend has a dedicated lowering:

Constant expressions

Not in the first patch. __has_constexpr_builtin returns 0, and a constant-expression context is diagnosed, including static-storage initializers.

Alternative: destination as a type argument

float  value  = __builtin_elementwise_convert_from_f8e4m3fn(encoded, float);
float4 packed = __builtin_elementwise_convert_from_f8e4m3fn(encoded4, float4);

The encoding stays in the name, so __has_builtin still tests the part that has no other query. The destination is a type argument: three names instead of nine, one per new encoding, typedefs and OpenCL half work directly, and f64 is a type-availability check rather than a new spelling.

Future directions

Narrowing, stochastic rounding, and scaled MX conversions are out of scope. They need IR changes: rounding on llvm.convert.to.arbitrary.fp is metadata and cannot hold a runtime seed or scale, and llvm.convert.from.arbitrary.fp has no scale operand. A later scaled family can add an argument or a second set of names. That is separate from dest-in-the-name here.

Target-specific notes

AArch64. __mfp8 is an extra scalar source: an opaque 8-bit container with no interpretation of its own. Neon vectors of __mfp8, and other Neon vector kinds, are rejected. Diagnostics for a bad operand should not mention __mfp8 on targets that lack the type. Widening a Neon vector while keeping the Neon kind can produce an invalid ABI type, which is why target-specific vector kinds are rejected in general.

AMDGPU. Target builtins put the destination first (__builtin_amdgcn_cvt_f32_fp8). Packed-word and lane-select forms stay on those builtins.

RISC-V / SVE. RVV and SVE sizeless vectors are rejected with the other unsupported vector kinds.

Open questions

  1. Destination in the name vs a type argument? Suffix form: 9 spellings, no typedef or OpenCL half destinations, __has_builtin silent on destination availability. Type-argument form: prototyped above. Neither needs a new mangling.
  2. Constant evaluation in the first patch? Proposed: no.

References

AI Disclosure: AI was used in writing up this post and for the reference implementation. It has gone through several manual review iterations in the process.

2

CCing folks that engaged on the initial PR:

@shiltian @MrSidims @arsenm @AaronBallman @jcranmer

3

As I’ve mentioned before, I think the most practical option is to have a family of builtins which includes both the source and destination element types in the builtin name. This allows writing a __has_builtin check. If we had a builtin that mapped directly to the intrinsic with an arbitrary format string or enum, we would have to invent a new test mechanism for compiler support for a particular format.

For me the main bike-shedding issue is how to spell the “f32” component.

4

How is this any different from __builtin_convertvector and using an one element vector instead?

5

Hello, thanks for engaging. The two builtins compute different things.

__builtin_convertvector performs a value-preserving numeric conversion between element types. Whereas, the proposed builtin reinterprets a bit pattern under a named floating-point encoding. On the same input byte they disagree by construction:

// 0x38 is Float8E4M3FN's encoding of 1.0 (sign 0, exp 0111 = bias, mantissa 000)
__builtin_convertvector((uchar1){0x38}, float1)
// 56.0f  -- uitofp i8 56 to float
__builtin_elementwise_convert_from_f8e4m3fn_f32(0x38) 
// 1.0f   -- reads the bits as E4M3FN

convertvector emits uitofp and there is no spelling of it that reads the operand as an FP8 encoding, because the element type it converts from is an 8-bit integer and C has no FP8 arithmetic type for it to convert from instead.

For a bit of AMDGPU context, the two select quite differently.

For instance, on a <4 x i8> operand the proposed builtin lowers to llvm.convert.from.arbitrary.fp and selects two v_cvt_pk_f32_fp8_e32 on targets that support it.

On the other hand __builtin_convertvector on the same operand selects integer-to-float conversions.

This is a different computation rather than a worse encoding of the same one. The existing way to get the packed instruction is the target-specific __builtin_amdgcn_cvt_pk_f32_fp8, which takes an instruction-shaped 32-bit word rather than a vector of elements. The motivation is to expose a more unified interface to HIP library consumers.

Note that if Clang had a first-class _Float8E4M3FN type, __builtin_convertvector would probably be the right tool and this proposal would be unnecessary.

6

Hi Matt. There’s a few ways we can go about this:

  1. f16 / bf16 / f32, as proposed. These are the exact tokens LLVM already uses for intrinsic mangling (llvm.convert.from.arbitrary.fp.f32.i8). They keep both halves of the name encoding-flavored which matches the f8e5m2 source token.

  2. _Float16 / __bf16 / float, the C type spellings. These read more “self-documenting”, but produces an awkward ..._f8e5m2___bf16 and mixes casing conventions.

  3. half / bfloat / float, the IR type names. But I feel like here we run into the ambiguity due to the distinction between HalfTy vs Float16Ty and such.