Add MX floating point types (f4E2M1FN, f6E2M3FN, f6E3M2FN, f8E8M0FNU) (…

…#2582) This PR is adding MX (microscaling) floating point types to StableHLO. The RFC is created in #2581 Bumps VHLO version to 1.8.0
openxla · Oct 24, 2024 · 5ae7716 · 5ae7716
1 parent 7c06dd2
commit 5ae7716
Show file tree

Hide file tree

Showing 17 changed files with 3,249 additions and 114 deletions.
diff --git a/docs/spec.md b/docs/spec.md
@@ -245,8 +245,9 @@ BooleanType ::= 'i1'
 IntegerType ::= SignedIntegerType | UnsignedIntegerType
 SignedIntegerType ::= 'si2' | 'si4' | 'si8' | 'si16' | 'si32' | 'si64'
 UnsignedIntegerType ::= 'ui2' | 'ui4' | 'ui8' | 'ui16' | 'ui32' | 'ui64'
-FloatType ::= 'f8E3M4' | 'f8E4M3' | 'f8E4M3FN' | 'f8E4M3FNUZ' | 'f8E4M3B11FNUZ'
-            | 'f8E5M2' | 'f8E5M2FNUZ' | 'bf16' | 'f16' | 'f32' | 'f64'
+FloatType ::= 'f4E2M1FN' | 'f6E2M3FN' | 'f6E3M2FN' | 'f8E3M4' | 'f8E4M3'
+            | 'f8E4M3FN' | 'f8E4M3FNUZ' | 'f8E4M3B11FNUZ' | 'f8E5M2'
+            | 'f8E5M2FNUZ' | 'f8E8M0FNU' | 'bf16' | 'f16' | 'f32' | 'f64'
 TensorFloat32 ::= 'tf32'
 ComplexType ::= 'complex' '<' ComplexElementType '>'
 ComplexElementType ::= 'f32' | 'f64'
@@ -284,6 +285,9 @@ values of type `tensor<T>`).
     [the IEEE 754 standard](https://ieeexplore.ieee.org/document/8766229).
   * `tf32` type corresponds to the [TensorFloat32 format](https://blogs.nvidia.com/blog/tensorfloat-32-precision-format/)
     and has limited support in StableHLO.
+  * `f4E2M1FN`, `f6E2M3FN`, `f6E3M2FN` and `f8E8M0FNU` MX (microscaling) types
+    described in
+    [OCP Microscaling Formats Specification](https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf).
 * **Complex types** represent complex values that have a **real part**
   and an **imaginary part** of the same **element type**. Supported complex
   types are `complex<f32>` (both parts are of type `f32`) and `complex<f64>`

diff --git a/stablehlo/dialect/Base.td b/stablehlo/dialect/Base.td
@@ -42,8 +42,9 @@ def HLO_SInt : SignlessIntOfWidths<[2, 4, 8, 16, 32, 64]>;
 def HLO_UInt : UnsignedIntOfWidths<[2, 4, 8, 16, 32, 64]>;
 def HLO_Int : AnyTypeOf<[HLO_SInt, HLO_UInt]>;
 
-def HLO_Float : AnyTypeOf<[F8E3M4, F8E4M3, F8E4M3FN, F8E4M3FNUZ, F8E4M3B11FNUZ,
-                           F8E5M2, F8E5M2FNUZ, F16, F32, F64, BF16]>;
+def HLO_Float : AnyTypeOf<[F4E2M1FN, F6E2M3FN, F6E3M2FN, F8E3M4, F8E4M3,
+                           F8E4M3FN, F8E4M3FNUZ, F8E4M3B11FNUZ, F8E5M2,
+                           F8E5M2FNUZ, F8E8M0FNU, F16, F32, F64, BF16]>;
 def HLO_Float32Or64 : AnyTypeOf<[F32, F64]>;
 
 def HLO_Complex : Complex<AnyTypeOf<[F32, F64]>>;

diff --git a/stablehlo/dialect/Version.h b/stablehlo/dialect/Version.h
@@ -38,7 +38,7 @@ class Version {
   static FailureOr<Version> fromString(llvm::StringRef versionRef);
 
   /// Return a Version representing the current VHLO dialect version.
-  static Version getCurrentVersion() { return Version(1, 7, 9); }
+  static Version getCurrentVersion() { return Version(1, 8, 0); }
 
   /// Return a Version representing the minimum supported VHLO dialect version.
   static Version getMinimumVersion() { return Version(0, 9, 0); }

diff --git a/stablehlo/dialect/VhloBytecode.cpp b/stablehlo/dialect/VhloBytecode.cpp
@@ -189,7 +189,7 @@ enum AttributeCode {
 /// location is updated.
 enum TypeCode {
   // TO ADD TYPE: Add an enum value with doc string for new type.
-  // Next available code: 37
+  // Next available code: 41
 
   ///   BooleanV1Type {
   ///   }
@@ -336,6 +336,18 @@ enum TypeCode {
   ///   }
   kWitnessV1Type = 26,
 
+  ///   FloatF4E2M1FNV1Type {
+  ///   }
+  kFloatF4E2M1FNV1Type = 37,
+
+  ///   FloatF6E2M3FNV1Type {
+  ///   }
+  kFloatF6E2M3FNV1Type = 38,
+
+  ///   FloatF6E3M2FNV1Type {
+  ///   }
+  kFloatF6E3M2FNV1Type = 39,
+
   ///   FloatF8E4M3FNUZV1Type {
   ///   }
   kFloatF8E4M3FNUZV1Type = 27,
@@ -348,6 +360,10 @@ enum TypeCode {
   ///   }
   kFloatF8E4M3B11FNUZV1Type = 29,
 
+  ///   FloatF8E8M0FNUV1Type {
+  ///   }
+  kFloatF8E8M0FNUV1Type = 40,
+
   ///   UniformQuantizedPerAxisV1Type {
   ///     flags: varint
   ///     storageType: Type
@@ -705,14 +721,18 @@ const llvm::fltSemantics &getFloatSemantics(Type type) {
   if (isa<FloatBF16V1Type>(type)) return APFloat::BFloat();
   if (isa<FloatF16V1Type>(type)) return APFloat::IEEEhalf();
   if (isa<FloatF32V1Type>(type)) return APFloat::IEEEsingle();
+  if (isa<FloatF4E2M1FNV1Type>(type)) return APFloat::Float4E2M1FN();
   if (isa<FloatF64V1Type>(type)) return APFloat::IEEEdouble();
+  if (isa<FloatF6E2M3FNV1Type>(type)) return APFloat::Float6E2M3FN();
+  if (isa<FloatF6E3M2FNV1Type>(type)) return APFloat::Float6E3M2FN();
   if (isa<FloatF8E3M4V1Type>(type)) return APFloat::Float8E3M4();
   if (isa<FloatF8E4M3FNUZV1Type>(type)) return APFloat::Float8E4M3FNUZ();
   if (isa<FloatF8E4M3B11FNUZV1Type>(type)) return APFloat::Float8E4M3B11FNUZ();
   if (isa<FloatF8E4M3FNV1Type>(type)) return APFloat::Float8E4M3FN();
   if (isa<FloatF8E4M3V1Type>(type)) return APFloat::Float8E4M3();
   if (isa<FloatF8E5M2FNUZV1Type>(type)) return APFloat::Float8E5M2FNUZ();
   if (isa<FloatF8E5M2V1Type>(type)) return APFloat::Float8E5M2();
+  if (isa<FloatF8E8M0FNUV1Type>(type)) return APFloat::Float8E8M0FNU();
   if (isa<FloatTF32V1Type>(type)) return APFloat::FloatTF32();
   llvm::report_fatal_error("unsupported floating-point type");
 }
@@ -974,8 +994,14 @@ Type VhloBytecodeInterface::readType(DialectBytecodeReader &reader) const {
       return FloatF16V1Type::get(getContext());
     case vhlo_encoding::kFloatF32V1Type:
       return FloatF32V1Type::get(getContext());
+    case vhlo_encoding::kFloatF4E2M1FNV1Type:
+      return FloatF4E2M1FNV1Type::get(getContext());
     case vhlo_encoding::kFloatF64V1Type:
       return FloatF64V1Type::get(getContext());
+    case vhlo_encoding::kFloatF6E2M3FNV1Type:
+      return FloatF6E2M3FNV1Type::get(getContext());
+    case vhlo_encoding::kFloatF6E3M2FNV1Type:
+      return FloatF6E3M2FNV1Type::get(getContext());
     case vhlo_encoding::kFloatF8E5M2V1Type:
       return FloatF8E5M2V1Type::get(getContext());
     case vhlo_encoding::kFloatF8E4M3V1Type:
@@ -990,6 +1016,8 @@ Type VhloBytecodeInterface::readType(DialectBytecodeReader &reader) const {
       return FloatF8E4M3B11FNUZV1Type::get(getContext());
     case vhlo_encoding::kFloatF8E3M4V1Type:
       return FloatF8E3M4V1Type::get(getContext());
+    case vhlo_encoding::kFloatF8E8M0FNUV1Type:
+      return FloatF8E8M0FNUV1Type::get(getContext());
     case vhlo_encoding::kFloatTF32V1Type:
       return FloatTF32V1Type::get(getContext());
     case vhlo_encoding::kFunctionV1Type:
@@ -1070,10 +1098,25 @@ LogicalResult VhloBytecodeInterface::writeType(
         LOG_WRITE_CALL;
         return writer.writeVarInt(vhlo_encoding::kFloatF32V1Type), success();
       })
+      .Case([&](FloatF4E2M1FNV1Type) {
+        LOG_WRITE_CALL;
+        return writer.writeVarInt(vhlo_encoding::kFloatF4E2M1FNV1Type),
+               success();
+      })
       .Case([&](FloatF64V1Type) {
         LOG_WRITE_CALL;
         return writer.writeVarInt(vhlo_encoding::kFloatF64V1Type), success();
       })
+      .Case([&](FloatF6E2M3FNV1Type) {
+        LOG_WRITE_CALL;
+        return writer.writeVarInt(vhlo_encoding::kFloatF6E2M3FNV1Type),
+               success();
+      })
+      .Case([&](FloatF6E3M2FNV1Type) {
+        LOG_WRITE_CALL;
+        return writer.writeVarInt(vhlo_encoding::kFloatF6E3M2FNV1Type),
+               success();
+      })
       .Case([&](FloatF8E3M4V1Type) {
         LOG_WRITE_CALL;
         return writer.writeVarInt(vhlo_encoding::kFloatF8E3M4V1Type), success();
@@ -1106,6 +1149,11 @@ LogicalResult VhloBytecodeInterface::writeType(
         return writer.writeVarInt(vhlo_encoding::kFloatF8E5M2FNUZV1Type),
                success();
       })
+      .Case([&](FloatF8E8M0FNUV1Type) {
+        LOG_WRITE_CALL;
+        return writer.writeVarInt(vhlo_encoding::kFloatF8E8M0FNUV1Type),
+               success();
+      })
       .Case([&](FloatTF32V1Type) {
         LOG_WRITE_CALL;
         return writer.writeVarInt(vhlo_encoding::kFloatTF32V1Type), success();

diff --git a/stablehlo/dialect/VhloDialect.td b/stablehlo/dialect/VhloDialect.td
@@ -46,6 +46,7 @@ def VHLO_Dialect : Dialect {
       1.5.0: Make collective ops (`all_reduce`, `all_gather`, `all_to_all`) variadic.
       1.6.0: Add DotAlgorithm specificaiton to `dot_general`.
       1.7.0: Introduce `f8E4M3` and `f8E3M4` types.
+      1.8.0: Introduce `f4E2M1FN`, `f6E2M3FN`, `f6E3M2FN` and `f8E8M0FNU` types.
   }];
 
   let useDefaultAttributePrinterParser = 0;

diff --git a/stablehlo/dialect/VhloTypes.cpp b/stablehlo/dialect/VhloTypes.cpp
@@ -84,6 +84,15 @@ void VhloTypeConverter::addBuiltinToVhloConversions() {
       [&](Float32Type type) { return FloatF32V1Type::get(type.getContext()); });
   addConversion(
       [&](Float64Type type) { return FloatF64V1Type::get(type.getContext()); });
+  addConversion([&](Float4E2M1FNType type) {
+    return FloatF4E2M1FNV1Type::get(type.getContext());
+  });
+  addConversion([&](Float6E2M3FNType type) {
+    return FloatF6E2M3FNV1Type::get(type.getContext());
+  });
+  addConversion([&](Float6E3M2FNType type) {
+    return FloatF6E3M2FNV1Type::get(type.getContext());
+  });
   addConversion([&](Float8E3M4Type type) {
     return FloatF8E3M4V1Type::get(type.getContext());
   });
@@ -105,6 +114,9 @@ void VhloTypeConverter::addBuiltinToVhloConversions() {
   addConversion([&](Float8E5M2FNUZType type) {
     return FloatF8E5M2FNUZV1Type::get(type.getContext());
   });
+  addConversion([&](Float8E8M0FNUType type) {
+    return FloatF8E8M0FNUV1Type::get(type.getContext());
+  });
   addConversion([&](FloatTF32Type type) {
     return FloatTF32V1Type::get(type.getContext());
   });
@@ -182,6 +194,15 @@ void VhloTypeConverter::addVhloToBuiltinConversions() {
       [&](FloatF32V1Type type) { return Float32Type::get(type.getContext()); });
   addConversion(
       [&](FloatF64V1Type type) { return Float64Type::get(type.getContext()); });
+  addConversion([&](FloatF4E2M1FNV1Type type) {
+    return Float4E2M1FNType::get(type.getContext());
+  });
+  addConversion([&](FloatF6E2M3FNV1Type type) {
+    return Float6E2M3FNType::get(type.getContext());
+  });
+  addConversion([&](FloatF6E3M2FNV1Type type) {
+    return Float6E3M2FNType::get(type.getContext());
+  });
   addConversion([&](FloatF8E3M4V1Type type) {
     return Float8E3M4Type::get(type.getContext());
   });
@@ -203,6 +224,9 @@ void VhloTypeConverter::addVhloToBuiltinConversions() {
   addConversion([&](FloatF8E5M2FNUZV1Type type) {
     return Float8E5M2FNUZType::get(type.getContext());
   });
+  addConversion([&](FloatF8E8M0FNUV1Type type) {
+    return Float8E8M0FNUType::get(type.getContext());
+  });
   addConversion([&](FloatTF32V1Type type) {
     return FloatTF32Type::get(type.getContext());
   });

diff --git a/stablehlo/dialect/VhloTypes.td b/stablehlo/dialect/VhloTypes.td
@@ -79,6 +79,15 @@ def VHLO_FloatF32V1 : VHLO_TypeDef<"FloatF32V1", "f32_v1", "0.9.0", "current">;
 // Corresponds to the 'f64' FloatType from the StableHLO spec.
 def VHLO_FloatF64V1 : VHLO_TypeDef<"FloatF64V1","f64_v1", "0.9.0", "current">;
 
+// Corresponds to the 'f4E2M1FN' FloatType from the StableHLO spec.
+def VHLO_FloatF4E2M1FNV1 : VHLO_TypeDef<"FloatF4E2M1FNV1", "f4E2M1FN_v1", "1.8.0", "current">;
+
+// Corresponds to the 'f6E2M3FN' FloatType from the StableHLO spec.
+def VHLO_FloatF6E2M3FNV1 : VHLO_TypeDef<"FloatF6E2M3FNV1", "f6E2M3FN_v1", "1.8.0", "current">;
+
+// Corresponds to the 'f6E3M2FN' FloatType from the StableHLO spec.
+def VHLO_FloatF6E3M2FNV1 : VHLO_TypeDef<"FloatF6E3M2FNV1", "f6E3M2FN_v1", "1.8.0", "current">;
+
 // Corresponds to the 'f8E3M4' FloatType from the StableHLO spec.
 def VHLO_FloatF8E3M4V1 : VHLO_TypeDef<"FloatF8E3M4V1", "f8E3M4_v1", "1.7.0", "current">;
 
@@ -100,6 +109,9 @@ def VHLO_FloatF8E4M3B11FNUZV1 : VHLO_TypeDef<"FloatF8E4M3B11FNUZV1", "f8E4M3B11F
 // Corresponds to the 'f8E5M2FNUZ' FloatType from the StableHLO spec.
 def VHLO_FloatF8E5M2FNUZV1 : VHLO_TypeDef<"FloatF8E5M2FNUZV1", "f8E5M2FNUZ_v1", "0.10.0", "current">;
 
+// Corresponds to the 'f8E8M0FNU' FloatType from the StableHLO spec.
+def VHLO_FloatF8E8M0FNUV1 : VHLO_TypeDef<"FloatF8E8M0FNUV1", "f8E8M0FNU_v1", "1.8.0", "current">;
+
 // Corresponds to the 'tf32' FloatType from the StableHLO spec.
 def VHLO_FloatTF32V1 : VHLO_TypeDef<"FloatTF32V1", "tf31_v1", "1.6.0", "current">;
 

diff --git a/stablehlo/reference/Tensor.cpp b/stablehlo/reference/Tensor.cpp
@@ -118,40 +118,13 @@ Element Tensor::get(const Index &index) const {
       getSizeInBytes(elementType) * flattenIndex(getShape(), index);
 
   // Handle floating-point types.
-  if (elementType.isFloat8E3M4()) {
+  if (isSupportedFloatType(elementType) &&
+      cast<FloatType>(elementType).getWidth() <= 8) {
     auto elementData = reinterpret_cast<const uint8_t *>(elementPtr);
-    return Element(elementType, APFloat(llvm::APFloatBase::Float8E3M4(),
-                                        APInt(8, *elementData)));
-  }
-  if (elementType.isFloat8E4M3B11FNUZ()) {
-    auto elementData = reinterpret_cast<const uint8_t *>(elementPtr);
-    return Element(elementType, APFloat(llvm::APFloatBase::Float8E4M3B11FNUZ(),
-                                        APInt(8, *elementData)));
-  }
-  if (elementType.isFloat8E4M3()) {
-    auto elementData = reinterpret_cast<const uint8_t *>(elementPtr);
-    return Element(elementType, APFloat(llvm::APFloatBase::Float8E4M3(),
-                                        APInt(8, *elementData)));
-  }
-  if (elementType.isFloat8E4M3FN()) {
-    auto elementData = reinterpret_cast<const uint8_t *>(elementPtr);
-    return Element(elementType, APFloat(llvm::APFloatBase::Float8E4M3FN(),
-                                        APInt(8, *elementData)));
-  }
-  if (elementType.isFloat8E4M3FNUZ()) {
-    auto elementData = reinterpret_cast<const uint8_t *>(elementPtr);
-    return Element(elementType, APFloat(llvm::APFloatBase::Float8E4M3FNUZ(),
-                                        APInt(8, *elementData)));
-  }
-  if (elementType.isFloat8E5M2()) {
-    auto elementData = reinterpret_cast<const uint8_t *>(elementPtr);
-    return Element(elementType, APFloat(llvm::APFloatBase::Float8E5M2(),
-                                        APInt(8, *elementData)));
-  }
-  if (elementType.isFloat8E5M2FNUZ()) {
-    auto elementData = reinterpret_cast<const uint8_t *>(elementPtr);
-    return Element(elementType, APFloat(llvm::APFloatBase::Float8E5M2FNUZ(),
-                                        APInt(8, *elementData)));
+    auto floatTy = cast<FloatType>(elementType);
+    return Element(elementType,
+                   APFloat(floatTy.getFloatSemantics(),
+                           APInt(floatTy.getWidth(), *elementData)));
   }
   if (elementType.isF16()) {
     auto elementData = reinterpret_cast<const uint16_t *>(elementPtr);
@@ -262,10 +235,8 @@ void Tensor::set(const Index &index, const Element &element) {
       getSizeInBytes(elementType) * flattenIndex(getShape(), index);
 
   // Handle floating-point types.
-  if (elementType.isFloat8E3M4() || elementType.isFloat8E4M3B11FNUZ() ||
-      elementType.isFloat8E4M3() || elementType.isFloat8E4M3FN() ||
-      elementType.isFloat8E4M3FNUZ() || elementType.isFloat8E5M2() ||
-      elementType.isFloat8E5M2FNUZ()) {
+  if (isSupportedFloatType(elementType) &&
+      cast<FloatType>(elementType).getWidth() <= 8) {
     auto elementData = reinterpret_cast<uint8_t *>(elementPtr);
     auto value = element.getFloatValue();
     *elementData = (uint8_t)value.bitcastToAPInt().getZExtValue();
@@ -457,18 +428,19 @@ Tensor makeTensor(DenseElementsAttr attr) {
   auto elementType = type.getElementType();
 
   // Handle floating-point types.
-  if (elementType.isFloat8E3M4() || elementType.isFloat8E4M3B11FNUZ() ||
-      elementType.isFloat8E4M3() || elementType.isFloat8E4M3FN() ||
-      elementType.isFloat8E4M3FNUZ() || elementType.isFloat8E5M2() ||
-      elementType.isFloat8E5M2FNUZ()) {
+  if (isSupportedFloatType(elementType) &&
+      cast<FloatType>(elementType).getWidth() <= 8) {
     auto floatValues = llvm::map_to_vector(
         attr.getValues<APFloat>(), [&](APFloat value) -> uint8_t {
           return value.bitcastToAPInt().getZExtValue();
         });
 
     // For f8E3M4, f8E4M3, f8E4M3FN, f8E4M3FNUZ, f8E4M3B11FNUZ, f8E5M2, and
-    // f8E5M2FNUZ floating-point types, we use uint8_t as their storage type
-    // because there are no builtin types for those.
+    // f8E5M2FNUZ, f8E8M0FNU floating-point types, we use uint8_t as their
+    // storage type because there are no builtin types for those.
+    // For f4E2M1FN, f6E2M3FN, and f6E3M2FN floating-point types, we still use
+    // uint8_t, even though the underlying types require less bits (similar
+    // to how ui2/ui4 types are handled).
     return Tensor(type, HeapAsmResourceBlob::allocateAndCopyInferAlign<uint8_t>(
                             floatValues));
   }

diff --git a/stablehlo/reference/Types.cpp b/stablehlo/reference/Types.cpp
@@ -48,10 +48,12 @@ bool isSupportedIntegerType(Type type) {
 }
 
 bool isSupportedFloatType(Type type) {
-  return type.isFloat8E3M4() || type.isFloat8E4M3B11FNUZ() ||
-         type.isFloat8E4M3() || type.isFloat8E4M3FN() ||
-         type.isFloat8E4M3FNUZ() || type.isFloat8E5M2() ||
-         type.isFloat8E5M2FNUZ() || type.isF16() || type.isBF16() ||
+  return type.isFloat4E2M1FN() || type.isFloat6E2M3FN() ||
+         type.isFloat6E3M2FN() || type.isFloat8E3M4() ||
+         type.isFloat8E4M3B11FNUZ() || type.isFloat8E4M3() ||
+         type.isFloat8E4M3FN() || type.isFloat8E4M3FNUZ() ||
+         type.isFloat8E5M2() || type.isFloat8E5M2FNUZ() ||
+         type.isFloat8E8M0FNU() || type.isF16() || type.isBF16() ||
          type.isF32() || type.isF64();
 }
 

diff --git a/stablehlo/tests/interpret/constant.mlir b/stablehlo/tests/interpret/constant.mlir
@@ -96,6 +96,30 @@ func.func @constant_op_test_ui64() {
 
 // -----
 
+func.func @constant_op_test_f4_e2m1fn() {
+  %0 = stablehlo.constant dense<[0.0, -0.0, 1.0, 0.125, 0.1, 3.1415, 0x07, 0x0F, 0x01, 0x09]> : tensor<10xf4E2M1FN>
+  check.expect_almost_eq_const %0, dense<[0.0, -0.0, 1.0, 0.0, 0.0, 3.0, 6.0, -6.0, 0.5, -0.5]> : tensor<10xf4E2M1FN>
+  func.return
+}
+
+// -----
+
+func.func @constant_op_test_f6_e2m3fn() {
+  %0 = stablehlo.constant dense<[0.0, -0.0, 1.0, 0.125, 0.1, 3.1415, 0x1F, 0x3F, 0x01, 0x21]> : tensor<10xf6E2M3FN>
+  check.expect_almost_eq_const %0, dense<[0.0, -0.0, 1.0, 0.125, 0.125, 3.25, 7.5, -7.5, 0.125, -0.125]> : tensor<10xf6E2M3FN>
+  func.return
+}
+
+// -----
+
+func.func @constant_op_test_f6_e3m2fn() {
+  %0 = stablehlo.constant dense<[0.0, -0.0, 1.0, 0.125, 0.1, 3.1415, 0x1F, 0x3F, 0x01, 0x21]> : tensor<10xf6E3M2FN>
+  check.expect_almost_eq_const %0, dense<[0.0, -0.0, 1.0, 0.125, 0.125, 3.0, 28.0, -28.0, 0.0625, -0.0625]> : tensor<10xf6E3M2FN>
+  func.return
+}
+
+// -----
+
 func.func @constant_op_test_f8_e3m4() {
   %0 = stablehlo.constant dense<[0.0, -0.0, 1.0, 0.125, 0.1, 3.1415, 0x7F, 0xFF, 0x01, 0x81]> : tensor<10xf8E3M4>
   check.expect_almost_eq_const %0, dense<[0.0, -0.0, 1.0, 0.125, 0.09375, 3.125, 0x7F, 0xFF, 0.015625, -0.015625]> : tensor<10xf8E3M4>
@@ -159,6 +183,14 @@ func.func @constant_op_test_f8_e5m2_fnuz() {
 
 // -----
 
+func.func @constant_op_test_f8_e8m0fnu() {
+  %0 = stablehlo.constant dense<[0.0, 1.0, 0.125, 0.1, 3.1415, 0x00, 0x80, 0xFF]> : tensor<8xf8E8M0FNU>
+  check.expect_almost_eq_const %0, dense<[0.0, 1.0, 0.125, 0.125, 4.0, 1.175490e-38, 2.0, 0xFF]> : tensor<8xf8E8M0FNU>
+  func.return
+}
+
+// -----
+
 func.func @constant_op_test_bf16() {
   %0 = stablehlo.constant dense<[0.0, -0.0, 1.0, 0.125, 0.1, 3.140630, 0x7F80, 0xFF80, 0x7FFF, 0x0001, 0x8001]> : tensor<11xbf16>
   check.expect_almost_eq_const %0, dense<[0.000000e+00, -0.000000e+00, 1.000000e+00, 1.250000e-01, 1.000980e-01, 3.140630e+00, 0x7F80, 0xFF80, 0x7FFF, 9.183550e-41, -9.183550e-41]> : tensor<11xbf16>