Import Tint changes from Dawn

Changes:
  - 8ce7183e686954e50f41df1e9da221f983a16ce0 tint: Remove public include of ast include by Ben Clayton <bclayton@google.com>
  - a5988a3058c02d028785551fc249aff664205324 tint/writer/hlsl: Support for F16 type, constructor, and ... by Zhaoming Jiang <zhaoming.jiang@intel.com>
GitOrigin-RevId: 8ce7183e686954e50f41df1e9da221f983a16ce0
Change-Id: I0e374527f9324928ae41a028fbd85734490dbe44
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/95988
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/include/tint/tint.h b/include/tint/tint.h
index c397cdc..42ceda4 100644
--- a/include/tint/tint.h
+++ b/include/tint/tint.h
@@ -18,7 +18,6 @@
 // TODO(tint:88): When implementing support for an install target, all of these
 //                headers will need to be moved to include/tint/.
 
-#include "src/tint/ast/pipeline_stage.h"
 #include "src/tint/demangler.h"
 #include "src/tint/diagnostic/printer.h"
 #include "src/tint/inspector/inspector.h"
diff --git a/src/tint/inspector/entry_point.h b/src/tint/inspector/entry_point.h
index b9ac4e4..493a93f 100644
--- a/src/tint/inspector/entry_point.h
+++ b/src/tint/inspector/entry_point.h
@@ -120,6 +120,9 @@
     bool is_numeric_id_specified = false;
 };
 
+/// The pipeline stage
+enum class PipelineStage { kVertex, kFragment, kCompute };
+
 /// Reflection data for an entry point in the shader.
 struct EntryPoint {
     /// Constructors
@@ -135,7 +138,7 @@
     /// Remapped entry point name in the backend
     std::string remapped_name;
     /// The entry point stage
-    ast::PipelineStage stage = ast::PipelineStage::kNone;
+    PipelineStage stage;
     /// The workgroup x size
     uint32_t workgroup_size_x = 0;
     /// The workgroup y size
diff --git a/src/tint/inspector/inspector.cc b/src/tint/inspector/inspector.cc
index c5b3eb7..4355487 100644
--- a/src/tint/inspector/inspector.cc
+++ b/src/tint/inspector/inspector.cc
@@ -145,7 +145,22 @@
         EntryPoint entry_point;
         entry_point.name = program_->Symbols().NameFor(func->symbol);
         entry_point.remapped_name = program_->Symbols().NameFor(func->symbol);
-        entry_point.stage = func->PipelineStage();
+
+        switch (func->PipelineStage()) {
+            case ast::PipelineStage::kCompute:
+                entry_point.stage = PipelineStage::kCompute;
+                break;
+            case ast::PipelineStage::kFragment:
+                entry_point.stage = PipelineStage::kFragment;
+                break;
+            case ast::PipelineStage::kVertex:
+                entry_point.stage = PipelineStage::kVertex;
+                break;
+            default:
+                TINT_UNREACHABLE(Inspector, diagnostics_)
+                    << "invalid pipeline stage for entry point '" << entry_point.name << "'";
+                break;
+        }
 
         auto wgsize = sem->WorkgroupSize();
         entry_point.workgroup_size_x = wgsize[0].value;
diff --git a/src/tint/inspector/inspector_test.cc b/src/tint/inspector/inspector_test.cc
index 4e369d5..2681f78 100644
--- a/src/tint/inspector/inspector_test.cc
+++ b/src/tint/inspector/inspector_test.cc
@@ -168,7 +168,7 @@
     ASSERT_EQ(1u, result.size());
     EXPECT_EQ("foo", result[0].name);
     EXPECT_EQ("foo", result[0].remapped_name);
-    EXPECT_EQ(ast::PipelineStage::kFragment, result[0].stage);
+    EXPECT_EQ(PipelineStage::kFragment, result[0].stage);
 }
 
 TEST_F(InspectorGetEntryPointTest, MultipleEntryPoints) {
@@ -191,10 +191,10 @@
     ASSERT_EQ(2u, result.size());
     EXPECT_EQ("foo", result[0].name);
     EXPECT_EQ("foo", result[0].remapped_name);
-    EXPECT_EQ(ast::PipelineStage::kFragment, result[0].stage);
+    EXPECT_EQ(PipelineStage::kFragment, result[0].stage);
     EXPECT_EQ("bar", result[1].name);
     EXPECT_EQ("bar", result[1].remapped_name);
-    EXPECT_EQ(ast::PipelineStage::kCompute, result[1].stage);
+    EXPECT_EQ(PipelineStage::kCompute, result[1].stage);
 }
 
 TEST_F(InspectorGetEntryPointTest, MixFunctionsAndEntryPoints) {
@@ -221,10 +221,10 @@
     ASSERT_EQ(2u, result.size());
     EXPECT_EQ("foo", result[0].name);
     EXPECT_EQ("foo", result[0].remapped_name);
-    EXPECT_EQ(ast::PipelineStage::kCompute, result[0].stage);
+    EXPECT_EQ(PipelineStage::kCompute, result[0].stage);
     EXPECT_EQ("bar", result[1].name);
     EXPECT_EQ("bar", result[1].remapped_name);
-    EXPECT_EQ(ast::PipelineStage::kFragment, result[1].stage);
+    EXPECT_EQ(PipelineStage::kFragment, result[1].stage);
 }
 
 TEST_F(InspectorGetEntryPointTest, DefaultWorkgroupSize) {
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index b307c8d..c9861a1 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -122,6 +122,16 @@
     }
 }
 
+bool PrintF16(std::ostream& out, float value) {
+    // Note: Currently inf and nan should not be constructable, don't emit them.
+    if (std::isinf(value) || std::isnan(value)) {
+        return false;
+    } else {
+        out << FloatToString(value) << "h";
+        return true;
+    }
+}
+
 // Helper for writing " : register(RX, spaceY)", where R is the register, X is
 // the binding point binding value, and Y is the binding point group value.
 struct RegisterAndSpace {
@@ -3122,6 +3132,13 @@
             PrintF32(out, constant->As<float>());
             return true;
         },
+        [&](const sem::F16*) {
+            // emit a f16 scalar with explicit float16_t type declaration.
+            out << "float16_t(";
+            bool valid = PrintF16(out, constant->As<float>());
+            out << ")";
+            return valid;
+        },
         [&](const sem::I32*) {
             out << constant->As<AInt>();
             return true;
@@ -3218,6 +3235,13 @@
             return true;
         },
         [&](const ast::FloatLiteralExpression* l) {
+            if (l->suffix == ast::FloatLiteralExpression::Suffix::kH) {
+                // Emit f16 literal with explicit float16_t type declaration.
+                out << "float16_t(";
+                bool valid = PrintF16(out, static_cast<float>(l->value));
+                out << ")";
+                return valid;
+            }
             PrintF32(out, static_cast<float>(l->value));
             return true;
         },
@@ -3251,6 +3275,10 @@
             out << value << ".0f";
             return true;
         },
+        [&](const sem::F16*) {
+            out << "float16_t(" << value << ".0h)";
+            return true;
+        },
         [&](const sem::I32*) {
             out << value;
             return true;
@@ -3723,15 +3751,23 @@
             return true;
         },
         [&](const sem::F16*) {
-            diagnostics_.add_error(diag::System::Writer,
-                                   "Type f16 is not completely implemented yet.");
-            return false;
+            out << "float16_t";
+            return true;
         },
         [&](const sem::I32*) {
             out << "int";
             return true;
         },
         [&](const sem::Matrix* mat) {
+            if (mat->type()->Is<sem::F16>()) {
+                // Use matrix<type, N, M> for f16 matrix
+                out << "matrix<";
+                if (!EmitType(out, mat->type(), storage_class, access, "")) {
+                    return false;
+                }
+                out << ", " << mat->columns() << ", " << mat->rows() << ">";
+                return true;
+            }
             if (!EmitType(out, mat->type(), storage_class, access, "")) {
                 return false;
             }
@@ -3847,6 +3883,7 @@
             } else if (vec->type()->Is<sem::Bool>() && width >= 1 && width <= 4) {
                 out << "bool" << width;
             } else {
+                // For example, use "vector<float16_t, N>" for f16 vector.
                 out << "vector<";
                 if (!EmitType(out, vec->type(), storage_class, access, "")) {
                     return false;
diff --git a/src/tint/writer/hlsl/generator_impl_constructor_test.cc b/src/tint/writer/hlsl/generator_impl_constructor_test.cc
index 6a1d696..f8b9357 100644
--- a/src/tint/writer/hlsl/generator_impl_constructor_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_constructor_test.cc
@@ -61,6 +61,18 @@
     EXPECT_THAT(gen.result(), HasSubstr("1073741824.0f"));
 }
 
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_F16) {
+    Enable(ast::Extension::kF16);
+
+    // Use a number close to 1<<16 but whose decimal representation ends in 0.
+    WrapInFunction(Expr(f16((1 << 15) - 8)));
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_THAT(gen.result(), HasSubstr("float16_t(32752.0h)"));
+}
+
 TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Float) {
     WrapInFunction(Construct<f32>(-1.2e-5_f));
 
@@ -70,6 +82,17 @@
     EXPECT_THAT(gen.result(), HasSubstr("-0.000012f"));
 }
 
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_F16) {
+    Enable(ast::Extension::kF16);
+
+    WrapInFunction(Construct<f16>(-1.2e-3_h));
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_THAT(gen.result(), HasSubstr("float16_t(-0.00119972229h)"));
+}
+
 TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Bool) {
     WrapInFunction(Construct<bool>(true));
 
@@ -97,7 +120,7 @@
     EXPECT_THAT(gen.result(), HasSubstr("12345u"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_F32) {
     WrapInFunction(vec3<f32>(1_f, 2_f, 3_f));
 
     GeneratorImpl& gen = Build();
@@ -106,7 +129,20 @@
     EXPECT_THAT(gen.result(), HasSubstr("float3(1.0f, 2.0f, 3.0f)"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_Empty) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_F16) {
+    Enable(ast::Extension::kF16);
+
+    WrapInFunction(vec3<f16>(1_h, 2_h, 3_h));
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_THAT(
+        gen.result(),
+        HasSubstr("vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h))"));
+}
+
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_Empty_F32) {
     WrapInFunction(vec3<f32>());
 
     GeneratorImpl& gen = Build();
@@ -115,7 +151,18 @@
     EXPECT_THAT(gen.result(), HasSubstr("0.0f).xxx"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_SingleScalar_Float_Literal) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_Empty_F16) {
+    Enable(ast::Extension::kF16);
+
+    WrapInFunction(vec3<f16>());
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_THAT(gen.result(), HasSubstr("(float16_t(0.0h)).xxx"));
+}
+
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_SingleScalar_F32_Literal) {
     WrapInFunction(vec3<f32>(2_f));
 
     GeneratorImpl& gen = Build();
@@ -124,7 +171,18 @@
     EXPECT_THAT(gen.result(), HasSubstr("2.0f).xxx"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_SingleScalar_Float_Var) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_SingleScalar_F16_Literal) {
+    Enable(ast::Extension::kF16);
+
+    WrapInFunction(vec3<f16>(2_h));
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_THAT(gen.result(), HasSubstr("(float16_t(2.0h)).xxx"));
+}
+
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_SingleScalar_F32_Var) {
     auto* var = Var("v", nullptr, Expr(2_f));
     auto* cast = vec3<f32>(var);
     WrapInFunction(var, cast);
@@ -136,6 +194,20 @@
   const float3 tint_symbol = float3((v).xxx);)"));
 }
 
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_SingleScalar_F16_Var) {
+    Enable(ast::Extension::kF16);
+
+    auto* var = Var("v", nullptr, Expr(2_h));
+    auto* cast = vec3<f16>(var);
+    WrapInFunction(var, cast);
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_THAT(gen.result(), HasSubstr(R"(float16_t v = float16_t(2.0h);
+  const vector<float16_t, 3> tint_symbol = vector<float16_t, 3>((v).xxx);)"));
+}
+
 TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Vec_SingleScalar_Bool_Literal) {
     WrapInFunction(vec3<bool>(true));
 
@@ -175,7 +247,7 @@
     EXPECT_THAT(gen.result(), HasSubstr("2u).xxx"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_F32) {
     WrapInFunction(mat2x3<f32>(vec3<f32>(1_f, 2_f, 3_f), vec3<f32>(3_f, 4_f, 5_f)));
 
     GeneratorImpl& gen = Build();
@@ -186,7 +258,22 @@
                 HasSubstr("float2x3(float3(1.0f, 2.0f, 3.0f), float3(3.0f, 4.0f, 5.0f))"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Complex) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_F16) {
+    Enable(ast::Extension::kF16);
+
+    WrapInFunction(mat2x3<f16>(vec3<f16>(1_h, 2_h, 3_h), vec3<f16>(3_h, 4_h, 5_h)));
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_THAT(gen.result(),
+                HasSubstr("matrix<float16_t, 2, 3>(vector<float16_t, 3>(float16_t(1.0h), "
+                          "float16_t(2.0h), float16_t(3.0h)), vector<float16_t, "
+                          "3>(float16_t(3.0h), float16_t(4.0h), float16_t(5.0h)))"));
+}
+
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Complex_F32) {
     // mat4x4<f32>(
     //     vec4<f32>(2.0f, 3.0f, 4.0f, 8.0f),
     //     vec4<f32>(),
@@ -213,7 +300,40 @@
                                         "(7.0f).xxxx, float4(42.0f, 21.0f, 6.0f, -5.0f))"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Empty) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Complex_F16) {
+    // mat4x4<f16>(
+    //     vec4<f16>(2.0h, 3.0h, 4.0h, 8.0h),
+    //     vec4<f16>(),
+    //     vec4<f16>(7.0h),
+    //     vec4<f16>(vec4<f16>(42.0h, 21.0h, 6.0h, -5.0h)),
+    //   );
+    Enable(ast::Extension::kF16);
+
+    auto* vector_literal =
+        vec4<f16>(Expr(f16(2.0)), Expr(f16(3.0)), Expr(f16(4.0)), Expr(f16(8.0)));
+    auto* vector_zero_ctor = vec4<f16>();
+    auto* vector_single_scalar_ctor = vec4<f16>(Expr(f16(7.0)));
+    auto* vector_identical_ctor =
+        vec4<f16>(vec4<f16>(Expr(f16(42.0)), Expr(f16(21.0)), Expr(f16(6.0)), Expr(f16(-5.0))));
+
+    auto* constructor = mat4x4<f16>(vector_literal, vector_zero_ctor, vector_single_scalar_ctor,
+                                    vector_identical_ctor);
+
+    WrapInFunction(constructor);
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_THAT(gen.result(), HasSubstr("matrix<float16_t, 4, 4>("
+                                        "vector<float16_t, 4>(float16_t(2.0h), float16_t(3.0h), "
+                                        "float16_t(4.0h), float16_t(8.0h)), "
+                                        "(float16_t(0.0h)).xxxx, (float16_t(7.0h)).xxxx, "
+                                        "vector<float16_t, 4>(float16_t(42.0h), float16_t(21.0h), "
+                                        "float16_t(6.0h), float16_t(-5.0h)))"));
+}
+
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Empty_F32) {
     WrapInFunction(mat2x3<f32>());
 
     GeneratorImpl& gen = Build();
@@ -223,7 +343,20 @@
     EXPECT_THAT(gen.result(), HasSubstr("float2x3 tint_symbol = float2x3((0.0f).xxx, (0.0f).xxx)"));
 }
 
-TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Identity) {
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Empty_F16) {
+    Enable(ast::Extension::kF16);
+
+    WrapInFunction(mat2x3<f16>());
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_THAT(gen.result(),
+                HasSubstr("matrix<float16_t, 2, 3>((float16_t(0.0h)).xxx, (float16_t(0.0h)).xxx)"));
+}
+
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Identity_F32) {
     // fn f() {
     //     var m_1: mat4x4<f32> = mat4x4<f32>();
     //     var m_2: mat4x4<f32> = mat4x4<f32>(m_1);
@@ -241,6 +374,27 @@
     EXPECT_THAT(gen.result(), HasSubstr("float4x4 m_2 = float4x4(m_1);"));
 }
 
+TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Mat_Identity_F16) {
+    // fn f() {
+    //     var m_1: mat4x4<f16> = mat4x4<f16>();
+    //     var m_2: mat4x4<f16> = mat4x4<f16>(m_1);
+    // }
+
+    Enable(ast::Extension::kF16);
+
+    auto* m_1 = Var("m_1", ty.mat4x4(ty.f16()), mat4x4<f16>());
+    auto* m_2 = Var("m_2", ty.mat4x4(ty.f16()), mat4x4<f16>(m_1));
+
+    WrapInFunction(m_1, m_2);
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_THAT(gen.result(),
+                HasSubstr("matrix<float16_t, 4, 4> m_2 = matrix<float16_t, 4, 4>(m_1);"));
+}
+
 TEST_F(HlslGeneratorImplTest_Constructor, EmitConstructor_Type_Array) {
     WrapInFunction(Construct(ty.array(ty.vec3<f32>(), 3_u), vec3<f32>(1_f, 2_f, 3_f),
                              vec3<f32>(4_f, 5_f, 6_f), vec3<f32>(7_f, 8_f, 9_f)));
diff --git a/src/tint/writer/hlsl/generator_impl_module_constant_test.cc b/src/tint/writer/hlsl/generator_impl_module_constant_test.cc
index 3a988a0..d5dee6f 100644
--- a/src/tint/writer/hlsl/generator_impl_module_constant_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_module_constant_test.cc
@@ -92,6 +92,22 @@
 )");
 }
 
+TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_GlobalConst_f16) {
+    Enable(ast::Extension::kF16);
+
+    auto* var = GlobalConst("G", nullptr, Expr(1_h));
+    Func("f", {}, ty.void_(), {Decl(Let("l", nullptr, Expr(var)))});
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_EQ(gen.result(), R"(void f() {
+  const float16_t l = float16_t(1.0h);
+}
+)");
+}
+
 TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_GlobalConst_vec3_AInt) {
     auto* var = GlobalConst("G", nullptr, Construct(ty.vec3(nullptr), 1_a, 2_a, 3_a));
     Func("f", {}, ty.void_(), {Decl(Let("l", nullptr, Expr(var)))});
@@ -134,6 +150,22 @@
 )");
 }
 
+TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_GlobalConst_vec3_f16) {
+    Enable(ast::Extension::kF16);
+
+    auto* var = GlobalConst("G", nullptr, vec3<f16>(1_h, 2_h, 3_h));
+    Func("f", {}, ty.void_(), {Decl(Let("l", nullptr, Expr(var)))});
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_EQ(gen.result(), R"(void f() {
+  const vector<float16_t, 3> l = vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h));
+}
+)");
+}
+
 TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_GlobalConst_mat2x3_AFloat) {
     auto* var = GlobalConst("G", nullptr,
                             Construct(ty.mat(nullptr, 2, 3), 1._a, 2._a, 3._a, 4._a, 5._a, 6._a));
@@ -163,6 +195,22 @@
 )");
 }
 
+TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_GlobalConst_mat2x3_f16) {
+    Enable(ast::Extension::kF16);
+
+    auto* var = GlobalConst("G", nullptr, mat2x3<f16>(1_h, 2_h, 3_h, 4_h, 5_h, 6_h));
+    Func("f", {}, ty.void_(), {Decl(Let("l", nullptr, Expr(var)))});
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_EQ(gen.result(), R"(void f() {
+  const matrix<float16_t, 2, 3> l = matrix<float16_t, 2, 3>(vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 3>(float16_t(4.0h), float16_t(5.0h), float16_t(6.0h)));
+}
+)");
+}
+
 TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_GlobalConst_arr_f32) {
     auto* var = GlobalConst("G", nullptr, Construct(ty.array<f32, 3>(), 1_f, 2_f, 3_f));
     Func("f", {}, ty.void_(), {Decl(Let("l", nullptr, Expr(var)))});
diff --git a/src/tint/writer/hlsl/generator_impl_type_test.cc b/src/tint/writer/hlsl/generator_impl_type_test.cc
index 614d8f1..991c5c0 100644
--- a/src/tint/writer/hlsl/generator_impl_type_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_type_test.cc
@@ -94,6 +94,17 @@
     EXPECT_EQ(out.str(), "bool");
 }
 
+TEST_F(HlslGeneratorImplTest_Type, EmitType_F16) {
+    auto* f16 = create<sem::F16>();
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitType(out, f16, ast::StorageClass::kNone, ast::Access::kReadWrite, ""))
+        << gen.error();
+    EXPECT_EQ(out.str(), "float16_t");
+}
+
 TEST_F(HlslGeneratorImplTest_Type, EmitType_F32) {
     auto* f32 = create<sem::F32>();
 
@@ -116,7 +127,20 @@
     EXPECT_EQ(out.str(), "int");
 }
 
-TEST_F(HlslGeneratorImplTest_Type, EmitType_Matrix) {
+TEST_F(HlslGeneratorImplTest_Type, EmitType_Matrix_F16) {
+    auto* f16 = create<sem::F16>();
+    auto* vec3 = create<sem::Vector>(f16, 3u);
+    auto* mat2x3 = create<sem::Matrix>(vec3, 2u);
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitType(out, mat2x3, ast::StorageClass::kNone, ast::Access::kReadWrite, ""))
+        << gen.error();
+    EXPECT_EQ(out.str(), "matrix<float16_t, 2, 3>");
+}
+
+TEST_F(HlslGeneratorImplTest_Type, EmitType_Matrix_F32) {
     auto* f32 = create<sem::F32>();
     auto* vec3 = create<sem::Vector>(f32, 3u);
     auto* mat2x3 = create<sem::Matrix>(vec3, 2u);
diff --git a/src/tint/writer/hlsl/generator_impl_variable_decl_statement_test.cc b/src/tint/writer/hlsl/generator_impl_variable_decl_statement_test.cc
index 1e111cd..9a741bc 100644
--- a/src/tint/writer/hlsl/generator_impl_variable_decl_statement_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_variable_decl_statement_test.cc
@@ -134,6 +134,22 @@
 )");
 }
 
+TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const_f16) {
+    Enable(ast::Extension::kF16);
+
+    auto* C = Const("C", nullptr, Expr(1_h));
+    Func("f", {}, ty.void_(), {Decl(C), Decl(Let("l", nullptr, Expr(C)))});
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_EQ(gen.result(), R"(void f() {
+  const float16_t l = float16_t(1.0h);
+}
+)");
+}
+
 TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const_vec3_AInt) {
     auto* C = Const("C", nullptr, Construct(ty.vec3(nullptr), 1_a, 2_a, 3_a));
     Func("f", {}, ty.void_(), {Decl(C), Decl(Let("l", nullptr, Expr(C)))});
@@ -176,6 +192,22 @@
 )");
 }
 
+TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const_vec3_f16) {
+    Enable(ast::Extension::kF16);
+
+    auto* C = Const("C", nullptr, vec3<f16>(1_h, 2_h, 3_h));
+    Func("f", {}, ty.void_(), {Decl(C), Decl(Let("l", nullptr, Expr(C)))});
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_EQ(gen.result(), R"(void f() {
+  const vector<float16_t, 3> l = vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h));
+}
+)");
+}
+
 TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const_mat2x3_AFloat) {
     auto* C =
         Const("C", nullptr, Construct(ty.mat(nullptr, 2, 3), 1._a, 2._a, 3._a, 4._a, 5._a, 6._a));
@@ -205,6 +237,22 @@
 )");
 }
 
+TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const_mat2x3_f16) {
+    Enable(ast::Extension::kF16);
+
+    auto* C = Const("C", nullptr, mat2x3<f16>(1_h, 2_h, 3_h, 4_h, 5_h, 6_h));
+    Func("f", {}, ty.void_(), {Decl(C), Decl(Let("l", nullptr, Expr(C)))});
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+
+    EXPECT_EQ(gen.result(), R"(void f() {
+  const matrix<float16_t, 2, 3> l = matrix<float16_t, 2, 3>(vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h)), vector<float16_t, 3>(float16_t(4.0h), float16_t(5.0h), float16_t(6.0h)));
+}
+)");
+}
+
 TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const_arr_f32) {
     auto* C = Const("C", nullptr, Construct(ty.array<f32, 3>(), 1_f, 2_f, 3_f));
     Func("f", {}, ty.void_(), {Decl(C), Decl(Let("l", nullptr, Expr(C)))});
@@ -263,7 +311,7 @@
     EXPECT_THAT(gen.result(), HasSubstr("  static float a = 0.0f;\n"));
 }
 
-TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Initializer_ZeroVec) {
+TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Initializer_ZeroVec_F32) {
     auto* var = Var("a", ty.vec3<f32>(), ast::StorageClass::kNone, vec3<f32>());
 
     auto* stmt = Decl(var);
@@ -276,7 +324,22 @@
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Initializer_ZeroMat) {
+TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Initializer_ZeroVec_F16) {
+    Enable(ast::Extension::kF16);
+
+    auto* var = Var("a", ty.vec3<f16>(), ast::StorageClass::kNone, vec3<f16>());
+
+    auto* stmt = Decl(var);
+    WrapInFunction(stmt);
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+    EXPECT_EQ(gen.result(), R"(vector<float16_t, 3> a = (float16_t(0.0h)).xxx;
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Initializer_ZeroMat_F32) {
     auto* var = Var("a", ty.mat2x3<f32>(), ast::StorageClass::kNone, mat2x3<f32>());
 
     auto* stmt = Decl(var);
@@ -290,5 +353,22 @@
 )");
 }
 
+TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Initializer_ZeroMat_F16) {
+    Enable(ast::Extension::kF16);
+
+    auto* var = Var("a", ty.mat2x3<f16>(), ast::StorageClass::kNone, mat2x3<f16>());
+
+    auto* stmt = Decl(var);
+    WrapInFunction(stmt);
+
+    GeneratorImpl& gen = Build();
+
+    ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
+    EXPECT_EQ(
+        gen.result(),
+        R"(matrix<float16_t, 2, 3> a = matrix<float16_t, 2, 3>((float16_t(0.0h)).xxx, (float16_t(0.0h)).xxx);
+)");
+}
+
 }  // namespace
 }  // namespace tint::writer::hlsl