diff --git a/mlir/include/mlir/IR/CommonTypeConstraints.td b/mlir/include/mlir/IR/CommonTypeConstraints.td index 6f52195c1d7c9..e752cdfb47fbb 100644 --- a/mlir/include/mlir/IR/CommonTypeConstraints.td +++ b/mlir/include/mlir/IR/CommonTypeConstraints.td @@ -329,31 +329,31 @@ def F64 : F<64>; def F80 : F<80>; def F128 : F<128>; -def BF16 : Type, "bfloat16 type">, +def BF16 : Type($_self)">, "bfloat16 type">, BuildableType<"$_builder.getType()">; -def TF32 : Type, "tf32 type">, +def TF32 : Type($_self)">, "tf32 type">, BuildableType<"$_builder.getType()">; -def F8E4M3FN : Type, "f8E4M3FN type">, +def F8E4M3FN : Type($_self)">, "f8E4M3FN type">, BuildableType<"$_builder.getType()">; -def F8E5M2 : Type, "f8E5M2 type">, +def F8E5M2 : Type($_self)">, "f8E5M2 type">, BuildableType<"$_builder.getType()">; -def F8E4M3 : Type, "f8E4M3 type">, +def F8E4M3 : Type($_self)">, "f8E4M3 type">, BuildableType<"$_builder.getType()">; -def F8E4M3FNUZ : Type, "f8E4M3FNUZ type">, +def F8E4M3FNUZ : Type($_self)">, "f8E4M3FNUZ type">, BuildableType<"$_builder.getType()">; -def F8E4M3B11FNUZ : Type, "f8E4M3B11FNUZ type">, +def F8E4M3B11FNUZ : Type($_self)">, "f8E4M3B11FNUZ type">, BuildableType<"$_builder.getType()">; -def F8E5M2FNUZ : Type, "f8E5M2FNUZ type">, +def F8E5M2FNUZ : Type($_self)">, "f8E5M2FNUZ type">, BuildableType<"$_builder.getType()">; -def F8E3M4 : Type, "f8E3M4 type">, +def F8E3M4 : Type($_self)">, "f8E3M4 type">, BuildableType<"$_builder.getType()">; -def F4E2M1FN : Type, "f4E2M1FN type">, +def F4E2M1FN : Type($_self)">, "f4E2M1FN type">, BuildableType<"$_builder.getType()">; -def F6E2M3FN : Type, "f6E2M3FN type">, +def F6E2M3FN : Type($_self)">, "f6E2M3FN type">, BuildableType<"$_builder.getType()">; -def F6E3M2FN : Type, "f6E3M2FN type">, +def F6E3M2FN : Type, "f6E3M2FN type">, BuildableType<"$_builder.getType()">; -def F8E8M0FNU : Type, "f8E8M0FNU type">, +def F8E8M0FNU : Type($_self)">, "f8E8M0FNU type">, BuildableType<"$_builder.getType()">; def AnyComplex : Type($_self)">, diff --git a/mlir/include/mlir/IR/Types.h b/mlir/include/mlir/IR/Types.h index acd0f894abbbe..0e82ad2be907a 100644 --- a/mlir/include/mlir/IR/Types.h +++ b/mlir/include/mlir/IR/Types.h @@ -125,17 +125,6 @@ class Type { // Convenience predicates. This is only for floating point types, // derived types should use isa/dyn_cast. bool isIndex() const; - bool isFloat4E2M1FN() const; - bool isFloat6E2M3FN() const; - bool isFloat6E3M2FN() const; - bool isFloat8E5M2() const; - bool isFloat8E4M3() const; - bool isFloat8E4M3FN() const; - bool isFloat8E5M2FNUZ() const; - bool isFloat8E4M3FNUZ() const; - bool isFloat8E4M3B11FNUZ() const; - bool isFloat8E3M4() const; - bool isFloat8E8M0FNU() const; bool isBF16() const; bool isF16() const; bool isTF32() const; diff --git a/mlir/lib/CAPI/IR/BuiltinTypes.cpp b/mlir/lib/CAPI/IR/BuiltinTypes.cpp index 250e4a6bbf8df..98ca9c3d23909 100644 --- a/mlir/lib/CAPI/IR/BuiltinTypes.cpp +++ b/mlir/lib/CAPI/IR/BuiltinTypes.cpp @@ -90,7 +90,7 @@ MlirTypeID mlirFloat4E2M1FNTypeGetTypeID() { } bool mlirTypeIsAFloat4E2M1FN(MlirType type) { - return unwrap(type).isFloat4E2M1FN(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat4E2M1FNTypeGet(MlirContext ctx) { @@ -102,7 +102,7 @@ MlirTypeID mlirFloat6E2M3FNTypeGetTypeID() { } bool mlirTypeIsAFloat6E2M3FN(MlirType type) { - return unwrap(type).isFloat6E2M3FN(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat6E2M3FNTypeGet(MlirContext ctx) { @@ -114,7 +114,7 @@ MlirTypeID mlirFloat6E3M2FNTypeGetTypeID() { } bool mlirTypeIsAFloat6E3M2FN(MlirType type) { - return unwrap(type).isFloat6E3M2FN(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat6E3M2FNTypeGet(MlirContext ctx) { @@ -126,7 +126,7 @@ MlirTypeID mlirFloat8E5M2TypeGetTypeID() { } bool mlirTypeIsAFloat8E5M2(MlirType type) { - return unwrap(type).isFloat8E5M2(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E5M2TypeGet(MlirContext ctx) { @@ -138,7 +138,7 @@ MlirTypeID mlirFloat8E4M3TypeGetTypeID() { } bool mlirTypeIsAFloat8E4M3(MlirType type) { - return unwrap(type).isFloat8E4M3(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E4M3TypeGet(MlirContext ctx) { @@ -150,7 +150,7 @@ MlirTypeID mlirFloat8E4M3FNTypeGetTypeID() { } bool mlirTypeIsAFloat8E4M3FN(MlirType type) { - return unwrap(type).isFloat8E4M3FN(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E4M3FNTypeGet(MlirContext ctx) { @@ -162,7 +162,7 @@ MlirTypeID mlirFloat8E5M2FNUZTypeGetTypeID() { } bool mlirTypeIsAFloat8E5M2FNUZ(MlirType type) { - return unwrap(type).isFloat8E5M2FNUZ(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E5M2FNUZTypeGet(MlirContext ctx) { @@ -174,7 +174,7 @@ MlirTypeID mlirFloat8E4M3FNUZTypeGetTypeID() { } bool mlirTypeIsAFloat8E4M3FNUZ(MlirType type) { - return unwrap(type).isFloat8E4M3FNUZ(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E4M3FNUZTypeGet(MlirContext ctx) { @@ -186,7 +186,7 @@ MlirTypeID mlirFloat8E4M3B11FNUZTypeGetTypeID() { } bool mlirTypeIsAFloat8E4M3B11FNUZ(MlirType type) { - return unwrap(type).isFloat8E4M3B11FNUZ(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E4M3B11FNUZTypeGet(MlirContext ctx) { @@ -198,7 +198,7 @@ MlirTypeID mlirFloat8E3M4TypeGetTypeID() { } bool mlirTypeIsAFloat8E3M4(MlirType type) { - return unwrap(type).isFloat8E3M4(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E3M4TypeGet(MlirContext ctx) { @@ -210,7 +210,7 @@ MlirTypeID mlirFloat8E8M0FNUTypeGetTypeID() { } bool mlirTypeIsAFloat8E8M0FNU(MlirType type) { - return unwrap(type).isFloat8E8M0FNU(); + return llvm::isa(unwrap(type)); } MlirType mlirFloat8E8M0FNUTypeGet(MlirContext ctx) { @@ -221,7 +221,9 @@ MlirTypeID mlirBFloat16TypeGetTypeID() { return wrap(BFloat16Type::getTypeID()); } -bool mlirTypeIsABF16(MlirType type) { return unwrap(type).isBF16(); } +bool mlirTypeIsABF16(MlirType type) { + return llvm::isa(unwrap(type)); +} MlirType mlirBF16TypeGet(MlirContext ctx) { return wrap(BFloat16Type::get(unwrap(ctx))); @@ -229,7 +231,9 @@ MlirType mlirBF16TypeGet(MlirContext ctx) { MlirTypeID mlirFloat16TypeGetTypeID() { return wrap(Float16Type::getTypeID()); } -bool mlirTypeIsAF16(MlirType type) { return unwrap(type).isF16(); } +bool mlirTypeIsAF16(MlirType type) { + return llvm::isa(unwrap(type)); +} MlirType mlirF16TypeGet(MlirContext ctx) { return wrap(Float16Type::get(unwrap(ctx))); @@ -239,7 +243,9 @@ MlirTypeID mlirFloatTF32TypeGetTypeID() { return wrap(FloatTF32Type::getTypeID()); } -bool mlirTypeIsATF32(MlirType type) { return unwrap(type).isTF32(); } +bool mlirTypeIsATF32(MlirType type) { + return llvm::isa(unwrap(type)); +} MlirType mlirTF32TypeGet(MlirContext ctx) { return wrap(FloatTF32Type::get(unwrap(ctx))); @@ -247,7 +253,9 @@ MlirType mlirTF32TypeGet(MlirContext ctx) { MlirTypeID mlirFloat32TypeGetTypeID() { return wrap(Float32Type::getTypeID()); } -bool mlirTypeIsAF32(MlirType type) { return unwrap(type).isF32(); } +bool mlirTypeIsAF32(MlirType type) { + return llvm::isa(unwrap(type)); +} MlirType mlirF32TypeGet(MlirContext ctx) { return wrap(Float32Type::get(unwrap(ctx))); @@ -255,7 +263,9 @@ MlirType mlirF32TypeGet(MlirContext ctx) { MlirTypeID mlirFloat64TypeGetTypeID() { return wrap(Float64Type::getTypeID()); } -bool mlirTypeIsAF64(MlirType type) { return unwrap(type).isF64(); } +bool mlirTypeIsAF64(MlirType type) { + return llvm::isa(unwrap(type)); +} MlirType mlirF64TypeGet(MlirContext ctx) { return wrap(Float64Type::get(unwrap(ctx))); diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp index 1564e417a7a48..5d09d6f1d6952 100644 --- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp +++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp @@ -564,38 +564,40 @@ static std::optional mfmaOpToIntrinsic(MFMAOp mfma, return ROCDL::mfma_f64_4x4x4f64::getOperationName(); } - if (sourceElem.isFloat8E5M2FNUZ() && destElem.isF32() && chipset >= kGfx940) { + if (isa(sourceElem) && destElem.isF32() && + chipset >= kGfx940) { // Known to be correct because there are no scalar f8 instructions and // because a length mismatch will have been caught by the verifier. Type sourceBElem = cast(mfma.getSourceB().getType()).getElementType(); if (m == 16 && n == 16 && k == 32 && b == 1) { - if (sourceBElem.isFloat8E5M2FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_16x16x32_bf8_bf8::getOperationName(); - if (sourceBElem.isFloat8E4M3FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_16x16x32_bf8_fp8::getOperationName(); } if (m == 32 && n == 32 && k == 16 && b == 1) { - if (sourceBElem.isFloat8E5M2FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_32x32x16_bf8_bf8::getOperationName(); - if (sourceBElem.isFloat8E4M3FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_32x32x16_bf8_fp8::getOperationName(); } } - if (sourceElem.isFloat8E4M3FNUZ() && destElem.isF32() && chipset >= kGfx940) { + if (isa(sourceElem) && destElem.isF32() && + chipset >= kGfx940) { Type sourceBElem = cast(mfma.getSourceB().getType()).getElementType(); if (m == 16 && n == 16 && k == 32 && b == 1) { - if (sourceBElem.isFloat8E5M2FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_16x16x32_fp8_bf8::getOperationName(); - if (sourceBElem.isFloat8E4M3FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_16x16x32_fp8_fp8::getOperationName(); } if (m == 32 && n == 32 && k == 16 && b == 1) { - if (sourceBElem.isFloat8E5M2FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_32x32x16_fp8_bf8::getOperationName(); - if (sourceBElem.isFloat8E4M3FNUZ()) + if (isa(sourceBElem)) return ROCDL::mfma_f32_32x32x16_fp8_fp8::getOperationName(); } } @@ -623,9 +625,9 @@ static std::optional wmmaOpToIntrinsic(WMMAOp wmma, return ROCDL::wmma_bf16_16x16x16_bf16::getOperationName(); if (elemSourceType.isInteger(8) && elemDestType.isInteger(32)) return ROCDL::wmma_i32_16x16x16_iu8::getOperationName(); - if (elemSourceType.isFloat8E4M3FN() && elemDestType.isF32()) + if (isa(elemSourceType) && elemDestType.isF32()) return ROCDL::wmma_f32_16x16x16_fp8::getOperationName(); - if (elemSourceType.isFloat8E5M2() && elemDestType.isF32()) + if (isa(elemSourceType) && elemDestType.isF32()) return ROCDL::wmma_f32_16x16x16_bf8::getOperationName(); return std::nullopt; } @@ -803,10 +805,10 @@ LogicalResult ExtPackedFp8OpLowering::matchAndRewrite( } Value i32Source = rewriter.create(loc, i32, source); Value wordSel = createI32Constant(rewriter, loc, op.getIndex()); - if (sourceElemType.isFloat8E5M2FNUZ()) { + if (isa(sourceElemType)) { rewriter.replaceOpWithNewOp(op, f32, i32Source, wordSel); - } else if (sourceElemType.isFloat8E4M3FNUZ()) { + } else if (isa(sourceElemType)) { rewriter.replaceOpWithNewOp(op, f32, i32Source, wordSel); } @@ -838,10 +840,10 @@ LogicalResult PackedTrunc2xFp8OpLowering::matchAndRewrite( Value wordSel = createI1Constant(rewriter, loc, op.getWordIndex()); Value result; - if (resultElemType.isFloat8E5M2FNUZ()) + if (isa(resultElemType)) result = rewriter.create(loc, i32, sourceA, sourceB, existing, wordSel); - else if (resultElemType.isFloat8E4M3FNUZ()) + else if (isa(resultElemType)) result = rewriter.create(loc, i32, sourceA, sourceB, existing, wordSel); @@ -873,10 +875,10 @@ LogicalResult PackedStochRoundFp8OpLowering::matchAndRewrite( Value byteSel = createI32Constant(rewriter, loc, op.getStoreIndex()); Value result; - if (resultElemType.isFloat8E5M2FNUZ()) + if (isa(resultElemType)) result = rewriter.create(loc, i32, source, stoch, existing, byteSel); - else if (resultElemType.isFloat8E4M3FNUZ()) + else if (isa(resultElemType)) result = rewriter.create(loc, i32, source, stoch, existing, byteSel); diff --git a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp index a8283023afc53..33370566996ee 100644 --- a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp +++ b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp @@ -86,7 +86,7 @@ LogicalResult ExtFOnFloat8RewritePattern::match(arith::ExtFOp op) const { return failure(); inType = inVecType.getElementType(); } - return success(inType.isFloat8E5M2FNUZ() || inType.isFloat8E4M3FNUZ()); + return success(isa(inType)); } void ExtFOnFloat8RewritePattern::rewrite(arith::ExtFOp op, @@ -216,7 +216,7 @@ LogicalResult TruncFToFloat8RewritePattern::match(arith::TruncFOp op) const { if (inType && inType.getWidth() <= 8 && saturateFP8) // Conversion between 8-bit floats is not supported with truncation enabled. return failure(); - return success(outType.isFloat8E5M2FNUZ() || outType.isFloat8E4M3FNUZ()); + return success(isa(outType)); } void TruncFToFloat8RewritePattern::rewrite(arith::TruncFOp op, diff --git a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp index 64bdb248dff43..247a8ab28a44b 100644 --- a/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp +++ b/mlir/lib/Conversion/LLVMCommon/TypeConverter.cpp @@ -299,11 +299,10 @@ Type LLVMTypeConverter::convertFloatType(FloatType type) const { return type; // F4, F6, F8 types are converted to integer types with the same bit width. - if (type.isFloat8E5M2() || type.isFloat8E4M3() || type.isFloat8E4M3FN() || - type.isFloat8E5M2FNUZ() || type.isFloat8E4M3FNUZ() || - type.isFloat8E4M3B11FNUZ() || type.isFloat8E3M4() || - type.isFloat4E2M1FN() || type.isFloat6E2M3FN() || type.isFloat6E3M2FN() || - type.isFloat8E8M0FNU()) + if (isa(type)) return IntegerType::get(&getContext(), type.getWidth()); // Other floating-point types: A custom type conversion rule must be diff --git a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp index 34a6b1d506540..7e97fb84434f8 100644 --- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp +++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp @@ -1254,8 +1254,8 @@ struct NVGPUWarpgroupMmaOpLowering wgmmaK = 8; } else if (inputElemType.isF16() || inputElemType.isBF16()) { wgmmaK = 16; - } else if (inputElemType.isFloat8E4M3FN() || - inputElemType.isFloat8E5M2() || inputElemType.isInteger(16)) { + } else if (isa(inputElemType) || + inputElemType.isInteger(16)) { wgmmaK = 32; } else if (inputElemType.isInteger(1)) { wgmmaK = 256; @@ -1276,9 +1276,9 @@ struct NVGPUWarpgroupMmaOpLowering return NVVM::WGMMATypes::f16; if (elemType.isBF16()) return NVVM::WGMMATypes::bf16; - if (elemType.isFloat8E4M3FN()) + if (isa(elemType)) return NVVM::WGMMATypes::e4m3; - if (elemType.isFloat8E5M2()) + if (isa(elemType)) return NVVM::WGMMATypes::e5m2; if (elemType.isInteger(1)) return NVVM::WGMMATypes::b1; diff --git a/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp b/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp index 492e4781f5781..5af0cb0c7ba1c 100644 --- a/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp +++ b/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp @@ -272,14 +272,14 @@ LogicalResult MFMAOp::verify() { } Type sourceBType = getSourceB().getType(); - if (sourceElem.isFloat8E5M2FNUZ() || sourceElem.isFloat8E4M3FNUZ()) { + if (isa(sourceElem)) { int64_t sourceBLen = 1; Type sourceBElem = sourceBType; if (auto sourceBVector = llvm::dyn_cast(sourceBType)) { sourceBLen = sourceBVector.getNumElements(); sourceBElem = sourceBVector.getElementType(); } - if (!sourceBElem.isFloat8E5M2FNUZ() && !sourceBElem.isFloat8E4M3FNUZ()) + if (!isa(sourceBElem)) return emitOpError("expected both source operands to have f8 elements"); if (sourceLen != sourceBLen) return emitOpError( diff --git a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp index a027350e8a5f7..47d1b8492e06e 100644 --- a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp +++ b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp @@ -525,8 +525,8 @@ LogicalResult isAllowedWGMMADataType(Type typeD, Type typeA, Type typeB) { return success(); // F16 += f8 + f8 // F32 += f8 + f8 - if ((typeA.isFloat8E5M2() || typeA.isFloat8E4M3FN()) && - (typeB.isFloat8E5M2() || typeB.isFloat8E4M3FN()) && + if (isa(typeA) && + isa(typeB) && (typeD.isF32() || typeD.isF16())) return success(); @@ -548,7 +548,7 @@ LogicalResult isAllowedSizeN(int sizeN, Type typeA) { 80, 96, 112, 128, 144, 160, 176, 192, 208, 224, 240, 256}; if (typeA.isBF16() || typeA.isF16() || typeA.isF32() || typeA.isTF32() || - typeA.isFloat8E4M3FN() || typeA.isFloat8E5M2()) + isa(typeA)) if (llvm::is_contained(allowedN, sizeN)) return success(); diff --git a/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp b/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp index 83cf4a9415fe6..75b7da708cbb5 100644 --- a/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp +++ b/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp @@ -306,8 +306,7 @@ static LogicalResult verifyConvOpModes(T op) { if (inputEType.isInteger(16) && !accType.isInteger(48)) return op.emitOpError("accumulator type for i16 tensor is not i48"); - if ((inputEType.isFloat8E5M2() || inputEType.isFloat8E4M3()) && - !accType.isF16()) + if (isa(inputEType) && !accType.isF16()) return op.emitOpError("accumulator type for f8 tensor is not f16"); if (inputEType.isF16() && !(accType.isF16() || accType.isF32())) diff --git a/mlir/lib/IR/Types.cpp b/mlir/lib/IR/Types.cpp index e190902b2e489..bca90de6f4a8a 100644 --- a/mlir/lib/IR/Types.cpp +++ b/mlir/lib/IR/Types.cpp @@ -34,25 +34,6 @@ Type AbstractType::replaceImmediateSubElements(Type type, MLIRContext *Type::getContext() const { return getDialect().getContext(); } -bool Type::isFloat4E2M1FN() const { return llvm::isa(*this); } -bool Type::isFloat6E2M3FN() const { return llvm::isa(*this); } -bool Type::isFloat6E3M2FN() const { return llvm::isa(*this); } -bool Type::isFloat8E5M2() const { return llvm::isa(*this); } -bool Type::isFloat8E4M3() const { return llvm::isa(*this); } -bool Type::isFloat8E4M3FN() const { return llvm::isa(*this); } -bool Type::isFloat8E5M2FNUZ() const { - return llvm::isa(*this); -} -bool Type::isFloat8E4M3FNUZ() const { - return llvm::isa(*this); -} -bool Type::isFloat8E4M3B11FNUZ() const { - return llvm::isa(*this); -} -bool Type::isFloat8E8M0FNU() const { - return llvm::isa(*this); -} -bool Type::isFloat8E3M4() const { return llvm::isa(*this); } bool Type::isBF16() const { return llvm::isa(*this); } bool Type::isF16() const { return llvm::isa(*this); } bool Type::isTF32() const { return llvm::isa(*this); }