writer/spirv: Fix [[override]] for zero-init defaults
These expressions were not emitted as spec-ops, failing new CTS tests.
Change-Id: I56e8f56a22de2b8dac9a8bd7a2d694d8d81dca35
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67480
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 5481a71..1ca479f 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -577,7 +577,7 @@
return GenerateCallExpression(c);
}
if (auto* c = expr->As<ast::ConstructorExpression>()) {
- return GenerateConstructorExpression(nullptr, c, false);
+ return GenerateConstructorExpression(nullptr, c);
}
if (auto* i = expr->As<ast::IdentifierExpression>()) {
return GenerateIdentifierExpression(i);
@@ -763,7 +763,7 @@
}
init_id = GenerateConstructorExpression(
- var, var->constructor->As<ast::ConstructorExpression>(), true);
+ var, var->constructor->As<ast::ConstructorExpression>());
if (init_id == 0) {
return false;
}
@@ -1268,13 +1268,12 @@
uint32_t Builder::GenerateConstructorExpression(
const ast::Variable* var,
- const ast::ConstructorExpression* expr,
- bool is_global_init) {
+ const ast::ConstructorExpression* expr) {
if (auto* scalar = expr->As<ast::ScalarConstructorExpression>()) {
return GenerateLiteralIfNeeded(var, scalar->literal);
}
if (auto* type = expr->As<ast::TypeConstructorExpression>()) {
- return GenerateTypeConstructorExpression(type, is_global_init);
+ return GenerateTypeConstructorExpression(var, type);
}
error_ = "unknown constructor expression";
@@ -1338,14 +1337,35 @@
}
uint32_t Builder::GenerateTypeConstructorExpression(
- const ast::TypeConstructorExpression* init,
- bool is_global_init) {
+ const ast::Variable* var,
+ const ast::TypeConstructorExpression* init) {
+ auto* global_var = builder_.Sem().Get<sem::GlobalVariable>(var);
+
auto& values = init->values;
auto* result_type = TypeOf(init);
// Generate the zero initializer if there are no values provided.
if (values.empty()) {
+ if (global_var && global_var->IsPipelineConstant()) {
+ auto constant_id = global_var->ConstantId();
+ if (result_type->Is<sem::I32>()) {
+ return GenerateConstantIfNeeded(
+ ScalarConstant::I32(0).AsSpecOp(constant_id));
+ }
+ if (result_type->Is<sem::U32>()) {
+ return GenerateConstantIfNeeded(
+ ScalarConstant::U32(0).AsSpecOp(constant_id));
+ }
+ if (result_type->Is<sem::F32>()) {
+ return GenerateConstantIfNeeded(
+ ScalarConstant::F32(0).AsSpecOp(constant_id));
+ }
+ if (result_type->Is<sem::Bool>()) {
+ return GenerateConstantIfNeeded(
+ ScalarConstant::Bool(false).AsSpecOp(constant_id));
+ }
+ }
return GenerateConstantNullIfNeeded(result_type->UnwrapRef());
}
@@ -1353,7 +1373,7 @@
out << "__const_" << init->type->FriendlyName(builder_.Symbols()) << "_";
result_type = result_type->UnwrapRef();
- bool constructor_is_const = is_constructor_const(init, is_global_init);
+ bool constructor_is_const = is_constructor_const(init, global_var);
if (has_error()) {
return 0;
}
@@ -1372,8 +1392,7 @@
}
if (can_cast_or_copy) {
- return GenerateCastOrCopyOrPassthrough(result_type, values[0],
- is_global_init);
+ return GenerateCastOrCopyOrPassthrough(result_type, values[0], global_var);
}
auto type_id = GenerateTypeIfNeeded(result_type);
@@ -1392,8 +1411,8 @@
for (auto* e : values) {
uint32_t id = 0;
if (constructor_is_const) {
- id = GenerateConstructorExpression(
- nullptr, e->As<ast::ConstructorExpression>(), is_global_init);
+ id = GenerateConstructorExpression(nullptr,
+ e->As<ast::ConstructorExpression>());
} else {
id = GenerateExpression(e);
id = GenerateLoadIfNeeded(TypeOf(e), id);
@@ -1417,8 +1436,7 @@
// Both scalars, but not the same type so we need to generate a conversion
// of the value.
if (value_type->is_scalar() && result_type->is_scalar()) {
- id = GenerateCastOrCopyOrPassthrough(result_type, values[0],
- is_global_init);
+ id = GenerateCastOrCopyOrPassthrough(result_type, values[0], global_var);
out << "_" << id;
ops.push_back(Operand::Int(id));
continue;
@@ -1445,7 +1463,7 @@
auto extract = result_op();
auto extract_id = extract.to_i();
- if (!is_global_init) {
+ if (!global_var) {
// A non-global initializer. Case 2.
if (!push_function_inst(spv::Op::OpCompositeExtract,
{Operand::Int(value_type_id), extract,
@@ -2816,7 +2834,7 @@
if (auto* array_index = arg(Usage::kArrayIndex)) {
// Array index needs to be appended to the coordinates.
auto* packed = AppendVector(&builder_, arg(Usage::kCoords), array_index);
- auto param = GenerateTypeConstructorExpression(packed, false);
+ auto param = GenerateTypeConstructorExpression(nullptr, packed);
if (param == 0) {
return false;
}
diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h
index 0aad2dd..77d878f 100644
--- a/src/writer/spirv/builder.h
+++ b/src/writer/spirv/builder.h
@@ -338,18 +338,17 @@
/// Generates a constructor expression
/// @param var the variable generated for, nullptr if no variable associated.
/// @param expr the expression to generate
- /// @param is_global_init set true if this is a global variable constructor
/// @returns the ID of the expression or 0 on failure.
- uint32_t GenerateConstructorExpression(const ast::Variable* var,
- const ast::ConstructorExpression* expr,
- bool is_global_init);
+ uint32_t GenerateConstructorExpression(
+ const ast::Variable* var,
+ const ast::ConstructorExpression* expr);
/// Generates a type constructor expression
+ /// @param var the variable generated for, nullptr if no variable associated.
/// @param init the expression to generate
- /// @param is_global_init set true if this is a global variable constructor
/// @returns the ID of the expression or 0 on failure.
uint32_t GenerateTypeConstructorExpression(
- const ast::TypeConstructorExpression* init,
- bool is_global_init);
+ const ast::Variable* var,
+ const ast::TypeConstructorExpression* init);
/// Generates a literal constant if needed
/// @param var the variable generated for, nullptr if no variable associated.
/// @param lit the literal to generate
diff --git a/src/writer/spirv/builder_constructor_expression_test.cc b/src/writer/spirv/builder_constructor_expression_test.cc
index 0f65afe..3488b2b 100644
--- a/src/writer/spirv/builder_constructor_expression_test.cc
+++ b/src/writer/spirv/builder_constructor_expression_test.cc
@@ -24,11 +24,11 @@
TEST_F(SpvBuilderConstructorTest, Const) {
auto* c = Expr(42.2f);
- WrapInFunction(c);
+ auto* g = Global("g", ty.f32(), c, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, c, true), 2u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, c), 2u);
ASSERT_FALSE(b.has_error()) << b.error();
EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32
@@ -55,7 +55,7 @@
spirv::Builder& b = Build();
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, t, true), 5u);
+ EXPECT_EQ(b.GenerateConstructorExpression(nullptr, t), 5u);
ASSERT_FALSE(b.has_error()) << b.error();
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
@@ -170,11 +170,11 @@
Expr(3.0f));
auto* t = vec2<f32>(1.0f, rel);
- WrapInFunction(t);
+ auto* g = Global("g", ty.vec2<f32>(), t, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, t, true), 0u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, t), 0u);
EXPECT_TRUE(b.has_error());
EXPECT_EQ(b.error(), R"(constructor must be a constant expression)");
}
@@ -670,12 +670,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec2_With_F32) {
auto* cast = vec2<f32>(2.0f);
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec2<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 4u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 4u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 2
@@ -740,12 +740,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_F32) {
auto* cast = vec3<f32>(2.0f);
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec3<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 4u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 4u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 3
@@ -756,12 +756,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_F32_Vec2) {
auto* cast = vec3<f32>(2.0f, vec2<f32>(2.0f, 2.0f));
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec3<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 3
@@ -779,12 +779,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec3_With_Vec2_F32) {
auto* cast = vec3<f32>(vec2<f32>(2.0f, 2.0f), 2.0f);
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec3<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 3
@@ -802,12 +802,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32) {
auto* cast = vec4<f32>(2.0f);
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec4<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 4u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 4u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 4
@@ -818,12 +818,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_F32_Vec2) {
auto* cast = vec4<f32>(2.0f, 2.0f, vec2<f32>(2.0f, 2.0f));
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec4<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 4
@@ -841,12 +841,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_Vec2_F32) {
auto* cast = vec4<f32>(2.0f, vec2<f32>(2.0f, 2.0f), 2.0f);
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec4<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 4
@@ -864,12 +864,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec2_F32_F32) {
auto* cast = vec4<f32>(vec2<f32>(2.0f, 2.0f), 2.0f, 2.0f);
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec4<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 11u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 11u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 4
@@ -887,12 +887,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec2_Vec2) {
auto* cast = vec4<f32>(vec2<f32>(2.0f, 2.0f), vec2<f32>(2.0f, 2.0f));
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec4<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 13u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 13u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 4
@@ -912,12 +912,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_F32_Vec3) {
auto* cast = vec4<f32>(2.0f, vec3<f32>(2.0f, 2.0f, 2.0f));
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec4<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 13u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 13u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 4
@@ -937,12 +937,12 @@
TEST_F(SpvBuilderConstructorTest, Type_ModuleScope_Vec4_With_Vec3_F32) {
auto* cast = vec4<f32>(vec3<f32>(2.0f, 2.0f, 2.0f), 2.0f);
- WrapInFunction(cast);
+ auto* g = Global("g", ty.vec4<f32>(), cast, ast::StorageClass::kPrivate);
spirv::Builder& b = Build();
b.push_function(Function{});
- EXPECT_EQ(b.GenerateConstructorExpression(nullptr, cast, true), 13u);
+ EXPECT_EQ(b.GenerateConstructorExpression(g, cast), 13u);
EXPECT_EQ(DumpInstructions(b.types()), R"(%2 = OpTypeFloat 32
%1 = OpTypeVector %2 4
@@ -1275,7 +1275,7 @@
}
TEST_F(SpvBuilderConstructorTest, Type_ZeroInit_F32) {
- auto* t = Construct(ty.f32());
+ auto* t = Construct<f32>();
WrapInFunction(t);
@@ -1326,7 +1326,7 @@
}
TEST_F(SpvBuilderConstructorTest, Type_ZeroInit_Bool) {
- auto* t = Construct(ty.bool_());
+ auto* t = Construct<bool>();
WrapInFunction(t);
diff --git a/src/writer/spirv/builder_global_variable_test.cc b/src/writer/spirv/builder_global_variable_test.cc
index f464bee..4160617 100644
--- a/src/writer/spirv/builder_global_variable_test.cc
+++ b/src/writer/spirv/builder_global_variable_test.cc
@@ -167,6 +167,24 @@
)");
}
+TEST_F(BuilderTest, GlobalVar_Override_Bool_ZeroValue) {
+ auto* v = GlobalConst("var", ty.bool_(), Construct<bool>(),
+ ast::DecorationList{
+ create<ast::OverrideDecoration>(1200),
+ });
+
+ spirv::Builder& b = Build();
+
+ EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var"
+)");
+ EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 1200
+)");
+ EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeBool
+%2 = OpSpecConstantFalse %1
+)");
+}
+
TEST_F(BuilderTest, GlobalVar_Override_Bool_NoConstructor) {
auto* v = GlobalConst("var", ty.bool_(), nullptr,
ast::DecorationList{
@@ -203,6 +221,24 @@
)");
}
+TEST_F(BuilderTest, GlobalVar_Override_Scalar_ZeroValue) {
+ auto* v = GlobalConst("var", ty.f32(), Construct<f32>(),
+ ast::DecorationList{
+ create<ast::OverrideDecoration>(0),
+ });
+
+ spirv::Builder& b = Build();
+
+ EXPECT_TRUE(b.GenerateGlobalVariable(v)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %2 "var"
+)");
+ EXPECT_EQ(DumpInstructions(b.annots()), R"(OpDecorate %2 SpecId 0
+)");
+ EXPECT_EQ(DumpInstructions(b.types()), R"(%1 = OpTypeFloat 32
+%2 = OpSpecConstant %1 0
+)");
+}
+
TEST_F(BuilderTest, GlobalVar_Override_Scalar_F32_NoConstructor) {
auto* v = GlobalConst("var", ty.f32(), nullptr,
ast::DecorationList{
diff --git a/src/writer/spirv/scalar_constant.h b/src/writer/spirv/scalar_constant.h
index 2911600..d8deece 100644
--- a/src/writer/spirv/scalar_constant.h
+++ b/src/writer/spirv/scalar_constant.h
@@ -83,6 +83,15 @@
return c;
}
+ /// @param value the value of the constant
+ /// @returns a new ScalarConstant with the provided value and kind Kind::kBool
+ static inline ScalarConstant Bool(bool value) {
+ ScalarConstant c;
+ c.value.b = value;
+ c.kind = Kind::kBool;
+ return c;
+ }
+
/// Equality operator
/// @param rhs the ScalarConstant to compare against
/// @returns true if this ScalarConstant is equal to `rhs`
@@ -98,6 +107,16 @@
return !(*this == rhs);
}
+ /// @returns this ScalarConstant as a specialization op with the given
+ /// specialization constant identifier
+ /// @param id the constant identifier
+ ScalarConstant AsSpecOp(uint32_t id) const {
+ auto ret = *this;
+ ret.is_spec_op = true;
+ ret.constant_id = id;
+ return ret;
+ }
+
/// The constant value
Value value;
/// The constant value kind
diff --git a/test/var/override/named/no_init/bool.wgsl b/test/var/override/named/no_init/bool.wgsl
new file mode 100644
index 0000000..4b58e68
--- /dev/null
+++ b/test/var/override/named/no_init/bool.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : bool;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/no_init/bool.wgsl.expected.hlsl b/test/var/override/named/no_init/bool.wgsl.expected.hlsl
new file mode 100644
index 0000000..977045c
--- /dev/null
+++ b/test/var/override/named/no_init/bool.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#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() {
+ return;
+}
+/tmp/tint_N41FiO:2:2: error: spec constant required for constant id 0
+#error spec constant required for constant id 0
+ ^
+
+
diff --git a/test/var/override/named/no_init/bool.wgsl.expected.msl b/test/var/override/named/no_init/bool.wgsl.expected.msl
new file mode 100644
index 0000000..95ce385
--- /dev/null
+++ b/test/var/override/named/no_init/bool.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant bool o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/no_init/bool.wgsl.expected.spvasm b/test/var/override/named/no_init/bool.wgsl.expected.spvasm
new file mode 100644
index 0000000..ff12e00
--- /dev/null
+++ b/test/var/override/named/no_init/bool.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/no_init/bool.wgsl.expected.wgsl b/test/var/override/named/no_init/bool.wgsl.expected.wgsl
new file mode 100644
index 0000000..8fcab9b
--- /dev/null
+++ b/test/var/override/named/no_init/bool.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : bool;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/no_init/f32.wgsl b/test/var/override/named/no_init/f32.wgsl
new file mode 100644
index 0000000..c09196d
--- /dev/null
+++ b/test/var/override/named/no_init/f32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/no_init/f32.wgsl.expected.hlsl b/test/var/override/named/no_init/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2ae175b
--- /dev/null
+++ b/test/var/override/named/no_init/f32.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#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() {
+ return;
+}
+/tmp/tint_gYmCvK:2:2: error: spec constant required for constant id 0
+#error spec constant required for constant id 0
+ ^
+
+
diff --git a/test/var/override/named/no_init/f32.wgsl.expected.msl b/test/var/override/named/no_init/f32.wgsl.expected.msl
new file mode 100644
index 0000000..2ed922c
--- /dev/null
+++ b/test/var/override/named/no_init/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant float o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/no_init/f32.wgsl.expected.spvasm b/test/var/override/named/no_init/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..7ae9060
--- /dev/null
+++ b/test/var/override/named/no_init/f32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/no_init/f32.wgsl.expected.wgsl b/test/var/override/named/no_init/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..5e60be6
--- /dev/null
+++ b/test/var/override/named/no_init/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/no_init/i32.wgsl b/test/var/override/named/no_init/i32.wgsl
new file mode 100644
index 0000000..790c56d
--- /dev/null
+++ b/test/var/override/named/no_init/i32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/no_init/i32.wgsl.expected.hlsl b/test/var/override/named/no_init/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..7685972
--- /dev/null
+++ b/test/var/override/named/no_init/i32.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#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() {
+ return;
+}
+/tmp/tint_kYjxDo:2:2: error: spec constant required for constant id 0
+#error spec constant required for constant id 0
+ ^
+
+
diff --git a/test/var/override/named/no_init/i32.wgsl.expected.msl b/test/var/override/named/no_init/i32.wgsl.expected.msl
new file mode 100644
index 0000000..b5b1fae
--- /dev/null
+++ b/test/var/override/named/no_init/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant int o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/no_init/i32.wgsl.expected.spvasm b/test/var/override/named/no_init/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..18c8840
--- /dev/null
+++ b/test/var/override/named/no_init/i32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/no_init/i32.wgsl.expected.wgsl b/test/var/override/named/no_init/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..56a81dd
--- /dev/null
+++ b/test/var/override/named/no_init/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/no_init/u32.wgsl b/test/var/override/named/no_init/u32.wgsl
new file mode 100644
index 0000000..b4826d8
--- /dev/null
+++ b/test/var/override/named/no_init/u32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/no_init/u32.wgsl.expected.hlsl b/test/var/override/named/no_init/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..4d5ed7a
--- /dev/null
+++ b/test/var/override/named/no_init/u32.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#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() {
+ return;
+}
+/tmp/tint_7FSY6m:2:2: error: spec constant required for constant id 0
+#error spec constant required for constant id 0
+ ^
+
+
diff --git a/test/var/override/named/no_init/u32.wgsl.expected.msl b/test/var/override/named/no_init/u32.wgsl.expected.msl
new file mode 100644
index 0000000..918348c
--- /dev/null
+++ b/test/var/override/named/no_init/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant uint o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/no_init/u32.wgsl.expected.spvasm b/test/var/override/named/no_init/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..e5486b0
--- /dev/null
+++ b/test/var/override/named/no_init/u32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/no_init/u32.wgsl.expected.wgsl b/test/var/override/named/no_init/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..ad44144
--- /dev/null
+++ b/test/var/override/named/no_init/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/bool.wgsl b/test/var/override/named/val_init/bool.wgsl
new file mode 100644
index 0000000..0cb9566
--- /dev/null
+++ b/test/var/override/named/val_init/bool.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : bool = true;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/bool.wgsl.expected.hlsl b/test/var/override/named/val_init/bool.wgsl.expected.hlsl
new file mode 100644
index 0000000..e4ba790
--- /dev/null
+++ b/test/var/override/named/val_init/bool.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/val_init/bool.wgsl.expected.msl b/test/var/override/named/val_init/bool.wgsl.expected.msl
new file mode 100644
index 0000000..95ce385
--- /dev/null
+++ b/test/var/override/named/val_init/bool.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant bool o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/val_init/bool.wgsl.expected.spvasm b/test/var/override/named/val_init/bool.wgsl.expected.spvasm
new file mode 100644
index 0000000..277782e
--- /dev/null
+++ b/test/var/override/named/val_init/bool.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/val_init/bool.wgsl.expected.wgsl b/test/var/override/named/val_init/bool.wgsl.expected.wgsl
new file mode 100644
index 0000000..8e9edf5
--- /dev/null
+++ b/test/var/override/named/val_init/bool.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : bool = true;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/f32.wgsl b/test/var/override/named/val_init/f32.wgsl
new file mode 100644
index 0000000..904a8ad
--- /dev/null
+++ b/test/var/override/named/val_init/f32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : f32 = 1.0;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/f32.wgsl.expected.hlsl b/test/var/override/named/val_init/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..9d0794c
--- /dev/null
+++ b/test/var/override/named/val_init/f32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/val_init/f32.wgsl.expected.msl b/test/var/override/named/val_init/f32.wgsl.expected.msl
new file mode 100644
index 0000000..2ed922c
--- /dev/null
+++ b/test/var/override/named/val_init/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant float o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/val_init/f32.wgsl.expected.spvasm b/test/var/override/named/val_init/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..d5ce59a
--- /dev/null
+++ b/test/var/override/named/val_init/f32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/val_init/f32.wgsl.expected.wgsl b/test/var/override/named/val_init/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..32660b8
--- /dev/null
+++ b/test/var/override/named/val_init/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : f32 = 1.0;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/i32.wgsl b/test/var/override/named/val_init/i32.wgsl
new file mode 100644
index 0000000..c94ee67
--- /dev/null
+++ b/test/var/override/named/val_init/i32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : i32 = 1;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/i32.wgsl.expected.hlsl b/test/var/override/named/val_init/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..fa2861c
--- /dev/null
+++ b/test/var/override/named/val_init/i32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/val_init/i32.wgsl.expected.msl b/test/var/override/named/val_init/i32.wgsl.expected.msl
new file mode 100644
index 0000000..b5b1fae
--- /dev/null
+++ b/test/var/override/named/val_init/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant int o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/val_init/i32.wgsl.expected.spvasm b/test/var/override/named/val_init/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..7d30d25
--- /dev/null
+++ b/test/var/override/named/val_init/i32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/val_init/i32.wgsl.expected.wgsl b/test/var/override/named/val_init/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..b897fd53
--- /dev/null
+++ b/test/var/override/named/val_init/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : i32 = 1;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/u32.wgsl b/test/var/override/named/val_init/u32.wgsl
new file mode 100644
index 0000000..8c8df63
--- /dev/null
+++ b/test/var/override/named/val_init/u32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : u32 = 1u;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/val_init/u32.wgsl.expected.hlsl b/test/var/override/named/val_init/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..9529784
--- /dev/null
+++ b/test/var/override/named/val_init/u32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/val_init/u32.wgsl.expected.msl b/test/var/override/named/val_init/u32.wgsl.expected.msl
new file mode 100644
index 0000000..918348c
--- /dev/null
+++ b/test/var/override/named/val_init/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant uint o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/val_init/u32.wgsl.expected.spvasm b/test/var/override/named/val_init/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..93718f1
--- /dev/null
+++ b/test/var/override/named/val_init/u32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/val_init/u32.wgsl.expected.wgsl b/test/var/override/named/val_init/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..3a62e22
--- /dev/null
+++ b/test/var/override/named/val_init/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : u32 = 1u;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/bool.wgsl b/test/var/override/named/zero_init/bool.wgsl
new file mode 100644
index 0000000..1e03a45
--- /dev/null
+++ b/test/var/override/named/zero_init/bool.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : bool = bool();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.hlsl b/test/var/override/named/zero_init/bool.wgsl.expected.hlsl
new file mode 100644
index 0000000..d14a339
--- /dev/null
+++ b/test/var/override/named/zero_init/bool.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.msl b/test/var/override/named/zero_init/bool.wgsl.expected.msl
new file mode 100644
index 0000000..95ce385
--- /dev/null
+++ b/test/var/override/named/zero_init/bool.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant bool o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.spvasm b/test/var/override/named/zero_init/bool.wgsl.expected.spvasm
new file mode 100644
index 0000000..ff12e00
--- /dev/null
+++ b/test/var/override/named/zero_init/bool.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/zero_init/bool.wgsl.expected.wgsl b/test/var/override/named/zero_init/bool.wgsl.expected.wgsl
new file mode 100644
index 0000000..511bd67
--- /dev/null
+++ b/test/var/override/named/zero_init/bool.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : bool = bool();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/f32.wgsl b/test/var/override/named/zero_init/f32.wgsl
new file mode 100644
index 0000000..a47d934
--- /dev/null
+++ b/test/var/override/named/zero_init/f32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : f32 = f32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.hlsl b/test/var/override/named/zero_init/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..5ae59a4
--- /dev/null
+++ b/test/var/override/named/zero_init/f32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.msl b/test/var/override/named/zero_init/f32.wgsl.expected.msl
new file mode 100644
index 0000000..2ed922c
--- /dev/null
+++ b/test/var/override/named/zero_init/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant float o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.spvasm b/test/var/override/named/zero_init/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..7ae9060
--- /dev/null
+++ b/test/var/override/named/zero_init/f32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/zero_init/f32.wgsl.expected.wgsl b/test/var/override/named/zero_init/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..fe1468f
--- /dev/null
+++ b/test/var/override/named/zero_init/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : f32 = f32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/i32.wgsl b/test/var/override/named/zero_init/i32.wgsl
new file mode 100644
index 0000000..6455227
--- /dev/null
+++ b/test/var/override/named/zero_init/i32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : i32 = i32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.hlsl b/test/var/override/named/zero_init/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..6af794b
--- /dev/null
+++ b/test/var/override/named/zero_init/i32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.msl b/test/var/override/named/zero_init/i32.wgsl.expected.msl
new file mode 100644
index 0000000..b5b1fae
--- /dev/null
+++ b/test/var/override/named/zero_init/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant int o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.spvasm b/test/var/override/named/zero_init/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..18c8840
--- /dev/null
+++ b/test/var/override/named/zero_init/i32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/zero_init/i32.wgsl.expected.wgsl b/test/var/override/named/zero_init/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..5cde670
--- /dev/null
+++ b/test/var/override/named/zero_init/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : i32 = i32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/u32.wgsl b/test/var/override/named/zero_init/u32.wgsl
new file mode 100644
index 0000000..9d1205e
--- /dev/null
+++ b/test/var/override/named/zero_init/u32.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : u32 = u32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.hlsl b/test/var/override/named/zero_init/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..17f5ddf
--- /dev/null
+++ b/test/var/override/named/zero_init/u32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#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() {
+ return;
+}
diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.msl b/test/var/override/named/zero_init/u32.wgsl.expected.msl
new file mode 100644
index 0000000..918348c
--- /dev/null
+++ b/test/var/override/named/zero_init/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant uint o [[function_constant(0)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.spvasm b/test/var/override/named/zero_init/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..e5486b0
--- /dev/null
+++ b/test/var/override/named/zero_init/u32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/named/zero_init/u32.wgsl.expected.wgsl b/test/var/override/named/zero_init/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..f9d5644
--- /dev/null
+++ b/test/var/override/named/zero_init/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override]] let o : u32 = u32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/bool.wgsl b/test/var/override/numbered/no_init/bool.wgsl
new file mode 100644
index 0000000..a8b81c7
--- /dev/null
+++ b/test/var/override/numbered/no_init/bool.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : bool;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.hlsl b/test/var/override/numbered/no_init/bool.wgsl.expected.hlsl
new file mode 100644
index 0000000..eca5977
--- /dev/null
+++ b/test/var/override/numbered/no_init/bool.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#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() {
+ return;
+}
+/tmp/tint_4uqZfq:2:2: error: spec constant required for constant id 1234
+#error spec constant required for constant id 1234
+ ^
+
+
diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.msl b/test/var/override/numbered/no_init/bool.wgsl.expected.msl
new file mode 100644
index 0000000..d6cbfe9
--- /dev/null
+++ b/test/var/override/numbered/no_init/bool.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant bool o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.spvasm b/test/var/override/numbered/no_init/bool.wgsl.expected.spvasm
new file mode 100644
index 0000000..b1169e6
--- /dev/null
+++ b/test/var/override/numbered/no_init/bool.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/no_init/bool.wgsl.expected.wgsl b/test/var/override/numbered/no_init/bool.wgsl.expected.wgsl
new file mode 100644
index 0000000..e8d6daf
--- /dev/null
+++ b/test/var/override/numbered/no_init/bool.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : bool;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/f32.wgsl b/test/var/override/numbered/no_init/f32.wgsl
new file mode 100644
index 0000000..95e9dfd
--- /dev/null
+++ b/test/var/override/numbered/no_init/f32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.hlsl b/test/var/override/numbered/no_init/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f5c5a75
--- /dev/null
+++ b/test/var/override/numbered/no_init/f32.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#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() {
+ return;
+}
+/tmp/tint_UepNPT:2:2: error: spec constant required for constant id 1234
+#error spec constant required for constant id 1234
+ ^
+
+
diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.msl b/test/var/override/numbered/no_init/f32.wgsl.expected.msl
new file mode 100644
index 0000000..41134e2
--- /dev/null
+++ b/test/var/override/numbered/no_init/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant float o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.spvasm b/test/var/override/numbered/no_init/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..a253f64
--- /dev/null
+++ b/test/var/override/numbered/no_init/f32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/no_init/f32.wgsl.expected.wgsl b/test/var/override/numbered/no_init/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..cabed6a
--- /dev/null
+++ b/test/var/override/numbered/no_init/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/i32.wgsl b/test/var/override/numbered/no_init/i32.wgsl
new file mode 100644
index 0000000..74846ef
--- /dev/null
+++ b/test/var/override/numbered/no_init/i32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.hlsl b/test/var/override/numbered/no_init/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..0b62c77
--- /dev/null
+++ b/test/var/override/numbered/no_init/i32.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#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() {
+ return;
+}
+/tmp/tint_mwzbqN:2:2: error: spec constant required for constant id 1234
+#error spec constant required for constant id 1234
+ ^
+
+
diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.msl b/test/var/override/numbered/no_init/i32.wgsl.expected.msl
new file mode 100644
index 0000000..812c466
--- /dev/null
+++ b/test/var/override/numbered/no_init/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant int o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.spvasm b/test/var/override/numbered/no_init/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..aca0b07
--- /dev/null
+++ b/test/var/override/numbered/no_init/i32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %int = OpTypeInt 32 1
+ %o = OpSpecConstant %int 0
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/no_init/i32.wgsl.expected.wgsl b/test/var/override/numbered/no_init/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..68fb678
--- /dev/null
+++ b/test/var/override/numbered/no_init/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/u32.wgsl b/test/var/override/numbered/no_init/u32.wgsl
new file mode 100644
index 0000000..152c709
--- /dev/null
+++ b/test/var/override/numbered/no_init/u32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.hlsl b/test/var/override/numbered/no_init/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..443a2b1
--- /dev/null
+++ b/test/var/override/numbered/no_init/u32.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+SKIP: FAILED
+
+#ifndef WGSL_SPEC_CONSTANT_1234
+#error spec constant required for constant id 1234
+#endif
+static const uint o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
+/tmp/tint_U0EZFZ:2:2: error: spec constant required for constant id 1234
+#error spec constant required for constant id 1234
+ ^
+
+
diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.msl b/test/var/override/numbered/no_init/u32.wgsl.expected.msl
new file mode 100644
index 0000000..799a682
--- /dev/null
+++ b/test/var/override/numbered/no_init/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant uint o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.spvasm b/test/var/override/numbered/no_init/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..ca38a9a
--- /dev/null
+++ b/test/var/override/numbered/no_init/u32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %uint = OpTypeInt 32 0
+ %o = OpSpecConstant %uint 0
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/no_init/u32.wgsl.expected.wgsl b/test/var/override/numbered/no_init/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..40cb6c0
--- /dev/null
+++ b/test/var/override/numbered/no_init/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/bool.wgsl b/test/var/override/numbered/val_init/bool.wgsl
new file mode 100644
index 0000000..c29ec30
--- /dev/null
+++ b/test/var/override/numbered/val_init/bool.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : bool = true;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.hlsl b/test/var/override/numbered/val_init/bool.wgsl.expected.hlsl
new file mode 100644
index 0000000..6e189ed
--- /dev/null
+++ b/test/var/override/numbered/val_init/bool.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 true
+#endif
+static const bool o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.msl b/test/var/override/numbered/val_init/bool.wgsl.expected.msl
new file mode 100644
index 0000000..d6cbfe9
--- /dev/null
+++ b/test/var/override/numbered/val_init/bool.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant bool o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.spvasm b/test/var/override/numbered/val_init/bool.wgsl.expected.spvasm
new file mode 100644
index 0000000..060f860
--- /dev/null
+++ b/test/var/override/numbered/val_init/bool.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %bool = OpTypeBool
+ %o = OpSpecConstantTrue %bool
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/val_init/bool.wgsl.expected.wgsl b/test/var/override/numbered/val_init/bool.wgsl.expected.wgsl
new file mode 100644
index 0000000..2ab9051
--- /dev/null
+++ b/test/var/override/numbered/val_init/bool.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : bool = true;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/f32.wgsl b/test/var/override/numbered/val_init/f32.wgsl
new file mode 100644
index 0000000..649d4e8
--- /dev/null
+++ b/test/var/override/numbered/val_init/f32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : f32 = 1.0;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.hlsl b/test/var/override/numbered/val_init/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..4671b12
--- /dev/null
+++ b/test/var/override/numbered/val_init/f32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 1.0f
+#endif
+static const float o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.msl b/test/var/override/numbered/val_init/f32.wgsl.expected.msl
new file mode 100644
index 0000000..41134e2
--- /dev/null
+++ b/test/var/override/numbered/val_init/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant float o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.spvasm b/test/var/override/numbered/val_init/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..c757557
--- /dev/null
+++ b/test/var/override/numbered/val_init/f32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %float = OpTypeFloat 32
+ %o = OpSpecConstant %float 1
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/val_init/f32.wgsl.expected.wgsl b/test/var/override/numbered/val_init/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..bb12693
--- /dev/null
+++ b/test/var/override/numbered/val_init/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : f32 = 1.0;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/i32.wgsl b/test/var/override/numbered/val_init/i32.wgsl
new file mode 100644
index 0000000..d867f6a
--- /dev/null
+++ b/test/var/override/numbered/val_init/i32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : i32 = 1;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.hlsl b/test/var/override/numbered/val_init/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..cfb2fd9
--- /dev/null
+++ b/test/var/override/numbered/val_init/i32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 1
+#endif
+static const int o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.msl b/test/var/override/numbered/val_init/i32.wgsl.expected.msl
new file mode 100644
index 0000000..812c466
--- /dev/null
+++ b/test/var/override/numbered/val_init/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant int o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.spvasm b/test/var/override/numbered/val_init/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..f66889a
--- /dev/null
+++ b/test/var/override/numbered/val_init/i32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %int = OpTypeInt 32 1
+ %o = OpSpecConstant %int 1
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/val_init/i32.wgsl.expected.wgsl b/test/var/override/numbered/val_init/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..592dd92
--- /dev/null
+++ b/test/var/override/numbered/val_init/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : i32 = 1;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/u32.wgsl b/test/var/override/numbered/val_init/u32.wgsl
new file mode 100644
index 0000000..22c55f2
--- /dev/null
+++ b/test/var/override/numbered/val_init/u32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : u32 = 1u;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.hlsl b/test/var/override/numbered/val_init/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..36437dd
--- /dev/null
+++ b/test/var/override/numbered/val_init/u32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 1u
+#endif
+static const uint o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.msl b/test/var/override/numbered/val_init/u32.wgsl.expected.msl
new file mode 100644
index 0000000..799a682
--- /dev/null
+++ b/test/var/override/numbered/val_init/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant uint o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.spvasm b/test/var/override/numbered/val_init/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..2b51d1c
--- /dev/null
+++ b/test/var/override/numbered/val_init/u32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %uint = OpTypeInt 32 0
+ %o = OpSpecConstant %uint 1
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/val_init/u32.wgsl.expected.wgsl b/test/var/override/numbered/val_init/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..18bb63b
--- /dev/null
+++ b/test/var/override/numbered/val_init/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : u32 = 1u;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/bool.wgsl b/test/var/override/numbered/zero_init/bool.wgsl
new file mode 100644
index 0000000..894b584
--- /dev/null
+++ b/test/var/override/numbered/zero_init/bool.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : bool = bool();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/bool.wgsl.expected.hlsl
new file mode 100644
index 0000000..86e9b98
--- /dev/null
+++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 false
+#endif
+static const bool o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.msl b/test/var/override/numbered/zero_init/bool.wgsl.expected.msl
new file mode 100644
index 0000000..d6cbfe9
--- /dev/null
+++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant bool o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/bool.wgsl.expected.spvasm
new file mode 100644
index 0000000..b1169e6
--- /dev/null
+++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/zero_init/bool.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/bool.wgsl.expected.wgsl
new file mode 100644
index 0000000..653b1da
--- /dev/null
+++ b/test/var/override/numbered/zero_init/bool.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : bool = bool();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/f32.wgsl b/test/var/override/numbered/zero_init/f32.wgsl
new file mode 100644
index 0000000..0a3908c
--- /dev/null
+++ b/test/var/override/numbered/zero_init/f32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : f32 = f32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..70f868f
--- /dev/null
+++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 0.0f
+#endif
+static const float o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.msl b/test/var/override/numbered/zero_init/f32.wgsl.expected.msl
new file mode 100644
index 0000000..41134e2
--- /dev/null
+++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant float o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..a253f64
--- /dev/null
+++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; 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
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/zero_init/f32.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..589751b
--- /dev/null
+++ b/test/var/override/numbered/zero_init/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : f32 = f32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/i32.wgsl b/test/var/override/numbered/zero_init/i32.wgsl
new file mode 100644
index 0000000..4dd63fd
--- /dev/null
+++ b/test/var/override/numbered/zero_init/i32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : i32 = i32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..654ffd6
--- /dev/null
+++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 0
+#endif
+static const int o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.msl b/test/var/override/numbered/zero_init/i32.wgsl.expected.msl
new file mode 100644
index 0000000..812c466
--- /dev/null
+++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant int o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..aca0b07
--- /dev/null
+++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %int = OpTypeInt 32 1
+ %o = OpSpecConstant %int 0
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/zero_init/i32.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..8c0adca
--- /dev/null
+++ b/test/var/override/numbered/zero_init/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : i32 = i32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/u32.wgsl b/test/var/override/numbered/zero_init/u32.wgsl
new file mode 100644
index 0000000..d525204
--- /dev/null
+++ b/test/var/override/numbered/zero_init/u32.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : u32 = u32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.hlsl b/test/var/override/numbered/zero_init/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2fe25d8
--- /dev/null
+++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+#ifndef WGSL_SPEC_CONSTANT_1234
+#define WGSL_SPEC_CONSTANT_1234 0u
+#endif
+static const uint o = WGSL_SPEC_CONSTANT_1234;
+
+[numthreads(1, 1, 1)]
+void main() {
+ return;
+}
diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.msl b/test/var/override/numbered/zero_init/u32.wgsl.expected.msl
new file mode 100644
index 0000000..799a682
--- /dev/null
+++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant uint o [[function_constant(1234)]];
+kernel void tint_symbol() {
+ return;
+}
+
diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.spvasm b/test/var/override/numbered/zero_init/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..ca38a9a
--- /dev/null
+++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.spvasm
@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %o "o"
+ OpName %main "main"
+ OpDecorate %o SpecId 1234
+ %uint = OpTypeInt 32 0
+ %o = OpSpecConstant %uint 0
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %6 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/var/override/numbered/zero_init/u32.wgsl.expected.wgsl b/test/var/override/numbered/zero_init/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..d45213b
--- /dev/null
+++ b/test/var/override/numbered/zero_init/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[override(1234)]] let o : u32 = u32();
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ _ = o;
+}
