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>
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
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;
