spirv-reader: support fconvert
Fixed: 402726010
Change-Id: I64d03a23fd20f008ab5564f448f49381a754d880
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/249574
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: David Neto <dneto@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 6cd964b..0735d32 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -370,6 +370,11 @@
case spv::Op::OpUConvert:
TINT_ICE() << "can't translate UConvert: WGSL does not have concrete "
"integer types of different widths";
+ case spv::Op::OpFConvert:
+ Emit(b_.Convert(Type(inst.type_id()),
+ Value(inst.GetSingleWordInOperand(1))),
+ inst.result_id());
+ break;
default:
TINT_ICE() << "Unknown spec constant operation: " << op;
}
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 86afef8..23472c2 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -1309,6 +1309,41 @@
)");
}
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_FConvert) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability Float16
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %f16 = OpTypeFloat 16
+ %one = OpConstant %f16 1.0
+ %myconst = OpSpecConstantOp %f32 FConvert %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %f32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:f32 = convert 1.0h
+ %myconst:f32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:f32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
// In the case of all literals, SPIR-V opt treats the `OpSpecConstantComposite` as an
// `OpConstantComposite` so it appears in the constant manager already. This then needs no handling
// on our side.
