Remove unrollConstEvalLoops flag from CTS runner
Issues on ARM Mac and apple silicon have been resolved.
Issues on compat (pixel6) have been added to expectations.txt
(these are pre-existing issues that have been exposed by this change
now that we are no longer loop unrolling these test cases)
I have investigated timings on Intel-Mac CTS runs and they are
equivalent before and after this change.
End2end test that verifies max compile/runtime:
https://dawn-review.googlesource.com/c/dawn/+/243854
Bug: 417519810
Change-Id: I03074e233ac266613d419ab20a95d82f45a24391
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/245254
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp
index dd87d12..7ab0247 100644
--- a/src/dawn/native/Toggles.cpp
+++ b/src/dawn/native/Toggles.cpp
@@ -573,8 +573,9 @@
{Toggle::ScalarizeMaxMinClamp,
{"scalarize_max_min_clamp", "Scalarize max, min, and clamp builtins.",
"https://crbug.com/422144514", ToggleStage::Device}},
- {Toggle::MetalEnableModuleConstant,
- {"metal_enable_module_constant_transform", "Enable the module constant transform.",
+ {Toggle::MetalDisableModuleConstantF16,
+ {"metal_disable_module_constant_f16",
+ "Disable module constant hoisting for values that contain f16 types.",
"https://crbug.com/419804339", ToggleStage::Device}},
{Toggle::EnableImmediateErrorHandling,
{"enable_immediate_error_handling",
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h
index e939459..4650c45 100644
--- a/src/dawn/native/Toggles.h
+++ b/src/dawn/native/Toggles.h
@@ -140,7 +140,7 @@
ExposeWGSLExperimentalFeatures,
DisablePolyfillsOnIntegerDivisonAndModulo,
ScalarizeMaxMinClamp,
- MetalEnableModuleConstant,
+ MetalDisableModuleConstantF16,
EnableImmediateErrorHandling,
VulkanUseStorageInputOutput16,
D3D12DontUseShaderModel66OrHigher,
diff --git a/src/dawn/native/metal/PhysicalDeviceMTL.mm b/src/dawn/native/metal/PhysicalDeviceMTL.mm
index 1fa71ae..73c9fef 100644
--- a/src/dawn/native/metal/PhysicalDeviceMTL.mm
+++ b/src/dawn/native/metal/PhysicalDeviceMTL.mm
@@ -442,13 +442,10 @@
deviceToggles->Default(Toggle::MetalRenderR8RG8UnormSmallMipToTempTexture, true);
}
- // chromium:419804339: Module constant hoisting is broadly available as a msl transform but
- // there are execution correction issues with f16 for non apple silicon (Intel/AMD). Therefore
- // we only enable for apple silicon for now.
- // chromium:417519810: Mutiple cts tests will fail on AMD if module scope hoisting is not
- // enabled on AMD. These failures will be internal compiler errors.
- if (gpu_info::IsApple(vendorId) || gpu_info::IsAMD(vendorId)) {
- deviceToggles->Default(Toggle::MetalEnableModuleConstant, true);
+ // chromium:419804339: Module constant hoisting is not supported for values containing f16
+ // types on Intel.
+ if (gpu_info::IsIntel(vendorId)) {
+ deviceToggles->Default(Toggle::MetalDisableModuleConstantF16, true);
}
// On some Intel GPUs vertex only render pipeline get wrong depth result if no fragment
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm
index ec35b39..d5af501 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -296,8 +296,8 @@
req.tintOptions.disable_polyfill_integer_div_mod =
device->IsToggleEnabled(Toggle::DisablePolyfillsOnIntegerDivisonAndModulo);
req.tintOptions.scalarize_max_min_clamp = device->IsToggleEnabled(Toggle::ScalarizeMaxMinClamp);
- req.tintOptions.enable_module_constant =
- device->IsToggleEnabled(Toggle::MetalEnableModuleConstant);
+ req.tintOptions.disable_module_constant_f16 =
+ device->IsToggleEnabled(Toggle::MetalDisableModuleConstantF16);
req.tintOptions.vertex_pulling_config = std::move(vertexPullingTransformConfig);
req.tintOptions.enable_integer_range_analysis =
device->IsToggleEnabled(Toggle::EnableIntegerRangeAnalysisInRobustness);
diff --git a/src/dawn/node/README.md b/src/dawn/node/README.md
index ed325b7..cea704d 100644
--- a/src/dawn/node/README.md
+++ b/src/dawn/node/README.md
@@ -272,7 +272,7 @@
loop:nested_loops:preventValueOptimizations=false'
<SNIP>
Running:
- Cmd: /home/user/src/dawn/third_party/node/node-linux-x64/bin/node -e "require('./out-node/common/runtime/cmdline.js');" -- placeholder-arg --gpu-provider /home/user/src/dawn/build-clang/cts.js --verbose --quiet --gpu-provider-flag verbose=1 --colors --unroll-const-eval-loops --gpu-provider-flag enable-dawn-features=allow_unsafe_apis "webgpu:shader,execution,flow_control,loop:nested_loops:preventValueOptimizations=false"
+ Cmd: /home/user/src/dawn/third_party/node/node-linux-x64/bin/node -e "require('./out-node/common/runtime/cmdline.js');" -- placeholder-arg --gpu-provider /home/user/src/dawn/build-clang/cts.js --verbose --quiet --gpu-provider-flag verbose=1 --colors --gpu-provider-flag enable-dawn-features=allow_unsafe_apis "webgpu:shader,execution,flow_control,loop:nested_loops:preventValueOptimizations=false"
Dir: /home/user/src/dawn/third_party/webgpu-cts
For VS Code launch.json:
@@ -289,7 +289,6 @@
"--gpu-provider-flag",
"verbose=1",
"--colors",
- "--unroll-const-eval-loops",
"--gpu-provider-flag",
"enable-dawn-features=allow_unsafe_apis",
"webgpu:shader,execution,flow_control,loop:nested_loops:preventValueOptimizations=false"
diff --git a/src/tint/lang/msl/writer/common/options.h b/src/tint/lang/msl/writer/common/options.h
index 925dcfb..551ee4c 100644
--- a/src/tint/lang/msl/writer/common/options.h
+++ b/src/tint/lang/msl/writer/common/options.h
@@ -166,8 +166,8 @@
/// Set to `true` to scalarize max min and clamp builtins.
bool scalarize_max_min_clamp = false;
- /// Set to `true` to enable the module constant transform
- bool enable_module_constant = false;
+ /// Set to `true` to disable the module constant transform for f16
+ bool disable_module_constant_f16 = false;
/// Emit argument buffers
bool use_argument_buffers = false;
@@ -204,7 +204,7 @@
emit_vertex_point_size,
disable_polyfill_integer_div_mod,
scalarize_max_min_clamp,
- enable_module_constant,
+ disable_module_constant_f16,
use_argument_buffers,
buffer_size_ubo_index,
fixed_sample_mask,
diff --git a/src/tint/lang/msl/writer/raise/module_constant.cc b/src/tint/lang/msl/writer/raise/module_constant.cc
index 6c9c8b6..3d31ae4 100644
--- a/src/tint/lang/msl/writer/raise/module_constant.cc
+++ b/src/tint/lang/msl/writer/raise/module_constant.cc
@@ -42,9 +42,32 @@
/// PIMPL state for the transform.
struct State {
core::ir::Module& ir;
+ const ModuleConstantConfig& config;
core::ir::Builder b{ir};
core::type::Manager& ty{ir.Types()};
+ /// @returns true if @p type is or contains a f16 type
+ bool ContainsF16Type(const core::type::Type* type) {
+ if (type->Is<core::type::F16>()) {
+ return true;
+ }
+
+ if (type->IsScalar()) {
+ return false;
+ }
+
+ if (const auto* str = type->As<core::type::Struct>()) {
+ for (auto* member : str->Members()) {
+ if (ContainsF16Type(member->Type())) {
+ return true;
+ }
+ }
+ return false;
+ }
+
+ return ContainsF16Type(type->DeepestElement());
+ }
+
void Process() {
Hashmap<core::ir::Value*, core::ir::Value*, 16> object_to_var;
for (auto* inst : ir.Instructions()) {
@@ -64,6 +87,11 @@
if (!curr_const->Type()->IsAnyOf<core::type::Array, core::type::Struct>()) {
continue;
}
+
+ if (config.disable_module_constant_f16 && ContainsF16Type(curr_const->Type())) {
+ continue;
+ }
+
// Declare a variable and copy the source object to it.
auto* var = object_to_var.GetOrAdd(source_object, [&] {
// If the source object is a constant we use a module-scope variable
@@ -78,13 +106,13 @@
} // namespace
-Result<SuccessType> ModuleConstant(core::ir::Module& ir) {
+Result<SuccessType> ModuleConstant(core::ir::Module& ir, const ModuleConstantConfig& config) {
auto result = ValidateAndDumpIfNeeded(ir, "msl.ModuleConstant", kModuleConstantCapabilities);
if (result != Success) {
return result;
}
- State{ir}.Process();
+ State{ir, config}.Process();
return Success;
}
diff --git a/src/tint/lang/msl/writer/raise/module_constant.h b/src/tint/lang/msl/writer/raise/module_constant.h
index cf8b5ab..33dfd6e 100644
--- a/src/tint/lang/msl/writer/raise/module_constant.h
+++ b/src/tint/lang/msl/writer/raise/module_constant.h
@@ -53,12 +53,18 @@
core::ir::Capability::kAllowModuleScopeLets,
};
+/// The set of polyfills that should be applied.
+struct ModuleConstantConfig {
+ // Set to true to disable module constant transform on constant data that has any f16.
+ bool disable_module_constant_f16 = false;
+};
+
/// ModuleConstant is a transform that moves all const data associated with access to a module scope
/// let. This transform is used to support 'program scope constants' in msl and thereby avoid the
/// potential for copying of large const in nested loops.
/// @param module the module to transform
/// @returns success or failure
-Result<SuccessType> ModuleConstant(core::ir::Module& module);
+Result<SuccessType> ModuleConstant(core::ir::Module& module, const ModuleConstantConfig& config);
} // namespace tint::msl::writer::raise
diff --git a/src/tint/lang/msl/writer/raise/module_constant_test.cc b/src/tint/lang/msl/writer/raise/module_constant_test.cc
index 5e56a28..2a1b783 100644
--- a/src/tint/lang/msl/writer/raise/module_constant_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_constant_test.cc
@@ -82,7 +82,8 @@
}
)";
- Run(ModuleConstant);
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
EXPECT_EQ(expect, str());
}
@@ -119,7 +120,8 @@
}
}
)";
- Run(ModuleConstant);
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
EXPECT_EQ(expect, str());
}
@@ -172,7 +174,8 @@
}
}
)";
- Run(ModuleConstant);
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
EXPECT_EQ(expect, str());
}
@@ -215,7 +218,8 @@
}
)";
- Run(ModuleConstant);
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
EXPECT_EQ(expect, str());
}
@@ -273,7 +277,109 @@
}
)";
- Run(ModuleConstant);
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstArrayStruct) {
+ auto* func = b.Function("foo", ty.u32());
+ auto* s = ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.u32()},
+ {mod.symbols.Register("b"), ty.f16()},
+ });
+
+ b.Append(func->Block(), [&] {
+ auto array_struct_type = ty.array(s, 2);
+ auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+ auto* index = b.Let(1_u);
+ auto* access = b.Access(ty.u32(), c, index, 0_u);
+ auto* r = b.Let("q", access);
+ b.Return(func, r);
+ });
+
+ auto* src = R"(
+S = struct @align(4) {
+ a:u32 @offset(0)
+ b:f16 @offset(4)
+}
+
+%foo = func():u32 {
+ $B1: {
+ %2:u32 = let 1u
+ %3:u32 = access array<S, 2>(S(1u), S(2u)), %2, 0u
+ %q:u32 = let %3
+ ret %q
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+ Run(ModuleConstant, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, EnableF16_ConstArrayStruct) {
+ auto* func = b.Function("foo", ty.u32());
+ auto* s = ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.u32()},
+ {mod.symbols.Register("b"), ty.f16()},
+ });
+
+ b.Append(func->Block(), [&] {
+ auto array_struct_type = ty.array(s, 2);
+ auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+ auto* index = b.Let(1_u);
+ auto* access = b.Access(ty.u32(), c, index, 0_u);
+ auto* r = b.Let("q", access);
+ b.Return(func, r);
+ });
+
+ auto* src = R"(
+S = struct @align(4) {
+ a:u32 @offset(0)
+ b:f16 @offset(4)
+}
+
+%foo = func():u32 {
+ $B1: {
+ %2:u32 = let 1u
+ %3:u32 = access array<S, 2>(S(1u), S(2u)), %2, 0u
+ %q:u32 = let %3
+ ret %q
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+S = struct @align(4) {
+ a:u32 @offset(0)
+ b:f16 @offset(4)
+}
+
+$B1: { # root
+ %1:array<S, 2> = let array<S, 2>(S(1u), S(2u))
+}
+
+%foo = func():u32 {
+ $B2: {
+ %3:u32 = let 1u
+ %4:u32 = access %1, %3, 0u
+ %q:u32 = let %4
+ ret %q
+ }
+}
+)";
+
+ // Note the disable f16 is false by default.
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
EXPECT_EQ(expect, str());
}
@@ -331,7 +437,49 @@
}
)";
- Run(ModuleConstant);
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstStructArray) {
+ auto* func = b.Function("foo", ty.u32());
+ auto* s =
+ ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.array(ty.u32(), 2)},
+ {mod.symbols.Register("b"), ty.array(ty.f16(), 2)},
+ });
+
+ b.Append(func->Block(), [&] {
+ auto* c = b.Splat(s, 1_u);
+ auto* index = b.Let(1_u);
+ auto* access = b.Access(ty.u32(), c, 0_u, index);
+ auto* r = b.Let("q", access);
+ b.Return(func, r);
+ });
+
+ auto* src = R"(
+S = struct @align(4) {
+ a:array<u32, 2> @offset(0)
+ b:array<f16, 2> @offset(8)
+}
+
+%foo = func():u32 {
+ $B1: {
+ %2:u32 = let 1u
+ %3:u32 = access S(1u), 0u, %2
+ %q:u32 = let %3
+ ret %q
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+ Run(ModuleConstant, cfg);
EXPECT_EQ(expect, str());
}
@@ -377,7 +525,211 @@
}
)";
- Run(ModuleConstant);
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstVecArrayStruct) {
+ auto* func = b.Function("foo", ty.f16());
+ auto* s = ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.vec2(ty.f16())},
+ {mod.symbols.Register("b"), ty.u32()},
+ });
+
+ b.Append(func->Block(), [&] {
+ auto array_struct_type = ty.array(s, 2);
+ auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+ auto* index = b.Let(1_u);
+ auto* access = b.Access(ty.f16(), c, index, 0_u, 0_u);
+ auto* r = b.Let("q", access);
+ b.Return(func, r);
+ });
+
+ auto* src = R"(
+S = struct @align(4) {
+ a:vec2<f16> @offset(0)
+ b:u32 @offset(4)
+}
+
+%foo = func():f16 {
+ $B1: {
+ %2:u32 = let 1u
+ %3:f16 = access array<S, 2>(S(1u), S(2u)), %2, 0u, 0u
+ %q:f16 = let %3
+ ret %q
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+ Run(ModuleConstant, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, EnableF16_ConstVecArrayStruct) {
+ auto* func = b.Function("foo", ty.f16());
+ auto* s = ty.Struct(mod.symbols.New("S"), {
+ {mod.symbols.Register("a"), ty.vec2(ty.f16())},
+ {mod.symbols.Register("b"), ty.f16()},
+ });
+
+ b.Append(func->Block(), [&] {
+ auto array_struct_type = ty.array(s, 2);
+ auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+ auto* index = b.Let(1_u);
+ auto* access = b.Access(ty.f16(), c, index, 0_u, 0_u);
+ auto* r = b.Let("q", access);
+ b.Return(func, r);
+ });
+
+ auto* src = R"(
+S = struct @align(4) {
+ a:vec2<f16> @offset(0)
+ b:f16 @offset(4)
+}
+
+%foo = func():f16 {
+ $B1: {
+ %2:u32 = let 1u
+ %3:f16 = access array<S, 2>(S(1u), S(2u)), %2, 0u, 0u
+ %q:f16 = let %3
+ ret %q
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+S = struct @align(4) {
+ a:vec2<f16> @offset(0)
+ b:f16 @offset(4)
+}
+
+$B1: { # root
+ %1:array<S, 2> = let array<S, 2>(S(1u), S(2u))
+}
+
+%foo = func():f16 {
+ $B2: {
+ %3:u32 = let 1u
+ %4:f16 = access %1, %3, 0u, 0u
+ %q:f16 = let %4
+ ret %q
+ }
+}
+)";
+
+ // Note the disable f16 is false by default.
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstMatStructArray) {
+ auto* func = b.Function("foo", ty.f16());
+ auto* s = ty.Struct(mod.symbols.New("S"),
+ {
+ {mod.symbols.Register("a"), ty.array(ty.u32(), 2)},
+ {mod.symbols.Register("b"), ty.array(ty.mat3x3(ty.f16()), 3)},
+ });
+
+ b.Append(func->Block(), [&] {
+ auto* c = b.Splat(s, 1_u);
+ auto* index = b.Let(1_u);
+ auto* access = b.Access(ty.f16(), c, 1_u, index, 0_u, 0_u);
+ auto* r = b.Let("q", access);
+ b.Return(func, r);
+ });
+
+ auto* src = R"(
+S = struct @align(8) {
+ a:array<u32, 2> @offset(0)
+ b:array<mat3x3<f16>, 3> @offset(8)
+}
+
+%foo = func():f16 {
+ $B1: {
+ %2:u32 = let 1u
+ %3:f16 = access S(1u), 1u, %2, 0u, 0u
+ %q:f16 = let %3
+ ret %q
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+ Run(ModuleConstant, cfg);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, EnabledF16_ConstMatStructArray) {
+ auto* func = b.Function("foo", ty.f16());
+ auto* s = ty.Struct(mod.symbols.New("S"),
+ {
+ {mod.symbols.Register("a"), ty.array(ty.u32(), 2)},
+ {mod.symbols.Register("b"), ty.array(ty.mat3x3(ty.f16()), 3)},
+ });
+
+ b.Append(func->Block(), [&] {
+ auto* c = b.Splat(s, 1_u);
+ auto* index = b.Let(1_u);
+ auto* access = b.Access(ty.f16(), c, 1_u, index, 0_u, 0_u);
+ auto* r = b.Let("q", access);
+ b.Return(func, r);
+ });
+
+ auto* src = R"(
+S = struct @align(8) {
+ a:array<u32, 2> @offset(0)
+ b:array<mat3x3<f16>, 3> @offset(8)
+}
+
+%foo = func():f16 {
+ $B1: {
+ %2:u32 = let 1u
+ %3:f16 = access S(1u), 1u, %2, 0u, 0u
+ %q:f16 = let %3
+ ret %q
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+S = struct @align(8) {
+ a:array<u32, 2> @offset(0)
+ b:array<mat3x3<f16>, 3> @offset(8)
+}
+
+$B1: { # root
+ %1:S = let S(1u)
+}
+
+%foo = func():f16 {
+ $B2: {
+ %3:u32 = let 1u
+ %4:f16 = access %1, 1u, %3, 0u, 0u
+ %q:f16 = let %4
+ ret %q
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+ // Note the disable f16 is false by default.
+ ModuleConstantConfig cfg;
+ Run(ModuleConstant, cfg);
EXPECT_EQ(expect, str());
}
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index 2a59e32..aead774 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -168,9 +168,8 @@
};
RUN_TRANSFORM(core::ir::transform::BuiltinScalarize, module, scalarize_config);
- if (options.enable_module_constant) {
- RUN_TRANSFORM(raise::ModuleConstant, module);
- }
+ raise::ModuleConstantConfig module_const_config{options.disable_module_constant_f16};
+ RUN_TRANSFORM(raise::ModuleConstant, module, module_const_config);
// These transforms need to be run last as various transforms introduce terminator arguments,
// naming conflicts, and expressions that need to be explicitly not inlined.
diff --git a/test/tint/bug/tint/1641.wgsl.expected.msl b/test/tint/bug/tint/1641.wgsl.expected.msl
index 1525b08..d0ac73b 100644
--- a/test/tint/bug/tint/1641.wgsl.expected.msl
+++ b/test/tint/bug/tint/1641.wgsl.expected.msl
@@ -20,13 +20,14 @@
struct main_outputs {
float4 tint_symbol [[position]];
};
+const constant tint_array<Normals, 1> v = tint_array<Normals, 1>{Normals{.f=float3(0.0f, 0.0f, 1.0f)}};
float4 main_inner() {
int const zero = 0;
- return float4(tint_array<Normals, 1>{Normals{.f=float3(0.0f, 0.0f, 1.0f)}}[min(uint(zero), 0u)].f, 1.0f);
+ return float4(v[min(uint(zero), 0u)].f, 1.0f);
}
-vertex main_outputs v() {
+vertex main_outputs v_1() {
main_outputs tint_wrapper_result = {};
tint_wrapper_result.tint_symbol = main_inner();
return tint_wrapper_result;
diff --git a/test/tint/bug/tint/1653.wgsl.expected.msl b/test/tint/bug/tint/1653.wgsl.expected.msl
index 3c39813..317c598 100644
--- a/test/tint/bug/tint/1653.wgsl.expected.msl
+++ b/test/tint/bug/tint/1653.wgsl.expected.msl
@@ -16,9 +16,10 @@
struct vs_main_outputs {
float4 tint_symbol [[position]];
};
+const constant tint_array<float4, 3> v = tint_array<float4, 3>{float4(0.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(1.0f, 1.0f, 0.0f, 1.0f)};
float4 vs_main_inner(uint in_vertex_index) {
- return tint_array<float4, 3>{float4(0.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(1.0f, 1.0f, 0.0f, 1.0f)}[min(in_vertex_index, 2u)];
+ return v[min(in_vertex_index, 2u)];
}
vertex vs_main_outputs vs_main(uint in_vertex_index [[vertex_id]]) {
diff --git a/test/tint/bug/tint/2237.wgsl.expected.msl b/test/tint/bug/tint/2237.wgsl.expected.msl
index f9846ac..3301641 100644
--- a/test/tint/bug/tint/2237.wgsl.expected.msl
+++ b/test/tint/bug/tint/2237.wgsl.expected.msl
@@ -1,10 +1,6 @@
#include <metal_stdlib>
using namespace metal;
-struct tint_module_vars_struct {
- device uint* tint_member;
-};
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -17,12 +13,17 @@
T elements[N];
};
+struct tint_module_vars_struct {
+ device uint* tint_member;
+};
+const constant tint_array<uint, 4> v_1 = tint_array<uint, 4>{0u, 1u, 2u, 4u};
+
uint foo(tint_module_vars_struct tint_module_vars) {
- return tint_array<uint, 4>{0u, 1u, 2u, 4u}[min((*tint_module_vars.tint_member), 3u)];
+ return v_1[min((*tint_module_vars.tint_member), 3u)];
}
-kernel void v_1(device uint* v_2 [[buffer(0)]]) {
- tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_member=v_2};
- uint const v = tint_array<uint, 4>{0u, 1u, 2u, 4u}[min((*tint_module_vars.tint_member), 3u)];
+kernel void v_2(device uint* v_3 [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_member=v_3};
+ uint const v = v_1[min((*tint_module_vars.tint_member), 3u)];
(*tint_module_vars.tint_member) = (v + foo(tint_module_vars));
}
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.msl b/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
index a798803..223f418 100644
--- a/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
@@ -21,15 +21,16 @@
struct tint_array_lengths_struct {
uint tint_array_length_0_0;
};
+const constant tint_array<half3x2, 2> v = tint_array<half3x2, 2>{half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h)), half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h))};
uint tint_f16_to_u32(half value) {
return uint(clamp(value, 0.0h, 65504.0h));
}
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint q = 0u;
- uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
- device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
- (*v_2) = tint_f16_to_u32(tint_array<half3x2, 2>{half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h)), half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h))}[min(q, 1u)][0u].x);
+ uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+ device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+ (*v_3) = tint_f16_to_u32(v[min(q, 1u)][0u].x);
}
diff --git a/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl b/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl
index cf141d3..8c998ac 100644
--- a/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl
+++ b/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl
@@ -21,11 +21,12 @@
struct tint_array_lengths_struct {
uint tint_array_length_0_0;
};
+const constant tint_array<tint_array<tint_array<uint, 2>, 2>, 2> v = tint_array<tint_array<tint_array<uint, 2>, 2>, 2>{tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{0u, 1u}, tint_array<uint, 2>{2u, 3u}}, tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{4u, 5u}, tint_array<uint, 2>{6u, 7u}}};
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint q = 0u;
- uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
- device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
- (*v_2) = tint_array<tint_array<tint_array<uint, 2>, 2>, 2>{tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{0u, 1u}, tint_array<uint, 2>{2u, 3u}}, tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{4u, 5u}, tint_array<uint, 2>{6u, 7u}}}[min(q, 1u)][min(q, 1u)][min(q, 1u)];
+ uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+ device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+ (*v_3) = v[min(q, 1u)][min(q, 1u)][min(q, 1u)];
}
diff --git a/test/tint/const/array/array_of_struct.wgsl.expected.msl b/test/tint/const/array/array_of_struct.wgsl.expected.msl
index d81bcdd..e8677fb 100644
--- a/test/tint/const/array/array_of_struct.wgsl.expected.msl
+++ b/test/tint/const/array/array_of_struct.wgsl.expected.msl
@@ -13,6 +13,11 @@
T elements[N];
};
+struct A {
+ uint2 b;
+ uint c;
+};
+
struct tint_module_vars_struct {
device tint_array<uint, 1>* s;
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
@@ -21,16 +26,12 @@
struct tint_array_lengths_struct {
uint tint_array_length_0_0;
};
+const constant tint_array<A, 2> v = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}};
-struct A {
- uint2 b;
- uint c;
-};
-
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint q = 0u;
- uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
- device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
- (*v_2) = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}}[min(q, 1u)].b.x;
+ uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+ device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+ (*v_3) = v[min(q, 1u)].b.x;
}
diff --git a/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl b/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl
index d81bcdd..e8677fb 100644
--- a/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl
+++ b/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl
@@ -13,6 +13,11 @@
T elements[N];
};
+struct A {
+ uint2 b;
+ uint c;
+};
+
struct tint_module_vars_struct {
device tint_array<uint, 1>* s;
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
@@ -21,16 +26,12 @@
struct tint_array_lengths_struct {
uint tint_array_length_0_0;
};
+const constant tint_array<A, 2> v = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}};
-struct A {
- uint2 b;
- uint c;
-};
-
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint q = 0u;
- uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
- device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
- (*v_2) = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}}[min(q, 1u)].b.x;
+ uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+ device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+ (*v_3) = v[min(q, 1u)].b.x;
}
diff --git a/test/tint/const/array/large_array.wgsl.expected.msl b/test/tint/const/array/large_array.wgsl.expected.msl
index ba6f9f3..c61cb31 100644
--- a/test/tint/const/array/large_array.wgsl.expected.msl
+++ b/test/tint/const/array/large_array.wgsl.expected.msl
@@ -21,11 +21,12 @@
struct tint_array_lengths_struct {
uint tint_array_length_0_0;
};
+const constant tint_array<uint, 2310> v = tint_array<uint, 2310>{0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u};
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint q = 0u;
- uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
- device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
- (*v_2) = tint_array<uint, 2310>{0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u}[min(q, 2309u)];
+ uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+ device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+ (*v_3) = v[min(q, 2309u)];
}
diff --git a/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl b/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl
index 6ba612c..407d525 100644
--- a/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl
+++ b/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl
@@ -21,11 +21,12 @@
struct tint_array_lengths_struct {
uint tint_array_length_0_0;
};
+const constant tint_array<uint2, 2> v = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)};
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint q = 0u;
- uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
- device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
- (*v_2) = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)}[min(q, 1u)].x;
+ uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+ device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+ (*v_3) = v[min(q, 1u)].x;
}
diff --git a/test/tint/const/struct/struct_of_array.wgsl.expected.msl b/test/tint/const/struct/struct_of_array.wgsl.expected.msl
index 6ba612c..407d525 100644
--- a/test/tint/const/struct/struct_of_array.wgsl.expected.msl
+++ b/test/tint/const/struct/struct_of_array.wgsl.expected.msl
@@ -21,11 +21,12 @@
struct tint_array_lengths_struct {
uint tint_array_length_0_0;
};
+const constant tint_array<uint2, 2> v = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)};
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
uint q = 0u;
- uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
- device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
- (*v_2) = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)}[min(q, 1u)].x;
+ uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+ device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+ (*v_3) = v[min(q, 1u)].x;
}
diff --git a/test/tint/samples/triangle.wgsl.expected.msl b/test/tint/samples/triangle.wgsl.expected.msl
index b8152ed..469449e 100644
--- a/test/tint/samples/triangle.wgsl.expected.msl
+++ b/test/tint/samples/triangle.wgsl.expected.msl
@@ -19,9 +19,10 @@
struct vtx_main_outputs {
float4 tint_symbol [[position]];
};
+const constant tint_array<float2, 3> v = tint_array<float2, 3>{float2(0.0f, 0.5f), float2(-0.5f), float2(0.5f, -0.5f)};
float4 vtx_main_inner(uint VertexIndex) {
- return float4(tint_array<float2, 3>{float2(0.0f, 0.5f), float2(-0.5f), float2(0.5f, -0.5f)}[min(VertexIndex, 2u)], 0.0f, 1.0f);
+ return float4(v[min(VertexIndex, 2u)], 0.0f, 1.0f);
}
vertex vtx_main_outputs vtx_main(uint VertexIndex [[vertex_id]]) {
diff --git a/tools/src/cmd/run-cts/node/cmd.go b/tools/src/cmd/run-cts/node/cmd.go
index f7fb281..960c6da 100644
--- a/tools/src/cmd/run-cts/node/cmd.go
+++ b/tools/src/cmd/run-cts/node/cmd.go
@@ -87,7 +87,7 @@
}
func (c *cmd) RegisterFlags(ctx context.Context, cfg common.Config) ([]string, error) {
- unrollConstEvalLoopsDefault := runtime.GOOS != "windows"
+ unrollConstEvalLoopsDefault := false
backendDefault := "default"
if vkIcdFilenames := cfg.OsWrapper.Getenv("VK_ICD_FILENAMES"); vkIcdFilenames != "" {
diff --git a/webgpu-cts/compat-expectations.txt b/webgpu-cts/compat-expectations.txt
index 463b443..9596069 100644
--- a/webgpu-cts/compat-expectations.txt
+++ b/webgpu-cts/compat-expectations.txt
@@ -610,6 +610,14 @@
crbug.com/373670504 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="uniform";type="bool";length=5;infer_type=false [ Failure ]
crbug.com/373670504 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="uniform";type="bool";length=5;infer_type=true [ Failure ]
+# Const bool arrays on Pixel 6
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=10;infer_type=false [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=10;infer_type=true [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=5;infer_type=false [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=5;infer_type=true [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,zero_value:array:type="bool";length=10 [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,zero_value:array:type="bool";length=5 [ Failure ]
+
# set_index_buffer_without_changing_buffer failing on Pixel 6
crbug.com/373664074 [ android-t arm ] webgpu:api,operation,command_buffer,render,state_tracking:set_index_buffer_without_changing_buffer: [ Failure ]
diff --git a/webgpu-cts/test_runner.js b/webgpu-cts/test_runner.js
index 5992de8..aa977ce 100644
--- a/webgpu-cts/test_runner.js
+++ b/webgpu-cts/test_runner.js
@@ -167,14 +167,9 @@
// simultaneously on a 32-bit system easily runs out of memory).
globalTestConfig.maxSubcasesInFlight = 100;
-// FXC is very slow to compile unrolled const-eval loops, where the metal shader
-// compiler (Intel GPU) is very slow to compile rolled loops. Intel drivers for
-// linux may also suffer the same performance issues, so unroll const-eval loops
-// if we're not running on Windows.
-const isWindows = navigator.userAgent.includes("Windows");
-if (!isWindows) {
- globalTestConfig.unrollConstEvalLoops = true;
-}
+
+
+
let lastOptionsKey, testWorker;
