https://github.com/chinmaydd updated https://github.com/llvm/llvm-project/pull/212647
>From dc8ea004d83533e486458898989eeb21843261fa Mon Sep 17 00:00:00 2001 From: Chinmay Deshpande <[email protected]> Date: Sun, 20 Sep 2026 18:40:37 -0700 Subject: [PATCH] [Clang] Add elementwise conversions from encoded FP8 values Add nine builtins converting Float8E5M2, Float8E4M3FN, and Float8E5M3FNU encodings to _Float16, __bf16, or float through llvm.convert.from.arbitrary.fp. Accept exactly 8-bit integer scalars and generic fixed-length vectors, preserving vector kinds and element counts without integer promotions. Also support scalar AArch64 __mfp8 containers. Apply destination type availability checks, including deferred offload diagnostics, and reject constant-expression use. Document the interface and add semantic, template, language-mode, target-specific, and IR-generation regression coverage. RFC: https://discourse.llvm.org/t/rfc-clang-elementwise-builtins-for-converting-encoded-floating-point-values/91644 --- clang/docs/LanguageExtensions.md | 85 +++++++ clang/docs/ReleaseNotes.md | 7 + clang/include/clang/Basic/Builtins.td | 54 ++++ .../clang/Basic/DiagnosticSemaKinds.td | 7 + clang/lib/CodeGen/CGBuiltin.cpp | 33 +++ clang/lib/Sema/SemaChecking.cpp | 87 +++++++ .../builtins-elementwise-convert-from-fp8.c | 27 ++ .../builtins-elementwise-convert-from-fp8.c | 236 ++++++++++++++++++ ...ins-elementwise-convert-from-fp8-aarch64.c | 54 ++++ ...ns-elementwise-convert-from-fp8-overflow.c | 49 ++++ ...ltins-elementwise-convert-from-fp8-riscv.c | 9 + ...tins-elementwise-convert-from-fp8-target.c | 35 +++ .../builtins-elementwise-convert-from-fp8.c | 166 ++++++++++++ .../builtins-elementwise-convert-from-fp8.cu | 78 ++++++ .../builtins-elementwise-convert-from-fp8.cpp | 123 +++++++++ .../builtins-elementwise-convert-from-fp8.cl | 46 ++++ 16 files changed, 1096 insertions(+) create mode 100644 clang/test/CodeGen/AArch64/builtins-elementwise-convert-from-fp8.c create mode 100644 clang/test/CodeGen/builtins-elementwise-convert-from-fp8.c create mode 100644 clang/test/Sema/builtins-elementwise-convert-from-fp8-aarch64.c create mode 100644 clang/test/Sema/builtins-elementwise-convert-from-fp8-overflow.c create mode 100644 clang/test/Sema/builtins-elementwise-convert-from-fp8-riscv.c create mode 100644 clang/test/Sema/builtins-elementwise-convert-from-fp8-target.c create mode 100644 clang/test/Sema/builtins-elementwise-convert-from-fp8.c create mode 100644 clang/test/SemaCUDA/builtins-elementwise-convert-from-fp8.cu create mode 100644 clang/test/SemaCXX/builtins-elementwise-convert-from-fp8.cpp create mode 100644 clang/test/SemaOpenCL/builtins-elementwise-convert-from-fp8.cl diff --git a/clang/docs/LanguageExtensions.md b/clang/docs/LanguageExtensions.md index bd47b18da24812..0d989d8e0f4fc1 100644 --- a/clang/docs/LanguageExtensions.md +++ b/clang/docs/LanguageExtensions.md @@ -805,6 +805,91 @@ of different sizes and signs is forbidden in binary and ternary builtins. | T \_\_builtin_elementwise_pext(T x, T m) | extract bits from x selected by the mask m, pack them contiguously into the least significant bits of the result, and zero the rest. | integer types | | T \_\_builtin_elementwise_pdep(T x, T m) | deposit the least significant bits of x at the positions where m has a 1-bit, and zero the rest. | integer types | +(langext-elementwise-encoded-fp-conversions)= + +*Conversions from encoded floating-point values* + +The `__builtin_elementwise_convert_from_<source_format>_<destination_type>` +builtins interpret integer bit patterns as floating-point encodings and convert +them to native floating-point values. They take one scalar or fixed-length +vector argument and are available in C, C++, OpenCL, CUDA, and HIP. + +Each builtin name selects one source encoding and one destination element type: + +| Source suffix | Encoding | Element width | +| ------------- | -------- | ------------- | +| `f8e5m2` | `Float8E5M2` | 8 bits | +| `f8e4m3fn` | `Float8E4M3FN` | 8 bits | +| `f8e5m3fnu` | `Float8E5M3FNU` | 8 bits | + +| Destination suffix | Element type | +| ------------------ | ------------ | +| `f16` | `_Float16` | +| `bf16` | `__bf16` | +| `f32` | `float` | + +All nine combinations are supported. The destination suffix has the same meaning +in every language mode: `f16` denotes `_Float16`, distinct from `__fp16` and +OpenCL `half`. The destination type must be supported by the target and language +mode, including the usual checks for offload code. + +The argument's element type must be an integer of exactly 8 bits, excluding +`bool` and enumerations. Signedness does not affect the interpretation. +`char`, `signed char`, `unsigned char`, and `_BitInt(8)` are accepted when they +have the required width. On AArch64, a scalar `__mfp8` argument is also accepted +as an uninterpreted 8-bit container. Integer promotions and the usual arithmetic +conversions do not apply to the argument. + +For vector arguments, the result has the same number of elements and the same +vector kind as the input, with the selected destination element type. GNU +`vector_size` and Clang/OpenCL `ext_vector_type` vectors are supported. Scalable, +sizeless, matrix, and target-specific vector types, including Neon vectors of +`__mfp8`, are not supported. + +```c +typedef unsigned char uchar4 __attribute__((ext_vector_type(4))); +typedef float float4 __attribute__((ext_vector_type(4))); + +float convert_scalar(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(bits); +} + +float4 convert_vector(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(bits); +} + +void convert_shifted(unsigned char bits) { + // The shift promotes bits to int, which is not an 8-bit element type. + __builtin_elementwise_convert_from_f8e5m2_f32(bits >> 1); // error + __builtin_elementwise_convert_from_f8e5m2_f32((unsigned char)(bits >> 1)); // OK +} +``` + +These builtins interpret the encoding rather than converting an integer's +numeric value. For example, converting `(unsigned char)0x38` from `f8e4m3fn` to +`f32` produces `1.0f`. + +Every defined input bit pattern produces a defined result. Finite values convert +exactly, except that the seven largest finite `Float8E5M3FNU` values, from 65536 +through 114688, convert to infinity in `_Float16`. All finite inputs are exactly +representable in `__bf16` and `float`. `Float8E4M3FN` and `Float8E5M3FNU` have no +infinity encoding, and `Float8E5M3FNU` has no sign bit. A NaN input produces a NaN, +but its sign, quiet or signaling status, and payload are unspecified. The +conversion does not depend on the dynamic rounding mode and has no +floating-point environment side effects. + +The builtins lower to the LLVM +[`llvm.convert.from.arbitrary.fp`](https://llvm.org/docs/LangRef.html#llvm-convert-from-arbitrary-fp-intrinsic) +intrinsic. They do not require a particular instruction or packing, and backend +support for lowering the intrinsic varies by target. + +Use `__has_builtin` with a complete spelling, such as +`__has_builtin(__builtin_elementwise_convert_from_f8e5m2_f32)`, to query whether +Clang recognizes the builtin. This does not imply that the destination type is +available or that the backend can lower the operation. These builtins cannot be +used in constant expressions, including C static-storage initializers, and +`__has_constexpr_builtin` returns 0 for them. + *Reduction Builtins* Each builtin returns a scalar equivalent to applying the specified diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index e5da258b9950a3..540d333cb3d900 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -247,6 +247,13 @@ features cannot lower the translation-unit ABI level; - Added support for the `__builtin_strlcat` and `__builtin_strlcpy` builtins. +- Added the + [`__builtin_elementwise_convert_from_<source_format>_<destination_type>`](langext-elementwise-encoded-fp-conversions) + builtins to convert 8-bit integer encodings of `Float8E5M2`, `Float8E4M3FN`, and + `Float8E5M3FNU` values to `_Float16`, `__bf16`, or `float`. They support scalar + and fixed-length GNU and Clang/OpenCL vector arguments in C, C++, OpenCL, CUDA, + and HIP. + ### New Compiler Flags - New option `-fdefined-pointer-subtraction` added to preserve stable semantics diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td index 80d09c089ae952..d352e2f372d0e4 100644 --- a/clang/include/clang/Basic/Builtins.td +++ b/clang/include/clang/Basic/Builtins.td @@ -1583,6 +1583,60 @@ def ElementwiseBitreverse : Builtin { let Prototype = "void(...)"; } +def ElementwiseConvertFromF8E5M2F16 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e5m2_f16"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E5M2BF16 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e5m2_bf16"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E5M2F32 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e5m2_f32"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E4M3FNF16 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e4m3fn_f16"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E4M3FNBF16 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e4m3fn_bf16"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E4M3FNF32 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e4m3fn_f32"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E5M3FNUF16 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e5m3fnu_f16"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E5M3FNUBF16 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e5m3fnu_bf16"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + +def ElementwiseConvertFromF8E5M3FNUF32 : Builtin { + let Spellings = ["__builtin_elementwise_convert_from_f8e5m3fnu_f32"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + def ElementwiseMax : Builtin { let Spellings = ["__builtin_elementwise_max"]; let Attributes = [NoThrow, Const, CustomTypeChecking, Constexpr]; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 36a18473f4d4cc..1a271f33e62553 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13417,6 +13417,13 @@ def err_builtin_invalid_arg_type: Error< def err_builtin_requires_double_type: Error< "%ordinal0 argument must be a scalar, vector, or matrix of double type (was %1)">; +def err_builtin_elementwise_convert_from_fp8_arg_type : Error< + "argument to %0 must be an 8-bit integer or a vector of 8-bit integers" + "%select{|, or '__mfp8'}1 (was %2)">; +def err_builtin_elementwise_convert_from_fp8_vector_type : Error< + "argument to %0 must be a scalar or a fixed-length vector declared with " + "'vector_size' or 'ext_vector_type' (was %1)">; + def err_bswapg_invalid_bit_width : Error< "_BitInt type %0 (%1 bits) must be a multiple of 16 bits for byte swapping">; diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index f5001625327777..586b09b351ab8c 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -615,6 +615,24 @@ static Value *EmitISOVolatileStore(CodeGenFunction &CGF, const CallExpr *E) { return Store; } +static Value *emitConvertFromArbitraryFPBuiltin(CodeGenFunction &CGF, + const CallExpr *E, + StringRef Format) { + Value *Bits = CGF.EmitScalarExpr(E->getArg(0)); + // The opaque scalar __mfp8 container is represented as <1 x i8> in IR. + if (E->getArg(0)->getType()->isMFloat8Type()) + Bits = CGF.Builder.CreateBitCast(Bits, CGF.Builder.getInt8Ty()); + + Function *F = + CGF.CGM.getIntrinsic(Intrinsic::convert_from_arbitrary_fp, + {CGF.ConvertType(E->getType()), Bits->getType()}); + Value *Interpretation = MetadataAsValue::get( + CGF.getLLVMContext(), MDString::get(CGF.getLLVMContext(), Format)); + // These conversions have no floating-point environment side effects, even + // when the enclosing function uses strict floating-point semantics. + return CGF.Builder.CreateCall(F, {Bits, Interpretation}); +} + // Emit a simple mangled intrinsic that has 1 argument and a return type // matching the argument type. Depending on mode, this may be a constrained // floating-point intrinsic. @@ -4364,6 +4382,21 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, case Builtin::BI__builtin_elementwise_bitreverse: return RValue::get(emitBuiltinWithOneOverloadedType<1>( *this, E, Intrinsic::bitreverse, "elt.bitreverse")); + case Builtin::BI__builtin_elementwise_convert_from_f8e5m2_f16: + case Builtin::BI__builtin_elementwise_convert_from_f8e5m2_bf16: + case Builtin::BI__builtin_elementwise_convert_from_f8e5m2_f32: + return RValue::get( + emitConvertFromArbitraryFPBuiltin(*this, E, "Float8E5M2")); + case Builtin::BI__builtin_elementwise_convert_from_f8e4m3fn_f16: + case Builtin::BI__builtin_elementwise_convert_from_f8e4m3fn_bf16: + case Builtin::BI__builtin_elementwise_convert_from_f8e4m3fn_f32: + return RValue::get( + emitConvertFromArbitraryFPBuiltin(*this, E, "Float8E4M3FN")); + case Builtin::BI__builtin_elementwise_convert_from_f8e5m3fnu_f16: + case Builtin::BI__builtin_elementwise_convert_from_f8e5m3fnu_bf16: + case Builtin::BI__builtin_elementwise_convert_from_f8e5m3fnu_f32: + return RValue::get( + emitConvertFromArbitraryFPBuiltin(*this, E, "Float8E5M3FNU")); case Builtin::BI__builtin_elementwise_popcount: return RValue::get(emitBuiltinWithOneOverloadedType<1>( *this, E, Intrinsic::ctpop, "elt.ctpop")); diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index ea8f3babcb0e01..442651b469f19e 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -3050,6 +3050,74 @@ static QualType getVectorElementType(ASTContext &Context, QualType VecTy) { return QualType(); } +static bool BuiltinElementwiseConvertFromFP8(Sema &S, CallExpr *TheCall, + QualType ResultEltTy) { + if (S.checkArgCount(TheCall, 1)) + return true; + + // The integer is a bit container. In particular, do not promote narrow + // integers or apply user-defined conversions to the operand. + ExprResult Arg = S.DefaultFunctionArrayLvalueConversion(TheCall->getArg(0)); + if (Arg.isInvalid()) + return true; + + ASTContext &Context = S.Context; + QualType ArgTy = Arg.get()->getType(); + const auto *VecTy = ArgTy->getAs<VectorType>(); + if (ArgTy->isSizelessType() || + (VecTy && VecTy->getVectorKind() != VectorKind::Generic)) { + S.Diag(Arg.get()->getBeginLoc(), + diag::err_builtin_elementwise_convert_from_fp8_vector_type) + << Context.BuiltinInfo.getQuotedName(TheCall->getBuiltinCallee()) + << ArgTy << Arg.get()->getSourceRange(); + return true; + } + + QualType EltTy = VecTy ? VecTy->getElementType() : ArgTy; + if (const auto *OBT = EltTy->getAs<OverflowBehaviorType>()) + EltTy = OBT->getUnderlyingType(); + bool IsInteger = EltTy->isIntegerType() && !EltTy->isBooleanType() && + !EltTy->isEnumeralType() && Context.getIntWidth(EltTy) == 8; + if (!IsInteger && !ArgTy->isMFloat8Type()) { + S.Diag(Arg.get()->getBeginLoc(), + diag::err_builtin_elementwise_convert_from_fp8_arg_type) + << Context.BuiltinInfo.getQuotedName(TheCall->getBuiltinCallee()) + << !Context.MFloat8Ty.isNull() << ArgTy << Arg.get()->getSourceRange(); + return true; + } + + // These result types are implicit, so apply the availability checks that + // would normally be performed when spelling them in a declaration. + const TargetInfo &TI = Context.getTargetInfo(); + const LangOptions &LangOpts = S.getLangOpts(); + bool IsOpenMPDevice = LangOpts.OpenMP && LangOpts.OpenMPIsTargetDevice; + if (ResultEltTy->isFloat16Type() && !TI.hasFloat16Type() && !LangOpts.CUDA && + !IsOpenMPDevice) { + S.Diag(TheCall->getBeginLoc(), diag::err_type_unsupported) << "_Float16"; + return true; + } + if (ResultEltTy->isBFloat16Type() && !TI.hasBFloat16Type() && + !IsOpenMPDevice && !LangOpts.SYCLIsDevice) { + S.Diag(TheCall->getBeginLoc(), diag::err_type_unsupported) << "__bf16"; + return true; + } + // Check the element type explicitly: checkTypeSupport does not look through + // vector types. It also handles deferred diagnostics in offloading modes. + S.checkTypeSupport(ResultEltTy, TheCall->getBeginLoc()); + + QualType ResultTy = ResultEltTy; + if (VecTy) { + if (isa<ExtVectorType>(VecTy)) + ResultTy = Context.getExtVectorType(ResultEltTy, VecTy->getNumElements()); + else + ResultTy = Context.getVectorType(ResultEltTy, VecTy->getNumElements(), + VecTy->getVectorKind()); + } + TheCall->setArg(0, Arg.get()); + TheCall->setType(ResultTy); + return false; +} + ExprResult Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID, CallExpr *TheCall) { @@ -3759,6 +3827,25 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID, break; } + case Builtin::BI__builtin_elementwise_convert_from_f8e5m2_f16: + case Builtin::BI__builtin_elementwise_convert_from_f8e4m3fn_f16: + case Builtin::BI__builtin_elementwise_convert_from_f8e5m3fnu_f16: + if (BuiltinElementwiseConvertFromFP8(*this, TheCall, Context.Float16Ty)) + return ExprError(); + break; + case Builtin::BI__builtin_elementwise_convert_from_f8e5m2_bf16: + case Builtin::BI__builtin_elementwise_convert_from_f8e4m3fn_bf16: + case Builtin::BI__builtin_elementwise_convert_from_f8e5m3fnu_bf16: + if (BuiltinElementwiseConvertFromFP8(*this, TheCall, Context.BFloat16Ty)) + return ExprError(); + break; + case Builtin::BI__builtin_elementwise_convert_from_f8e5m2_f32: + case Builtin::BI__builtin_elementwise_convert_from_f8e4m3fn_f32: + case Builtin::BI__builtin_elementwise_convert_from_f8e5m3fnu_f32: + if (BuiltinElementwiseConvertFromFP8(*this, TheCall, Context.FloatTy)) + return ExprError(); + break; + // __builtin_elementwise_abs restricts the element type to signed integers or // floating point types only. case Builtin::BI__builtin_elementwise_abs: diff --git a/clang/test/CodeGen/AArch64/builtins-elementwise-convert-from-fp8.c b/clang/test/CodeGen/AArch64/builtins-elementwise-convert-from-fp8.c new file mode 100644 index 00000000000000..8dccb96beec919 --- /dev/null +++ b/clang/test/CodeGen/AArch64/builtins-elementwise-convert-from-fp8.c @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 -triple aarch64-unknown-linux-gnu -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s +// REQUIRES: aarch64-registered-target + +// __mfp8 is a scalar source even though its IR representation is <1 x i8>. +// CHECK-LABEL: define {{.*}}@mfp8_f16( +// CHECK: [[BITS:%.*]] = load <1 x i8>, ptr +// CHECK-NEXT: [[SCALAR:%.*]] = bitcast <1 x i8> [[BITS]] to i8 +// CHECK-NEXT: call half @llvm.convert.from.arbitrary.fp.f16.i8(i8 [[SCALAR]], metadata !"Float8E5M2") +_Float16 mfp8_f16(__mfp8 bits) { + return __builtin_elementwise_convert_from_f8e5m2_f16(bits); +} + +// CHECK-LABEL: define {{.*}}@mfp8_bf16( +// CHECK: [[BITS:%.*]] = load <1 x i8>, ptr +// CHECK-NEXT: [[SCALAR:%.*]] = bitcast <1 x i8> [[BITS]] to i8 +// CHECK-NEXT: call bfloat @llvm.convert.from.arbitrary.fp.bf16.i8(i8 [[SCALAR]], metadata !"Float8E4M3FN") +__bf16 mfp8_bf16(__mfp8 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@mfp8_f32( +// CHECK: [[BITS:%.*]] = load <1 x i8>, ptr +// CHECK-NEXT: [[SCALAR:%.*]] = bitcast <1 x i8> [[BITS]] to i8 +// CHECK-NEXT: call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 [[SCALAR]], metadata !"Float8E5M3FNU") +float mfp8_f32(__mfp8 bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} diff --git a/clang/test/CodeGen/builtins-elementwise-convert-from-fp8.c b/clang/test/CodeGen/builtins-elementwise-convert-from-fp8.c new file mode 100644 index 00000000000000..a9f5cb33abf35c --- /dev/null +++ b/clang/test/CodeGen/builtins-elementwise-convert-from-fp8.c @@ -0,0 +1,236 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -disable-llvm-passes -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -disable-llvm-passes -frounding-math -ffp-exception-behavior=strict -DTEST_STRICT_FP -o - %s | FileCheck %s --check-prefixes=CHECK,STRICT + +typedef unsigned char uchar4 __attribute__((ext_vector_type(4))); +typedef _Float16 half4 __attribute__((ext_vector_type(4))); +typedef __bf16 bfloat4 __attribute__((ext_vector_type(4))); +typedef float float4 __attribute__((ext_vector_type(4))); + +// CHECK-LABEL: define {{.*}}@scalar_f8e5m2_f16( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call half @llvm.convert.from.arbitrary.fp.f16.i8(i8 [[BITS]], metadata !"Float8E5M2") +_Float16 scalar_f8e5m2_f16(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m2_f16(bits); +} + +// CHECK: declare half @llvm.convert.from.arbitrary.fp.f16.i8(i8, metadata) #[[FP8_ATTRS:[0-9]+]] + +// CHECK-LABEL: define {{.*}}@scalar_f8e5m2_bf16( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call bfloat @llvm.convert.from.arbitrary.fp.bf16.i8(i8 [[BITS]], metadata !"Float8E5M2") +__bf16 scalar_f8e5m2_bf16(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m2_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@scalar_f8e5m2_f32( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 [[BITS]], metadata !"Float8E5M2") +float scalar_f8e5m2_f32(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m2_f32(bits); +} + +// CHECK-LABEL: define {{.*}}@scalar_f8e4m3fn_f16( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call half @llvm.convert.from.arbitrary.fp.f16.i8(i8 [[BITS]], metadata !"Float8E4M3FN") +_Float16 scalar_f8e4m3fn_f16(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f16(bits); +} + +// CHECK-LABEL: define {{.*}}@scalar_f8e4m3fn_bf16( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call bfloat @llvm.convert.from.arbitrary.fp.bf16.i8(i8 [[BITS]], metadata !"Float8E4M3FN") +__bf16 scalar_f8e4m3fn_bf16(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@scalar_f8e4m3fn_f32( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 [[BITS]], metadata !"Float8E4M3FN") +float scalar_f8e4m3fn_f32(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(bits); +} + +// CHECK-LABEL: define {{.*}}@scalar_f8e5m3fnu_f16( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call half @llvm.convert.from.arbitrary.fp.f16.i8(i8 [[BITS]], metadata !"Float8E5M3FNU") +_Float16 scalar_f8e5m3fnu_f16(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f16(bits); +} + +// CHECK-LABEL: define {{.*}}@scalar_f8e5m3fnu_bf16( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call bfloat @llvm.convert.from.arbitrary.fp.bf16.i8(i8 [[BITS]], metadata !"Float8E5M3FNU") +__bf16 scalar_f8e5m3fnu_bf16(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@scalar_f8e5m3fnu_f32( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 [[BITS]], metadata !"Float8E5M3FNU") +float scalar_f8e5m3fnu_f32(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e5m2_f16( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x half> @llvm.convert.from.arbitrary.fp.v4f16.v4i8(<4 x i8> [[BITS]], metadata !"Float8E5M2") +half4 vector_f8e5m2_f16(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e5m2_f16(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e5m2_bf16( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x bfloat> @llvm.convert.from.arbitrary.fp.v4bf16.v4i8(<4 x i8> [[BITS]], metadata !"Float8E5M2") +bfloat4 vector_f8e5m2_bf16(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e5m2_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e5m2_f32( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> [[BITS]], metadata !"Float8E5M2") +float4 vector_f8e5m2_f32(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e5m2_f32(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e4m3fn_f16( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x half> @llvm.convert.from.arbitrary.fp.v4f16.v4i8(<4 x i8> [[BITS]], metadata !"Float8E4M3FN") +half4 vector_f8e4m3fn_f16(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f16(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e4m3fn_bf16( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x bfloat> @llvm.convert.from.arbitrary.fp.v4bf16.v4i8(<4 x i8> [[BITS]], metadata !"Float8E4M3FN") +bfloat4 vector_f8e4m3fn_bf16(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e4m3fn_f32( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> [[BITS]], metadata !"Float8E4M3FN") +float4 vector_f8e4m3fn_f32(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e5m3fnu_f16( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x half> @llvm.convert.from.arbitrary.fp.v4f16.v4i8(<4 x i8> [[BITS]], metadata !"Float8E5M3FNU") +half4 vector_f8e5m3fnu_f16(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f16(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e5m3fnu_bf16( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x bfloat> @llvm.convert.from.arbitrary.fp.v4bf16.v4i8(<4 x i8> [[BITS]], metadata !"Float8E5M3FNU") +bfloat4 vector_f8e5m3fnu_bf16(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@vector_f8e5m3fnu_f32( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> [[BITS]], metadata !"Float8E5M3FNU") +float4 vector_f8e5m3fnu_f32(uchar4 bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} + +// Signedness does not change the interpretation of the bits. +// CHECK-LABEL: define {{.*}}@signed_bits( +// CHECK: call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 -68, metadata !"Float8E5M2") +float signed_bits(void) { + return __builtin_elementwise_convert_from_f8e5m2_f32((signed char)0xbc); +} + +// CHECK-LABEL: define {{.*}}@plain_char_bits( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 [[BITS]], metadata !"Float8E4M3FN") +float plain_char_bits(char bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(bits); +} + +// CHECK-LABEL: define {{.*}}@signed_bitint_bits( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 [[BITS]], metadata !"Float8E5M3FNU") +float signed_bitint_bits(_BitInt(8) bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} + +// CHECK-LABEL: define {{.*}}@unsigned_bitint_bits( +// CHECK: [[BITS:%.*]] = load i8, ptr +// CHECK-NEXT: call half @llvm.convert.from.arbitrary.fp.f16.i8(i8 [[BITS]], metadata !"Float8E4M3FN") +_Float16 unsigned_bitint_bits(unsigned _BitInt(8) bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f16(bits); +} + +typedef _BitInt(8) bitint4 __attribute__((ext_vector_type(4))); + +// CHECK-LABEL: define {{.*}}@signed_bitint_vector( +// CHECK: [[BITS:%.*]] = load <4 x i8>, ptr %bits.addr +// CHECK-NEXT: call <4 x bfloat> @llvm.convert.from.arbitrary.fp.v4bf16.v4i8(<4 x i8> [[BITS]], metadata !"Float8E5M3FNU") +bfloat4 signed_bitint_vector(bitint4 bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_bf16(bits); +} + +typedef signed char schar3 __attribute__((ext_vector_type(3))); +typedef float float3 __attribute__((ext_vector_type(3))); + +// CHECK-LABEL: define {{.*}}@three_lanes( +// CHECK: call <3 x float> @llvm.convert.from.arbitrary.fp.v3f32.v3i8(<3 x i8> {{.*}}, metadata !"Float8E4M3FN") +float3 three_lanes(schar3 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(bits); +} + +typedef unsigned char uchar1 __attribute__((ext_vector_type(1))); +typedef _Float16 half1 __attribute__((ext_vector_type(1))); + +// CHECK-LABEL: define {{.*}}@one_lane( +// CHECK: call <1 x half> @llvm.convert.from.arbitrary.fp.v1f16.v1i8(<1 x i8> {{.*}}, metadata !"Float8E5M2") +half1 one_lane(uchar1 bits) { + return __builtin_elementwise_convert_from_f8e5m2_f16(bits); +} + +typedef signed char gnu_char4 __attribute__((vector_size(4))); +typedef _Float16 gnu_half4 __attribute__((vector_size(8))); +typedef __bf16 gnu_bfloat4 __attribute__((vector_size(8))); +typedef float gnu_float4 __attribute__((vector_size(16))); + +// CHECK-LABEL: define {{.*}}@gnu_vector_f16( +// CHECK: call <4 x half> @llvm.convert.from.arbitrary.fp.v4f16.v4i8(<4 x i8> {{.*}}, metadata !"Float8E5M3FNU") +gnu_half4 gnu_vector_f16(gnu_char4 bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f16(bits); +} + +// CHECK-LABEL: define {{.*}}@gnu_vector_bf16( +// CHECK: call <4 x bfloat> @llvm.convert.from.arbitrary.fp.v4bf16.v4i8(<4 x i8> {{.*}}, metadata !"Float8E4M3FN") +gnu_bfloat4 gnu_vector_bf16(gnu_char4 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); +} + +// CHECK-LABEL: define {{.*}}@gnu_vector_f32( +// CHECK: call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> {{.*}}, metadata !"Float8E5M2") +gnu_float4 gnu_vector_f32(gnu_char4 bits) { + return __builtin_elementwise_convert_from_f8e5m2_f32(bits); +} + +unsigned char next_bits(void); + +// CHECK-LABEL: define {{.*}}@evaluate_once( +// CHECK: [[BITS:%.*]] = call {{.*}}i8 @next_bits() +// CHECK-NEXT: [[RESULT:%.*]] = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 [[BITS]], metadata !"Float8E5M2") +// CHECK-NEXT: ret float [[RESULT]] +float evaluate_once(void) { + return __builtin_elementwise_convert_from_f8e5m2_f32(next_bits()); +} + +#ifdef TEST_STRICT_FP +// The conversion ignores the floating-point environment while the addition +// respects the enclosing function's strict floating-point semantics. +// STRICT-LABEL: define {{.*}}@strict_convert( +// STRICT: [[CONVERTED:%.*]] = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 {{.*}}, metadata !"Float8E5M2") +// STRICT-NEXT: call float @llvm.experimental.constrained.fadd.f32(float [[CONVERTED]], float 1.000000e+00, metadata !"round.dynamic", metadata !"fpexcept.strict") +float strict_convert(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m2_f32(bits) + 1.0f; +} +#endif + +// CHECK: attributes #[[FP8_ATTRS]] = { {{.*}}nounwind{{.*}}speculatable{{.*}}memory(none){{.*}} } diff --git a/clang/test/Sema/builtins-elementwise-convert-from-fp8-aarch64.c b/clang/test/Sema/builtins-elementwise-convert-from-fp8-aarch64.c new file mode 100644 index 00000000000000..86424630fee896 --- /dev/null +++ b/clang/test/Sema/builtins-elementwise-convert-from-fp8-aarch64.c @@ -0,0 +1,54 @@ +// RUN: %clang_cc1 -triple aarch64-none-linux-gnu -target-feature +neon -target-feature +sve -target-feature -fp8 -mvscale-min=1 -mvscale-max=1 -fsyntax-only -verify %s + +// Scalar __mfp8 is an opaque bit container, available without the FP8 extension. +_Float16 mfp8_to_f16(__mfp8 bits) { + return __builtin_elementwise_convert_from_f8e5m2_f16(bits); +} + +__bf16 mfp8_to_bf16(__mfp8 bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); +} + +float mfp8_to_f32(__mfp8 bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} + +void invalid_scalar(int bits) { + (void)__builtin_elementwise_convert_from_f8e5m2_f32(bits); // expected-error {{must be an 8-bit integer or a vector of 8-bit integers, or '__mfp8'}} +} + +typedef signed char neon_int8x8 __attribute__((neon_vector_type(8))); +typedef unsigned char neon_uint8x16 __attribute__((neon_vector_type(16))); +typedef unsigned char neon_poly8x8 __attribute__((neon_polyvector_type(8))); +typedef __mfp8 neon_mfp8x8 __attribute__((neon_vector_type(8))); +typedef __mfp8 neon_mfp8x16 __attribute__((neon_vector_type(16))); + +// Widening a Neon vector while retaining its vector kind could create an +// invalid ABI type. All Neon kinds are rejected, including integer vectors. +void neon_vectors(neon_int8x8 a, neon_uint8x16 b, neon_poly8x8 c, + neon_mfp8x8 d, neon_mfp8x16 e) { + (void)__builtin_elementwise_convert_from_f8e5m2_f32(a); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m2_f16(b); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m2_bf16(c); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e4m3fn_f32(d); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m3fnu_f32(e); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} +} + +typedef __SVUint8_t fixed_sve_uint8 __attribute__((arm_sve_vector_bits(128))); +typedef __SVBool_t fixed_sve_bool __attribute__((arm_sve_vector_bits(128))); + +void sve_vectors(__SVUint8_t a, __SVBool_t b, fixed_sve_uint8 c, + fixed_sve_bool d) { + (void)__builtin_elementwise_convert_from_f8e5m2_f32(a); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m2_f32(b); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m2_f32(c); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m2_f32(d); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} +} + +typedef unsigned char uchar8 __attribute__((ext_vector_type(8))); +typedef float float8 __attribute__((ext_vector_type(8))); + +// An ordinary vector can produce a result larger than a Neon register. +float8 ordinary_vector(uchar8 bits) { + return __builtin_elementwise_convert_from_f8e5m2_f32(bits); +} diff --git a/clang/test/Sema/builtins-elementwise-convert-from-fp8-overflow.c b/clang/test/Sema/builtins-elementwise-convert-from-fp8-overflow.c new file mode 100644 index 00000000000000..c14ca3ca17c119 --- /dev/null +++ b/clang/test/Sema/builtins-elementwise-convert-from-fp8-overflow.c @@ -0,0 +1,49 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -std=c23 -fexperimental-overflow-behavior-types -verify %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x c++ -std=c++20 -fexperimental-overflow-behavior-types -verify %s + +typedef unsigned char __attribute__((overflow_behavior(wrap))) wrap_uchar; +typedef signed _BitInt(8) __attribute__((overflow_behavior(trap))) trap_int8; +typedef unsigned _BitInt(7) __attribute__((overflow_behavior(wrap))) wrap_uint7; +typedef unsigned _BitInt(1) __attribute__((overflow_behavior(trap))) trap_uint1; +typedef bool __attribute__((overflow_behavior(wrap))) wrap_bool; +enum byte_enum : unsigned char { byte_zero }; +typedef enum byte_enum __attribute__((overflow_behavior(trap))) trap_enum; + +typedef const volatile wrap_uchar qualified_byte; +typedef const volatile wrap_uint7 qualified_uint7; +typedef qualified_byte byte_alias; +typedef qualified_uint7 uint7_alias; + +void valid_inputs(byte_alias *c, const trap_int8 i) { + // Overflow behavior does not change the source encoding or result type. + _Float16 h = __builtin_elementwise_convert_from_f8e5m2_f16(*c); + __bf16 b = __builtin_elementwise_convert_from_f8e4m3fn_bf16(i); + float f = __builtin_elementwise_convert_from_f8e5m3fnu_f32(i); +} + +void invalid_inputs(uint7_alias *u7, const trap_uint1 u1, + const wrap_bool b, const trap_enum e) { + // Check the underlying integer's width, not the wrapper's storage size. + __builtin_elementwise_convert_from_f8e5m2_f32(*u7); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(u1); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + + // Wrapping a boolean or enumeration must not hide its type category. + __builtin_elementwise_convert_from_f8e4m3fn_f16(b); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m3fnu_bf16(e); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} +} + +#ifdef __cplusplus +template <class T> +concept can_convert = requires(T bits) { + __builtin_elementwise_convert_from_f8e5m2_f32(bits); +}; +static_assert(can_convert<wrap_uchar>); +static_assert(can_convert<const trap_int8 &>); +static_assert(!can_convert<qualified_uint7>); +static_assert(!can_convert<wrap_bool>); +static_assert(!can_convert<trap_enum>); +#endif diff --git a/clang/test/Sema/builtins-elementwise-convert-from-fp8-riscv.c b/clang/test/Sema/builtins-elementwise-convert-from-fp8-riscv.c new file mode 100644 index 00000000000000..d823e0d8d3b0dd --- /dev/null +++ b/clang/test/Sema/builtins-elementwise-convert-from-fp8-riscv.c @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 -triple riscv64 -target-feature +v -target-feature +f -target-feature +d -mvscale-min=2 -mvscale-max=2 -fsyntax-only -verify %s + +typedef __rvv_uint8m1_t fixed_rvv_uint8 __attribute__((riscv_rvv_vector_bits(128))); + +void rvv_vectors(__rvv_uint8m1_t a, __rvv_bool8_t b, fixed_rvv_uint8 c) { + (void)__builtin_elementwise_convert_from_f8e5m2_f32(a); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m2_f32(b); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} + (void)__builtin_elementwise_convert_from_f8e5m2_f32(c); // expected-error {{must be a scalar or a fixed-length vector declared with 'vector_size' or 'ext_vector_type'}} +} diff --git a/clang/test/Sema/builtins-elementwise-convert-from-fp8-target.c b/clang/test/Sema/builtins-elementwise-convert-from-fp8-target.c new file mode 100644 index 00000000000000..917822cf664f12 --- /dev/null +++ b/clang/test/Sema/builtins-elementwise-convert-from-fp8-target.c @@ -0,0 +1,35 @@ +// RUN: %clang_cc1 -triple i386-unknown-unknown -target-feature -sse2 -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple i386-unknown-unknown -target-feature +sse2 -fsyntax-only -verify=supported %s +// supported-no-diagnostics + +// Recognition of the spelling does not imply support for the destination type. +#if !__has_builtin(__builtin_elementwise_convert_from_f8e5m2_f16) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e5m2_bf16) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e5m2_f32) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e4m3fn_f16) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e4m3fn_bf16) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e4m3fn_f32) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e5m3fnu_f16) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e5m3fnu_bf16) || \ + !__has_builtin(__builtin_elementwise_convert_from_f8e5m3fnu_f32) +#error missing encoded floating-point conversion builtin +#endif + +typedef unsigned char uchar4 __attribute__((ext_vector_type(4))); + +void scalar(unsigned char bits) { + (void)__builtin_elementwise_convert_from_f8e5m2_f16(bits); // expected-error {{_Float16 is not supported on this target}} + (void)__builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); // expected-error {{__bf16 is not supported on this target}} + (void)__builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} + +void vector(uchar4 bits) { + (void)__builtin_elementwise_convert_from_f8e5m2_f16(bits); // expected-error {{_Float16 is not supported on this target}} + (void)__builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); // expected-error {{__bf16 is not supported on this target}} + (void)__builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} + +void unevaluated(unsigned char bits, uchar4 packed) { + (void)sizeof(__builtin_elementwise_convert_from_f8e5m2_f16(bits)); // expected-error {{_Float16 is not supported on this target}} + (void)sizeof(__builtin_elementwise_convert_from_f8e4m3fn_bf16(packed)); // expected-error {{__bf16 is not supported on this target}} +} diff --git a/clang/test/Sema/builtins-elementwise-convert-from-fp8.c b/clang/test/Sema/builtins-elementwise-convert-from-fp8.c new file mode 100644 index 00000000000000..7a722306bca78b --- /dev/null +++ b/clang/test/Sema/builtins-elementwise-convert-from-fp8.c @@ -0,0 +1,166 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -std=c23 -fenable-matrix -verify %s + +#define CHECK_BUILTIN(NAME) \ + static_assert(__has_builtin(NAME), #NAME); \ + static_assert(!__has_constexpr_builtin(NAME), #NAME) + +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e5m2_f16); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e5m2_bf16); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e5m2_f32); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e4m3fn_f16); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e4m3fn_bf16); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e4m3fn_f32); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e5m3fnu_f16); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e5m3fnu_bf16); +CHECK_BUILTIN(__builtin_elementwise_convert_from_f8e5m3fnu_f32); + +typedef unsigned char uchar; +typedef signed _BitInt(8) int8; +typedef unsigned _BitInt(8) uint8; + +typedef char char1 __attribute__((ext_vector_type(1))); +typedef signed char schar2 __attribute__((ext_vector_type(2))); +typedef uchar uchar3 __attribute__((ext_vector_type(3))); +typedef uchar uchar4 __attribute__((ext_vector_type(4))); +typedef int8 int8x4 __attribute__((ext_vector_type(4))); +typedef uint8 uint8x4 __attribute__((ext_vector_type(4))); +typedef _Float16 half2 __attribute__((ext_vector_type(2))); +typedef _Float16 half4 __attribute__((ext_vector_type(4))); +typedef __bf16 bfloat3 __attribute__((ext_vector_type(3))); +typedef __bf16 bfloat4 __attribute__((ext_vector_type(4))); +typedef float float1 __attribute__((ext_vector_type(1))); +typedef float float4 __attribute__((ext_vector_type(4))); + +typedef uchar gnu_uchar4 __attribute__((vector_size(4))); +typedef int8 gnu_int8x4 __attribute__((vector_size(4))); +typedef _Float16 gnu_half4 __attribute__((vector_size(8))); +typedef __bf16 gnu_bfloat4 __attribute__((vector_size(8))); +typedef float gnu_float4 __attribute__((vector_size(16))); + +#define CHECK_TYPE(EXPR, TYPE) \ + static_assert(__builtin_types_compatible_p(__typeof__(EXPR), TYPE), "incorrect result type") + +#define CHECK_FORMAT(FORMAT, BITS, HALF, BFLOAT, FLOAT) \ + CHECK_TYPE(__builtin_elementwise_convert_from_##FORMAT##_f16(BITS), HALF); \ + CHECK_TYPE(__builtin_elementwise_convert_from_##FORMAT##_bf16(BITS), BFLOAT); \ + CHECK_TYPE(__builtin_elementwise_convert_from_##FORMAT##_f32(BITS), FLOAT) + +void scalar_types(char c, signed char sc, uchar uc, int8 i8, uint8 u8, + const uchar cc, volatile uchar vc) { + CHECK_FORMAT(f8e5m2, c, _Float16, __bf16, float); + CHECK_FORMAT(f8e4m3fn, sc, _Float16, __bf16, float); + CHECK_FORMAT(f8e5m3fnu, uc, _Float16, __bf16, float); + CHECK_FORMAT(f8e5m2, i8, _Float16, __bf16, float); + CHECK_FORMAT(f8e5m2, u8, _Float16, __bf16, float); + CHECK_FORMAT(f8e5m2, cc, _Float16, __bf16, float); + CHECK_FORMAT(f8e5m2, vc, _Float16, __bf16, float); + CHECK_FORMAT(f8e5m2, (uchar)(uc >> 1), _Float16, __bf16, float); +} + +void vector_types(char1 v1, schar2 v2, uchar3 v3, uchar4 v4, + int8x4 i8, uint8x4 u8, const uchar4 cv, + gnu_uchar4 gv, gnu_int8x4 gi8) { + CHECK_FORMAT(f8e5m2, v4, half4, bfloat4, float4); + CHECK_FORMAT(f8e4m3fn, v4, half4, bfloat4, float4); + CHECK_FORMAT(f8e5m3fnu, v4, half4, bfloat4, float4); + CHECK_FORMAT(f8e5m2, i8, half4, bfloat4, float4); + CHECK_FORMAT(f8e5m2, u8, half4, bfloat4, float4); + CHECK_FORMAT(f8e5m2, cv, half4, bfloat4, float4); + CHECK_FORMAT(f8e5m2, gv, gnu_half4, gnu_bfloat4, gnu_float4); + CHECK_FORMAT(f8e4m3fn, gv, gnu_half4, gnu_bfloat4, gnu_float4); + CHECK_FORMAT(f8e5m3fnu, gv, gnu_half4, gnu_bfloat4, gnu_float4); + CHECK_FORMAT(f8e5m2, gi8, gnu_half4, gnu_bfloat4, gnu_float4); + + // Element counts are preserved, including one and non-power-of-two counts. + CHECK_TYPE(__builtin_elementwise_convert_from_f8e5m2_f32(v1), float1); + CHECK_TYPE(__builtin_elementwise_convert_from_f8e4m3fn_f16(v2), half2); + CHECK_TYPE(__builtin_elementwise_convert_from_f8e5m3fnu_bf16(v3), bfloat3); + CHECK_TYPE(__builtin_elementwise_convert_from_f8e5m2_f32(v4.x), float); +} + +typedef short short4 __attribute__((ext_vector_type(4))); +typedef bool bool8 __attribute__((ext_vector_type(8))); +typedef unsigned _BitInt(4) uint4x4 __attribute__((ext_vector_type(4))); +typedef uchar uchar2x2 __attribute__((matrix_type(2, 2))); +enum byte_enum : unsigned char { byte_zero }; +struct bytes { + uchar value; +}; + +void invalid_types(bool b, enum byte_enum e, short s, int i, long l, + _BitInt(7) i7, _BitInt(9) i9, unsigned _BitInt(1) u1, + float f, __bf16 bf, _Complex float complex, + uchar *p, struct bytes aggregate, + short4 sv, float4 fv, bool8 bv, uint4x4 u4v, uchar2x2 m) { + uchar a[4]; + __builtin_elementwise_convert_from_f8e5m2_f32(b); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(e); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(s); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(i); + // expected-error@-1 {{argument to '__builtin_elementwise_convert_from_f8e5m2_f32' must be an 8-bit integer or a vector of 8-bit integers (was 'int')}} + __builtin_elementwise_convert_from_f8e5m2_f32(l); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(i7); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(i9); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(u1); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(f); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(bf); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(complex); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(p); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(a); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(aggregate); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(sv); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(fv); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(bv); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(u4v); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(m); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} +} + +void promotions(uchar bits) { + __builtin_elementwise_convert_from_f8e5m2_f32(0x38); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(bits >> 1); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(+bits); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} +} + +void arity(uchar bits) { + __builtin_elementwise_convert_from_f8e5m2_f16(); + // expected-error@-1 {{too few arguments to function call, expected 1, have 0}} + __builtin_elementwise_convert_from_f8e4m3fn_bf16(bits, bits); + // expected-error@-1 {{too many arguments to function call, expected 1, have 2}} + __builtin_elementwise_convert_from_f8e5m3fnu_f32(bits, bits, bits); + // expected-error@-1 {{too many arguments to function call, expected 1, have 3}} +} + +// These calls cannot be used in constant expressions, including initializers +// for objects with static storage duration in C. +_Float16 global_half = __builtin_elementwise_convert_from_f8e5m2_f16((uchar)0); +// expected-error@-1 {{initializer element is not a compile-time constant}} +__bf16 global_bfloat = __builtin_elementwise_convert_from_f8e4m3fn_bf16((uchar)0); +// expected-error@-1 {{initializer element is not a compile-time constant}} +float global_float = __builtin_elementwise_convert_from_f8e5m3fnu_f32((uchar)0); +// expected-error@-1 {{initializer element is not a compile-time constant}} + +void static_initializer(void) { + static float4 values = __builtin_elementwise_convert_from_f8e5m2_f32((uchar4){0}); + // expected-error@-1 {{initializer element is not a compile-time constant}} +} diff --git a/clang/test/SemaCUDA/builtins-elementwise-convert-from-fp8.cu b/clang/test/SemaCUDA/builtins-elementwise-convert-from-fp8.cu new file mode 100644 index 00000000000000..c8bdc0002eab7b --- /dev/null +++ b/clang/test/SemaCUDA/builtins-elementwise-convert-from-fp8.cu @@ -0,0 +1,78 @@ +// RUN: %clang_cc1 -x cuda -triple x86_64-unknown-linux-gnu -aux-triple nvptx64-nvidia-cuda -fsyntax-only -verify=supported %s +// RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -fsyntax-only -verify=supported %s +// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -fsyntax-only -verify=supported %s +// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -fsyntax-only -verify=supported %s +// RUN: %clang_cc1 -x cuda -triple i386-unknown-linux-gnu -aux-triple x86_64-unknown-linux-gnu -target-feature -sse2 -fcuda-is-device -DUNSUPPORTED_TARGET -fsyntax-only -verify=unsupported %s +// RUN: %clang_cc1 -x hip -triple i386-unknown-linux-gnu -aux-triple x86_64-unknown-linux-gnu -target-feature -sse2 -fcuda-is-device -DUNSUPPORTED_TARGET -fsyntax-only -verify=unsupported %s +// supported-no-diagnostics + +#define __host__ __attribute__((host)) +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) + +typedef unsigned char uchar3 __attribute__((ext_vector_type(3))); + +#ifdef UNSUPPORTED_TARGET + +// Unsupported result types in host-only or unused inline host/device functions +// must not be diagnosed during device compilation. +__host__ void host_only(unsigned char bits, uchar3 packed) { + (void)__builtin_elementwise_convert_from_f8e5m2_f16(bits); + (void)__builtin_elementwise_convert_from_f8e5m2_f16(packed); +} + +__host__ __device__ inline void unused(unsigned char bits, uchar3 packed) { + (void)__builtin_elementwise_convert_from_f8e5m2_f16(bits); + (void)__builtin_elementwise_convert_from_f8e5m2_f16(packed); +} + +__host__ __device__ inline void used(unsigned char bits, uchar3 packed) { + // CUDA permits spelling _Float16 even when one compilation target does not + // support it, and unevaluated operands do not require device type support. + static_assert(__is_same(decltype(__builtin_elementwise_convert_from_f8e5m2_f16(bits)), + _Float16), ""); + (void)__builtin_elementwise_convert_from_f8e5m2_f16(bits); // unsupported-error {{expression requires 16 bit size '_Float16' type support, but target 'i386-unknown-linux-gnu' does not support it}} + (void)__builtin_elementwise_convert_from_f8e5m2_f16(packed); // unsupported-error {{expression requires 16 bit size '_Float16' type support, but target 'i386-unknown-linux-gnu' does not support it}} +} + +__global__ void kernel(unsigned char bits, uchar3 packed) { + used(bits, packed); // unsupported-note {{called by 'kernel'}} +} + +#else + +typedef _Float16 float16_3 __attribute__((ext_vector_type(3))); +typedef __bf16 bfloat16_3 __attribute__((ext_vector_type(3))); +typedef float float3 __attribute__((ext_vector_type(3))); + +__host__ __device__ void results(unsigned char bits, uchar3 packed) { + static_assert(__is_same(decltype(__builtin_elementwise_convert_from_f8e5m2_f16(bits)), + _Float16), ""); + static_assert(__is_same(decltype(__builtin_elementwise_convert_from_f8e4m3fn_bf16(bits)), + __bf16), ""); + static_assert(__is_same(decltype(__builtin_elementwise_convert_from_f8e5m3fnu_f32(bits)), + float), ""); + static_assert(__is_same(decltype(__builtin_elementwise_convert_from_f8e5m2_f16(packed)), + float16_3), ""); + static_assert(__is_same(decltype(__builtin_elementwise_convert_from_f8e4m3fn_bf16(packed)), + bfloat16_3), ""); + static_assert(__is_same(decltype(__builtin_elementwise_convert_from_f8e5m3fnu_f32(packed)), + float3), ""); + + (void)__builtin_elementwise_convert_from_f8e5m2_f16(bits); + (void)__builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); + (void)__builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); + (void)__builtin_elementwise_convert_from_f8e5m2_f16(packed); + (void)__builtin_elementwise_convert_from_f8e4m3fn_bf16(packed); + (void)__builtin_elementwise_convert_from_f8e5m3fnu_f32(packed); +} + +__host__ void host(unsigned char bits, uchar3 packed) { + results(bits, packed); +} + +__global__ void kernel(unsigned char bits, uchar3 packed) { + results(bits, packed); +} + +#endif diff --git a/clang/test/SemaCXX/builtins-elementwise-convert-from-fp8.cpp b/clang/test/SemaCXX/builtins-elementwise-convert-from-fp8.cpp new file mode 100644 index 00000000000000..21ee58a2fbd6ed --- /dev/null +++ b/clang/test/SemaCXX/builtins-elementwise-convert-from-fp8.cpp @@ -0,0 +1,123 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -std=c++20 -verify %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -std=c++20 -fexperimental-new-constant-interpreter -verify %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -std=c++20 -DAST_DUMP -ast-dump -ast-dump-filter call_expr %s | FileCheck %s --check-prefix=AST + +#ifdef AST_DUMP + +// The builtin remains an ordinary CallExpr. Its operand undergoes an +// lvalue-to-rvalue conversion without an integer promotion. +float call_expr(unsigned char bits) { + return __builtin_elementwise_convert_from_f8e5m2_f32(bits); +} + +// AST-LABEL: FunctionDecl {{.*}} call_expr 'float (unsigned char)' +// AST: ReturnStmt +// AST-NEXT: CallExpr {{.*}} 'float' +// AST: ImplicitCastExpr {{.*}} 'unsigned char' <LValueToRValue> +// AST-NEXT: DeclRefExpr {{.*}} 'unsigned char' lvalue ParmVar {{.*}} 'bits' 'unsigned char' + +#else + +using uchar = unsigned char; +using uchar4 = uchar __attribute__((ext_vector_type(4))); +using half4 = _Float16 __attribute__((ext_vector_type(4))); +using bfloat4 = __bf16 __attribute__((ext_vector_type(4))); +using float4 = float __attribute__((ext_vector_type(4))); +using gnu_uchar4 = uchar __attribute__((vector_size(4))); +using gnu_half4 = _Float16 __attribute__((vector_size(8))); +using gnu_bfloat4 = __bf16 __attribute__((vector_size(8))); +using gnu_float4 = float __attribute__((vector_size(16))); + +#define CHECK_TYPE(EXPR, TYPE) static_assert(__is_same(decltype(EXPR), TYPE)) + +void cv_qualified(const uchar &c, volatile uchar &v, const volatile uchar &cv, + const uchar4 &vc, volatile uchar4 &vv) { + CHECK_TYPE(__builtin_elementwise_convert_from_f8e5m2_f16(c), _Float16); + CHECK_TYPE(__builtin_elementwise_convert_from_f8e4m3fn_bf16(v), __bf16); + CHECK_TYPE(__builtin_elementwise_convert_from_f8e5m3fnu_f32(cv), float); + CHECK_TYPE(__builtin_elementwise_convert_from_f8e5m2_f16(vc), half4); + CHECK_TYPE(__builtin_elementwise_convert_from_f8e4m3fn_bf16(vv), bfloat4); +} + +template <class T> auto convert_half(T bits) { + return __builtin_elementwise_convert_from_f8e5m2_f16(bits); +} + +template <class T> auto convert_bfloat(T bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); +} + +template <class T> auto convert_float(T bits) { + return __builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); +} + +void dependent_result_types(uchar scalar, uchar4 ext, gnu_uchar4 generic) { + CHECK_TYPE(convert_half(scalar), _Float16); + CHECK_TYPE(convert_bfloat(scalar), __bf16); + CHECK_TYPE(convert_float(scalar), float); + CHECK_TYPE(convert_half(ext), half4); + CHECK_TYPE(convert_bfloat(ext), bfloat4); + CHECK_TYPE(convert_float(ext), float4); + CHECK_TYPE(convert_half(generic), gnu_half4); + CHECK_TYPE(convert_bfloat(generic), gnu_bfloat4); + CHECK_TYPE(convert_float(generic), gnu_float4); +} + +// Value dependence does not change the argument's known 8-bit integer type. +template <uchar Bits> auto convert_value() { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(Bits); +} +CHECK_TYPE(convert_value<0x38>(), float); + +enum byte_enum : unsigned char { byte_zero }; +enum class scoped_byte_enum : unsigned char { zero }; +struct convertible { + operator unsigned char() const; +}; + +void invalid_types(bool b, byte_enum e, scoped_byte_enum se, convertible c) { + __builtin_elementwise_convert_from_f8e5m2_f32(b); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(e); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(se); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} + __builtin_elementwise_convert_from_f8e5m2_f32(c); + // expected-error@-1 {{must be an 8-bit integer or a vector of 8-bit integers}} +} + +template <class T> auto invalid_dependent(T bits) { + return __builtin_elementwise_convert_from_f8e5m2_f32(bits); + // expected-error@-1 3 {{must be an 8-bit integer or a vector of 8-bit integers}} +} + +template <class... T> auto dependent_arity(T... bits) { + return __builtin_elementwise_convert_from_f8e4m3fn_f32(bits...); + // expected-error@-1 {{too few arguments to function call, expected 1, have 0}} + // expected-error@-2 {{too many arguments to function call, expected 1, have 2}} +} + +void instantiate_errors() { + invalid_dependent(0); + // expected-note@-1 {{in instantiation of function template specialization 'invalid_dependent<int>' requested here}} + invalid_dependent(false); + // expected-note@-1 {{in instantiation of function template specialization 'invalid_dependent<bool>' requested here}} + invalid_dependent(byte_zero); + // expected-note@-1 {{in instantiation of function template specialization 'invalid_dependent<byte_enum>' requested here}} + dependent_arity(); + // expected-note@-1 {{in instantiation of function template specialization 'dependent_arity<>' requested here}} + dependent_arity(uchar{}, uchar{}); + // expected-note@-1 {{in instantiation of function template specialization 'dependent_arity<unsigned char, unsigned char>' requested here}} +} + +constexpr _Float16 constant_half = __builtin_elementwise_convert_from_f8e5m2_f16(uchar{}); +// expected-error@-1 {{constexpr variable 'constant_half' must be initialized by a constant expression}} +constexpr __bf16 constant_bfloat = __builtin_elementwise_convert_from_f8e4m3fn_bf16(uchar{}); +// expected-error@-1 {{constexpr variable 'constant_bfloat' must be initialized by a constant expression}} +constexpr float constant_float = __builtin_elementwise_convert_from_f8e5m3fnu_f32(uchar{}); +// expected-error@-1 {{constexpr variable 'constant_float' must be initialized by a constant expression}} + +static_assert(__builtin_elementwise_convert_from_f8e5m2_f32(uchar{}) == 0.0f); +// expected-error@-1 {{static assertion expression is not an integral constant expression}} + +#endif diff --git a/clang/test/SemaOpenCL/builtins-elementwise-convert-from-fp8.cl b/clang/test/SemaOpenCL/builtins-elementwise-convert-from-fp8.cl new file mode 100644 index 00000000000000..cc5b5fcc90fa8c --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-elementwise-convert-from-fp8.cl @@ -0,0 +1,46 @@ +// RUN: %clang_cc1 -triple spir64-unknown-unknown -cl-std=CL1.2 -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple spir64-unknown-unknown -cl-std=CL3.0 -cl-ext=-cl_khr_fp16 -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -fsyntax-only -verify %s +// expected-no-diagnostics + +typedef unsigned char uchar3 __attribute__((ext_vector_type(3))); +typedef _Float16 float16_3 __attribute__((ext_vector_type(3))); +typedef __bf16 bfloat16_3 __attribute__((ext_vector_type(3))); +typedef float float3 __attribute__((ext_vector_type(3))); +typedef half half3 __attribute__((ext_vector_type(3))); + +// The f16 suffix always produces _Float16, including when cl_khr_fp16 is +// disabled. OpenCL half is a distinct type. +void results(unsigned char bits, uchar3 packed) { + _Static_assert(__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e5m2_f16(bits)), + _Float16), ""); + _Static_assert(__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e4m3fn_bf16(bits)), + __bf16), ""); + _Static_assert(__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e5m3fnu_f32(bits)), + float), ""); + _Static_assert(__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e5m2_f16(packed)), + float16_3), ""); + _Static_assert(__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e4m3fn_bf16(packed)), + bfloat16_3), ""); + _Static_assert(__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e5m3fnu_f32(packed)), + float3), ""); + _Static_assert(!__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e5m2_f16(bits)), + half), ""); + _Static_assert(!__builtin_types_compatible_p( + __typeof__(__builtin_elementwise_convert_from_f8e5m2_f16(packed)), + half3), ""); + + (void)__builtin_elementwise_convert_from_f8e5m2_f16(bits); + (void)__builtin_elementwise_convert_from_f8e4m3fn_bf16(bits); + (void)__builtin_elementwise_convert_from_f8e5m3fnu_f32(bits); + (void)__builtin_elementwise_convert_from_f8e5m2_f16(packed); + (void)__builtin_elementwise_convert_from_f8e4m3fn_bf16(packed); + (void)__builtin_elementwise_convert_from_f8e5m3fnu_f32(packed); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
