From 256f1116b836aa88d8704e7f4ee2f6b32ed44dff Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Fri, 22 Jul 2022 16:05:06 +0000 Subject: [PATCH] Add transform to substitute overrides with const expressions. This CL adds a SubstituteOverride transform which will convert an `override` into a `const`. The transform is provided a map of (string, double) which matches what the WebGPU API accepts as data for overrides. Bug: tint:1582 Change-Id: I6e6bf51b98ce4d4746f8de55128666c36735e585 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/96760 Reviewed-by: Ben Clayton Commit-Queue: Dan Sinclair Kokoro: Kokoro --- include/tint/tint.h | 1 + src/tint/BUILD.gn | 6 +- src/tint/CMakeLists.txt | 4 + src/tint/ast/override.cc | 7 + src/tint/ast/override.h | 8 + src/tint/ast/override_test.cc | 35 +++ src/tint/cmd/main.cc | 47 +++- src/tint/transform/substitute_override.cc | 96 +++++++ src/tint/transform/substitute_override.h | 89 +++++++ .../transform/substitute_override_test.cc | 243 ++++++++++++++++++ src/tint/val/hlsl.cc | 31 +-- src/tint/val/val.h | 9 +- .../tint/var/override/named/no_init/bool.wgsl | 7 +- .../named/no_init/bool.wgsl.expected.glsl | 20 +- .../named/no_init/bool.wgsl.expected.hlsl | 11 +- .../named/no_init/bool.wgsl.expected.msl | 4 +- .../named/no_init/bool.wgsl.expected.spvasm | 26 +- .../named/no_init/bool.wgsl.expected.wgsl | 8 +- test/tint/var/override/named/no_init/f32.wgsl | 6 +- .../named/no_init/f32.wgsl.expected.glsl | 16 +- .../named/no_init/f32.wgsl.expected.hlsl | 7 +- .../named/no_init/f32.wgsl.expected.msl | 4 +- .../named/no_init/f32.wgsl.expected.spvasm | 21 +- .../named/no_init/f32.wgsl.expected.wgsl | 6 +- test/tint/var/override/named/no_init/i32.wgsl | 6 +- .../named/no_init/i32.wgsl.expected.glsl | 16 +- .../named/no_init/i32.wgsl.expected.hlsl | 7 +- .../named/no_init/i32.wgsl.expected.msl | 4 +- .../named/no_init/i32.wgsl.expected.spvasm | 22 +- .../named/no_init/i32.wgsl.expected.wgsl | 6 +- test/tint/var/override/named/no_init/u32.wgsl | 6 +- .../named/no_init/u32.wgsl.expected.glsl | 16 +- .../named/no_init/u32.wgsl.expected.hlsl | 7 +- .../named/no_init/u32.wgsl.expected.msl | 4 +- .../named/no_init/u32.wgsl.expected.spvasm | 22 +- .../named/no_init/u32.wgsl.expected.wgsl | 6 +- .../var/override/named/val_init/bool.wgsl | 6 +- .../named/val_init/bool.wgsl.expected.glsl | 10 +- .../named/val_init/bool.wgsl.expected.hlsl | 11 +- .../named/val_init/bool.wgsl.expected.msl | 4 +- .../named/val_init/bool.wgsl.expected.spvasm | 26 +- .../named/val_init/bool.wgsl.expected.wgsl | 8 +- .../tint/var/override/named/val_init/f32.wgsl | 5 +- .../named/val_init/f32.wgsl.expected.glsl | 6 +- .../named/val_init/f32.wgsl.expected.hlsl | 7 +- .../named/val_init/f32.wgsl.expected.msl | 4 +- .../named/val_init/f32.wgsl.expected.spvasm | 21 +- .../named/val_init/f32.wgsl.expected.wgsl | 6 +- .../tint/var/override/named/val_init/i32.wgsl | 5 +- .../named/val_init/i32.wgsl.expected.glsl | 6 +- .../named/val_init/i32.wgsl.expected.hlsl | 7 +- .../named/val_init/i32.wgsl.expected.msl | 4 +- .../named/val_init/i32.wgsl.expected.spvasm | 22 +- .../named/val_init/i32.wgsl.expected.wgsl | 6 +- .../tint/var/override/named/val_init/u32.wgsl | 5 +- .../named/val_init/u32.wgsl.expected.glsl | 6 +- .../named/val_init/u32.wgsl.expected.hlsl | 7 +- .../named/val_init/u32.wgsl.expected.msl | 4 +- .../named/val_init/u32.wgsl.expected.spvasm | 22 +- .../named/val_init/u32.wgsl.expected.wgsl | 6 +- .../var/override/named/zero_init/bool.wgsl | 6 +- .../named/zero_init/bool.wgsl.expected.glsl | 10 +- .../named/zero_init/bool.wgsl.expected.hlsl | 11 +- .../named/zero_init/bool.wgsl.expected.msl | 4 +- .../named/zero_init/bool.wgsl.expected.spvasm | 26 +- .../named/zero_init/bool.wgsl.expected.wgsl | 8 +- .../var/override/named/zero_init/f32.wgsl | 5 +- .../named/zero_init/f32.wgsl.expected.glsl | 6 +- .../named/zero_init/f32.wgsl.expected.hlsl | 7 +- .../named/zero_init/f32.wgsl.expected.msl | 4 +- .../named/zero_init/f32.wgsl.expected.spvasm | 21 +- .../named/zero_init/f32.wgsl.expected.wgsl | 6 +- .../var/override/named/zero_init/i32.wgsl | 5 +- .../named/zero_init/i32.wgsl.expected.glsl | 6 +- .../named/zero_init/i32.wgsl.expected.hlsl | 7 +- .../named/zero_init/i32.wgsl.expected.msl | 4 +- .../named/zero_init/i32.wgsl.expected.spvasm | 22 +- .../named/zero_init/i32.wgsl.expected.wgsl | 6 +- .../var/override/named/zero_init/u32.wgsl | 5 +- .../named/zero_init/u32.wgsl.expected.glsl | 6 +- .../named/zero_init/u32.wgsl.expected.hlsl | 7 +- .../named/zero_init/u32.wgsl.expected.msl | 4 +- .../named/zero_init/u32.wgsl.expected.spvasm | 22 +- .../named/zero_init/u32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/no_init/bool.wgsl | 6 +- .../numbered/no_init/bool.wgsl.expected.glsl | 16 +- .../numbered/no_init/bool.wgsl.expected.hlsl | 7 +- .../numbered/no_init/bool.wgsl.expected.msl | 4 +- .../no_init/bool.wgsl.expected.spvasm | 19 +- .../numbered/no_init/bool.wgsl.expected.wgsl | 6 +- .../var/override/numbered/no_init/f32.wgsl | 6 +- .../numbered/no_init/f32.wgsl.expected.glsl | 16 +- .../numbered/no_init/f32.wgsl.expected.hlsl | 7 +- .../numbered/no_init/f32.wgsl.expected.msl | 4 +- .../numbered/no_init/f32.wgsl.expected.spvasm | 21 +- .../numbered/no_init/f32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/no_init/i32.wgsl | 6 +- .../numbered/no_init/i32.wgsl.expected.glsl | 16 +- .../numbered/no_init/i32.wgsl.expected.hlsl | 7 +- .../numbered/no_init/i32.wgsl.expected.msl | 4 +- .../numbered/no_init/i32.wgsl.expected.spvasm | 21 +- .../numbered/no_init/i32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/no_init/u32.wgsl | 6 +- .../numbered/no_init/u32.wgsl.expected.glsl | 16 +- .../numbered/no_init/u32.wgsl.expected.hlsl | 7 +- .../numbered/no_init/u32.wgsl.expected.msl | 4 +- .../numbered/no_init/u32.wgsl.expected.spvasm | 22 +- .../numbered/no_init/u32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/val_init/bool.wgsl | 5 +- .../numbered/val_init/bool.wgsl.expected.glsl | 6 +- .../numbered/val_init/bool.wgsl.expected.hlsl | 7 +- .../numbered/val_init/bool.wgsl.expected.msl | 4 +- .../val_init/bool.wgsl.expected.spvasm | 19 +- .../numbered/val_init/bool.wgsl.expected.wgsl | 6 +- .../var/override/numbered/val_init/f32.wgsl | 5 +- .../numbered/val_init/f32.wgsl.expected.glsl | 6 +- .../numbered/val_init/f32.wgsl.expected.hlsl | 7 +- .../numbered/val_init/f32.wgsl.expected.msl | 4 +- .../val_init/f32.wgsl.expected.spvasm | 22 +- .../numbered/val_init/f32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/val_init/i32.wgsl | 5 +- .../numbered/val_init/i32.wgsl.expected.glsl | 6 +- .../numbered/val_init/i32.wgsl.expected.hlsl | 7 +- .../numbered/val_init/i32.wgsl.expected.msl | 4 +- .../val_init/i32.wgsl.expected.spvasm | 22 +- .../numbered/val_init/i32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/val_init/u32.wgsl | 5 +- .../numbered/val_init/u32.wgsl.expected.glsl | 6 +- .../numbered/val_init/u32.wgsl.expected.hlsl | 7 +- .../numbered/val_init/u32.wgsl.expected.msl | 4 +- .../val_init/u32.wgsl.expected.spvasm | 22 +- .../numbered/val_init/u32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/zero_init/bool.wgsl | 5 +- .../zero_init/bool.wgsl.expected.glsl | 6 +- .../zero_init/bool.wgsl.expected.hlsl | 7 +- .../numbered/zero_init/bool.wgsl.expected.msl | 4 +- .../zero_init/bool.wgsl.expected.spvasm | 19 +- .../zero_init/bool.wgsl.expected.wgsl | 6 +- .../var/override/numbered/zero_init/f32.wgsl | 5 +- .../numbered/zero_init/f32.wgsl.expected.glsl | 6 +- .../numbered/zero_init/f32.wgsl.expected.hlsl | 7 +- .../numbered/zero_init/f32.wgsl.expected.msl | 4 +- .../zero_init/f32.wgsl.expected.spvasm | 21 +- .../numbered/zero_init/f32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/zero_init/i32.wgsl | 5 +- .../numbered/zero_init/i32.wgsl.expected.glsl | 6 +- .../numbered/zero_init/i32.wgsl.expected.hlsl | 7 +- .../numbered/zero_init/i32.wgsl.expected.msl | 4 +- .../zero_init/i32.wgsl.expected.spvasm | 22 +- .../numbered/zero_init/i32.wgsl.expected.wgsl | 6 +- .../var/override/numbered/zero_init/u32.wgsl | 5 +- .../numbered/zero_init/u32.wgsl.expected.glsl | 6 +- .../numbered/zero_init/u32.wgsl.expected.hlsl | 7 +- .../numbered/zero_init/u32.wgsl.expected.msl | 4 +- .../zero_init/u32.wgsl.expected.spvasm | 22 +- .../numbered/zero_init/u32.wgsl.expected.wgsl | 6 +- 156 files changed, 1235 insertions(+), 659 deletions(-) create mode 100644 src/tint/ast/override_test.cc create mode 100644 src/tint/transform/substitute_override.cc create mode 100644 src/tint/transform/substitute_override.h create mode 100644 src/tint/transform/substitute_override_test.cc diff --git a/include/tint/tint.h b/include/tint/tint.h index 42ceda48e8..a6ab6caf47 100644 --- a/include/tint/tint.h +++ b/include/tint/tint.h @@ -31,6 +31,7 @@ #include "src/tint/transform/renamer.h" #include "src/tint/transform/robustness.h" #include "src/tint/transform/single_entry_point.h" +#include "src/tint/transform/substitute_override.h" #include "src/tint/transform/vertex_pulling.h" #include "src/tint/writer/flatten_bindings.h" #include "src/tint/writer/writer.h" diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 03f27e5183..20134a908b 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -287,9 +287,9 @@ libtint_source_set("libtint_core_all_src") { "ast/module.h", "ast/multisampled_texture.cc", "ast/multisampled_texture.h", - "ast/node_id.h", "ast/node.cc", "ast/node.h", + "ast/node_id.h", "ast/override.cc", "ast/override.h", "ast/parameter.cc", @@ -529,6 +529,8 @@ libtint_source_set("libtint_core_all_src") { "transform/single_entry_point.h", "transform/spirv_atomic.cc", "transform/spirv_atomic.h", + "transform/substitute_override.cc", + "transform/substitute_override.h", "transform/transform.cc", "transform/transform.h", "transform/unshadow.cc", @@ -1034,6 +1036,7 @@ if (tint_build_unittests) { "ast/module_clone_test.cc", "ast/module_test.cc", "ast/multisampled_texture_test.cc", + "ast/override_test.cc", "ast/phony_expression_test.cc", "ast/pointer_test.cc", "ast/return_statement_test.cc", @@ -1197,6 +1200,7 @@ if (tint_build_unittests) { "transform/simplify_pointers_test.cc", "transform/single_entry_point_test.cc", "transform/spirv_atomic_test.cc", + "transform/substitute_override_test.cc", "transform/test_helper.h", "transform/transform_test.cc", "transform/unshadow_test.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index 4eb97a2a1d..312aa5218c 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -441,6 +441,8 @@ set(TINT_LIB_SRCS transform/single_entry_point.h transform/spirv_atomic.cc transform/spirv_atomic.h + transform/substitute_override.cc + transform/substitute_override.h transform/transform.cc transform/transform.h transform/unshadow.cc @@ -730,6 +732,7 @@ if(TINT_BUILD_TESTS) ast/module_clone_test.cc ast/module_test.cc ast/multisampled_texture_test.cc + ast/override_test.cc ast/phony_expression_test.cc ast/pointer_test.cc ast/return_statement_test.cc @@ -1109,6 +1112,7 @@ if(TINT_BUILD_TESTS) transform/simplify_pointers_test.cc transform/single_entry_point_test.cc transform/spirv_atomic_test.cc + transform/substitute_override_test.cc transform/test_helper.h transform/unshadow_test.cc transform/unwind_discard_functions_test.cc diff --git a/src/tint/ast/override.cc b/src/tint/ast/override.cc index b7ee064103..2876920678 100644 --- a/src/tint/ast/override.cc +++ b/src/tint/ast/override.cc @@ -46,4 +46,11 @@ const Override* Override::Clone(CloneContext* ctx) const { return ctx->dst->create(src, sym, ty, ctor, attrs); } +std::string Override::Identifier(const SymbolTable& symbols) const { + if (auto* id = ast::GetAttribute(attributes)) { + return std::to_string(id->value); + } + return symbols.NameFor(symbol); +} + } // namespace tint::ast diff --git a/src/tint/ast/override.h b/src/tint/ast/override.h index e97f14c9c0..00c36fcb1e 100644 --- a/src/tint/ast/override.h +++ b/src/tint/ast/override.h @@ -15,6 +15,8 @@ #ifndef SRC_TINT_AST_OVERRIDE_H_ #define SRC_TINT_AST_OVERRIDE_H_ +#include + #include "src/tint/ast/variable.h" namespace tint::ast { @@ -60,6 +62,12 @@ class Override final : public Castable { /// @param ctx the clone context /// @return the newly cloned node const Override* Clone(CloneContext* ctx) const override; + + /// @param symbols the symbol table to retrieve the name from + /// @returns the identifier string for the override. If the override has + /// an ID attribute, the string is the id-stringified. Otherwise, the ID + /// is the symbol. + std::string Identifier(const SymbolTable& symbols) const; }; } // namespace tint::ast diff --git a/src/tint/ast/override_test.cc b/src/tint/ast/override_test.cc new file mode 100644 index 0000000000..b4179e5bc1 --- /dev/null +++ b/src/tint/ast/override_test.cc @@ -0,0 +1,35 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/ast/override.h" + +#include "src/tint/ast/test_helper.h" + +namespace tint::ast { +namespace { + +using OverrideTest = TestHelper; + +TEST_F(OverrideTest, Identifier_NoId) { + auto* o = Override("o", nullptr, Expr(f32(1.0))); + EXPECT_EQ(std::string("o"), o->Identifier(Symbols())); +} + +TEST_F(OverrideTest, Identifier_WithId) { + auto* o = Override("o", nullptr, Expr(f32(1.0)), {Id(4u)}); + EXPECT_EQ(std::string("4"), o->Identifier(Symbols())); +} + +} // namespace +} // namespace tint::ast diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc index f6029ac925..941b399456 100644 --- a/src/tint/cmd/main.cc +++ b/src/tint/cmd/main.cc @@ -20,6 +20,7 @@ #include #include #include +#include #include #if TINT_BUILD_GLSL_WRITER @@ -86,7 +87,7 @@ struct Options { bool use_fxc = false; std::string dxc_path; std::string xcrun_path; - std::vector overrides; + std::unordered_map overrides; std::optional hlsl_root_constant_binding_point; }; @@ -126,7 +127,8 @@ ${transforms} When specified, automatically enables --validate --xcrun -- Path to xcrun executable, used to validate MSL output. When specified, automatically enables --validate - --overrides -- Pipeline overrides as NAME=VALUE, comma-separated.)"; + --overrides -- Override values as IDENTIFIER=VALUE, comma-separated. +)"; Format parse_format(const std::string& fmt) { (void)fmt; @@ -215,18 +217,26 @@ Format infer_format(const std::string& filename) { return Format::kNone; } -std::vector split_on_comma(std::string list) { +std::vector split_on_char(std::string list, char c) { std::vector res; std::stringstream str(list); while (str.good()) { std::string substr; - getline(str, substr, ','); + getline(str, substr, c); res.push_back(substr); } return res; } +std::vector split_on_comma(std::string list) { + return split_on_char(list, ','); +} + +std::vector split_on_equal(std::string list) { + return split_on_char(list, '='); +} + std::optional parse_unsigned_number(std::string number) { for (char c : number) { if (!std::isdigit(c)) { @@ -429,7 +439,10 @@ bool ParseArgs(const std::vector& args, Options* opts) { std::cerr << "Missing value for " << arg << std::endl; return false; } - opts->overrides = split_on_comma(args[i]); + for (const auto& o : split_on_comma(args[i])) { + auto parts = split_on_equal(o); + opts->overrides.insert({parts[0], std::stod(parts[1])}); + } } else if (arg == "--hlsl-root-constant-binding-point") { ++i; if (i >= args.size()) { @@ -792,7 +805,7 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) { tint::val::Result res; if (options.use_fxc) { #ifdef _WIN32 - res = tint::val::HlslUsingFXC(result.hlsl, result.entry_points, options.overrides); + res = tint::val::HlslUsingFXC(result.hlsl, result.entry_points); #else res.failed = true; res.output = "FXC can only be used on Windows. Sorry :X"; @@ -801,8 +814,7 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) { auto dxc = tint::utils::Command::LookPath(options.dxc_path.empty() ? "dxc" : options.dxc_path); if (dxc.Found()) { - res = tint::val::HlslUsingDXC(dxc.Path(), result.hlsl, result.entry_points, - options.overrides); + res = tint::val::HlslUsingDXC(dxc.Path(), result.hlsl, result.entry_points); } else { res.failed = true; res.output = "DXC executable not found. Cannot validate"; @@ -947,6 +959,14 @@ int main(int argc, const char** argv) { tint::transform::DataMap&) { m.Add(); }}, {"robustness", [](tint::transform::Manager& m, tint::transform::DataMap&) { m.Add(); }}, + {"substitute_override", + [&](tint::transform::Manager& m, tint::transform::DataMap& i) { + tint::transform::SubstituteOverride::Config cfg; + cfg.map = options.overrides; + + i.Add(cfg); + m.Add(); + }}, }; auto transform_names = [&] { std::stringstream names; @@ -1078,6 +1098,17 @@ int main(int argc, const char** argv) { tint::transform::Manager transform_manager; tint::transform::DataMap transform_inputs; + + // If overrides are provided, add the SubstituteOverride transform. + if (!options.overrides.empty()) { + for (auto& t : transforms) { + if (t.name == std::string("substitute_override")) { + t.make(transform_manager, transform_inputs); + break; + } + } + } + for (const auto& name : options.transforms) { // TODO(dsinclair): The vertex pulling transform requires setup code to // be run that needs user input. Should we find a way to support that here diff --git a/src/tint/transform/substitute_override.cc b/src/tint/transform/substitute_override.cc new file mode 100644 index 0000000000..aa7abff58a --- /dev/null +++ b/src/tint/transform/substitute_override.cc @@ -0,0 +1,96 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/transform/substitute_override.h" + +#include + +#include "src/tint/program_builder.h" +#include "src/tint/sem/variable.h" + +TINT_INSTANTIATE_TYPEINFO(tint::transform::SubstituteOverride); +TINT_INSTANTIATE_TYPEINFO(tint::transform::SubstituteOverride::Config); + +namespace tint::transform { + +SubstituteOverride::SubstituteOverride() = default; + +SubstituteOverride::~SubstituteOverride() = default; + +bool SubstituteOverride::ShouldRun(const Program* program, const DataMap&) const { + for (auto* node : program->AST().GlobalVariables()) { + if (node->Is()) { + return true; + } + } + return false; +} + +void SubstituteOverride::Run(CloneContext& ctx, const DataMap& config, DataMap&) const { + const auto* data = config.Get(); + if (!data) { + ctx.dst->Diagnostics().add_error(diag::System::Transform, + "Missing override substitution data"); + return; + } + + ctx.ReplaceAll([&](const ast::Override* w) -> const ast::Const* { + auto ident = w->Identifier(ctx.src->Symbols()); + + auto src = ctx.Clone(w->source); + auto sym = ctx.Clone(w->symbol); + auto* ty = ctx.Clone(w->type); + + // No replacement provided, just clone the override node as a const. + auto iter = data->map.find(ident); + if (iter == data->map.end()) { + if (!w->constructor) { + ctx.dst->Diagnostics().add_error( + diag::System::Transform, + "Initializer not provided for override, and override not overridden."); + return nullptr; + } + return ctx.dst->Const(src, sym, ty, ctx.Clone(w->constructor)); + } + + auto value = iter->second; + auto* ctor = Switch( + ctx.src->Sem().Get(w)->Type(), + [&](const sem::Bool*) { return ctx.dst->Expr(!std::equal_to()(value, 0.0)); }, + [&](const sem::I32*) { return ctx.dst->Expr(i32(value)); }, + [&](const sem::U32*) { return ctx.dst->Expr(u32(value)); }, + [&](const sem::F32*) { return ctx.dst->Expr(f32(value)); }, + [&](const sem::F16*) { return ctx.dst->Expr(f16(value)); }); + + if (!ctor) { + ctx.dst->Diagnostics().add_error(diag::System::Transform, + "Failed to create override expression"); + return nullptr; + } + + return ctx.dst->Const(src, sym, ty, ctor); + }); + + ctx.Clone(); +} + +SubstituteOverride::Config::Config() = default; + +SubstituteOverride::Config::Config(const Config&) = default; + +SubstituteOverride::Config::~Config() = default; + +SubstituteOverride::Config& SubstituteOverride::Config::operator=(const Config&) = default; + +} // namespace tint::transform diff --git a/src/tint/transform/substitute_override.h b/src/tint/transform/substitute_override.h new file mode 100644 index 0000000000..ab4b38a842 --- /dev/null +++ b/src/tint/transform/substitute_override.h @@ -0,0 +1,89 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef SRC_TINT_TRANSFORM_SUBSTITUTE_OVERRIDE_H_ +#define SRC_TINT_TRANSFORM_SUBSTITUTE_OVERRIDE_H_ + +#include +#include + +#include "src/tint/transform/transform.h" + +namespace tint::transform { + +/// A transform that replaces overrides with the constant values provided. +/// +/// # Example +/// ``` +/// override width: f32; +/// @id(1) override height: i32 = 4; +/// override depth = 1i; +/// ``` +/// +/// When transformed with `width` -> 1, `1` -> 22, `depth` -> 42 +/// +/// ``` +/// const width: f32 = 1f; +/// const height: i32 = 22i; +/// const depth = 42i; +/// ``` +/// +/// @see crbug.com/tint/1582 +class SubstituteOverride final : public Castable { + public: + /// Configuration options for the transform + struct Config final : public Castable { + /// Constructor + Config(); + + /// Copy constructor + Config(const Config&); + + /// Destructor + ~Config() override; + + /// Assignment operator + /// @returns this Config + Config& operator=(const Config&); + + /// The map of override name (either the identifier or id) to value. + /// The value is always a double coming into the transform and will be + /// converted to the correct type through and initializer. + std::unordered_map map; + }; + + /// Constructor + SubstituteOverride(); + + /// Destructor + ~SubstituteOverride() override; + + /// @param program the program to inspect + /// @param data optional extra transform-specific input data + /// @returns true if this transform should be run for the given program + bool ShouldRun(const Program* program, const DataMap& data = {}) const override; + + protected: + /// Runs the transform using the CloneContext built for transforming a + /// program. Run() is responsible for calling Clone() on the CloneContext. + /// @param ctx the CloneContext primed with the input program and + /// ProgramBuilder + /// @param inputs optional extra transform-specific input data + /// @param outputs optional extra transform-specific output data + void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) const override; +}; + +} // namespace tint::transform + +#endif // SRC_TINT_TRANSFORM_SUBSTITUTE_OVERRIDE_H_ diff --git a/src/tint/transform/substitute_override_test.cc b/src/tint/transform/substitute_override_test.cc new file mode 100644 index 0000000000..70ac67514f --- /dev/null +++ b/src/tint/transform/substitute_override_test.cc @@ -0,0 +1,243 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/transform/substitute_override.h" + +#include "src/tint/transform/test_helper.h" + +namespace tint::transform { +namespace { + +using SubstituteOverrideTest = TransformTest; + +TEST_F(SubstituteOverrideTest, Error_NoData) { + auto* src = R"( +override width: i32; +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + auto* expect = "error: Missing override substitution data"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(SubstituteOverrideTest, Error_NoOverrideValue) { + auto* src = R"( +override width: i32; +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + auto* expect = "error: Initializer not provided for override, and override not overridden."; + + SubstituteOverride::Config cfg; + DataMap data; + data.Add(cfg); + + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(SubstituteOverrideTest, Module_NoOverrides) { + auto* src = R"( +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + auto* expect = R"( +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + SubstituteOverride::Config cfg; + + DataMap data; + data.Add(cfg); + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(SubstituteOverrideTest, Identifier) { + auto* src = R"( +override i_width: i32; +override i_height = 1i; + +override f_width: f32; +override f_height = 1.f; + +// TODO(crbug.com/tint/1473) +// override h_width: f16; +// override h_height = 1.h; + +override b_width: bool; +override b_height = true; + +override o_width = 2i; + +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + auto* expect = R"( +const i_width : i32 = 42i; + +const i_height = 11i; + +const f_width : f32 = 22.299999237f; + +const f_height = 12.399999619f; + +const b_width : bool = true; + +const b_height = false; + +const o_width = 2i; + +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + SubstituteOverride::Config cfg; + cfg.map.insert({"i_width", 42.0}); + cfg.map.insert({"i_height", 11.0}); + cfg.map.insert({"f_width", 22.3}); + cfg.map.insert({"f_height", 12.4}); + cfg.map.insert({"h_width", 9.4}); + cfg.map.insert({"h_height", 3.4}); + cfg.map.insert({"b_width", 1.0}); + cfg.map.insert({"b_height", 0.0}); + + DataMap data; + data.Add(cfg); + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(SubstituteOverrideTest, Id) { + auto* src = R"( +enable f16; + +@id(0) override i_width: i32; +@id(10) override i_height = 1i; + +@id(1) override f_width: f32; +@id(9) override f_height = 1.f; + +// TODO(crbug.com/tint/1473) +// @id(2) override h_width: f16; +// @id(8) override h_height = 1.h; + +@id(3) override b_width: bool; +@id(7) override b_height = true; + +@id(5) override o_width = 2i; + +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + auto* expect = R"( +enable f16; + +const i_width : i32 = 42i; + +const i_height = 11i; + +const f_width : f32 = 22.299999237f; + +const f_height = 12.399999619f; + +const b_width : bool = true; + +const b_height = false; + +const o_width = 2i; + +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + SubstituteOverride::Config cfg; + cfg.map.insert({"0", 42.0}); + cfg.map.insert({"10", 11.0}); + cfg.map.insert({"1", 22.3}); + cfg.map.insert({"9", 12.4}); + cfg.map.insert({"2", 9.4}); + cfg.map.insert({"8", 3.4}); + cfg.map.insert({"3", 1.0}); + cfg.map.insert({"7", 0.0}); + // No effect because an @id is set for o_width + cfg.map.insert({"o_width", 13}); + + DataMap data; + data.Add(cfg); + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(SubstituteOverrideTest, Identifier_Expression) { + auto* src = R"( +override i_height = ~2i; + +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + auto* expect = R"( +const i_height = 11i; + +@vertex +fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + + SubstituteOverride::Config cfg; + cfg.map.insert({"i_height", 11.0}); + + DataMap data; + data.Add(cfg); + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + +} // namespace +} // namespace tint::transform diff --git a/src/tint/val/hlsl.cc b/src/tint/val/hlsl.cc index 49c38501a9..164152d8f8 100644 --- a/src/tint/val/hlsl.cc +++ b/src/tint/val/hlsl.cc @@ -30,8 +30,7 @@ namespace tint::val { Result HlslUsingDXC(const std::string& dxc_path, const std::string& source, - const EntryPointList& entry_points, - const std::vector& overrides) { + const EntryPointList& entry_points) { Result result; auto dxc = utils::Command(dxc_path); @@ -70,13 +69,7 @@ Result HlslUsingDXC(const std::string& dxc_path, "/Zpr " // D3DCOMPILE_PACK_MATRIX_ROW_MAJOR "/Gis"; // D3DCOMPILE_IEEE_STRICTNESS - std::string defs; - defs.reserve(overrides.size() * 20); - for (auto& o : overrides) { - defs += "/D" + o + " "; - } - - auto res = dxc(profile, "-E " + ep.first, compileFlags, file.Path(), defs); + auto res = dxc(profile, "-E " + ep.first, compileFlags, file.Path()); if (!res.out.empty()) { if (!result.output.empty()) { result.output += "\n"; @@ -102,9 +95,7 @@ Result HlslUsingDXC(const std::string& dxc_path, } #ifdef _WIN32 -Result HlslUsingFXC(const std::string& source, - const EntryPointList& entry_points, - const std::vector& overrides) { +Result HlslUsingFXC(const std::string& source, const EntryPointList& entry_points) { Result result; // This library leaks if an error happens in this function, but it is ok @@ -148,26 +139,12 @@ Result HlslUsingFXC(const std::string& source, UINT compileFlags = D3DCOMPILE_OPTIMIZATION_LEVEL0 | D3DCOMPILE_PACK_MATRIX_ROW_MAJOR | D3DCOMPILE_IEEE_STRICTNESS; - auto overrides_copy = overrides; // Copy so that we can replace '=' with '\0' - std::vector macros; - macros.reserve(overrides_copy.size() * 2); - for (auto& o : overrides_copy) { - if (auto sep = o.find_first_of('='); sep != std::string::npos) { - // Replace '=' with '\0' so we can point directly into the allocated string buffer - o[sep] = '\0'; - macros.push_back(D3D_SHADER_MACRO{&o[0], &o[sep + 1]}); - } else { - macros.emplace_back(D3D_SHADER_MACRO{o.c_str(), NULL}); - } - } - macros.emplace_back(D3D_SHADER_MACRO{NULL, NULL}); - ComPtr compiledShader; ComPtr errors; HRESULT cr = d3dCompile(source.c_str(), // pSrcData source.length(), // SrcDataSize nullptr, // pSourceName - macros.data(), // pDefines + nullptr, // pDefines nullptr, // pInclude ep.first.c_str(), // pEntrypoint profile, // pTarget diff --git a/src/tint/val/val.h b/src/tint/val/val.h index c869efb7aa..4b3fff9cb2 100644 --- a/src/tint/val/val.h +++ b/src/tint/val/val.h @@ -43,23 +43,18 @@ struct Result { /// @param dxc_path path to DXC /// @param source the generated HLSL source /// @param entry_points the list of entry points to validate -/// @param overrides optional list of pipeline overrides /// @return the result of the compile Result HlslUsingDXC(const std::string& dxc_path, const std::string& source, - const EntryPointList& entry_points, - const std::vector& overrides); + const EntryPointList& entry_points); #ifdef _WIN32 /// Hlsl attempts to compile the shader with FXC, verifying that the shader /// compiles successfully. /// @param source the generated HLSL source /// @param entry_points the list of entry points to validate -/// @param overrides optional list of pipeline overrides /// @return the result of the compile -Result HlslUsingFXC(const std::string& source, - const EntryPointList& entry_points, - const std::vector& overrides); +Result HlslUsingFXC(const std::string& source, const EntryPointList& entry_points); #endif // _WIN32 /// Msl attempts to compile the shader with the Metal Shader Compiler, diff --git a/test/tint/var/override/named/no_init/bool.wgsl b/test/tint/var/override/named/no_init/bool.wgsl index de58af9ed7..1044195686 100644 --- a/test/tint/var/override/named/no_init/bool.wgsl +++ b/test/tint/var/override/named/no_init/bool.wgsl @@ -1,7 +1,10 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_0=0 +// flags: --overrides o=0,j=1 override o : bool; +override j : bool; @compute @workgroup_size(1) fn main() { - _ = o; + if o && j { + _ = 1; + } } diff --git a/test/tint/var/override/named/no_init/bool.wgsl.expected.glsl b/test/tint/var/override/named/no_init/bool.wgsl.expected.glsl index 4df044d579..677dd8d573 100644 --- a/test/tint/var/override/named/no_init/bool.wgsl.expected.glsl +++ b/test/tint/var/override/named/no_init/bool.wgsl.expected.glsl @@ -1,12 +1,12 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -const bool o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + bool tint_tmp = false; + if (tint_tmp) { + tint_tmp = true; + } + if ((tint_tmp)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +14,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 0 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl index 3ea88ba0b2..fba945b8e2 100644 --- a/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl +++ b/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl @@ -1,9 +1,10 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -static const bool o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + bool tint_tmp = false; + if (tint_tmp) { + tint_tmp = true; + } + if ((tint_tmp)) { + } return; } diff --git a/test/tint/var/override/named/no_init/bool.wgsl.expected.msl b/test/tint/var/override/named/no_init/bool.wgsl.expected.msl index 014e3b1fd3..c502139d39 100644 --- a/test/tint/var/override/named/no_init/bool.wgsl.expected.msl +++ b/test/tint/var/override/named/no_init/bool.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant bool o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((false && true)) { + } return; } diff --git a/test/tint/var/override/named/no_init/bool.wgsl.expected.spvasm b/test/tint/var/override/named/no_init/bool.wgsl.expected.spvasm index ff12e004f7..3c8eb1b1e5 100644 --- a/test/tint/var/override/named/no_init/bool.wgsl.expected.spvasm +++ b/test/tint/var/override/named/no_init/bool.wgsl.expected.spvasm @@ -1,20 +1,30 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 13 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %bool = OpTypeBool - %o = OpSpecConstantFalse %bool %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %bool = OpTypeBool + %6 = OpConstantNull %bool + %true = OpConstantTrue %bool + %main = OpFunction %void None %1 + %4 = OpLabel + OpSelectionMerge %7 None + OpBranchConditional %6 %8 %7 + %8 = OpLabel + OpBranch %7 + %7 = OpLabel + %10 = OpPhi %bool %6 %4 %true %8 + OpSelectionMerge %11 None + OpBranchConditional %10 %12 %11 + %12 = OpLabel + OpBranch %11 + %11 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/no_init/bool.wgsl.expected.wgsl b/test/tint/var/override/named/no_init/bool.wgsl.expected.wgsl index 6ee52bb2b6..532159899c 100644 --- a/test/tint/var/override/named/no_init/bool.wgsl.expected.wgsl +++ b/test/tint/var/override/named/no_init/bool.wgsl.expected.wgsl @@ -1,6 +1,10 @@ -override o : bool; +const o : bool = false; + +const j : bool = true; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o && j)) { + _ = 1; + } } diff --git a/test/tint/var/override/named/no_init/f32.wgsl b/test/tint/var/override/named/no_init/f32.wgsl index ddc14507a7..fa827f9c80 100644 --- a/test/tint/var/override/named/no_init/f32.wgsl +++ b/test/tint/var/override/named/no_init/f32.wgsl @@ -1,7 +1,9 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_0=0 +// flags: --overrides o=0 override o : f32; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 0.0 { + _ = 1; + } } diff --git a/test/tint/var/override/named/no_init/f32.wgsl.expected.glsl b/test/tint/var/override/named/no_init/f32.wgsl.expected.glsl index bb100addb6..3d082fd3e3 100644 --- a/test/tint/var/override/named/no_init/f32.wgsl.expected.glsl +++ b/test/tint/var/override/named/no_init/f32.wgsl.expected.glsl @@ -1,12 +1,8 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -const float o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0.0f == 0.0f)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +10,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 0 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl index 34bec48685..e42b8fb027 100644 --- a/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -static const float o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/named/no_init/f32.wgsl.expected.msl b/test/tint/var/override/named/no_init/f32.wgsl.expected.msl index ecf89c4c91..ecda7f6a64 100644 --- a/test/tint/var/override/named/no_init/f32.wgsl.expected.msl +++ b/test/tint/var/override/named/no_init/f32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant float o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/named/no_init/f32.wgsl.expected.spvasm b/test/tint/var/override/named/no_init/f32.wgsl.expected.spvasm index 7ae9060492..9190d7bd41 100644 --- a/test/tint/var/override/named/no_init/f32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/no_init/f32.wgsl.expected.spvasm @@ -1,20 +1,25 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 11 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %float = OpTypeFloat 32 - %o = OpSpecConstant %float 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %6 = OpConstantNull %float + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %7 = OpFOrdEqual %bool %6 %6 + OpSelectionMerge %9 None + OpBranchConditional %7 %10 %9 + %10 = OpLabel + OpBranch %9 + %9 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/no_init/f32.wgsl.expected.wgsl b/test/tint/var/override/named/no_init/f32.wgsl.expected.wgsl index bc62e28b2d..362776f333 100644 --- a/test/tint/var/override/named/no_init/f32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/no_init/f32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : f32; +const o : f32 = 0.0f; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 0.0)) { + _ = 1; + } } diff --git a/test/tint/var/override/named/no_init/i32.wgsl b/test/tint/var/override/named/no_init/i32.wgsl index 1a951c97db..d78b6e7c66 100644 --- a/test/tint/var/override/named/no_init/i32.wgsl +++ b/test/tint/var/override/named/no_init/i32.wgsl @@ -1,7 +1,9 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_0=0 +// flags: --overrides o=0 override o : i32; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/named/no_init/i32.wgsl.expected.glsl b/test/tint/var/override/named/no_init/i32.wgsl.expected.glsl index e443a44f4f..aadd0384bb 100644 --- a/test/tint/var/override/named/no_init/i32.wgsl.expected.glsl +++ b/test/tint/var/override/named/no_init/i32.wgsl.expected.glsl @@ -1,12 +1,8 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -const int o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0 == 1)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +10,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 0 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl index 034990e973..d384517be6 100644 --- a/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -static const int o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0 == 1)) { + } return; } diff --git a/test/tint/var/override/named/no_init/i32.wgsl.expected.msl b/test/tint/var/override/named/no_init/i32.wgsl.expected.msl index 6df39feed4..4bcde5e5ec 100644 --- a/test/tint/var/override/named/no_init/i32.wgsl.expected.msl +++ b/test/tint/var/override/named/no_init/i32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant int o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0 == 1)) { + } return; } diff --git a/test/tint/var/override/named/no_init/i32.wgsl.expected.spvasm b/test/tint/var/override/named/no_init/i32.wgsl.expected.spvasm index 18c8840440..74678ef974 100644 --- a/test/tint/var/override/named/no_init/i32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/no_init/i32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %int = OpTypeInt 32 1 - %o = OpSpecConstant %int 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %6 = OpConstantNull %int + %int_1 = OpConstant %int 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %int_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/no_init/i32.wgsl.expected.wgsl b/test/tint/var/override/named/no_init/i32.wgsl.expected.wgsl index e3f13ae612..88b26b31eb 100644 --- a/test/tint/var/override/named/no_init/i32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/no_init/i32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : i32; +const o : i32 = 0i; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } } diff --git a/test/tint/var/override/named/no_init/u32.wgsl b/test/tint/var/override/named/no_init/u32.wgsl index 6440d75e35..e69b91b73e 100644 --- a/test/tint/var/override/named/no_init/u32.wgsl +++ b/test/tint/var/override/named/no_init/u32.wgsl @@ -1,7 +1,9 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_0=0 +// flags: --overrides o=0 override o : u32; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/named/no_init/u32.wgsl.expected.glsl b/test/tint/var/override/named/no_init/u32.wgsl.expected.glsl index ba68f323bd..368da7abbf 100644 --- a/test/tint/var/override/named/no_init/u32.wgsl.expected.glsl +++ b/test/tint/var/override/named/no_init/u32.wgsl.expected.glsl @@ -1,12 +1,8 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -const uint o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0u == 1u)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +10,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 0 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl index fb48bbafcf..fe084141f7 100644 --- a/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#error spec constant required for constant id 0 -#endif -static const uint o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/named/no_init/u32.wgsl.expected.msl b/test/tint/var/override/named/no_init/u32.wgsl.expected.msl index 8254bed60e..641561bdc7 100644 --- a/test/tint/var/override/named/no_init/u32.wgsl.expected.msl +++ b/test/tint/var/override/named/no_init/u32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant uint o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/named/no_init/u32.wgsl.expected.spvasm b/test/tint/var/override/named/no_init/u32.wgsl.expected.spvasm index e5486b0dfb..fa0711ab91 100644 --- a/test/tint/var/override/named/no_init/u32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/no_init/u32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %uint = OpTypeInt 32 0 - %o = OpSpecConstant %uint 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %6 = OpConstantNull %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %uint_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/no_init/u32.wgsl.expected.wgsl b/test/tint/var/override/named/no_init/u32.wgsl.expected.wgsl index 785862494d..b7ac534cc5 100644 --- a/test/tint/var/override/named/no_init/u32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/no_init/u32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : u32; +const o : u32 = 0u; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/bool.wgsl b/test/tint/var/override/named/val_init/bool.wgsl index 839f0ae5c5..7d35a70c03 100644 --- a/test/tint/var/override/named/val_init/bool.wgsl +++ b/test/tint/var/override/named/val_init/bool.wgsl @@ -1,6 +1,10 @@ +// flags: --overrides o=0,j=1 override o : bool = true; +override j : bool = false; @compute @workgroup_size(1) fn main() { - _ = o; + if o && j { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/bool.wgsl.expected.glsl b/test/tint/var/override/named/val_init/bool.wgsl.expected.glsl index 06d171d4ba..677dd8d573 100644 --- a/test/tint/var/override/named/val_init/bool.wgsl.expected.glsl +++ b/test/tint/var/override/named/val_init/bool.wgsl.expected.glsl @@ -1,10 +1,12 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 true -#endif -const bool o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + bool tint_tmp = false; + if (tint_tmp) { + tint_tmp = true; + } + if ((tint_tmp)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/val_init/bool.wgsl.expected.hlsl b/test/tint/var/override/named/val_init/bool.wgsl.expected.hlsl index e4ba79005c..fba945b8e2 100644 --- a/test/tint/var/override/named/val_init/bool.wgsl.expected.hlsl +++ b/test/tint/var/override/named/val_init/bool.wgsl.expected.hlsl @@ -1,9 +1,10 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 true -#endif -static const bool o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + bool tint_tmp = false; + if (tint_tmp) { + tint_tmp = true; + } + if ((tint_tmp)) { + } return; } diff --git a/test/tint/var/override/named/val_init/bool.wgsl.expected.msl b/test/tint/var/override/named/val_init/bool.wgsl.expected.msl index 014e3b1fd3..c502139d39 100644 --- a/test/tint/var/override/named/val_init/bool.wgsl.expected.msl +++ b/test/tint/var/override/named/val_init/bool.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant bool o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((false && true)) { + } return; } diff --git a/test/tint/var/override/named/val_init/bool.wgsl.expected.spvasm b/test/tint/var/override/named/val_init/bool.wgsl.expected.spvasm index 277782e3ef..3c8eb1b1e5 100644 --- a/test/tint/var/override/named/val_init/bool.wgsl.expected.spvasm +++ b/test/tint/var/override/named/val_init/bool.wgsl.expected.spvasm @@ -1,20 +1,30 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 13 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %bool = OpTypeBool - %o = OpSpecConstantTrue %bool %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %bool = OpTypeBool + %6 = OpConstantNull %bool + %true = OpConstantTrue %bool + %main = OpFunction %void None %1 + %4 = OpLabel + OpSelectionMerge %7 None + OpBranchConditional %6 %8 %7 + %8 = OpLabel + OpBranch %7 + %7 = OpLabel + %10 = OpPhi %bool %6 %4 %true %8 + OpSelectionMerge %11 None + OpBranchConditional %10 %12 %11 + %12 = OpLabel + OpBranch %11 + %11 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/val_init/bool.wgsl.expected.wgsl b/test/tint/var/override/named/val_init/bool.wgsl.expected.wgsl index 46c64f0fe0..c6e8c19d7c 100644 --- a/test/tint/var/override/named/val_init/bool.wgsl.expected.wgsl +++ b/test/tint/var/override/named/val_init/bool.wgsl.expected.wgsl @@ -1,6 +1,10 @@ -override o : bool = true; +const o : bool = false; + +const j : bool = true; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o && j)) { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/f32.wgsl b/test/tint/var/override/named/val_init/f32.wgsl index 7b981109ad..4a7f3db11f 100644 --- a/test/tint/var/override/named/val_init/f32.wgsl +++ b/test/tint/var/override/named/val_init/f32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides o=0 override o : f32 = 1.0; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 0.0 { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/f32.wgsl.expected.glsl b/test/tint/var/override/named/val_init/f32.wgsl.expected.glsl index 62ee1b4236..3d082fd3e3 100644 --- a/test/tint/var/override/named/val_init/f32.wgsl.expected.glsl +++ b/test/tint/var/override/named/val_init/f32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 1.0f -#endif -const float o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0.0f == 0.0f)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/val_init/f32.wgsl.expected.hlsl b/test/tint/var/override/named/val_init/f32.wgsl.expected.hlsl index 9d0794c1af..e42b8fb027 100644 --- a/test/tint/var/override/named/val_init/f32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/val_init/f32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 1.0f -#endif -static const float o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/named/val_init/f32.wgsl.expected.msl b/test/tint/var/override/named/val_init/f32.wgsl.expected.msl index ecf89c4c91..ecda7f6a64 100644 --- a/test/tint/var/override/named/val_init/f32.wgsl.expected.msl +++ b/test/tint/var/override/named/val_init/f32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant float o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/named/val_init/f32.wgsl.expected.spvasm b/test/tint/var/override/named/val_init/f32.wgsl.expected.spvasm index d5ce59ae6e..9190d7bd41 100644 --- a/test/tint/var/override/named/val_init/f32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/val_init/f32.wgsl.expected.spvasm @@ -1,20 +1,25 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 11 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %float = OpTypeFloat 32 - %o = OpSpecConstant %float 1 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %6 = OpConstantNull %float + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %7 = OpFOrdEqual %bool %6 %6 + OpSelectionMerge %9 None + OpBranchConditional %7 %10 %9 + %10 = OpLabel + OpBranch %9 + %9 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/val_init/f32.wgsl.expected.wgsl b/test/tint/var/override/named/val_init/f32.wgsl.expected.wgsl index b297f1a55c..80171943a9 100644 --- a/test/tint/var/override/named/val_init/f32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/val_init/f32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : f32 = 1.0; +const o : f32 = 0.0f; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 0.0)) { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/i32.wgsl b/test/tint/var/override/named/val_init/i32.wgsl index bcc63d5892..1acc5a1cbd 100644 --- a/test/tint/var/override/named/val_init/i32.wgsl +++ b/test/tint/var/override/named/val_init/i32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides o=0 override o : i32 = 1; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 2 { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/i32.wgsl.expected.glsl b/test/tint/var/override/named/val_init/i32.wgsl.expected.glsl index 299dbaef3b..e6f817c290 100644 --- a/test/tint/var/override/named/val_init/i32.wgsl.expected.glsl +++ b/test/tint/var/override/named/val_init/i32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 1 -#endif -const int o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0 == 2)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/val_init/i32.wgsl.expected.hlsl b/test/tint/var/override/named/val_init/i32.wgsl.expected.hlsl index fa2861c472..21744ce5ca 100644 --- a/test/tint/var/override/named/val_init/i32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/val_init/i32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 1 -#endif -static const int o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0 == 2)) { + } return; } diff --git a/test/tint/var/override/named/val_init/i32.wgsl.expected.msl b/test/tint/var/override/named/val_init/i32.wgsl.expected.msl index 6df39feed4..4563b8d153 100644 --- a/test/tint/var/override/named/val_init/i32.wgsl.expected.msl +++ b/test/tint/var/override/named/val_init/i32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant int o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0 == 2)) { + } return; } diff --git a/test/tint/var/override/named/val_init/i32.wgsl.expected.spvasm b/test/tint/var/override/named/val_init/i32.wgsl.expected.spvasm index 7d30d25b6e..5e40713526 100644 --- a/test/tint/var/override/named/val_init/i32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/val_init/i32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %int = OpTypeInt 32 1 - %o = OpSpecConstant %int 1 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %6 = OpConstantNull %int + %int_2 = OpConstant %int 2 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %int_2 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/val_init/i32.wgsl.expected.wgsl b/test/tint/var/override/named/val_init/i32.wgsl.expected.wgsl index 52b23b936f..6675efccb5 100644 --- a/test/tint/var/override/named/val_init/i32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/val_init/i32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : i32 = 1; +const o : i32 = 0i; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 2)) { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/u32.wgsl b/test/tint/var/override/named/val_init/u32.wgsl index 0ae52cc4f7..a4e4f65d96 100644 --- a/test/tint/var/override/named/val_init/u32.wgsl +++ b/test/tint/var/override/named/val_init/u32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides o=0 override o : u32 = 1u; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 2 { + _ = o; + } } diff --git a/test/tint/var/override/named/val_init/u32.wgsl.expected.glsl b/test/tint/var/override/named/val_init/u32.wgsl.expected.glsl index e1067f0375..b19260b440 100644 --- a/test/tint/var/override/named/val_init/u32.wgsl.expected.glsl +++ b/test/tint/var/override/named/val_init/u32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 1u -#endif -const uint o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0u == 2u)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/val_init/u32.wgsl.expected.hlsl b/test/tint/var/override/named/val_init/u32.wgsl.expected.hlsl index 9529784468..b51955d36d 100644 --- a/test/tint/var/override/named/val_init/u32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/val_init/u32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 1u -#endif -static const uint o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0u == 2u)) { + } return; } diff --git a/test/tint/var/override/named/val_init/u32.wgsl.expected.msl b/test/tint/var/override/named/val_init/u32.wgsl.expected.msl index 8254bed60e..d5f3f223da 100644 --- a/test/tint/var/override/named/val_init/u32.wgsl.expected.msl +++ b/test/tint/var/override/named/val_init/u32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant uint o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0u == 2u)) { + } return; } diff --git a/test/tint/var/override/named/val_init/u32.wgsl.expected.spvasm b/test/tint/var/override/named/val_init/u32.wgsl.expected.spvasm index 93718f11d9..4212d43c39 100644 --- a/test/tint/var/override/named/val_init/u32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/val_init/u32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %uint = OpTypeInt 32 0 - %o = OpSpecConstant %uint 1 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %6 = OpConstantNull %uint + %uint_2 = OpConstant %uint 2 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %uint_2 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/val_init/u32.wgsl.expected.wgsl b/test/tint/var/override/named/val_init/u32.wgsl.expected.wgsl index ab5c07a1bf..020319ae6b 100644 --- a/test/tint/var/override/named/val_init/u32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/val_init/u32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : u32 = 1u; +const o : u32 = 0u; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 2)) { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/bool.wgsl b/test/tint/var/override/named/zero_init/bool.wgsl index b9c57b395e..a29e4ab86f 100644 --- a/test/tint/var/override/named/zero_init/bool.wgsl +++ b/test/tint/var/override/named/zero_init/bool.wgsl @@ -1,6 +1,10 @@ +// flags: --overrides o=0,j=1 override o : bool = bool(); +override j : bool = bool(); @compute @workgroup_size(1) fn main() { - _ = o; + if o && j { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/bool.wgsl.expected.glsl b/test/tint/var/override/named/zero_init/bool.wgsl.expected.glsl index 1d946b4672..677dd8d573 100644 --- a/test/tint/var/override/named/zero_init/bool.wgsl.expected.glsl +++ b/test/tint/var/override/named/zero_init/bool.wgsl.expected.glsl @@ -1,10 +1,12 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 false -#endif -const bool o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + bool tint_tmp = false; + if (tint_tmp) { + tint_tmp = true; + } + if ((tint_tmp)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/zero_init/bool.wgsl.expected.hlsl b/test/tint/var/override/named/zero_init/bool.wgsl.expected.hlsl index d14a339663..fba945b8e2 100644 --- a/test/tint/var/override/named/zero_init/bool.wgsl.expected.hlsl +++ b/test/tint/var/override/named/zero_init/bool.wgsl.expected.hlsl @@ -1,9 +1,10 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 false -#endif -static const bool o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + bool tint_tmp = false; + if (tint_tmp) { + tint_tmp = true; + } + if ((tint_tmp)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/bool.wgsl.expected.msl b/test/tint/var/override/named/zero_init/bool.wgsl.expected.msl index 014e3b1fd3..c502139d39 100644 --- a/test/tint/var/override/named/zero_init/bool.wgsl.expected.msl +++ b/test/tint/var/override/named/zero_init/bool.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant bool o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((false && true)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/bool.wgsl.expected.spvasm b/test/tint/var/override/named/zero_init/bool.wgsl.expected.spvasm index ff12e004f7..3c8eb1b1e5 100644 --- a/test/tint/var/override/named/zero_init/bool.wgsl.expected.spvasm +++ b/test/tint/var/override/named/zero_init/bool.wgsl.expected.spvasm @@ -1,20 +1,30 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 13 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %bool = OpTypeBool - %o = OpSpecConstantFalse %bool %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %bool = OpTypeBool + %6 = OpConstantNull %bool + %true = OpConstantTrue %bool + %main = OpFunction %void None %1 + %4 = OpLabel + OpSelectionMerge %7 None + OpBranchConditional %6 %8 %7 + %8 = OpLabel + OpBranch %7 + %7 = OpLabel + %10 = OpPhi %bool %6 %4 %true %8 + OpSelectionMerge %11 None + OpBranchConditional %10 %12 %11 + %12 = OpLabel + OpBranch %11 + %11 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/zero_init/bool.wgsl.expected.wgsl b/test/tint/var/override/named/zero_init/bool.wgsl.expected.wgsl index 283ad4ebfa..c6e8c19d7c 100644 --- a/test/tint/var/override/named/zero_init/bool.wgsl.expected.wgsl +++ b/test/tint/var/override/named/zero_init/bool.wgsl.expected.wgsl @@ -1,6 +1,10 @@ -override o : bool = bool(); +const o : bool = false; + +const j : bool = true; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o && j)) { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/f32.wgsl b/test/tint/var/override/named/zero_init/f32.wgsl index a49a8a392e..7c41ef3c99 100644 --- a/test/tint/var/override/named/zero_init/f32.wgsl +++ b/test/tint/var/override/named/zero_init/f32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides o=0 override o : f32 = f32(); @compute @workgroup_size(1) fn main() { - _ = o; + if o == 0.0 { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/f32.wgsl.expected.glsl b/test/tint/var/override/named/zero_init/f32.wgsl.expected.glsl index 73ed75be79..3d082fd3e3 100644 --- a/test/tint/var/override/named/zero_init/f32.wgsl.expected.glsl +++ b/test/tint/var/override/named/zero_init/f32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 0.0f -#endif -const float o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0.0f == 0.0f)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/zero_init/f32.wgsl.expected.hlsl b/test/tint/var/override/named/zero_init/f32.wgsl.expected.hlsl index 5ae59a4ccf..e42b8fb027 100644 --- a/test/tint/var/override/named/zero_init/f32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/zero_init/f32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 0.0f -#endif -static const float o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/f32.wgsl.expected.msl b/test/tint/var/override/named/zero_init/f32.wgsl.expected.msl index ecf89c4c91..ecda7f6a64 100644 --- a/test/tint/var/override/named/zero_init/f32.wgsl.expected.msl +++ b/test/tint/var/override/named/zero_init/f32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant float o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/f32.wgsl.expected.spvasm b/test/tint/var/override/named/zero_init/f32.wgsl.expected.spvasm index 7ae9060492..9190d7bd41 100644 --- a/test/tint/var/override/named/zero_init/f32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/zero_init/f32.wgsl.expected.spvasm @@ -1,20 +1,25 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 11 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %float = OpTypeFloat 32 - %o = OpSpecConstant %float 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %6 = OpConstantNull %float + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %7 = OpFOrdEqual %bool %6 %6 + OpSelectionMerge %9 None + OpBranchConditional %7 %10 %9 + %10 = OpLabel + OpBranch %9 + %9 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/zero_init/f32.wgsl.expected.wgsl b/test/tint/var/override/named/zero_init/f32.wgsl.expected.wgsl index b3b26c1145..80171943a9 100644 --- a/test/tint/var/override/named/zero_init/f32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/zero_init/f32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : f32 = f32(); +const o : f32 = 0.0f; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 0.0)) { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/i32.wgsl b/test/tint/var/override/named/zero_init/i32.wgsl index 21cb5cbeb2..6a8fc527ce 100644 --- a/test/tint/var/override/named/zero_init/i32.wgsl +++ b/test/tint/var/override/named/zero_init/i32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides o=0 override o : i32 = i32(); @compute @workgroup_size(1) fn main() { - _ = o; + if o == 2 { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/i32.wgsl.expected.glsl b/test/tint/var/override/named/zero_init/i32.wgsl.expected.glsl index d222d52e23..e6f817c290 100644 --- a/test/tint/var/override/named/zero_init/i32.wgsl.expected.glsl +++ b/test/tint/var/override/named/zero_init/i32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 0 -#endif -const int o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0 == 2)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/zero_init/i32.wgsl.expected.hlsl b/test/tint/var/override/named/zero_init/i32.wgsl.expected.hlsl index 6af794b4cb..21744ce5ca 100644 --- a/test/tint/var/override/named/zero_init/i32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/zero_init/i32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 0 -#endif -static const int o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0 == 2)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/i32.wgsl.expected.msl b/test/tint/var/override/named/zero_init/i32.wgsl.expected.msl index 6df39feed4..4563b8d153 100644 --- a/test/tint/var/override/named/zero_init/i32.wgsl.expected.msl +++ b/test/tint/var/override/named/zero_init/i32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant int o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0 == 2)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/i32.wgsl.expected.spvasm b/test/tint/var/override/named/zero_init/i32.wgsl.expected.spvasm index 18c8840440..5e40713526 100644 --- a/test/tint/var/override/named/zero_init/i32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/zero_init/i32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %int = OpTypeInt 32 1 - %o = OpSpecConstant %int 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %6 = OpConstantNull %int + %int_2 = OpConstant %int 2 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %int_2 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/zero_init/i32.wgsl.expected.wgsl b/test/tint/var/override/named/zero_init/i32.wgsl.expected.wgsl index b699ac1847..6675efccb5 100644 --- a/test/tint/var/override/named/zero_init/i32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/zero_init/i32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : i32 = i32(); +const o : i32 = 0i; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 2)) { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/u32.wgsl b/test/tint/var/override/named/zero_init/u32.wgsl index 137b5f8167..84f53a44b1 100644 --- a/test/tint/var/override/named/zero_init/u32.wgsl +++ b/test/tint/var/override/named/zero_init/u32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides o=0 override o : u32 = u32(); @compute @workgroup_size(1) fn main() { - _ = o; + if o == 2 { + _ = o; + } } diff --git a/test/tint/var/override/named/zero_init/u32.wgsl.expected.glsl b/test/tint/var/override/named/zero_init/u32.wgsl.expected.glsl index 76638469b9..b19260b440 100644 --- a/test/tint/var/override/named/zero_init/u32.wgsl.expected.glsl +++ b/test/tint/var/override/named/zero_init/u32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 0u -#endif -const uint o = WGSL_SPEC_CONSTANT_0; void tint_symbol() { + if ((0u == 2u)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/named/zero_init/u32.wgsl.expected.hlsl b/test/tint/var/override/named/zero_init/u32.wgsl.expected.hlsl index 17f5ddf612..b51955d36d 100644 --- a/test/tint/var/override/named/zero_init/u32.wgsl.expected.hlsl +++ b/test/tint/var/override/named/zero_init/u32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_0 -#define WGSL_SPEC_CONSTANT_0 0u -#endif -static const uint o = WGSL_SPEC_CONSTANT_0; - [numthreads(1, 1, 1)] void main() { + if ((0u == 2u)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/u32.wgsl.expected.msl b/test/tint/var/override/named/zero_init/u32.wgsl.expected.msl index 8254bed60e..d5f3f223da 100644 --- a/test/tint/var/override/named/zero_init/u32.wgsl.expected.msl +++ b/test/tint/var/override/named/zero_init/u32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant uint o [[function_constant(0)]]; - kernel void tint_symbol() { + if ((0u == 2u)) { + } return; } diff --git a/test/tint/var/override/named/zero_init/u32.wgsl.expected.spvasm b/test/tint/var/override/named/zero_init/u32.wgsl.expected.spvasm index e5486b0dfb..4212d43c39 100644 --- a/test/tint/var/override/named/zero_init/u32.wgsl.expected.spvasm +++ b/test/tint/var/override/named/zero_init/u32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 0 - %uint = OpTypeInt 32 0 - %o = OpSpecConstant %uint 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %6 = OpConstantNull %uint + %uint_2 = OpConstant %uint 2 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %uint_2 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/named/zero_init/u32.wgsl.expected.wgsl b/test/tint/var/override/named/zero_init/u32.wgsl.expected.wgsl index 01620769ca..020319ae6b 100644 --- a/test/tint/var/override/named/zero_init/u32.wgsl.expected.wgsl +++ b/test/tint/var/override/named/zero_init/u32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -override o : u32 = u32(); +const o : u32 = 0u; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 2)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl b/test/tint/var/override/numbered/no_init/bool.wgsl index 955c29f1e6..ded2aa4e55 100644 --- a/test/tint/var/override/numbered/no_init/bool.wgsl +++ b/test/tint/var/override/numbered/no_init/bool.wgsl @@ -1,7 +1,9 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_1234=0 +// flags: --overrides 1234=0 @id(1234) override o : bool; @compute @workgroup_size(1) fn main() { - _ = o; + if o { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.glsl b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.glsl index 2e5beb9985..94d2bb42ca 100644 --- a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.glsl @@ -1,12 +1,8 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -const bool o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if (false) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +10,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 1234 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl index fd25726f1a..af95aed96e 100644 --- a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -static const bool o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if (false) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.msl b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.msl index 28e456b41b..6fa26c20be 100644 --- a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.msl +++ b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant bool o [[function_constant(1234)]]; - kernel void tint_symbol() { + if (false) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.spvasm b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.spvasm index b1169e6faa..14d2a40e27 100644 --- a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.spvasm @@ -1,20 +1,23 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 9 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %bool = OpTypeBool - %o = OpSpecConstantFalse %bool %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %bool = OpTypeBool + %6 = OpConstantNull %bool + %main = OpFunction %void None %1 + %4 = OpLabel + OpSelectionMerge %7 None + OpBranchConditional %6 %8 %7 + %8 = OpLabel + OpBranch %7 + %7 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.wgsl b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.wgsl index fac817d292..f662a7e29d 100644 --- a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : bool; +const o : bool = false; @compute @workgroup_size(1) fn main() { - _ = o; + if (o) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl b/test/tint/var/override/numbered/no_init/f32.wgsl index ecef6e18bd..52af3f1244 100644 --- a/test/tint/var/override/numbered/no_init/f32.wgsl +++ b/test/tint/var/override/numbered/no_init/f32.wgsl @@ -1,7 +1,9 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_1234=0 +// flags: --overrides 1234=0 @id(1234) override o : f32; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 0.0 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.glsl b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.glsl index ef6bc78b2c..3d082fd3e3 100644 --- a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.glsl @@ -1,12 +1,8 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -const float o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0.0f == 0.0f)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +10,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 1234 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl index 2ae70421ea..e42b8fb027 100644 --- a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -static const float o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.msl b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.msl index ce09e6ef83..ecda7f6a64 100644 --- a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant float o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.spvasm b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.spvasm index a253f64b1c..9190d7bd41 100644 --- a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.spvasm @@ -1,20 +1,25 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 11 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %float = OpTypeFloat 32 - %o = OpSpecConstant %float 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %6 = OpConstantNull %float + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %7 = OpFOrdEqual %bool %6 %6 + OpSelectionMerge %9 None + OpBranchConditional %7 %10 %9 + %10 = OpLabel + OpBranch %9 + %9 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.wgsl b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.wgsl index b7cf516305..80171943a9 100644 --- a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : f32; +const o : f32 = 0.0f; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 0.0)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl b/test/tint/var/override/numbered/no_init/i32.wgsl index a347fc6313..2296dece78 100644 --- a/test/tint/var/override/numbered/no_init/i32.wgsl +++ b/test/tint/var/override/numbered/no_init/i32.wgsl @@ -1,7 +1,9 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_1234=0 +// flags: --overrides 1234=0 @id(1234) override o : i32; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 0 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.glsl b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.glsl index 6dabe7ff6c..9edb6dda05 100644 --- a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.glsl @@ -1,12 +1,8 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -const int o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0 == 0)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +10,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 1234 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl index a02d224038..95e256f3a8 100644 --- a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -static const int o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0 == 0)) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.msl b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.msl index 3e8475293b..abeb66026c 100644 --- a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant int o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0 == 0)) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.spvasm b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.spvasm index aca0b07388..a26f8edc2b 100644 --- a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.spvasm @@ -1,20 +1,25 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 11 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %int = OpTypeInt 32 1 - %o = OpSpecConstant %int 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %6 = OpConstantNull %int + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %7 = OpIEqual %bool %6 %6 + OpSelectionMerge %9 None + OpBranchConditional %7 %10 %9 + %10 = OpLabel + OpBranch %9 + %9 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.wgsl b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.wgsl index 8538212210..04d8f90c48 100644 --- a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : i32; +const o : i32 = 0i; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 0)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl b/test/tint/var/override/numbered/no_init/u32.wgsl index 0fe7b1293d..c00557b200 100644 --- a/test/tint/var/override/numbered/no_init/u32.wgsl +++ b/test/tint/var/override/numbered/no_init/u32.wgsl @@ -1,7 +1,9 @@ -// flags: --overrides WGSL_SPEC_CONSTANT_1234=0 +// flags: --overrides 1234=0 @id(1234) override o : u32; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.glsl b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.glsl index fcf949f115..368da7abbf 100644 --- a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.glsl @@ -1,12 +1,8 @@ -SKIP: FAILED - #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -const uint o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0u == 1u)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; @@ -14,11 +10,3 @@ void main() { tint_symbol(); return; } -Error parsing GLSL shader: -ERROR: 0:4: '#error' : spec constant required for constant id 1234 -ERROR: 0:5: '' : missing #endif -ERROR: 0:5: '' : compilation terminated -ERROR: 3 compilation errors. No code generated. - - - diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl index 4a3a2e1cfc..fe084141f7 100644 --- a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#error spec constant required for constant id 1234 -#endif -static const uint o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.msl b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.msl index 51e2d724d4..641561bdc7 100644 --- a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant uint o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.spvasm b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.spvasm index ca38a9a94c..fa0711ab91 100644 --- a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %uint = OpTypeInt 32 0 - %o = OpSpecConstant %uint 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %6 = OpConstantNull %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %uint_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.wgsl b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.wgsl index a187373791..b7ac534cc5 100644 --- a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : u32; +const o : u32 = 0u; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/bool.wgsl b/test/tint/var/override/numbered/val_init/bool.wgsl index 17983bbf01..0d6a548b05 100644 --- a/test/tint/var/override/numbered/val_init/bool.wgsl +++ b/test/tint/var/override/numbered/val_init/bool.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : bool = true; @compute @workgroup_size(1) fn main() { - _ = o; + if o { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.glsl b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.glsl index 3b5cc68b4a..94d2bb42ca 100644 --- a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 true -#endif -const bool o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if (false) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.hlsl b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.hlsl index 6e189ede17..af95aed96e 100644 --- a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 true -#endif -static const bool o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if (false) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.msl b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.msl index 28e456b41b..6fa26c20be 100644 --- a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.msl +++ b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant bool o [[function_constant(1234)]]; - kernel void tint_symbol() { + if (false) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.spvasm b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.spvasm index 060f8603f6..14d2a40e27 100644 --- a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.spvasm @@ -1,20 +1,23 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 9 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %bool = OpTypeBool - %o = OpSpecConstantTrue %bool %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %bool = OpTypeBool + %6 = OpConstantNull %bool + %main = OpFunction %void None %1 + %4 = OpLabel + OpSelectionMerge %7 None + OpBranchConditional %6 %8 %7 + %8 = OpLabel + OpBranch %7 + %7 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.wgsl b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.wgsl index ff5606e60b..f662a7e29d 100644 --- a/test/tint/var/override/numbered/val_init/bool.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/val_init/bool.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : bool = true; +const o : bool = false; @compute @workgroup_size(1) fn main() { - _ = o; + if (o) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/f32.wgsl b/test/tint/var/override/numbered/val_init/f32.wgsl index 55635abf4f..f76de51133 100644 --- a/test/tint/var/override/numbered/val_init/f32.wgsl +++ b/test/tint/var/override/numbered/val_init/f32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : f32 = 1.0; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.glsl b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.glsl index 5d67c5f1dc..7540a29c95 100644 --- a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 1.0f -#endif -const float o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0.0f == 1.0f)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.hlsl b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.hlsl index 4671b12b3b..0d9f804f9d 100644 --- a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 1.0f -#endif -static const float o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0.0f == 1.0f)) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.msl b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.msl index ce09e6ef83..5936a883f6 100644 --- a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant float o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0.0f == 1.0f)) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.spvasm b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.spvasm index c75755700c..704864876a 100644 --- a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %float = OpTypeFloat 32 - %o = OpSpecConstant %float 1 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %6 = OpConstantNull %float + %float_1 = OpConstant %float 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpFOrdEqual %bool %6 %float_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.wgsl b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.wgsl index 7c7a9378ad..c97a28d237 100644 --- a/test/tint/var/override/numbered/val_init/f32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/val_init/f32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : f32 = 1.0; +const o : f32 = 0.0f; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/i32.wgsl b/test/tint/var/override/numbered/val_init/i32.wgsl index d763d36595..6220760ced 100644 --- a/test/tint/var/override/numbered/val_init/i32.wgsl +++ b/test/tint/var/override/numbered/val_init/i32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : i32 = 1; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.glsl b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.glsl index 034f0daa88..aadd0384bb 100644 --- a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 1 -#endif -const int o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0 == 1)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.hlsl b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.hlsl index cfb2fd9371..d384517be6 100644 --- a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 1 -#endif -static const int o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0 == 1)) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.msl b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.msl index 3e8475293b..4bcde5e5ec 100644 --- a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant int o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0 == 1)) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.spvasm b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.spvasm index f66889a42e..74678ef974 100644 --- a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %int = OpTypeInt 32 1 - %o = OpSpecConstant %int 1 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %6 = OpConstantNull %int + %int_1 = OpConstant %int 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %int_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.wgsl b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.wgsl index 84151e88d4..88b26b31eb 100644 --- a/test/tint/var/override/numbered/val_init/i32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/val_init/i32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : i32 = 1; +const o : i32 = 0i; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/u32.wgsl b/test/tint/var/override/numbered/val_init/u32.wgsl index 107bb310d0..e28df648b0 100644 --- a/test/tint/var/override/numbered/val_init/u32.wgsl +++ b/test/tint/var/override/numbered/val_init/u32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : u32 = 1u; @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.glsl b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.glsl index 2b50d92d67..368da7abbf 100644 --- a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 1u -#endif -const uint o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0u == 1u)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.hlsl b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.hlsl index 36437dd9e2..fe084141f7 100644 --- a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 1u -#endif -static const uint o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.msl b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.msl index 51e2d724d4..641561bdc7 100644 --- a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant uint o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.spvasm b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.spvasm index 2b51d1ce62..fa0711ab91 100644 --- a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %uint = OpTypeInt 32 0 - %o = OpSpecConstant %uint 1 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %6 = OpConstantNull %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %uint_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.wgsl b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.wgsl index 9ea424d3a8..b7ac534cc5 100644 --- a/test/tint/var/override/numbered/val_init/u32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/val_init/u32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : u32 = 1u; +const o : u32 = 0u; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/bool.wgsl b/test/tint/var/override/numbered/zero_init/bool.wgsl index 531d595cf9..929b2e1df3 100644 --- a/test/tint/var/override/numbered/zero_init/bool.wgsl +++ b/test/tint/var/override/numbered/zero_init/bool.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : bool = bool(); @compute @workgroup_size(1) fn main() { - _ = o; + if o { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.glsl b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.glsl index 072eb92f2c..94d2bb42ca 100644 --- a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 false -#endif -const bool o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if (false) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.hlsl b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.hlsl index 86e9b98f87..af95aed96e 100644 --- a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 false -#endif -static const bool o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if (false) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.msl b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.msl index 28e456b41b..6fa26c20be 100644 --- a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.msl +++ b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant bool o [[function_constant(1234)]]; - kernel void tint_symbol() { + if (false) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.spvasm b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.spvasm index b1169e6faa..14d2a40e27 100644 --- a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.spvasm @@ -1,20 +1,23 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 9 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %bool = OpTypeBool - %o = OpSpecConstantFalse %bool %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %bool = OpTypeBool + %6 = OpConstantNull %bool + %main = OpFunction %void None %1 + %4 = OpLabel + OpSelectionMerge %7 None + OpBranchConditional %6 %8 %7 + %8 = OpLabel + OpBranch %7 + %7 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.wgsl b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.wgsl index 1f7f861313..f662a7e29d 100644 --- a/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/zero_init/bool.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : bool = bool(); +const o : bool = false; @compute @workgroup_size(1) fn main() { - _ = o; + if (o) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/f32.wgsl b/test/tint/var/override/numbered/zero_init/f32.wgsl index 40361fcf08..6308a76c3e 100644 --- a/test/tint/var/override/numbered/zero_init/f32.wgsl +++ b/test/tint/var/override/numbered/zero_init/f32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : f32 = f32(); @compute @workgroup_size(1) fn main() { - _ = o; + if o == 0.0 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.glsl b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.glsl index fe17e5b528..3d082fd3e3 100644 --- a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 0.0f -#endif -const float o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0.0f == 0.0f)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.hlsl b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.hlsl index 70f868fce9..e42b8fb027 100644 --- a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 0.0f -#endif -static const float o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.msl b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.msl index ce09e6ef83..ecda7f6a64 100644 --- a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant float o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0.0f == 0.0f)) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.spvasm b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.spvasm index a253f64b1c..9190d7bd41 100644 --- a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.spvasm @@ -1,20 +1,25 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 11 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %float = OpTypeFloat 32 - %o = OpSpecConstant %float 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %float = OpTypeFloat 32 + %6 = OpConstantNull %float + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %7 = OpFOrdEqual %bool %6 %6 + OpSelectionMerge %9 None + OpBranchConditional %7 %10 %9 + %10 = OpLabel + OpBranch %9 + %9 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.wgsl b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.wgsl index af740ac3e3..80171943a9 100644 --- a/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/zero_init/f32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : f32 = f32(); +const o : f32 = 0.0f; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 0.0)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/i32.wgsl b/test/tint/var/override/numbered/zero_init/i32.wgsl index c1e8fe712f..621161e679 100644 --- a/test/tint/var/override/numbered/zero_init/i32.wgsl +++ b/test/tint/var/override/numbered/zero_init/i32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : i32 = i32(); @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.glsl b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.glsl index a748cb70df..aadd0384bb 100644 --- a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 0 -#endif -const int o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0 == 1)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.hlsl b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.hlsl index 654ffd6295..d384517be6 100644 --- a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 0 -#endif -static const int o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0 == 1)) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.msl b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.msl index 3e8475293b..4bcde5e5ec 100644 --- a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant int o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0 == 1)) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.spvasm b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.spvasm index aca0b07388..74678ef974 100644 --- a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %int = OpTypeInt 32 1 - %o = OpSpecConstant %int 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %6 = OpConstantNull %int + %int_1 = OpConstant %int 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %int_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.wgsl b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.wgsl index ac2516f0ac..88b26b31eb 100644 --- a/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/zero_init/i32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : i32 = i32(); +const o : i32 = 0i; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/u32.wgsl b/test/tint/var/override/numbered/zero_init/u32.wgsl index 183862ff93..fbaf795ed1 100644 --- a/test/tint/var/override/numbered/zero_init/u32.wgsl +++ b/test/tint/var/override/numbered/zero_init/u32.wgsl @@ -1,6 +1,9 @@ +// flags: --overrides 1234=0 @id(1234) override o : u32 = u32(); @compute @workgroup_size(1) fn main() { - _ = o; + if o == 1 { + _ = o; + } } diff --git a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.glsl b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.glsl index 85057c1473..368da7abbf 100644 --- a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.glsl +++ b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.glsl @@ -1,10 +1,8 @@ #version 310 es -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 0u -#endif -const uint o = WGSL_SPEC_CONSTANT_1234; void tint_symbol() { + if ((0u == 1u)) { + } } layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.hlsl b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.hlsl index 2fe25d834a..fe084141f7 100644 --- a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.hlsl +++ b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.hlsl @@ -1,9 +1,6 @@ -#ifndef WGSL_SPEC_CONSTANT_1234 -#define WGSL_SPEC_CONSTANT_1234 0u -#endif -static const uint o = WGSL_SPEC_CONSTANT_1234; - [numthreads(1, 1, 1)] void main() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.msl b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.msl index 51e2d724d4..641561bdc7 100644 --- a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.msl +++ b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.msl @@ -1,9 +1,9 @@ #include using namespace metal; -constant uint o [[function_constant(1234)]]; - kernel void tint_symbol() { + if ((0u == 1u)) { + } return; } diff --git a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.spvasm b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.spvasm index ca38a9a94c..fa0711ab91 100644 --- a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.spvasm +++ b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.spvasm @@ -1,20 +1,26 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 7 +; Bound: 12 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %main "main" OpExecutionMode %main LocalSize 1 1 1 - OpName %o "o" OpName %main "main" - OpDecorate %o SpecId 1234 - %uint = OpTypeInt 32 0 - %o = OpSpecConstant %uint 0 %void = OpTypeVoid - %3 = OpTypeFunction %void - %main = OpFunction %void None %3 - %6 = OpLabel + %1 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %6 = OpConstantNull %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %main = OpFunction %void None %1 + %4 = OpLabel + %8 = OpIEqual %bool %6 %uint_1 + OpSelectionMerge %10 None + OpBranchConditional %8 %11 %10 + %11 = OpLabel + OpBranch %10 + %10 = OpLabel OpReturn OpFunctionEnd diff --git a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.wgsl b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.wgsl index b4551272d5..b7ac534cc5 100644 --- a/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.wgsl +++ b/test/tint/var/override/numbered/zero_init/u32.wgsl.expected.wgsl @@ -1,6 +1,8 @@ -@id(1234) override o : u32 = u32(); +const o : u32 = 0u; @compute @workgroup_size(1) fn main() { - _ = o; + if ((o == 1)) { + _ = o; + } }