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