add support for FConvert

Bug: 401002063
Change-Id: Ie81b54ba61b5a890817a2b1a7fa4d528a532f519
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/229534
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: dan sinclair <dsinclair@chromium.org>

src/tint/lang/spirv/reader/parser/convert_test.cc

diff --git a/src/tint/lang/spirv/reader/parser/convert_test.cc b/src/tint/lang/spirv/reader/parser/convert_test.cc
index 68aae98..98aa2fa 100644
--- a/src/tint/lang/spirv/reader/parser/convert_test.cc
+++ b/src/tint/lang/spirv/reader/parser/convert_test.cc
@@ -443,5 +443,127 @@
 )");
 }
 
+TEST_F(SpirvParserTest, FConvert_ScalarF32ToF16) {
+    EXPECT_IR(R"(
+             OpCapability Shader
+             OpCapability Float16
+             OpMemoryModel Logical GLSL450
+             OpEntryPoint GLCompute %main "main"
+             OpExecutionMode %main LocalSize 1 1 1
+     %void    = OpTypeVoid
+     %float16 = OpTypeFloat 16
+     %float32 = OpTypeFloat 32
+     %two     = OpConstant %float32 2
+     %void_fn = OpTypeFunction %void
+
+     %main    = OpFunction %void None %void_fn
+     %main_start = OpLabel
+             %1 = OpFConvert %float16 %two
+             OpReturn
+     OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f16 = convert 2.0f
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, FConvert_ScalarF16ToF32) {
+    EXPECT_IR(R"(
+             OpCapability Shader
+             OpCapability Float16
+             OpMemoryModel Logical GLSL450
+             OpEntryPoint GLCompute %main "main"
+             OpExecutionMode %main LocalSize 1 1 1
+     %void    = OpTypeVoid
+     %float16 = OpTypeFloat 16
+     %float32 = OpTypeFloat 32
+     %two     = OpConstant %float16 2
+     %void_fn = OpTypeFunction %void
+
+     %main    = OpFunction %void None %void_fn
+     %main_start = OpLabel
+             %1 = OpFConvert %float32 %two
+             OpReturn
+     OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = convert 2.0h
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, FConvert_VectorF32ToF16) {
+    EXPECT_IR(R"(
+             OpCapability Shader
+             OpCapability Float16
+             OpMemoryModel Logical GLSL450
+             OpEntryPoint GLCompute %main "main"
+             OpExecutionMode %main LocalSize 1 1 1
+     %void       = OpTypeVoid
+     %float16    = OpTypeFloat 16
+     %float32    = OpTypeFloat 32
+     %v2float16  = OpTypeVector %float16 2
+     %v2float32  = OpTypeVector %float32 2
+     %two        = OpConstant %float32 2
+     %v2_two     = OpConstantComposite %v2float32 %two %two
+     %void_fn    = OpTypeFunction %void
+
+     %main       = OpFunction %void None %void_fn
+     %main_start = OpLabel
+             %1 = OpFConvert %v2float16 %v2_two
+             OpReturn
+     OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f16> = convert vec2<f32>(2.0f)
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, FConvert_VectorF16ToF32) {
+    EXPECT_IR(R"(
+             OpCapability Shader
+             OpCapability Float16
+             OpMemoryModel Logical GLSL450
+             OpEntryPoint GLCompute %main "main"
+             OpExecutionMode %main LocalSize 1 1 1
+     %void       = OpTypeVoid
+     %float16    = OpTypeFloat 16
+     %float32    = OpTypeFloat 32
+     %v2float16  = OpTypeVector %float16 2
+     %v2float32  = OpTypeVector %float32 2
+     %two       = OpConstant %float16 2
+     %v2_two    = OpConstantComposite %v2float16 %two %two
+     %void_fn    = OpTypeFunction %void
+
+     %main       = OpFunction %void None %void_fn
+     %main_start = OpLabel
+             %1 = OpFConvert %v2float32 %v2_two
+             OpReturn
+     OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<f32> = convert vec2<f16>(2.0h)
+    ret
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::spirv::reader

src/tint/lang/spirv/reader/parser/parser.cc

diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index da22740..bd7a8a9 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -752,6 +752,10 @@
                 case spv::Op::OpConvertUToF:
                     EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kConvertUToF);
                     break;
+                case spv::Op::OpFConvert:
+                    Emit(b_.Convert(Type(inst.type_id()), Value(inst.GetSingleWordOperand(2))),
+                         inst.result_id());
+                    break;
                 case spv::Op::OpBitwiseAnd:
                     EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseAnd);
                     break;