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>

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)

src/tint/reader/wgsl/parser_impl.cc

@@ -641,7 +641,7 @@ Maybe<const ast::Variable*> ParserImpl::global_constant_decl(ast::AttributeList&
 
 // 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 @@ Maybe<ParserImpl::VarDeclInfo> ParserImpl::variable_decl(bool allow_inferred) {
         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 @@ Expect<ast::StructMember*> ParserImpl::expect_struct_member() {
         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::ParameterList> ParserImpl::expect_param_list() {
 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 @@ Maybe<const ast::ReturnStatement*> ParserImpl::return_stmt() {
 //   | 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 @@ Maybe<const ast::VariableDeclStatement*> ParserImpl::variable_stmt() {
         return create<ast::VariableDeclStatement>(decl->source, let);
     }
 
-    auto decl = variable_decl(/*allow_inferred = */ true);
+    auto decl = variable_decl();
     if (decl.errored) {
         return Failure::kErrored;
     }

src/tint/reader/wgsl/parser_impl.h

@@ -393,18 +393,15 @@ class ParserImpl {
     /// @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,
+    Expect<TypedIdentifier> expect_variable_ident_decl(std::string_view use, bool allow_inferred);
-                                                       bool allow_inferred = false);
     /// Parses a `variable_qualifier` grammar element
     /// @returns the variable qualifier information
     Maybe<VariableQualifier> variable_qualifier();

src/tint/reader/wgsl/parser_impl_global_decl_test.cc

@@ -33,13 +33,20 @@ TEST_F(ParserImplTest, GlobalDecl_GlobalVariable) {
 
     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());
+    ASSERT_FALSE(p->has_error()) << p->error();
-    EXPECT_EQ(p->error(), "1:16: expected ':' for variable declaration");
+
+    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) {

src/tint/reader/wgsl/parser_impl_variable_decl_test.cc

@@ -51,7 +51,7 @@ TEST_F(ParserImplTest, VariableDecl_Unicode_Parses) {
 
 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 @@ TEST_F(ParserImplTest, VariableDecl_MissingVar) {
     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();

src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc

@@ -19,7 +19,7 @@ namespace {
 
 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 @@ TEST_F(ParserImplTest, VariableIdentDecl_Parses) {
     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_Inferred_Parses) {
 
 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) {
+TEST_F(ParserImplTest, VariableIdentDecl_MissingIdent_AllowInferredType) {
-    auto p = parser("my_var f32");
+    auto p = parser(": f32");
-    auto decl = p->expect_variable_ident_decl("test");
+    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_MissingType) {
 
 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");

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

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

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

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

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

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