tint/reader: Allow module-scope 'var' type inferencing
Fixed: tint:1584
Change-Id: I193ad2c00faa4ae2001d981bb38a55d4d6a4c269
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94687
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
diff --git a/docs/tint/origin-trial-changes.md b/docs/tint/origin-trial-changes.md
index 6d04c01..3736506 100644
--- a/docs/tint/origin-trial-changes.md
+++ b/docs/tint/origin-trial-changes.md
@@ -2,6 +2,10 @@
## Changes for M105
+### New features
+
+* Module-scope `var<private>` can now infer the storage type, like function-scope `var`. [tint:1584](crbug.com/tint/1584)
+
### Breaking changes
* The `smoothStep()` builtin has been removed (use `smoothstep` instead). [tint:1483](crbug.com/tint/1483)
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index c39df85..35360e4 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -641,7 +641,7 @@
// variable_decl
// : VAR variable_qualifier? variable_ident_decl
-Maybe<ParserImpl::VarDeclInfo> ParserImpl::variable_decl(bool allow_inferred) {
+Maybe<ParserImpl::VarDeclInfo> ParserImpl::variable_decl() {
Source source;
if (!match(Token::Type::kVar, &source)) {
return Failure::kNoMatch;
@@ -656,7 +656,8 @@
vq = explicit_vq.value;
}
- auto decl = expect_variable_ident_decl("variable declaration", allow_inferred);
+ auto decl = expect_variable_ident_decl("variable declaration",
+ /*allow_inferred = */ true);
if (decl.errored) {
return Failure::kErrored;
}
@@ -1379,7 +1380,8 @@
return Failure::kErrored;
}
- auto decl = expect_variable_ident_decl("struct member");
+ auto decl = expect_variable_ident_decl("struct member",
+ /*allow_inferred = */ false);
if (decl.errored) {
return Failure::kErrored;
}
@@ -1515,7 +1517,8 @@
Expect<ast::Parameter*> ParserImpl::expect_param() {
auto attrs = attribute_list();
- auto decl = expect_variable_ident_decl("parameter");
+ auto decl = expect_variable_ident_decl("parameter",
+ /*allow_inferred = */ false);
if (decl.errored) {
return Failure::kErrored;
}
@@ -1798,7 +1801,7 @@
// | variable_decl EQUAL logical_or_expression
// | CONST variable_ident_decl EQUAL logical_or_expression
Maybe<const ast::VariableDeclStatement*> ParserImpl::variable_stmt() {
- if (const_enabled && match(Token::Type::kConst)) {
+ if (match(Token::Type::kConst)) {
auto decl = expect_variable_ident_decl("'const' declaration",
/*allow_inferred = */ true);
if (decl.errored) {
@@ -1854,7 +1857,7 @@
return create<ast::VariableDeclStatement>(decl->source, let);
}
- auto decl = variable_decl(/*allow_inferred = */ true);
+ auto decl = variable_decl();
if (decl.errored) {
return Failure::kErrored;
}
diff --git a/src/tint/reader/wgsl/parser_impl.h b/src/tint/reader/wgsl/parser_impl.h
index 73c5829..4ae4eef 100644
--- a/src/tint/reader/wgsl/parser_impl.h
+++ b/src/tint/reader/wgsl/parser_impl.h
@@ -393,18 +393,15 @@
/// @param attrs the list of attributes for the constant declaration.
Maybe<const ast::Variable*> global_constant_decl(ast::AttributeList& attrs);
/// Parses a `variable_decl` grammar element
- /// @param allow_inferred if true, do not fail if variable decl does not
- /// specify type
/// @returns the parsed variable declaration info
- Maybe<VarDeclInfo> variable_decl(bool allow_inferred = false);
+ Maybe<VarDeclInfo> variable_decl();
/// Parses a `variable_ident_decl` grammar element, erroring on parse
/// failure.
/// @param use a description of what was being parsed if an error was raised.
/// @param allow_inferred if true, do not fail if variable decl does not
/// specify type
/// @returns the identifier and type parsed or empty otherwise
- Expect<TypedIdentifier> expect_variable_ident_decl(std::string_view use,
- bool allow_inferred = false);
+ Expect<TypedIdentifier> expect_variable_ident_decl(std::string_view use, bool allow_inferred);
/// Parses a `variable_qualifier` grammar element
/// @returns the variable qualifier information
Maybe<VariableQualifier> variable_qualifier();
diff --git a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
index a09486d..3b48b5a 100644
--- a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
@@ -33,13 +33,20 @@
auto* v = program.AST().GlobalVariables()[0];
EXPECT_EQ(v->symbol, program.Symbols().Get("a"));
+ EXPECT_TRUE(Is<ast::Vector>(v->type));
}
-TEST_F(ParserImplTest, GlobalDecl_GlobalVariable_Inferred_Invalid) {
+TEST_F(ParserImplTest, GlobalDecl_GlobalVariable_Inferred) {
auto p = parser("var<private> a = vec2<i32>(1, 2);");
p->global_decl();
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:16: expected ':' for variable declaration");
+ ASSERT_FALSE(p->has_error()) << p->error();
+
+ auto program = p->program();
+ ASSERT_EQ(program.AST().GlobalVariables().size(), 1u);
+
+ auto* v = program.AST().GlobalVariables()[0];
+ EXPECT_EQ(v->symbol, program.Symbols().Get("a"));
+ EXPECT_EQ(v->type, nullptr);
}
TEST_F(ParserImplTest, GlobalDecl_GlobalVariable_MissingSemicolon) {
diff --git a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
index 70ec394..39f8b24 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
@@ -51,7 +51,7 @@
TEST_F(ParserImplTest, VariableDecl_Inferred_Parses) {
auto p = parser("var my_var = 1.0");
- auto v = p->variable_decl(/*allow_inferred = */ true);
+ auto v = p->variable_decl();
EXPECT_FALSE(p->has_error());
EXPECT_TRUE(v.matched);
EXPECT_FALSE(v.errored);
@@ -72,15 +72,6 @@
ASSERT_TRUE(t.IsIdentifier());
}
-TEST_F(ParserImplTest, VariableDecl_InvalidIdentDecl) {
- auto p = parser("var my_var f32");
- auto v = p->variable_decl();
- EXPECT_FALSE(v.matched);
- EXPECT_TRUE(v.errored);
- EXPECT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:12: expected ':' for variable declaration");
-}
-
TEST_F(ParserImplTest, VariableDecl_WithStorageClass) {
auto p = parser("var<private> my_var : f32");
auto v = p->variable_decl();
diff --git a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
index 3ffb29e..a811a82 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
@@ -19,7 +19,7 @@
TEST_F(ParserImplTest, VariableIdentDecl_Parses) {
auto p = parser("my_var : f32");
- auto decl = p->expect_variable_ident_decl("test");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ false);
ASSERT_FALSE(p->has_error()) << p->error();
ASSERT_FALSE(decl.errored);
ASSERT_EQ(decl->name, "my_var");
@@ -30,7 +30,27 @@
EXPECT_EQ(decl->type->source.range, (Source::Range{{1u, 10u}, {1u, 13u}}));
}
-TEST_F(ParserImplTest, VariableIdentDecl_Inferred_Parses) {
+TEST_F(ParserImplTest, VariableIdentDecl_Parses_AllowInferredType) {
+ auto p = parser("my_var : f32");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ true);
+ ASSERT_FALSE(p->has_error()) << p->error();
+ ASSERT_FALSE(decl.errored);
+ ASSERT_EQ(decl->name, "my_var");
+ ASSERT_NE(decl->type, nullptr);
+ ASSERT_TRUE(decl->type->Is<ast::F32>());
+
+ EXPECT_EQ(decl->source.range, (Source::Range{{1u, 1u}, {1u, 7u}}));
+ EXPECT_EQ(decl->type->source.range, (Source::Range{{1u, 10u}, {1u, 13u}}));
+}
+
+TEST_F(ParserImplTest, VariableIdentDecl_Inferred_Parse_Failure) {
+ auto p = parser("my_var = 1.0");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ false);
+ ASSERT_TRUE(p->has_error());
+ ASSERT_EQ(p->error(), "1:8: expected ':' for test");
+}
+
+TEST_F(ParserImplTest, VariableIdentDecl_Inferred_Parses_AllowInferredType) {
auto p = parser("my_var = 1.0");
auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ true);
ASSERT_FALSE(p->has_error()) << p->error();
@@ -43,23 +63,31 @@
TEST_F(ParserImplTest, VariableIdentDecl_MissingIdent) {
auto p = parser(": f32");
- auto decl = p->expect_variable_ident_decl("test");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ false);
ASSERT_TRUE(p->has_error());
ASSERT_TRUE(decl.errored);
ASSERT_EQ(p->error(), "1:1: expected identifier for test");
}
-TEST_F(ParserImplTest, VariableIdentDecl_MissingColon) {
- auto p = parser("my_var f32");
- auto decl = p->expect_variable_ident_decl("test");
+TEST_F(ParserImplTest, VariableIdentDecl_MissingIdent_AllowInferredType) {
+ auto p = parser(": f32");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ true);
ASSERT_TRUE(p->has_error());
ASSERT_TRUE(decl.errored);
- ASSERT_EQ(p->error(), "1:8: expected ':' for test");
+ ASSERT_EQ(p->error(), "1:1: expected identifier for test");
}
TEST_F(ParserImplTest, VariableIdentDecl_MissingType) {
auto p = parser("my_var :");
- auto decl = p->expect_variable_ident_decl("test");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ false);
+ ASSERT_TRUE(p->has_error());
+ ASSERT_TRUE(decl.errored);
+ ASSERT_EQ(p->error(), "1:9: invalid type for test");
+}
+
+TEST_F(ParserImplTest, VariableIdentDecl_MissingType_AllowInferredType) {
+ auto p = parser("my_var :");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ true);
ASSERT_TRUE(p->has_error());
ASSERT_TRUE(decl.errored);
ASSERT_EQ(p->error(), "1:9: invalid type for test");
@@ -67,7 +95,15 @@
TEST_F(ParserImplTest, VariableIdentDecl_InvalidIdent) {
auto p = parser("123 : f32");
- auto decl = p->expect_variable_ident_decl("test");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ false);
+ ASSERT_TRUE(p->has_error());
+ ASSERT_TRUE(decl.errored);
+ ASSERT_EQ(p->error(), "1:1: expected identifier for test");
+}
+
+TEST_F(ParserImplTest, VariableIdentDecl_InvalidIdent_AllowInferredType) {
+ auto p = parser("123 : f32");
+ auto decl = p->expect_variable_ident_decl("test", /*allow_inferred = */ true);
ASSERT_TRUE(p->has_error());
ASSERT_TRUE(decl.errored);
ASSERT_EQ(p->error(), "1:1: expected identifier for test");
diff --git a/test/tint/var/inferred/global.wgsl b/test/tint/var/inferred/global.wgsl
new file mode 100644
index 0000000..7ddb3e0
--- /dev/null
+++ b/test/tint/var/inferred/global.wgsl
@@ -0,0 +1,46 @@
+struct MyStruct {
+ f1 : f32,
+};
+
+type MyArray = array<f32, 10>;
+
+var<private> v1 = 1;
+var<private> v2 = 1u;
+var<private> v3 = 1.0;
+
+var<private> v4 = vec3<i32>(1, 1, 1);
+var<private> v5 = vec3<u32>(1u, 2u, 3u);
+var<private> v6 = vec3<f32>(1.0, 2.0, 3.0);
+
+var<private> v7 = MyStruct(1.0);
+var<private> v8 = MyArray();
+
+var<private> v9 = i32();
+var<private> v10 = u32();
+var<private> v11 = f32();
+var<private> v12 = MyStruct();
+var<private> v13 = MyStruct();
+var<private> v14 = MyArray();
+
+var<private> v15 = vec3(1, 2, 3);
+var<private> v16 = vec3(1.0, 2.0, 3.0);
+
+@compute @workgroup_size(1)
+fn f() {
+ let l1 = v1;
+ let l2 = v2;
+ let l3 = v3;
+ let l4 = v4;
+ let l5 = v5;
+ let l6 = v6;
+ let l7 = v7;
+ let l8 = v8;
+ let l9 = v9;
+ let l10 = v10;
+ let l11 = v11;
+ let l12 = v12;
+ let l13 = v13;
+ let l14 = v14;
+ let l15 = v15;
+ let l16 = v16;
+}
diff --git a/test/tint/var/inferred/global.wgsl.expected.glsl b/test/tint/var/inferred/global.wgsl.expected.glsl
new file mode 100644
index 0000000..0c7fe37
--- /dev/null
+++ b/test/tint/var/inferred/global.wgsl.expected.glsl
@@ -0,0 +1,46 @@
+#version 310 es
+
+struct MyStruct {
+ float f1;
+};
+
+int v1 = 1;
+uint v2 = 1u;
+float v3 = 1.0f;
+ivec3 v4 = ivec3(1);
+uvec3 v5 = uvec3(1u, 2u, 3u);
+vec3 v6 = vec3(1.0f, 2.0f, 3.0f);
+MyStruct v7 = MyStruct(1.0f);
+float v8[10] = float[10](0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+int v9 = 0;
+uint v10 = 0u;
+float v11 = 0.0f;
+MyStruct v12 = MyStruct(0.0f);
+MyStruct v13 = MyStruct(0.0f);
+float v14[10] = float[10](0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ivec3 v15 = ivec3(1, 2, 3);
+vec3 v16 = vec3(1.0f, 2.0f, 3.0f);
+void f() {
+ int l1 = v1;
+ uint l2 = v2;
+ float l3 = v3;
+ ivec3 l4 = v4;
+ uvec3 l5 = v5;
+ vec3 l6 = v6;
+ MyStruct l7 = v7;
+ float l8[10] = v8;
+ int l9 = v9;
+ uint l10 = v10;
+ float l11 = v11;
+ MyStruct l12 = v12;
+ MyStruct l13 = v13;
+ float l14[10] = v14;
+ ivec3 l15 = v15;
+ vec3 l16 = v16;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/var/inferred/global.wgsl.expected.hlsl b/test/tint/var/inferred/global.wgsl.expected.hlsl
new file mode 100644
index 0000000..09a7a21
--- /dev/null
+++ b/test/tint/var/inferred/global.wgsl.expected.hlsl
@@ -0,0 +1,41 @@
+struct MyStruct {
+ float f1;
+};
+
+static int v1 = 1;
+static uint v2 = 1u;
+static float v3 = 1.0f;
+static int3 v4 = (1).xxx;
+static uint3 v5 = uint3(1u, 2u, 3u);
+static float3 v6 = float3(1.0f, 2.0f, 3.0f);
+static MyStruct v7 = {1.0f};
+static float v8[10] = (float[10])0;
+static int v9 = 0;
+static uint v10 = 0u;
+static float v11 = 0.0f;
+static MyStruct v12 = (MyStruct)0;
+static MyStruct v13 = (MyStruct)0;
+static float v14[10] = (float[10])0;
+static int3 v15 = int3(1, 2, 3);
+static float3 v16 = float3(1.0f, 2.0f, 3.0f);
+
+[numthreads(1, 1, 1)]
+void f() {
+ const int l1 = v1;
+ const uint l2 = v2;
+ const float l3 = v3;
+ const int3 l4 = v4;
+ const uint3 l5 = v5;
+ const float3 l6 = v6;
+ const MyStruct l7 = v7;
+ const float l8[10] = v8;
+ const int l9 = v9;
+ const uint l10 = v10;
+ const float l11 = v11;
+ const MyStruct l12 = v12;
+ const MyStruct l13 = v13;
+ const float l14[10] = v14;
+ const int3 l15 = v15;
+ const float3 l16 = v16;
+ return;
+}
diff --git a/test/tint/var/inferred/global.wgsl.expected.msl b/test/tint/var/inferred/global.wgsl.expected.msl
new file mode 100644
index 0000000..50c5a7e
--- /dev/null
+++ b/test/tint/var/inferred/global.wgsl.expected.msl
@@ -0,0 +1,56 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct MyStruct {
+ float f1;
+};
+
+kernel void f() {
+ thread int tint_symbol = 1;
+ thread uint tint_symbol_1 = 1u;
+ thread float tint_symbol_2 = 1.0f;
+ thread int3 tint_symbol_3 = int3(1);
+ thread uint3 tint_symbol_4 = uint3(1u, 2u, 3u);
+ thread float3 tint_symbol_5 = float3(1.0f, 2.0f, 3.0f);
+ thread MyStruct tint_symbol_6 = {.f1=1.0f};
+ thread tint_array<float, 10> tint_symbol_7 = tint_array<float, 10>{};
+ thread int tint_symbol_8 = 0;
+ thread uint tint_symbol_9 = 0u;
+ thread float tint_symbol_10 = 0.0f;
+ thread MyStruct tint_symbol_11 = {};
+ thread MyStruct tint_symbol_12 = {};
+ thread tint_array<float, 10> tint_symbol_13 = tint_array<float, 10>{};
+ thread int3 tint_symbol_14 = int3(1, 2, 3);
+ thread float3 tint_symbol_15 = float3(1.0f, 2.0f, 3.0f);
+ int const l1 = tint_symbol;
+ uint const l2 = tint_symbol_1;
+ float const l3 = tint_symbol_2;
+ int3 const l4 = tint_symbol_3;
+ uint3 const l5 = tint_symbol_4;
+ float3 const l6 = tint_symbol_5;
+ MyStruct const l7 = tint_symbol_6;
+ tint_array<float, 10> const l8 = tint_symbol_7;
+ int const l9 = tint_symbol_8;
+ uint const l10 = tint_symbol_9;
+ float const l11 = tint_symbol_10;
+ MyStruct const l12 = tint_symbol_11;
+ MyStruct const l13 = tint_symbol_12;
+ tint_array<float, 10> const l14 = tint_symbol_13;
+ int3 const l15 = tint_symbol_14;
+ float3 const l16 = tint_symbol_15;
+ return;
+}
+
diff --git a/test/tint/var/inferred/global.wgsl.expected.spvasm b/test/tint/var/inferred/global.wgsl.expected.spvasm
new file mode 100644
index 0000000..e4bae60
--- /dev/null
+++ b/test/tint/var/inferred/global.wgsl.expected.spvasm
@@ -0,0 +1,104 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 73
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %v1 "v1"
+ OpName %v2 "v2"
+ OpName %v3 "v3"
+ OpName %v4 "v4"
+ OpName %v5 "v5"
+ OpName %v6 "v6"
+ OpName %MyStruct "MyStruct"
+ OpMemberName %MyStruct 0 "f1"
+ OpName %v7 "v7"
+ OpName %v8 "v8"
+ OpName %v9 "v9"
+ OpName %v10 "v10"
+ OpName %v11 "v11"
+ OpName %v12 "v12"
+ OpName %v13 "v13"
+ OpName %v14 "v14"
+ OpName %v15 "v15"
+ OpName %v16 "v16"
+ OpName %f "f"
+ OpMemberDecorate %MyStruct 0 Offset 0
+ OpDecorate %_arr_float_uint_10 ArrayStride 4
+ %int = OpTypeInt 32 1
+ %int_1 = OpConstant %int 1
+%_ptr_Private_int = OpTypePointer Private %int
+ %v1 = OpVariable %_ptr_Private_int Private %int_1
+ %uint = OpTypeInt 32 0
+ %uint_1 = OpConstant %uint 1
+%_ptr_Private_uint = OpTypePointer Private %uint
+ %v2 = OpVariable %_ptr_Private_uint Private %uint_1
+ %float = OpTypeFloat 32
+ %float_1 = OpConstant %float 1
+%_ptr_Private_float = OpTypePointer Private %float
+ %v3 = OpVariable %_ptr_Private_float Private %float_1
+ %v3int = OpTypeVector %int 3
+ %14 = OpConstantComposite %v3int %int_1 %int_1 %int_1
+%_ptr_Private_v3int = OpTypePointer Private %v3int
+ %v4 = OpVariable %_ptr_Private_v3int Private %14
+ %v3uint = OpTypeVector %uint 3
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+ %20 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+%_ptr_Private_v3uint = OpTypePointer Private %v3uint
+ %v5 = OpVariable %_ptr_Private_v3uint Private %20
+ %v3float = OpTypeVector %float 3
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %26 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+%_ptr_Private_v3float = OpTypePointer Private %v3float
+ %v6 = OpVariable %_ptr_Private_v3float Private %26
+ %MyStruct = OpTypeStruct %float
+ %30 = OpConstantComposite %MyStruct %float_1
+%_ptr_Private_MyStruct = OpTypePointer Private %MyStruct
+ %v7 = OpVariable %_ptr_Private_MyStruct Private %30
+ %uint_10 = OpConstant %uint 10
+%_arr_float_uint_10 = OpTypeArray %float %uint_10
+ %35 = OpConstantNull %_arr_float_uint_10
+%_ptr_Private__arr_float_uint_10 = OpTypePointer Private %_arr_float_uint_10
+ %v8 = OpVariable %_ptr_Private__arr_float_uint_10 Private %35
+ %38 = OpConstantNull %int
+ %v9 = OpVariable %_ptr_Private_int Private %38
+ %40 = OpConstantNull %uint
+ %v10 = OpVariable %_ptr_Private_uint Private %40
+ %42 = OpConstantNull %float
+ %v11 = OpVariable %_ptr_Private_float Private %42
+ %44 = OpConstantNull %MyStruct
+ %v12 = OpVariable %_ptr_Private_MyStruct Private %44
+ %v13 = OpVariable %_ptr_Private_MyStruct Private %44
+ %v14 = OpVariable %_ptr_Private__arr_float_uint_10 Private %35
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %50 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+ %v15 = OpVariable %_ptr_Private_v3int Private %50
+ %v16 = OpVariable %_ptr_Private_v3float Private %26
+ %void = OpTypeVoid
+ %53 = OpTypeFunction %void
+ %f = OpFunction %void None %53
+ %56 = OpLabel
+ %57 = OpLoad %int %v1
+ %58 = OpLoad %uint %v2
+ %59 = OpLoad %float %v3
+ %60 = OpLoad %v3int %v4
+ %61 = OpLoad %v3uint %v5
+ %62 = OpLoad %v3float %v6
+ %63 = OpLoad %MyStruct %v7
+ %64 = OpLoad %_arr_float_uint_10 %v8
+ %65 = OpLoad %int %v9
+ %66 = OpLoad %uint %v10
+ %67 = OpLoad %float %v11
+ %68 = OpLoad %MyStruct %v12
+ %69 = OpLoad %MyStruct %v13
+ %70 = OpLoad %_arr_float_uint_10 %v14
+ %71 = OpLoad %v3int %v15
+ %72 = OpLoad %v3float %v16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/var/inferred/global.wgsl.expected.wgsl b/test/tint/var/inferred/global.wgsl.expected.wgsl
new file mode 100644
index 0000000..bc7132f
--- /dev/null
+++ b/test/tint/var/inferred/global.wgsl.expected.wgsl
@@ -0,0 +1,57 @@
+struct MyStruct {
+ f1 : f32,
+}
+
+type MyArray = array<f32, 10>;
+
+var<private> v1 = 1;
+
+var<private> v2 = 1u;
+
+var<private> v3 = 1.0;
+
+var<private> v4 = vec3<i32>(1, 1, 1);
+
+var<private> v5 = vec3<u32>(1u, 2u, 3u);
+
+var<private> v6 = vec3<f32>(1.0, 2.0, 3.0);
+
+var<private> v7 = MyStruct(1.0);
+
+var<private> v8 = MyArray();
+
+var<private> v9 = i32();
+
+var<private> v10 = u32();
+
+var<private> v11 = f32();
+
+var<private> v12 = MyStruct();
+
+var<private> v13 = MyStruct();
+
+var<private> v14 = MyArray();
+
+var<private> v15 = vec3(1, 2, 3);
+
+var<private> v16 = vec3(1.0, 2.0, 3.0);
+
+@compute @workgroup_size(1)
+fn f() {
+ let l1 = v1;
+ let l2 = v2;
+ let l3 = v3;
+ let l4 = v4;
+ let l5 = v5;
+ let l6 = v6;
+ let l7 = v7;
+ let l8 = v8;
+ let l9 = v9;
+ let l10 = v10;
+ let l11 = v11;
+ let l12 = v12;
+ let l13 = v13;
+ let l14 = v14;
+ let l15 = v15;
+ let l16 = v16;
+}
