diff --git a/src/tint/transform/module_scope_var_to_entry_point_param.cc b/src/tint/transform/module_scope_var_to_entry_point_param.cc index 5a4d222544..e61c56e239 100644 --- a/src/tint/transform/module_scope_var_to_entry_point_param.cc +++ b/src/tint/transform/module_scope_var_to_entry_point_param.cc @@ -23,6 +23,7 @@ #include "src/tint/program_builder.h" #include "src/tint/sem/call.h" #include "src/tint/sem/function.h" +#include "src/tint/sem/module.h" #include "src/tint/sem/statement.h" #include "src/tint/sem/variable.h" @@ -282,7 +283,12 @@ struct ModuleScopeVarToEntryPointParam::State { std::vector functions_to_process; // Build a list of functions that transitively reference any module-scope variables. - for (auto* func_ast : ctx.src->AST().Functions()) { + for (auto* decl : ctx.src->Sem().Module()->DependencyOrderedDeclarations()) { + auto* func_ast = decl->As(); + if (!func_ast) { + continue; + } + auto* func_sem = ctx.src->Sem().Get(func_ast); bool needs_processing = false; diff --git a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc index 9e81d318bd..6a8d5eca28 100644 --- a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc +++ b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc @@ -168,9 +168,9 @@ var w : f32; auto* expect = R"( @compute @workgroup_size(1) fn main() { - @internal(disable_validation__ignore_storage_class) var tint_symbol : f32; - @internal(disable_validation__ignore_storage_class) var tint_symbol_1 : f32; - foo(1.0, &(tint_symbol), &(tint_symbol_1)); + @internal(disable_validation__ignore_storage_class) var tint_symbol_4 : f32; + @internal(disable_validation__ignore_storage_class) var tint_symbol_5 : f32; + foo(1.0, &(tint_symbol_4), &(tint_symbol_5)); } fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr) { @@ -182,9 +182,9 @@ fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(di fn no_uses() { } -fn bar(a : f32, b : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_4 : ptr, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_5 : ptr) { - *(tint_symbol_4) = a; - *(tint_symbol_5) = b; +fn bar(a : f32, b : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr) { + *(tint_symbol) = a; + *(tint_symbol_1) = b; } )"; @@ -365,12 +365,12 @@ var v : f32; auto* expect = R"( @compute @workgroup_size(1) fn main() { - @internal(disable_validation__ignore_storage_class) var tint_symbol : f32; - foo(&(tint_symbol)); + @internal(disable_validation__ignore_storage_class) var tint_symbol_1 : f32; + foo(&(tint_symbol_1)); } -fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr) { - bar(tint_symbol_1); +fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr) { + bar(tint_symbol); } fn bar(p : ptr) { @@ -556,17 +556,17 @@ fn foo() { )"; auto* expect = R"( -struct tint_symbol_1 { +struct tint_symbol_2 { arr : array, } @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr) { - foo(&((*(tint_symbol)).arr)); +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr) { + foo(&((*(tint_symbol_1)).arr)); } -fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr>) { - _ = (*(tint_symbol_2))[0]; +fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr>) { + _ = (*(tint_symbol))[0]; } )"; @@ -802,8 +802,8 @@ var s : S; auto* expect = R"( @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr) { - foo(1.0, tint_symbol, tint_symbol_1); +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_4 : ptr, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_5 : ptr) { + foo(1.0, tint_symbol_4, tint_symbol_5); } fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr) { @@ -816,9 +816,9 @@ fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(di fn no_uses() { } -fn bar(a : f32, b : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_4 : ptr, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_5 : ptr) { - _ = *(tint_symbol_4); - _ = *(tint_symbol_5); +fn bar(a : f32, b : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr) { + _ = *(tint_symbol); + _ = *(tint_symbol_1); } struct S { @@ -937,8 +937,8 @@ fn bar(a : f32, b : f32) { auto* expect = R"( @compute @workgroup_size(1) -fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) tint_symbol : texture_2d, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) tint_symbol_1 : sampler) { - foo(1.0, tint_symbol, tint_symbol_1); +fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) tint_symbol_4 : texture_2d, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) tint_symbol_5 : sampler) { + foo(1.0, tint_symbol_4, tint_symbol_5); } fn foo(a : f32, tint_symbol_2 : texture_2d, tint_symbol_3 : sampler) { @@ -951,9 +951,9 @@ fn foo(a : f32, tint_symbol_2 : texture_2d, tint_symbol_3 : sampler) { fn no_uses() { } -fn bar(a : f32, b : f32, tint_symbol_4 : texture_2d, tint_symbol_5 : sampler) { - _ = tint_symbol_4; - _ = tint_symbol_5; +fn bar(a : f32, b : f32, tint_symbol : texture_2d, tint_symbol_1 : sampler) { + _ = tint_symbol; + _ = tint_symbol_1; } )";