[spirv-reader][ir] Add Not support for spec constants
Add support for the `OpNot` command to the spec constant instructions.
Bug: 402726637
Change-Id: I03211a6fd3cc59a96f7d3e73f89c4b47be68b436
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/237396
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index ac934bd..80440d9 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -8501,6 +8501,53 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvReader_BuiltinsTest, SpecConstantOp_Not) {
+ auto* ep = b.ComputeFunction("foo");
+
+ capabilities = core::ir::Capabilities{core::ir::Capability::kAllowOverrides};
+
+ b.Append(b.ir.root_block, [&] {
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kNot, Vector{ty.i32()},
+ 1_i);
+ });
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kNot, Vector{ty.i32()},
+ 1_i);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+$B1: { # root
+ %1:i32 = spirv.not<i32> 1i
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:i32 = spirv.not<i32> 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+$B1: { # root
+ %1:i32 = complement 1i
+}
+
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:i32 = complement 1i
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(SpirvReader_BuiltinsTest, Not_Scalar_Signed_Signed) {
auto* ep = b.ComputeFunction("foo");
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 080f67b..18bf5a8 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -208,6 +208,9 @@
case spv::Op::OpLogicalNotEqual:
EmitBinary(inst, core::BinaryOp::kNotEqual, 3);
break;
+ case spv::Op::OpNot:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kNot, 3);
+ break;
default:
TINT_ICE() << "Unknown spec constant operation: " << op;
}
@@ -1871,10 +1874,12 @@
Emit(b_.Call(Type(inst.type_id()), fn, Args(inst, 2)), inst.result_id());
}
- void EmitSpirvExplicitBuiltinCall(const spvtools::opt::Instruction& inst, spirv::BuiltinFn fn) {
+ void EmitSpirvExplicitBuiltinCall(const spvtools::opt::Instruction& inst,
+ spirv::BuiltinFn fn,
+ uint32_t first_operand_idx = 2) {
Emit(b_.CallExplicit<spirv::ir::BuiltinCall>(Type(inst.type_id()), fn,
Vector{Type(inst.type_id())->DeepestElement()},
- Args(inst, 2)),
+ Args(inst, first_operand_idx)),
inst.result_id());
}
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index c0c5c65..27c30b9 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -904,5 +904,41 @@
)");
}
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_Not) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %cond "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %i32 = OpTypeInt 32 1
+ %false = OpSpecConstantFalse %bool
+ %true = OpSpecConstantTrue %bool
+ %one = OpConstant %i32 1
+ %cond = OpSpecConstantOp %i32 Not %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpNot %i32 %cond
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.not<i32> 1i
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = spirv.not<i32> %myconst
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
