From 62bfd318aefae4b5a1f50813f90bf0b8b821849b Mon Sep 17 00:00:00 2001 From: Zhaoming Jiang Date: Fri, 13 May 2022 12:01:11 +0000 Subject: [PATCH] tint: Implement `f16` keyword in Tint frontend This patch: 1. Add the `f16` WGSL extension. 2. Add `f16` as keyword, and remove it from reserved word list. 3. Add ast::f16 and sem::f16, and implement validation that using `f16` type must be with `f16` extension enabled. 4. Add `Number` for f16 literal and constant, and add `ast::FloatLiteralExpression::Suffix::kH`. 5. Add placeholder in all writer which report error when try to emit f16 type. Bugs: tint:1473, tint:1502 Change-Id: Ifb363beeb2699ed7cac57e07227d1b2cfa8050b4 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/89922 Reviewed-by: Ben Clayton Kokoro: Kokoro Commit-Queue: Zhaoming Jiang --- src/tint/BUILD.gn | 7 ++- src/tint/CMakeLists.txt | 6 ++ src/tint/ast/enable.cc | 6 +- src/tint/ast/enable.h | 2 + src/tint/ast/f16.cc | 38 ++++++++++++ src/tint/ast/f16.h | 48 +++++++++++++++ src/tint/ast/f16_test.cc | 30 ++++++++++ src/tint/ast/float_literal_expression.h | 2 + src/tint/inspector/inspector.cc | 1 + src/tint/number.h | 56 ++++++++++++++++++ src/tint/program_builder.h | 30 ++++++++++ src/tint/reader/wgsl/lexer.cc | 2 + src/tint/reader/wgsl/parser_impl.cc | 15 ++++- .../wgsl/parser_impl_reserved_keyword_test.cc | 1 - .../reader/wgsl/parser_impl_type_decl_test.cc | 11 ++++ src/tint/reader/wgsl/token.cc | 2 + src/tint/reader/wgsl/token.h | 2 + src/tint/resolver/dependency_graph.cc | 2 +- src/tint/resolver/resolver.cc | 8 +++ src/tint/resolver/resolver_constants.cc | 9 +++ src/tint/resolver/type_validation_test.cc | 18 ++++++ src/tint/resolver/validator.cc | 6 +- src/tint/sem/constant.h | 7 +++ src/tint/sem/f16.cc | 55 ++++++++++++++++++ src/tint/sem/f16.h | 58 +++++++++++++++++++ src/tint/sem/f16_test.cc | 48 +++++++++++++++ src/tint/sem/type.cc | 7 ++- src/tint/transform/transform.cc | 3 + src/tint/writer/glsl/generator_impl.cc | 3 + src/tint/writer/hlsl/generator_impl.cc | 5 ++ src/tint/writer/msl/generator_impl.cc | 11 ++++ src/tint/writer/spirv/builder.cc | 7 +++ src/tint/writer/wgsl/generator_impl.cc | 5 ++ test/tint/BUILD.gn | 2 + 34 files changed, 502 insertions(+), 11 deletions(-) create mode 100644 src/tint/ast/f16.cc create mode 100644 src/tint/ast/f16.h create mode 100644 src/tint/ast/f16_test.cc create mode 100644 src/tint/sem/f16.cc create mode 100644 src/tint/sem/f16.h create mode 100644 src/tint/sem/f16_test.cc diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 4a22a67f72..9a78107b26 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -228,6 +228,8 @@ libtint_source_set("libtint_core_all_src") { "ast/expression.h", "ast/external_texture.cc", "ast/external_texture.h", + "ast/f16.cc", + "ast/f16.h", "ast/f32.cc", "ast/f32.h", "ast/fallthrough_statement.cc", @@ -342,7 +344,6 @@ libtint_source_set("libtint_core_all_src") { "debug.h", "demangler.cc", "demangler.h", - "number.h", "diagnostic/diagnostic.cc", "diagnostic/diagnostic.h", "diagnostic/formatter.cc", @@ -357,6 +358,7 @@ libtint_source_set("libtint_core_all_src") { "inspector/resource_binding.h", "inspector/scalar.cc", "inspector/scalar.h", + "number.h", "program.cc", "program.h", "program_builder.cc", @@ -394,6 +396,7 @@ libtint_source_set("libtint_core_all_src") { "sem/depth_texture.h", "sem/expression.h", "sem/external_texture.h", + "sem/f16.h", "sem/f32.h", "sem/for_loop_statement.h", "sem/i32.h", @@ -591,6 +594,8 @@ libtint_source_set("libtint_sem_src") { "sem/expression.h", "sem/external_texture.cc", "sem/external_texture.h", + "sem/f16.cc", + "sem/f16.h", "sem/f32.cc", "sem/f32.h", "sem/for_loop_statement.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index d2273fdd18..d49259538d 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -116,6 +116,8 @@ set(TINT_LIB_SRCS ast/expression.h ast/external_texture.cc ast/external_texture.h + ast/f16.cc + ast/f16.h ast/f32.cc ast/f32.h ast/fallthrough_statement.cc @@ -403,6 +405,8 @@ set(TINT_LIB_SRCS sem/depth_texture.h sem/external_texture.cc sem/external_texture.h + sem/f16.cc + sem/f16.h sem/f32.cc sem/f32.h sem/for_loop_statement.cc @@ -686,6 +690,7 @@ if(TINT_BUILD_TESTS) ast/discard_statement_test.cc ast/enable_test.cc ast/external_texture_test.cc + ast/f16_test.cc ast/f32_test.cc ast/fallthrough_statement_test.cc ast/float_literal_expression_test.cc @@ -790,6 +795,7 @@ if(TINT_BUILD_TESTS) sem/depth_multisampled_texture_test.cc sem/depth_texture_test.cc sem/external_texture_test.cc + sem/f16_test.cc sem/f32_test.cc sem/i32_test.cc sem/matrix_test.cc diff --git a/src/tint/ast/enable.cc b/src/tint/ast/enable.cc index a611cfbe19..200c2bec00 100644 --- a/src/tint/ast/enable.cc +++ b/src/tint/ast/enable.cc @@ -28,6 +28,9 @@ Enable::ExtensionKind Enable::NameToKind(const std::string& name) { if (name == "chromium_disable_uniformity_analysis") { return Enable::ExtensionKind::kChromiumDisableUniformityAnalysis; } + if (name == "f16") { + return Enable::ExtensionKind::kF16; + } // The reserved internal extension name for testing if (name == "InternalExtensionForTesting") { @@ -43,7 +46,8 @@ std::string Enable::KindToName(ExtensionKind kind) { return "chromium_experimental_dp4a"; case ExtensionKind::kChromiumDisableUniformityAnalysis: return "chromium_disable_uniformity_analysis"; - + case ExtensionKind::kF16: + return "f16"; // The reserved internal extension for testing case ExtensionKind::kInternalExtensionForTesting: return "InternalExtensionForTesting"; diff --git a/src/tint/ast/enable.h b/src/tint/ast/enable.h index 69622e0033..f190b0a792 100644 --- a/src/tint/ast/enable.h +++ b/src/tint/ast/enable.h @@ -35,6 +35,8 @@ class Enable : public Castable { /// An internal reserved extension for test, named /// "InternalExtensionForTesting". kInternalExtensionForTesting, + /// WGSL Extension "f16" + kF16, /// An extension for the experimental feature /// "chromium_experimental_dp4a". diff --git a/src/tint/ast/f16.cc b/src/tint/ast/f16.cc new file mode 100644 index 0000000000..0eb1be5fdd --- /dev/null +++ b/src/tint/ast/f16.cc @@ -0,0 +1,38 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/ast/f16.h" + +#include "src/tint/program_builder.h" + +TINT_INSTANTIATE_TYPEINFO(tint::ast::F16); + +namespace tint::ast { + +F16::F16(ProgramID pid, const Source& src) : Base(pid, src) {} + +F16::F16(F16&&) = default; + +F16::~F16() = default; + +std::string F16::FriendlyName(const SymbolTable&) const { + return "f16"; +} + +const F16* F16::Clone(CloneContext* ctx) const { + auto src = ctx->Clone(source); + return ctx->dst->create(src); +} + +} // namespace tint::ast diff --git a/src/tint/ast/f16.h b/src/tint/ast/f16.h new file mode 100644 index 0000000000..1b84f09529 --- /dev/null +++ b/src/tint/ast/f16.h @@ -0,0 +1,48 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef SRC_TINT_AST_F16_H_ +#define SRC_TINT_AST_F16_H_ + +#include + +#include "src/tint/ast/type.h" + +namespace tint::ast { + +/// A float 16 type +class F16 : public Castable { + public: + /// Constructor + /// @param pid the identifier of the program that owns this node + /// @param src the source of this node + F16(ProgramID pid, const Source& src); + /// Move constructor + F16(F16&&); + ~F16() override; + + /// @param symbols the program's symbol table + /// @returns the name for this type that closely resembles how it would be + /// declared in WGSL. + std::string FriendlyName(const SymbolTable& symbols) const override; + + /// Clones this type and all transitive types using the `CloneContext` `ctx`. + /// @param ctx the clone context + /// @return the newly cloned type + const F16* Clone(CloneContext* ctx) const override; +}; + +} // namespace tint::ast + +#endif // SRC_TINT_AST_F16_H_ diff --git a/src/tint/ast/f16_test.cc b/src/tint/ast/f16_test.cc new file mode 100644 index 0000000000..48ab284f55 --- /dev/null +++ b/src/tint/ast/f16_test.cc @@ -0,0 +1,30 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/ast/f16.h" + +#include "src/tint/ast/test_helper.h" + +namespace tint::ast { +namespace { + +using AstF16Test = TestHelper; + +TEST_F(AstF16Test, FriendlyName) { + auto* f = create(); + EXPECT_EQ(f->FriendlyName(Symbols()), "f16"); +} + +} // namespace +} // namespace tint::ast diff --git a/src/tint/ast/float_literal_expression.h b/src/tint/ast/float_literal_expression.h index 321efc891c..72a395f8d4 100644 --- a/src/tint/ast/float_literal_expression.h +++ b/src/tint/ast/float_literal_expression.h @@ -30,6 +30,8 @@ class FloatLiteralExpression final : public Castable #include +namespace tint::detail { +/// An empty structure used as a unique template type for Number when +/// specializing for the f16 type. +struct NumberKindF16 {}; +} // namespace tint::detail + namespace tint { /// Number wraps a integer or floating point number, enforcing explicit casting. @@ -72,6 +78,43 @@ bool operator==(A a, Number b) { return Number(a) == b; } +/// The partial specification of Number for f16 type, storing the f16 value as float, +/// and enforcing proper explicit casting. +template <> +struct Number { + /// Constructor. The value is zero-initialized. + Number() = default; + + /// Constructor. + /// @param v the value to initialize this Number to + template + explicit Number(U v) : value(static_cast(v)) {} + + /// Constructor. + /// @param v the value to initialize this Number to + template + explicit Number(Number v) : value(static_cast(v.value)) {} + + /// Conversion operator + /// @returns the value as the internal representation type of F16 + operator float() const { return value; } + + /// Negation operator + /// @returns the negative value of the number + Number operator-() const { return Number(-value); } + + /// Assignment operator with parameter as native floating point type + /// @param v the new value + /// @returns this Number so calls can be chained + Number& operator=(float v) { + value = v; + return *this; + } + + /// The number value, stored as float + float value = {}; +}; + /// `AInt` is a type alias to `Number`. using AInt = Number; /// `AFloat` is a type alias to `Number`. @@ -83,6 +126,9 @@ using i32 = Number; using u32 = Number; /// `f32` is a type alias to `Number` using f32 = Number; +/// `f16` is a type alias to `Number`, which should be IEEE 754 binary16. +/// However since C++ don't have native binary16 type, the value is stored as float. +using f16 = Number; } // namespace tint @@ -118,6 +164,16 @@ inline f32 operator"" _f(unsigned long long int value) { // NOLINT return f32(static_cast(value)); } +/// Literal suffix for f16 literals +inline f16 operator"" _h(long double value) { // NOLINT + return f16(static_cast(value)); +} + +/// Literal suffix for f16 literals +inline f16 operator"" _h(unsigned long long int value) { // NOLINT + return f16(static_cast(value)); +} + } // namespace tint::number_suffixes #endif // SRC_TINT_NUMBER_H_ diff --git a/src/tint/program_builder.h b/src/tint/program_builder.h index 7dd2c814b3..a91d87c82e 100644 --- a/src/tint/program_builder.h +++ b/src/tint/program_builder.h @@ -40,6 +40,7 @@ #include "src/tint/ast/discard_statement.h" #include "src/tint/ast/enable.h" #include "src/tint/ast/external_texture.h" +#include "src/tint/ast/f16.h" #include "src/tint/ast/f32.h" #include "src/tint/ast/fallthrough_statement.h" #include "src/tint/ast/float_literal_expression.h" @@ -82,6 +83,7 @@ #include "src/tint/sem/bool.h" #include "src/tint/sem/depth_texture.h" #include "src/tint/sem/external_texture.h" +#include "src/tint/sem/f16.h" #include "src/tint/sem/f32.h" #include "src/tint/sem/i32.h" #include "src/tint/sem/matrix.h" @@ -385,6 +387,15 @@ class ProgramBuilder { return builder->create(source); } + /// @returns a f16 type + const ast::F16* f16() const { return builder->create(); } + + /// @param source the Source of the node + /// @returns a f16 type + const ast::F16* f16(const Source& source) const { + return builder->create(source); + } + /// @returns a f32 type const ast::F32* f32() const { return builder->create(); } @@ -1004,6 +1015,21 @@ class ProgramBuilder { ast::FloatLiteralExpression::Suffix::kF); } + /// @param source the source information + /// @param value the float value + /// @return a 'h'-suffixed FloatLiteralExpression for the f16 value + const ast::FloatLiteralExpression* Expr(const Source& source, f16 value) { + return create(source, static_cast(value.value), + ast::FloatLiteralExpression::Suffix::kH); + } + + /// @param value the float value + /// @return a 'h'-suffixed FloatLiteralExpression for the f16 value + const ast::FloatLiteralExpression* Expr(f16 value) { + return create(static_cast(value.value), + ast::FloatLiteralExpression::Suffix::kH); + } + /// @param source the source information /// @param value the integer value /// @return an unsuffixed IntLiteralExpression for the AInt value @@ -2675,6 +2701,10 @@ struct ProgramBuilder::TypesBuilder::CToAST { static const ast::Type* get(const ProgramBuilder::TypesBuilder* t) { return t->f32(); } }; template <> +struct ProgramBuilder::TypesBuilder::CToAST { + static const ast::Type* get(const ProgramBuilder::TypesBuilder* t) { return t->f16(); } +}; +template <> struct ProgramBuilder::TypesBuilder::CToAST { static const ast::Type* get(const ProgramBuilder::TypesBuilder* t) { return t->bool_(); } }; diff --git a/src/tint/reader/wgsl/lexer.cc b/src/tint/reader/wgsl/lexer.cc index db75613fab..4b5f955dbc 100644 --- a/src/tint/reader/wgsl/lexer.cc +++ b/src/tint/reader/wgsl/lexer.cc @@ -1067,6 +1067,8 @@ Token Lexer::check_keyword(const Source& source, std::string_view str) { return {Token::Type::kElse, source, "else"}; if (str == "enable") return {Token::Type::kEnable, source, "enable"}; + if (str == "f16") + return {Token::Type::kF16, source, "f16"}; if (str == "f32") return {Token::Type::kF32, source, "f32"}; if (str == "fallthrough") diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc index e684790dcd..abd97e50ca 100644 --- a/src/tint/reader/wgsl/parser_impl.cc +++ b/src/tint/reader/wgsl/parser_impl.cc @@ -124,8 +124,8 @@ const char kWorkgroupSizeAttribute[] = "workgroup_size"; // https://gpuweb.github.io/gpuweb/wgsl.html#reserved-keywords bool is_reserved(Token t) { - return t == "asm" || t == "bf16" || t == "const" || t == "do" || t == "enum" || t == "f16" || - t == "f64" || t == "handle" || t == "i8" || t == "i16" || t == "i64" || t == "mat" || + return t == "asm" || t == "bf16" || t == "const" || t == "do" || t == "enum" || t == "f64" || + t == "handle" || t == "i8" || t == "i16" || t == "i64" || t == "mat" || t == "premerge" || t == "regardless" || t == "typedef" || t == "u8" || t == "u16" || t == "u64" || t == "unless" || t == "using" || t == "vec" || t == "void" || t == "while"; } @@ -310,6 +310,9 @@ void ParserImpl::translation_unit() { if (after_global_decl) { add_error(p, "enable directives must come before all global declarations"); } + } else if (ed.errored) { + // Found a invalid enable directive. + continue; } else { auto gd = global_decl(); @@ -345,6 +348,11 @@ Maybe ParserImpl::enable_directive() { synchronized_ = true; next(); name = {t.to_str(), t.source()}; + } else if (t.Is(Token::Type::kF16)) { + // `f16` is a valid extension name and also a keyword + synchronized_ = true; + next(); + name = {"f16", t.source()}; } else if (handle_error(t)) { // The token might itself be an error. return Failure::kErrored; @@ -976,6 +984,9 @@ Maybe ParserImpl::type_decl() { if (match(Token::Type::kBool, &source)) return builder_.ty.bool_(source); + if (match(Token::Type::kF16, &source)) + return builder_.ty.f16(source); + if (match(Token::Type::kF32, &source)) return builder_.ty.f32(source); diff --git a/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc b/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc index 3226ef5e53..e840af7c0f 100644 --- a/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc +++ b/src/tint/reader/wgsl/parser_impl_reserved_keyword_test.cc @@ -88,7 +88,6 @@ INSTANTIATE_TEST_SUITE_P(ParserImplReservedKeywordTest, "const", "do", "enum", - "f16", "f64", "handle", "i8", diff --git a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc index a721d77ec1..bc2bfe5d66 100644 --- a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc +++ b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc @@ -55,6 +55,17 @@ TEST_F(ParserImplTest, TypeDecl_Bool) { EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 1u}, {1u, 5u}})); } +TEST_F(ParserImplTest, TypeDecl_F16) { + auto p = parser("f16"); + + auto t = p->type_decl(); + EXPECT_TRUE(t.matched); + EXPECT_FALSE(t.errored); + ASSERT_NE(t.value, nullptr) << p->error(); + ASSERT_TRUE(t.value->Is()); + EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 1u}, {1u, 4u}})); +} + TEST_F(ParserImplTest, TypeDecl_F32) { auto p = parser("f32"); diff --git a/src/tint/reader/wgsl/token.cc b/src/tint/reader/wgsl/token.cc index 06eb6fb18b..4680eee0d7 100644 --- a/src/tint/reader/wgsl/token.cc +++ b/src/tint/reader/wgsl/token.cc @@ -151,6 +151,8 @@ std::string_view Token::TypeToName(Type type) { return "else"; case Token::Type::kEnable: return "enable"; + case Token::Type::kF16: + return "f16"; case Token::Type::kF32: return "f32"; case Token::Type::kFallthrough: diff --git a/src/tint/reader/wgsl/token.h b/src/tint/reader/wgsl/token.h index 82b3fabda3..0a68f9b055 100644 --- a/src/tint/reader/wgsl/token.h +++ b/src/tint/reader/wgsl/token.h @@ -162,6 +162,8 @@ class Token { kElse, /// A 'enable' kEnable, + /// A 'f16' + kF16, /// A 'f32' kF32, /// A 'fallthrough' diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc index 94efafb3e7..7e668997c6 100644 --- a/src/tint/resolver/dependency_graph.cc +++ b/src/tint/resolver/dependency_graph.cc @@ -342,7 +342,7 @@ class DependencyScanner { TraverseType(tex->type); }, [&](Default) { - if (!ty->IsAnyOfIsAnyOf()) { UnhandledNode(diagnostics_, ty); diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc index 962dd02a05..38a502587d 100644 --- a/src/tint/resolver/resolver.cc +++ b/src/tint/resolver/resolver.cc @@ -172,6 +172,14 @@ sem::Type* Resolver::Type(const ast::Type* ty) { [&](const ast::Bool*) { return builder_->create(); }, [&](const ast::I32*) { return builder_->create(); }, [&](const ast::U32*) { return builder_->create(); }, + [&](const ast::F16* t) -> sem::F16* { + // Validate if f16 type is allowed. + if (builder_->AST().Extensions().count(ast::Enable::ExtensionKind::kF16) == 0) { + AddError("f16 used without 'f16' extension enabled", t->source); + return nullptr; + } + return builder_->create(); + }, [&](const ast::F32*) { return builder_->create(); }, [&](const ast::Vector* t) -> sem::Vector* { if (!t->type) { diff --git a/src/tint/resolver/resolver_constants.cc b/src/tint/resolver/resolver_constants.cc index 2a8a1d1086..3ec6f94a65 100644 --- a/src/tint/resolver/resolver_constants.cc +++ b/src/tint/resolver/resolver_constants.cc @@ -70,6 +70,10 @@ sem::Constant Resolver::EvaluateConstantValue(const ast::CallExpression* call, if (elem_type->Is()) { return sem::Constant(type, sem::Constant::Scalars(result_size, 0_u)); } + // Add f16 zero scalar here + if (elem_type->Is()) { + return sem::Constant(type, sem::Constant::Scalars(result_size, f16{0.f})); + } if (elem_type->Is()) { return sem::Constant(type, sem::Constant::Scalars(result_size, 0_f)); } @@ -120,6 +124,11 @@ sem::Constant Resolver::ConstantCast(const sem::Constant& value, return u32(static_cast(s)); }); }, + [&](const sem::F16*) { + return value.WithScalarAt(i, [](auto&& s) { // + return f16{static_cast(s)}; + }); + }, [&](const sem::F32*) { return value.WithScalarAt(i, [](auto&& s) { // return static_cast(s); diff --git a/src/tint/resolver/type_validation_test.cc b/src/tint/resolver/type_validation_test.cc index 90e18e15a0..4c2f56d1ef 100644 --- a/src/tint/resolver/type_validation_test.cc +++ b/src/tint/resolver/type_validation_test.cc @@ -662,6 +662,24 @@ TEST_F(ResolverTypeValidationTest, BuiltinAsType) { EXPECT_EQ(r()->error(), "error: cannot use builtin 'max' as type"); } +TEST_F(ResolverTypeValidationTest, F16TypeUsedWithExtension) { + // enable f16; + // var v : f16; + auto* ext = create("f16"); + AST().AddEnable(ext); + Global("v", ty.f16(), ast::StorageClass::kPrivate); + + EXPECT_TRUE(r()->Resolve()) << r()->error(); +} + +TEST_F(ResolverTypeValidationTest, F16TypeUsedWithoutExtension) { + // var v : f16; + Global("v", ty.f16(), ast::StorageClass::kPrivate); + + EXPECT_FALSE(r()->Resolve()); + EXPECT_EQ(r()->error(), "error: f16 used without 'f16' extension enabled"); +} + namespace GetCanonicalTests { struct Params { builder::ast_type_func_ptr create_ast_type; diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc index 7a66ec4f61..df9c7c08e0 100644 --- a/src/tint/resolver/validator.cc +++ b/src/tint/resolver/validator.cc @@ -194,7 +194,7 @@ bool Validator::IsFixedFootprint(const sem::Type* type) const { // https://gpuweb.github.io/gpuweb/wgsl.html#host-shareable-types bool Validator::IsHostShareable(const sem::Type* type) const { - if (type->IsAnyOf()) { + if (type->IsAnyOf()) { return true; } return Switch( @@ -1890,11 +1890,13 @@ bool Validator::ScalarConstructorOrCast(const ast::CallExpression* ctor, using Bool = sem::Bool; using I32 = sem::I32; using U32 = sem::U32; + using F16 = sem::F16; using F32 = sem::F32; const bool is_valid = (ty->Is() && value_ty->is_scalar()) || (ty->Is() && value_ty->is_scalar()) || - (ty->Is() && value_ty->is_scalar()) || (ty->Is() && value_ty->is_scalar()); + (ty->Is() && value_ty->is_scalar()) || (ty->Is() && value_ty->is_scalar()) || + (ty->Is() && value_ty->is_scalar()); if (!is_valid) { AddError("cannot construct '" + sem_.TypeNameOf(ty) + "' with a value of type '" + sem_.TypeNameOf(value_ty) + "'", diff --git a/src/tint/sem/constant.h b/src/tint/sem/constant.h index 673446f3cc..2e8258dd1c 100644 --- a/src/tint/sem/constant.h +++ b/src/tint/sem/constant.h @@ -35,6 +35,8 @@ class Constant { tint::u32 u32; /// The scalar value as a f32 tint::f32 f32; + /// The scalar value as a f16, internally stored as float + tint::f16 f16; /// The scalar value as a bool bool bool_; @@ -50,6 +52,10 @@ class Constant { /// @param v the value of the Scalar Scalar(tint::f32 v) : f32(v) {} // NOLINT + /// Constructs the scalar with the f16 value `v` + /// @param v the value of the Scalar + Scalar(tint::f16 v) : f16({v}) {} // NOLINT + /// Constructs the scalar with the bool value `v` /// @param v the value of the Scalar Scalar(bool v) : bool_(v) {} // NOLINT @@ -106,6 +112,7 @@ class Constant { ElementType(), // [&](const I32*) { return func(elems_[index].i32); }, [&](const U32*) { return func(elems_[index].u32); }, + [&](const F16*) { return func(elems_[index].f16); }, [&](const F32*) { return func(elems_[index].f32); }, [&](const Bool*) { return func(elems_[index].bool_); }, [&](Default) { diff --git a/src/tint/sem/f16.cc b/src/tint/sem/f16.cc new file mode 100644 index 0000000000..7da65fa91d --- /dev/null +++ b/src/tint/sem/f16.cc @@ -0,0 +1,55 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/sem/f16.h" + +#include "src/tint/program_builder.h" + +TINT_INSTANTIATE_TYPEINFO(tint::sem::F16); + +namespace tint { +namespace sem { + +F16::F16() = default; + +F16::F16(F16&&) = default; + +F16::~F16() = default; + +size_t F16::Hash() const { + return static_cast(TypeInfo::Of().full_hashcode); +} + +bool F16::Equals(const Type& other) const { + return other.Is(); +} + +std::string F16::FriendlyName(const SymbolTable&) const { + return "f16"; +} + +bool F16::IsConstructible() const { + return true; +} + +uint32_t F16::Size() const { + return 2; +} + +uint32_t F16::Align() const { + return 2; +} + +} // namespace sem +} // namespace tint diff --git a/src/tint/sem/f16.h b/src/tint/sem/f16.h new file mode 100644 index 0000000000..72984c13e7 --- /dev/null +++ b/src/tint/sem/f16.h @@ -0,0 +1,58 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef SRC_TINT_SEM_F16_H_ +#define SRC_TINT_SEM_F16_H_ + +#include + +#include "src/tint/sem/type.h" + +namespace tint::sem { + +/// A float 16 type +class F16 : public Castable { + public: + /// Constructor + F16(); + /// Move constructor + F16(F16&&); + ~F16() override; + + /// @returns a hash of the type. + size_t Hash() const override; + + /// @param other the other type to compare against + /// @returns true if the this type is equal to the given type + bool Equals(const Type& other) const override; + + /// @param symbols the program's symbol table + /// @returns the name for this type that closely resembles how it would be + /// declared in WGSL. + std::string FriendlyName(const SymbolTable& symbols) const override; + + /// @returns true if constructible as per + /// https://gpuweb.github.io/gpuweb/wgsl/#constructible-types + bool IsConstructible() const override; + + /// @returns the size in bytes of the type. + uint32_t Size() const override; + + /// @returns the alignment in bytes of the type. + uint32_t Align() const override; +}; + +} // namespace tint::sem + +#endif // SRC_TINT_SEM_F16_H_ diff --git a/src/tint/sem/f16_test.cc b/src/tint/sem/f16_test.cc new file mode 100644 index 0000000000..28fd0dab6b --- /dev/null +++ b/src/tint/sem/f16_test.cc @@ -0,0 +1,48 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/sem/test_helper.h" +#include "src/tint/sem/texture.h" + +namespace tint::sem { +namespace { + +using F16Test = TestHelper; + +TEST_F(F16Test, Creation) { + auto* a = create(); + auto* b = create(); + EXPECT_EQ(a, b); +} + +TEST_F(F16Test, Hash) { + auto* a = create(); + auto* b = create(); + EXPECT_EQ(a->Hash(), b->Hash()); +} + +TEST_F(F16Test, Equals) { + auto* a = create(); + auto* b = create(); + EXPECT_TRUE(a->Equals(*b)); + EXPECT_FALSE(a->Equals(Void{})); +} + +TEST_F(F16Test, FriendlyName) { + F16 f; + EXPECT_EQ(f.FriendlyName(Symbols()), "f16"); +} + +} // namespace +} // namespace tint::sem diff --git a/src/tint/sem/type.cc b/src/tint/sem/type.cc index 5776ff4fd2..0ee138da9c 100644 --- a/src/tint/sem/type.cc +++ b/src/tint/sem/type.cc @@ -15,6 +15,7 @@ #include "src/tint/sem/type.h" #include "src/tint/sem/bool.h" +#include "src/tint/sem/f16.h" #include "src/tint/sem/f32.h" #include "src/tint/sem/i32.h" #include "src/tint/sem/matrix.h" @@ -64,15 +65,15 @@ bool Type::IsConstructible() const { } bool Type::is_scalar() const { - return IsAnyOf(); + return IsAnyOf(); } bool Type::is_numeric_scalar() const { - return IsAnyOf(); + return IsAnyOf(); } bool Type::is_float_scalar() const { - return Is(); + return IsAnyOf(); } bool Type::is_float_matrix() const { diff --git a/src/tint/transform/transform.cc b/src/tint/transform/transform.cc index e3d0ea9475..f3ab173673 100644 --- a/src/tint/transform/transform.cc +++ b/src/tint/transform/transform.cc @@ -87,6 +87,9 @@ const ast::Type* Transform::CreateASTTypeFor(CloneContext& ctx, const sem::Type* if (ty->Is()) { return ctx.dst->create(); } + if (ty->Is()) { + return ctx.dst->create(); + } if (ty->Is()) { return ctx.dst->create(); } diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc index 98066d2efc..97450796aa 100644 --- a/src/tint/writer/glsl/generator_impl.cc +++ b/src/tint/writer/glsl/generator_impl.cc @@ -2565,6 +2565,9 @@ bool GeneratorImpl::EmitType(std::ostream& out, out << "bool"; } else if (type->Is()) { out << "float"; + } else if (type->Is()) { + diagnostics_.add_error(diag::System::Writer, "Type f16 is not completely implemented yet."); + return false; } else if (type->Is()) { out << "int"; } else if (auto* mat = type->As()) { diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc index 1a16079a20..fb4efb6c31 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc @@ -3513,6 +3513,11 @@ bool GeneratorImpl::EmitType(std::ostream& out, out << "float"; return true; }, + [&](const sem::F16*) { + diagnostics_.add_error(diag::System::Writer, + "Type f16 is not completely implemented yet."); + return false; + }, [&](const sem::I32*) { out << "int"; return true; diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index 8e8dc7aaa3..2f87cb257a 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -1467,6 +1467,12 @@ bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) { out << "false"; return true; }, + [&](const sem::F16*) { + // Placeholder for emitting f16 zero value + diagnostics_.add_error(diag::System::Writer, + "Type f16 is not completely implemented yet"); + return false; + }, [&](const sem::F32*) { out << "0.0f"; return true; @@ -2239,6 +2245,11 @@ bool GeneratorImpl::EmitType(std::ostream& out, out << "bool"; return true; }, + [&](const sem::F16*) { + diagnostics_.add_error(diag::System::Writer, + "Type f16 is not completely implemented yet"); + return false; + }, [&](const sem::F32*) { out << "float"; return true; diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc index 7d232c079b..aaa64cec57 100644 --- a/src/tint/writer/spirv/builder.cc +++ b/src/tint/writer/spirv/builder.cc @@ -1628,6 +1628,8 @@ uint32_t Builder::GenerateLiteralIfNeeded(const ast::Variable* var, constant.kind = ScalarConstant::Kind::kF32; constant.value.f32 = static_cast(f->value); return; + case ast::FloatLiteralExpression::Suffix::kH: + error_ = "Type f16 is not completely implemented yet"; } }, [&](Default) { error_ = "unknown literal type"; }); @@ -3672,6 +3674,11 @@ uint32_t Builder::GenerateTypeIfNeeded(const sem::Type* type) { push_type(spv::Op::OpTypeFloat, {result, Operand(32u)}); return true; }, + [&](const sem::F16*) { + // Should be `push_type(spv::Op::OpTypeFloat, {result, Operand(16u)});` + error_ = "Type f16 is not completely implemented yet."; + return false; + }, [&](const sem::I32*) { push_type(spv::Op::OpTypeInt, {result, Operand(32u), Operand(1u)}); return true; diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc index 3ef6269448..121d90425d 100644 --- a/src/tint/writer/wgsl/generator_impl.cc +++ b/src/tint/writer/wgsl/generator_impl.cc @@ -405,6 +405,11 @@ bool GeneratorImpl::EmitType(std::ostream& out, const ast::Type* ty) { out << "f32"; return true; }, + [&](const ast::F16*) { + diagnostics_.add_error(diag::System::Writer, + "Type f16 is not completely implemented yet."); + return false; + }, [&](const ast::I32*) { out << "i32"; return true; diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn index e17c6c7b48..87e26f2538 100644 --- a/test/tint/BUILD.gn +++ b/test/tint/BUILD.gn @@ -165,6 +165,7 @@ tint_unittests_source_set("tint_unittests_ast_src") { "../../src/tint/ast/discard_statement_test.cc", "../../src/tint/ast/enable_test.cc", "../../src/tint/ast/external_texture_test.cc", + "../../src/tint/ast/f16_test.cc", "../../src/tint/ast/f32_test.cc", "../../src/tint/ast/fallthrough_statement_test.cc", "../../src/tint/ast/float_literal_expression_test.cc", @@ -289,6 +290,7 @@ tint_unittests_source_set("tint_unittests_sem_src") { "../../src/tint/sem/depth_multisampled_texture_test.cc", "../../src/tint/sem/depth_texture_test.cc", "../../src/tint/sem/external_texture_test.cc", + "../../src/tint/sem/f16_test.cc", "../../src/tint/sem/f32_test.cc", "../../src/tint/sem/i32_test.cc", "../../src/tint/sem/matrix_test.cc",