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 <bclayton@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
dan sinclair 2022-07-22 16:05:06 +00:00 committed by Dawn LUCI CQ
parent 9ec7893ad4
commit 256f1116b8
156 changed files with 1235 additions and 659 deletions

View File

@ -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"

View File

@ -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",

View File

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

View File

@ -46,4 +46,11 @@ const Override* Override::Clone(CloneContext* ctx) const {
return ctx->dst->create<Override>(src, sym, ty, ctor, attrs);
}
std::string Override::Identifier(const SymbolTable& symbols) const {
if (auto* id = ast::GetAttribute<ast::IdAttribute>(attributes)) {
return std::to_string(id->value);
}
return symbols.NameFor(symbol);
}
} // namespace tint::ast

View File

@ -15,6 +15,8 @@
#ifndef SRC_TINT_AST_OVERRIDE_H_
#define SRC_TINT_AST_OVERRIDE_H_
#include <string>
#include "src/tint/ast/variable.h"
namespace tint::ast {
@ -60,6 +62,12 @@ class Override final : public Castable<Override, Variable> {
/// @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

View File

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

View File

@ -20,6 +20,7 @@
#include <optional>
#include <sstream>
#include <string>
#include <unordered_map>
#include <vector>
#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<std::string> overrides;
std::unordered_map<std::string, double> overrides;
std::optional<tint::sem::BindingPoint> 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<std::string> split_on_comma(std::string list) {
std::vector<std::string> split_on_char(std::string list, char c) {
std::vector<std::string> 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<std::string> split_on_comma(std::string list) {
return split_on_char(list, ',');
}
std::vector<std::string> split_on_equal(std::string list) {
return split_on_char(list, '=');
}
std::optional<uint64_t> parse_unsigned_number(std::string number) {
for (char c : number) {
if (!std::isdigit(c)) {
@ -429,7 +439,10 @@ bool ParseArgs(const std::vector<std::string>& 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<tint::transform::Renamer>(); }},
{"robustness", [](tint::transform::Manager& m,
tint::transform::DataMap&) { m.Add<tint::transform::Robustness>(); }},
{"substitute_override",
[&](tint::transform::Manager& m, tint::transform::DataMap& i) {
tint::transform::SubstituteOverride::Config cfg;
cfg.map = options.overrides;
i.Add<tint::transform::SubstituteOverride::Config>(cfg);
m.Add<tint::transform::SubstituteOverride>();
}},
};
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

View File

@ -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 <functional>
#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<ast::Override>()) {
return true;
}
}
return false;
}
void SubstituteOverride::Run(CloneContext& ctx, const DataMap& config, DataMap&) const {
const auto* data = config.Get<Config>();
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<double>()(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

View File

@ -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 <string>
#include <unordered_map>
#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<SubstituteOverride, Transform> {
public:
/// Configuration options for the transform
struct Config final : public Castable<Config, Data> {
/// 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<std::string, double> 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_

View File

@ -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<f32> {
return vec4<f32>();
}
)";
auto* expect = "error: Missing override substitution data";
DataMap data;
auto got = Run<SubstituteOverride>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SubstituteOverrideTest, Error_NoOverrideValue) {
auto* src = R"(
override width: i32;
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
)";
auto* expect = "error: Initializer not provided for override, and override not overridden.";
SubstituteOverride::Config cfg;
DataMap data;
data.Add<SubstituteOverride::Config>(cfg);
auto got = Run<SubstituteOverride>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SubstituteOverrideTest, Module_NoOverrides) {
auto* src = R"(
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
)";
auto* expect = R"(
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
)";
SubstituteOverride::Config cfg;
DataMap data;
data.Add<SubstituteOverride::Config>(cfg);
auto got = Run<SubstituteOverride>(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<f32> {
return vec4<f32>();
}
)";
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<f32> {
return vec4<f32>();
}
)";
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<SubstituteOverride::Config>(cfg);
auto got = Run<SubstituteOverride>(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<f32> {
return vec4<f32>();
}
)";
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<f32> {
return vec4<f32>();
}
)";
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<SubstituteOverride::Config>(cfg);
auto got = Run<SubstituteOverride>(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<f32> {
return vec4<f32>();
}
)";
auto* expect = R"(
const i_height = 11i;
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
)";
SubstituteOverride::Config cfg;
cfg.map.insert({"i_height", 11.0});
DataMap data;
data.Add<SubstituteOverride::Config>(cfg);
auto got = Run<SubstituteOverride>(src, data);
EXPECT_EQ(expect, str(got));
}
} // namespace
} // namespace tint::transform

View File

@ -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<std::string>& 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<std::string>& 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<D3D_SHADER_MACRO> 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<ID3DBlob> compiledShader;
ComPtr<ID3DBlob> 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

View File

@ -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<std::string>& 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<std::string>& overrides);
Result HlslUsingFXC(const std::string& source, const EntryPointList& entry_points);
#endif // _WIN32
/// Msl attempts to compile the shader with the Metal Shader Compiler,

View File

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

View File

@ -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.

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(0)]];
kernel void tint_symbol() {
if ((false && true)) {
}
return;
}

View File

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

View File

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

View File

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

View File

@ -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.

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0.0f == 0.0f)) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
override o : f32;
const o : f32 = 0.0f;
@compute @workgroup_size(1)
fn main() {
_ = o;
if ((o == 0.0)) {
_ = 1;
}
}

View File

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

View File

@ -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.

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0 == 1)) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
override o : i32;
const o : i32 = 0i;
@compute @workgroup_size(1)
fn main() {
_ = o;
if ((o == 1)) {
_ = o;
}
}

View File

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

View File

@ -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.

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant uint o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0u == 1u)) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
override o : u32;
const o : u32 = 0u;
@compute @workgroup_size(1)
fn main() {
_ = o;
if ((o == 1)) {
_ = o;
}
}

View File

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

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(0)]];
kernel void tint_symbol() {
if ((false && true)) {
}
return;
}

View File

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

View File

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

View File

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

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0.0f == 0.0f)) {
}
return;
}

View File

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

View File

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

View File

@ -1,6 +1,9 @@
// flags: --overrides o=0
override o : i32 = 1;
@compute @workgroup_size(1)
fn main() {
_ = o;
if o == 2 {
_ = o;
}
}

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0 == 2)) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
override o : i32 = 1;
const o : i32 = 0i;
@compute @workgroup_size(1)
fn main() {
_ = o;
if ((o == 2)) {
_ = o;
}
}

View File

@ -1,6 +1,9 @@
// flags: --overrides o=0
override o : u32 = 1u;
@compute @workgroup_size(1)
fn main() {
_ = o;
if o == 2 {
_ = o;
}
}

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant uint o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0u == 2u)) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
override o : u32 = 1u;
const o : u32 = 0u;
@compute @workgroup_size(1)
fn main() {
_ = o;
if ((o == 2)) {
_ = o;
}
}

View File

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

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(0)]];
kernel void tint_symbol() {
if ((false && true)) {
}
return;
}

View File

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

View File

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

View File

@ -1,6 +1,9 @@
// flags: --overrides o=0
override o : f32 = f32();
@compute @workgroup_size(1)
fn main() {
_ = o;
if o == 0.0 {
_ = o;
}
}

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0.0f == 0.0f)) {
}
return;
}

View File

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

View File

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

View File

@ -1,6 +1,9 @@
// flags: --overrides o=0
override o : i32 = i32();
@compute @workgroup_size(1)
fn main() {
_ = o;
if o == 2 {
_ = o;
}
}

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0 == 2)) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
override o : i32 = i32();
const o : i32 = 0i;
@compute @workgroup_size(1)
fn main() {
_ = o;
if ((o == 2)) {
_ = o;
}
}

View File

@ -1,6 +1,9 @@
// flags: --overrides o=0
override o : u32 = u32();
@compute @workgroup_size(1)
fn main() {
_ = o;
if o == 2 {
_ = o;
}
}

View File

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

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant uint o [[function_constant(0)]];
kernel void tint_symbol() {
if ((0u == 2u)) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
override o : u32 = u32();
const o : u32 = 0u;
@compute @workgroup_size(1)
fn main() {
_ = o;
if ((o == 2)) {
_ = o;
}
}

View File

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

View File

@ -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.

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant bool o [[function_constant(1234)]];
kernel void tint_symbol() {
if (false) {
}
return;
}

View File

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

View File

@ -1,6 +1,8 @@
@id(1234) override o : bool;
const o : bool = false;
@compute @workgroup_size(1)
fn main() {
_ = o;
if (o) {
_ = o;
}
}

View File

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

View File

@ -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.

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant float o [[function_constant(1234)]];
kernel void tint_symbol() {
if ((0.0f == 0.0f)) {
}
return;
}

View File

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

View File

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

View File

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

View File

@ -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.

View File

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

View File

@ -1,9 +1,9 @@
#include <metal_stdlib>
using namespace metal;
constant int o [[function_constant(1234)]];
kernel void tint_symbol() {
if ((0 == 0)) {
}
return;
}

Some files were not shown because too many files have changed in this diff Show More