[spirv-reader][ir] Add `OpSpecConstant{True,False}` support.
Adds supports for boolean spec constants.
Bug: 398007970
Change-Id: Ib900c767f308794b10f1daceed02b5cfc7af1798
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/230014
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/core/ir/transform/remove_terminator_args.h b/src/tint/lang/core/ir/transform/remove_terminator_args.h
index 96a1690..77e06ec 100644
--- a/src/tint/lang/core/ir/transform/remove_terminator_args.h
+++ b/src/tint/lang/core/ir/transform/remove_terminator_args.h
@@ -42,6 +42,7 @@
const core::ir::Capabilities kRemoveTerminatorArgsCapabilities{
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllow64BitIntegers,
+ core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowPointersAndHandlesInStructures,
core::ir::Capability::kAllowVectorElementPointer,
core::ir::Capability::kAllowHandleVarsWithoutBindings,
diff --git a/src/tint/lang/core/ir/transform/rename_conflicts.h b/src/tint/lang/core/ir/transform/rename_conflicts.h
index 1b608a0..f4b3337 100644
--- a/src/tint/lang/core/ir/transform/rename_conflicts.h
+++ b/src/tint/lang/core/ir/transform/rename_conflicts.h
@@ -42,6 +42,7 @@
const core::ir::Capabilities kRenameConflictsCapabilities{
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllow64BitIntegers,
+ core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowPointersAndHandlesInStructures,
core::ir::Capability::kAllowVectorElementPointer,
core::ir::Capability::kAllowHandleVarsWithoutBindings,
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 7c1cd54..5fe41de 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -1099,7 +1099,10 @@
} // namespace
Result<SuccessType> Builtins(core::ir::Module& ir) {
- auto result = ValidateAndDumpIfNeeded(ir, "spirv.Builtins");
+ auto result = ValidateAndDumpIfNeeded(ir, "spirv.Builtins",
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowOverrides,
+ });
if (result != Success) {
return result.Failure();
}
diff --git a/src/tint/lang/spirv/reader/lower/lower.cc b/src/tint/lang/spirv/reader/lower/lower.cc
index 55c1659..86d7241 100644
--- a/src/tint/lang/spirv/reader/lower/lower.cc
+++ b/src/tint/lang/spirv/reader/lower/lower.cc
@@ -56,7 +56,11 @@
// `||` statements.
RUN_TRANSFORM(core::ir::transform::RemoveTerminatorArgs, mod);
- if (auto res = core::ir::ValidateAndDumpIfNeeded(mod, "spirv.Lower"); res != Success) {
+ auto res = core::ir::ValidateAndDumpIfNeeded(mod, "spirv.Lower",
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowOverrides,
+ });
+ if (res != Success) {
return res.Failure();
}
diff --git a/src/tint/lang/spirv/reader/lower/shader_io.cc b/src/tint/lang/spirv/reader/lower/shader_io.cc
index 4e64e64..38faf72 100644
--- a/src/tint/lang/spirv/reader/lower/shader_io.cc
+++ b/src/tint/lang/spirv/reader/lower/shader_io.cc
@@ -464,7 +464,10 @@
} // namespace
Result<SuccessType> ShaderIO(core::ir::Module& ir) {
- auto result = ValidateAndDumpIfNeeded(ir, "spirv.ShaderIO");
+ auto result = ValidateAndDumpIfNeeded(ir, "spirv.ShaderIO",
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowOverrides,
+ });
if (result != Success) {
return result.Failure();
}
diff --git a/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
index f1f14cc..8db8634 100644
--- a/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
+++ b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
@@ -155,6 +155,7 @@
Result<SuccessType> VectorElementPointer(core::ir::Module& ir) {
auto result = ValidateAndDumpIfNeeded(ir, "spirv.VectorElementPointer",
core::ir::Capabilities{
+ core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowVectorElementPointer,
});
if (result != Success) {
diff --git a/src/tint/lang/spirv/reader/parser/helper_test.h b/src/tint/lang/spirv/reader/parser/helper_test.h
index 6b7b9f8..a3da582 100644
--- a/src/tint/lang/spirv/reader/parser/helper_test.h
+++ b/src/tint/lang/spirv/reader/parser/helper_test.h
@@ -75,6 +75,7 @@
// Validate the IR module against the capabilities supported by the SPIR-V dialect.
auto validated =
core::ir::Validate(parsed.Get(), core::ir::Capabilities{
+ core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowVectorElementPointer,
});
if (validated != Success) {
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 9327aa8..c665958 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -114,6 +114,7 @@
id_stack_.emplace_back();
{
TINT_SCOPED_ASSIGNMENT(current_block_, ir_.root_block);
+ EmitSpecConstants();
EmitModuleScopeVariables();
}
@@ -125,6 +126,47 @@
return std::move(ir_);
}
+ // Generate a module-scope const declaration for each instruction
+ // that is OpSpecConstantTrue, OpSpecConstantFalse, or OpSpecConstant.
+ void EmitSpecConstants() {
+ for (auto& inst : spirv_context_->types_values()) {
+ core::ir::Override* override_ = nullptr;
+ switch (inst.opcode()) {
+ case spv::Op::OpSpecConstantTrue:
+ case spv::Op::OpSpecConstantFalse:
+ override_ = b_.Override(Type(inst.type_id()));
+ override_->SetInitializer(
+ b_.Value(inst.opcode() == spv::Op::OpSpecConstantTrue));
+ break;
+ default:
+ break;
+ }
+ if (!override_) {
+ continue;
+ }
+
+ Emit(override_, inst.result_id());
+
+ Symbol name = GetSymbolFor(inst.result_id());
+ if (name.IsValid()) {
+ ir_.SetName(override_, name);
+ }
+
+ auto decos =
+ spirv_context_->get_decoration_mgr()->GetDecorationsFor(inst.result_id(), true);
+ for (const auto* deco_inst : decos) {
+ TINT_ASSERT(deco_inst->opcode() == spv::Op::OpDecorate);
+
+ if (deco_inst->GetSingleWordInOperand(1) ==
+ static_cast<uint32_t>(spv::Decoration::SpecId)) {
+ const uint16_t id = static_cast<uint16_t>(deco_inst->GetSingleWordInOperand(2));
+ override_->SetOverrideId(OverrideId{id});
+ break;
+ }
+ }
+ }
+ }
+
void RegisterNames() {
// Register names from OpName
for (const auto& inst : spirv_context_->debugs2()) {
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 3de62a5..88c141a 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -608,5 +608,73 @@
"var undef @location(6) @interpolate(linear, centroid)",
}));
+TEST_F(SpirvParserTest, Var_OpSpecConstantTrue) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %c "myconst"
+ OpDecorate %c SpecId 12
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %c = OpSpecConstantTrue %bool
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %b = OpLogicalAnd %bool %c %c
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %myconst:bool = override true @id(12)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:bool = and %myconst, %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantFalse) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %c "myconst"
+ OpDecorate %c SpecId 12
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %c = OpSpecConstantFalse %bool
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %b = OpLogicalAnd %bool %c %c
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %myconst:bool = override false @id(12)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:bool = and %myconst, %myconst
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index bc414fe..868c5ad 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -58,6 +58,7 @@
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/lang/core/ir/next_iteration.h"
+#include "src/tint/lang/core/ir/override.h"
#include "src/tint/lang/core/ir/phony.h"
#include "src/tint/lang/core/ir/return.h"
#include "src/tint/lang/core/ir/store.h"
@@ -114,6 +115,7 @@
Program Run(const ProgramOptions& options) {
core::ir::Capabilities caps{core::ir::Capability::kAllowRefTypes,
+ core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowPhonyInstructions};
if (auto res = core::ir::Validate(mod, caps); res != Success) {
// IR module failed validation.
@@ -199,8 +201,9 @@
void RootBlock(const core::ir::Block* root) {
for (auto* inst : *root) {
tint::Switch(
- inst, //
- [&](const core::ir::Var* var) { Var(var); }, //
+ inst, //
+ [&](const core::ir::Var* var) { Var(var); }, //
+ [&](const core::ir::Override* override_) { Override(override_); }, //
TINT_ICE_ON_NO_MATCH);
}
}
@@ -614,6 +617,22 @@
}
}
+ void Override(const core::ir::Override* override_) {
+ auto* val = override_->Result(0);
+ Symbol name = NameFor(override_->Result(0));
+ Bind(override_->Result(0), name);
+
+ Vector<const ast::Attribute*, 4> attrs;
+ attrs.Push(b.Id(override_->OverrideId()));
+
+ auto ty = Type(val->Type());
+ const ast::Expression* init = nullptr;
+ if (override_->Initializer()) {
+ init = Expr(override_->Initializer());
+ }
+ b.Override(name, ty, init, attrs);
+ }
+
void Let(const core::ir::Let* let) {
auto* result = let->Result(0);
Symbol name = NameFor(result);
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
index 33645fc..7d8e678 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
@@ -2690,6 +2690,59 @@
}
////////////////////////////////////////////////////////////////////////////////
+// Override Construct
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRToProgramTest, Override_DefaultId) {
+ core::ir::Override* o;
+ b.Append(b.ir.root_block, [&] { o = b.Override("o", false); });
+
+ auto* fn = b.Function("f", ty.bool_());
+ b.Append(fn->Block(), [&] { b.Return(fn, o); });
+
+ EXPECT_WGSL(R"(
+@id(0) override o : bool = false;
+
+fn f() -> bool {
+ return o;
+}
+)");
+}
+
+TEST_F(IRToProgramTest, Override_Id) {
+ core::ir::Override* o;
+ b.Append(b.ir.root_block, [&] { o = b.Override("o", true); });
+ o->SetOverrideId(OverrideId{10});
+
+ auto* fn = b.Function("f", ty.bool_());
+ b.Append(fn->Block(), [&] { b.Return(fn, o); });
+
+ EXPECT_WGSL(R"(
+@id(10) override o : bool = true;
+
+fn f() -> bool {
+ return o;
+}
+)");
+}
+
+TEST_F(IRToProgramTest, Override_NoInit) {
+ core::ir::Override* o;
+ b.Append(b.ir.root_block, [&] { o = b.Override("o", ty.i32()); });
+ o->SetOverrideId(OverrideId{10});
+
+ auto* fn = b.Function("f", ty.i32());
+ b.Append(fn->Block(), [&] { b.Return(fn, o); });
+
+ EXPECT_WGSL(R"(
+@id(10) override o : i32;
+
+fn f() -> i32 {
+ return o;
+}
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
// chromium_internal_graphite
////////////////////////////////////////////////////////////////////////////////
TEST_F(IRToProgramTest, Enable_ChromiumInternalGraphite_SubgroupBallot) {
diff --git a/src/tint/lang/wgsl/writer/raise/value_to_let.cc b/src/tint/lang/wgsl/writer/raise/value_to_let.cc
index 32ef545..4c08108 100644
--- a/src/tint/lang/wgsl/writer/raise/value_to_let.cc
+++ b/src/tint/lang/wgsl/writer/raise/value_to_let.cc
@@ -193,7 +193,12 @@
} // namespace
Result<SuccessType> ValueToLet(core::ir::Module& ir) {
- auto result = core::ir::ValidateAndDumpIfNeeded(ir, "wgsl.ValueToLet");
+ auto result = core::ir::ValidateAndDumpIfNeeded(ir, "wgsl.ValueToLet",
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowOverrides,
+ }
+
+ );
if (result != Success) {
return result;
}
