tint: Implement uniformity analaysis
This implements the uniformity analysis as currently described in the
WGSL specification. Uniformity issues are presented as warnings, and
will be switched to errors in a future release.
A follow-up patch will improve the error messages, which currently
just show the point at which a uniformity was detected.
In a future release, once we have obtained initial feedback from
users, uniformity issues will become errors.
Bug: tint:880
Change-Id: I7d0b3080932c786c5d50b55720fec6d19f00d356
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/88368
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Alan Baker <alanbaker@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/docs/tint/origin-trial-changes.md b/docs/tint/origin-trial-changes.md
index 31f91ce..3876816 100644
--- a/docs/tint/origin-trial-changes.md
+++ b/docs/tint/origin-trial-changes.md
@@ -1,5 +1,12 @@
# Tint changes during Origin Trial
+## Changes for M103
+
+### New features
+
+* Produce warnings for when calling barriers, textureSample, and derivative
+builtins in non-uniform control flow [tint:880](crbug.com/tint/880)
+
## Changes for M102
### New Features
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 843f7fd..4a22a67 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -375,6 +375,8 @@
"resolver/resolver_constants.cc",
"resolver/sem_helper.cc",
"resolver/sem_helper.h",
+ "resolver/uniformity.cc",
+ "resolver/uniformity.h",
"resolver/validator.cc",
"resolver/validator.h",
"scope_stack.h",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 9635bef..d2273fd 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -256,6 +256,8 @@
resolver/resolver.h
resolver/sem_helper.cc
resolver/sem_helper.h
+ resolver/uniformity.cc
+ resolver/uniformity.h
resolver/validator.cc
resolver/validator.h
scope_stack.h
@@ -833,6 +835,13 @@
writer/text_generator_test.cc
)
+ # Uniformity analysis tests depend on WGSL reader
+ if(${TINT_BUILD_WGSL_READER})
+ list(APPEND TINT_TEST_SRCS
+ resolver/uniformity_test.cc
+ )
+ endif()
+
# Inspector tests depend on WGSL reader
if(${TINT_BUILD_WGSL_READER})
list(APPEND TINT_TEST_SRCS
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index fc7a34d..962dd02 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -50,6 +50,7 @@
#include "src/tint/ast/variable_decl_statement.h"
#include "src/tint/ast/vector.h"
#include "src/tint/ast/workgroup_attribute.h"
+#include "src/tint/resolver/uniformity.h"
#include "src/tint/sem/array.h"
#include "src/tint/sem/atomic.h"
#include "src/tint/sem/call.h"
@@ -145,6 +146,10 @@
return false;
}
+ if (!AnalyzeUniformity(builder_, dependencies_)) {
+ // TODO(jrprice): Reject programs that fail uniformity analysis.
+ }
+
bool result = true;
for (auto* node : builder_->ASTNodes().Objects()) {
if (marked_.count(node) == 0) {
diff --git a/src/tint/resolver/uniformity.cc b/src/tint/resolver/uniformity.cc
new file mode 100644
index 0000000..381a6da
--- /dev/null
+++ b/src/tint/resolver/uniformity.cc
@@ -0,0 +1,1275 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/resolver/uniformity.h"
+
+#include <limits>
+#include <string>
+#include <unordered_map>
+#include <unordered_set>
+#include <utility>
+#include <vector>
+
+#include "src/tint/program_builder.h"
+#include "src/tint/resolver/dependency_graph.h"
+#include "src/tint/scope_stack.h"
+#include "src/tint/sem/block_statement.h"
+#include "src/tint/sem/for_loop_statement.h"
+#include "src/tint/sem/function.h"
+#include "src/tint/sem/if_statement.h"
+#include "src/tint/sem/info.h"
+#include "src/tint/sem/loop_statement.h"
+#include "src/tint/sem/statement.h"
+#include "src/tint/sem/switch_statement.h"
+#include "src/tint/sem/type_constructor.h"
+#include "src/tint/sem/type_conversion.h"
+#include "src/tint/sem/variable.h"
+#include "src/tint/utils/block_allocator.h"
+#include "src/tint/utils/map.h"
+#include "src/tint/utils/unique_vector.h"
+
+// Set to `1` to dump the uniformity graph for each function in graphviz format.
+#define TINT_DUMP_UNIFORMITY_GRAPH 0
+
+namespace tint::resolver {
+
+namespace {
+
+/// CallSiteTag describes the uniformity requirements on the call sites of a function.
+enum CallSiteTag {
+ CallSiteRequiredToBeUniform,
+ CallSiteNoRestriction,
+};
+
+/// FunctionTag describes a functions effects on uniformity.
+enum FunctionTag {
+ SubsequentControlFlowMayBeNonUniform,
+ ReturnValueMayBeNonUniform,
+ NoRestriction,
+};
+
+/// ParameterTag describes the uniformity requirements of values passed to a function parameter.
+enum ParameterTag {
+ ParameterRequiredToBeUniform,
+ ParameterRequiredToBeUniformForSubsequentControlFlow,
+ ParameterRequiredToBeUniformForReturnValue,
+ ParameterNoRestriction,
+};
+
+/// ParameterInfo holds information about the uniformity requirements and effects for a particular
+/// function parameter.
+struct ParameterInfo {
+ /// The parameter's uniformity requirements.
+ ParameterTag tag = ParameterNoRestriction;
+ /// Will be `true` if this function may cause the contents of this pointer parameter to become
+ /// non-uniform.
+ bool pointer_may_become_non_uniform = false;
+ /// The parameters that are required to be uniform for the contents of this pointer parameter to
+ /// be uniform at function exit.
+ std::vector<const sem::Parameter*> pointer_param_output_sources;
+};
+
+/// FunctionInfo holds information about the uniformity requirements and effects for a particular
+/// function.
+struct FunctionInfo {
+ /// The call site uniformity requirements.
+ CallSiteTag callsite_tag;
+ /// The function's uniformity effects.
+ FunctionTag function_tag;
+ /// The uniformity requirements of the function's parameters.
+ std::vector<ParameterInfo> parameters;
+};
+
+/// Node represents a node in the graph of control flow and value nodes within the analysis of a
+/// single function.
+struct Node {
+ /// Constructor
+ /// @param t the node tag (used for debugging)
+ Node([[maybe_unused]] std::string t, const ast::Node* a)
+ :
+#if TINT_DUMP_UNIFORMITY_GRAPH
+ tag(t),
+#endif
+ ast(a) {
+ }
+
+#if TINT_DUMP_UNIFORMITY_GRAPH
+ /// The node tag.
+ const std::string tag;
+#endif
+
+ /// The corresponding AST node, or nullptr.
+ const ast::Node* ast = nullptr;
+
+ /// The function call argument index, or UINT32_MAX.
+ uint32_t arg_index = std::numeric_limits<uint32_t>::max();
+
+ /// The set of edges from this node to other nodes in the graph.
+ utils::UniqueVector<Node*> edges;
+
+ /// The node that this node was visited from, or nullptr if not visited.
+ Node* visited_from = nullptr;
+
+ /// Add a edge to the `to` node.
+ /// @param to the destination node
+ void AddEdge(Node* to) { edges.add(to); }
+};
+
+/// UniformityGraph is used to analyze the uniformity requirements and effects of functions in a
+/// module.
+class UniformityGraph {
+ public:
+ /// Constructor.
+ /// @param builder the program to analyze
+ explicit UniformityGraph(ProgramBuilder* builder)
+ : builder_(builder), sem_(builder->Sem()), diagnostics_(builder->Diagnostics()) {}
+
+ /// Destructor.
+ ~UniformityGraph() {}
+
+ /// Build and analyze the graph to determine whether the program satisfies the uniformity
+ /// constraints of WGSL.
+ /// @param dependency_graph the dependency-ordered module-scope declarations
+ /// @returns true if all uniformity constraints are satisfied, otherise false
+ bool Build(const DependencyGraph& dependency_graph) {
+#if TINT_DUMP_UNIFORMITY_GRAPH
+ std::cout << "digraph G {\n";
+ std::cout << "rankdir=BT\n";
+#endif
+
+ // Process all functions in the module.
+ bool success = true;
+ for (auto* decl : dependency_graph.ordered_globals) {
+ if (auto* func = decl->As<ast::Function>()) {
+ if (!ProcessFunction(func)) {
+ success = false;
+ break;
+ }
+ }
+ }
+
+#if TINT_DUMP_UNIFORMITY_GRAPH
+ std::cout << "\n}\n";
+#endif
+
+ return success;
+ }
+
+ private:
+ const ProgramBuilder* builder_;
+ const sem::Info& sem_;
+ diag::List& diagnostics_;
+
+ /// Map of analyzed function results.
+ std::unordered_map<const ast::Function*, FunctionInfo> functions_;
+
+ /// Allocator for nodes.
+ utils::BlockAllocator<Node> nodes_;
+
+ /// Name of the function currently being analyzed.
+ std::string current_function_;
+
+ /// Special `RequiredToBeUniform` node.
+ Node* required_to_be_uniform_;
+ /// Special `MayBeNonUniform` node.
+ Node* may_be_non_uniform_;
+ /// Special `CF_return` node.
+ Node* cf_return_;
+ /// Special `Value_return` node.
+ Node* value_return_;
+ /// Special `{param}_return` nodes for pointer parameters.
+ std::unordered_map<const sem::Parameter*, Node*> pointer_param_returns_;
+
+ /// Map from variables to their value nodes in the graph, scoped with respect to control flow.
+ ScopeStack<const sem::Variable*, Node*> variables_;
+
+ /// The set of a local read-write vars that are in scope at any given point in the process.
+ /// Includes pointer parameters.
+ std::unordered_set<const sem::Variable*> local_var_decls_;
+
+ /// LoopSwitchInfo tracks information about the value of variables for a control flow construct.
+ struct LoopSwitchInfo {
+ /// The type of this control flow construct.
+ std::string type;
+ /// The input values for local variables at the start of this construct.
+ std::unordered_map<const sem::Variable*, Node*> var_in_nodes;
+ /// The exit values for local variables at the end of this construct.
+ std::unordered_map<const sem::Variable*, Node*> var_exit_nodes;
+ };
+
+ /// Map from control flow statements to the corresponding LoopSwitchInfo structure.
+ std::unordered_map<const sem::Statement*, LoopSwitchInfo> loop_switch_infos_;
+
+ /// A list of tags that have already been used within the current function.
+ std::unordered_set<std::string> tags_;
+
+ /// Create a new node.
+ /// @param tag a tag used to identify the node for debugging purposes.
+ /// @param ast the optional AST node that this node corresponds to
+ /// @returns the new node
+ Node* CreateNode(std::string tag, const ast::Node* ast = nullptr) {
+ std::string unique_tag = tag;
+
+#if TINT_DUMP_UNIFORMITY_GRAPH
+ // Make the tag unique.
+ // This only matters if we're dumping the graph.
+ int suffix = 0;
+ while (tags_.count(unique_tag)) {
+ unique_tag = tag + "_$" + std::to_string(++suffix);
+ }
+ tags_.insert(unique_tag);
+#endif
+
+ return nodes_.Create(current_function_ + "." + unique_tag, ast);
+ }
+
+ /// Process a function.
+ /// @param func the function to process
+ /// @returns true if there are no uniformity issues, false otherwise
+ bool ProcessFunction(const ast::Function* func) {
+ nodes_.Reset();
+ variables_.Clear();
+ pointer_param_returns_.clear();
+ tags_.clear();
+
+ current_function_ = builder_->Symbols().NameFor(func->symbol);
+
+ // Create special nodes.
+ Node* cf_start = CreateNode("CF_start");
+ required_to_be_uniform_ = CreateNode("RequiredToBeUniform");
+ may_be_non_uniform_ = CreateNode("MayBeNonUniform");
+ cf_return_ = CreateNode("CF_return");
+ if (func->return_type) {
+ value_return_ = CreateNode("Value_return");
+ }
+
+ // Create nodes for parameters.
+ std::vector<Node*> param_init_values(func->params.size(), nullptr);
+ for (size_t i = 0; i < func->params.size(); i++) {
+ auto* param = func->params[i];
+ auto name = builder_->Symbols().NameFor(param->symbol);
+ auto sem = sem_.Get<sem::Parameter>(param);
+
+ Node* node_init;
+ if (sem->Type()->Is<sem::Pointer>()) {
+ node_init = CreateNode("ptrparam_" + name + "_init");
+ pointer_param_returns_[sem] = CreateNode("ptrparam_" + name + "_return");
+ local_var_decls_.insert(sem);
+ } else {
+ node_init = CreateNode("param_" + name);
+ }
+ param_init_values[i] = node_init;
+ variables_.Set(sem, node_init);
+ }
+
+ // Process function body.
+ if (func->body) {
+ ProcessStatement(cf_start, func->body);
+ }
+
+#if TINT_DUMP_UNIFORMITY_GRAPH
+ // Dump the graph for this function as a subgraph.
+ std::cout << "\nsubgraph cluster_" << current_function_ << " {\n";
+ std::cout << " label=" << current_function_ << ";";
+ for (auto* node : nodes_.Objects()) {
+ std::cout << "\n \"" << node->tag << "\";";
+ for (auto* edge : node->edges) {
+ std::cout << "\n \"" << node->tag << "\" -> \"" << edge->tag << "\";";
+ }
+ }
+ std::cout << "\n}\n";
+#endif
+
+ FunctionInfo& info = functions_[func];
+ info.callsite_tag = CallSiteNoRestriction;
+ info.function_tag = NoRestriction;
+ info.parameters.resize(func->params.size());
+
+ // Look at which nodes are reachable from "RequiredToBeUniform".
+ {
+ utils::UniqueVector<Node*> reachable;
+ Traverse(required_to_be_uniform_, reachable);
+ if (reachable.contains(may_be_non_uniform_)) {
+ MakeError();
+ return false;
+ }
+ if (reachable.contains(cf_start)) {
+ info.callsite_tag = CallSiteRequiredToBeUniform;
+ }
+
+ // Set the parameter tag to ParameterRequiredToBeUniform for each parameter node that
+ // was reachable.
+ for (size_t i = 0; i < func->params.size(); i++) {
+ auto* param = func->params[i];
+ if (reachable.contains(variables_.Get(sem_.Get(param)))) {
+ info.parameters[i].tag = ParameterRequiredToBeUniform;
+ }
+ }
+ }
+
+ // Look at which nodes are reachable from "CF_return"
+ {
+ utils::UniqueVector<Node*> reachable;
+ Traverse(cf_return_, reachable);
+ if (reachable.contains(may_be_non_uniform_)) {
+ info.function_tag = SubsequentControlFlowMayBeNonUniform;
+ }
+
+ // Set the parameter tag to ParameterRequiredToBeUniformForSubsequentControlFlow for
+ // each parameter node that was reachable.
+ for (size_t i = 0; i < func->params.size(); i++) {
+ auto* param = func->params[i];
+ if (reachable.contains(variables_.Get(sem_.Get(param)))) {
+ info.parameters[i].tag = ParameterRequiredToBeUniformForSubsequentControlFlow;
+ }
+ }
+ }
+
+ // If "Value_return" exists, look at which nodes are reachable from it
+ if (value_return_) {
+ utils::UniqueVector<Node*> reachable;
+ Traverse(value_return_, reachable);
+ if (reachable.contains(may_be_non_uniform_)) {
+ info.function_tag = ReturnValueMayBeNonUniform;
+ }
+
+ // Set the parameter tag to ParameterRequiredToBeUniformForReturnValue for each
+ // parameter node that was reachable.
+ for (size_t i = 0; i < func->params.size(); i++) {
+ auto* param = func->params[i];
+ if (reachable.contains(variables_.Get(sem_.Get(param)))) {
+ info.parameters[i].tag = ParameterRequiredToBeUniformForReturnValue;
+ }
+ }
+ }
+
+ // Traverse the graph for each pointer parameter.
+ for (size_t i = 0; i < func->params.size(); i++) {
+ auto* param_dest = sem_.Get<sem::Parameter>(func->params[i]);
+ if (!param_dest->Type()->Is<sem::Pointer>()) {
+ continue;
+ }
+
+ // Reset "visited" state for all nodes.
+ for (auto* node : nodes_.Objects()) {
+ node->visited_from = nullptr;
+ }
+
+ utils::UniqueVector<Node*> reachable;
+ Traverse(pointer_param_returns_[param_dest], reachable);
+ if (reachable.contains(may_be_non_uniform_)) {
+ info.parameters[i].pointer_may_become_non_uniform = true;
+ }
+
+ // Check every other parameter to see if they feed into this parameter's final value.
+ for (size_t j = 0; j < func->params.size(); j++) {
+ auto* param_source = sem_.Get<sem::Parameter>(func->params[j]);
+ if (reachable.contains(param_init_values[j])) {
+ info.parameters[i].pointer_param_output_sources.push_back(param_source);
+ }
+ }
+ }
+
+ return true;
+ }
+
+ /// Process a statement, returning the new control flow node.
+ /// @param cf the input control flow node
+ /// @param stmt the statement to process d
+ /// @returns the new control flow node
+ Node* ProcessStatement(Node* cf, const ast::Statement* stmt) {
+ return Switch(
+ stmt,
+
+ [&](const ast::AssignmentStatement* a) {
+ auto [cf1, v1] = ProcessExpression(cf, a->rhs);
+ if (a->lhs->Is<ast::PhonyExpression>()) {
+ return cf1;
+ } else {
+ auto [cf2, l2] = ProcessLValueExpression(cf1, a->lhs);
+ l2->AddEdge(v1);
+ return cf2;
+ }
+ },
+
+ [&](const ast::BlockStatement* b) {
+ std::unordered_map<const sem::Variable*, Node*> scoped_assignments;
+ {
+ // Push a new scope for variable assignments in the block.
+ variables_.Push();
+ TINT_DEFER(variables_.Pop());
+
+ for (auto* s : b->statements) {
+ cf = ProcessStatement(cf, s);
+ if (!sem_.Get(s)->Behaviors().Contains(sem::Behavior::kNext)) {
+ break;
+ }
+ }
+
+ if (sem_.Get<sem::FunctionBlockStatement>(b)) {
+ // We've reached the end of the function body.
+ // Add edges from pointer parameter outputs to their current value.
+ for (auto param : pointer_param_returns_) {
+ param.second->AddEdge(variables_.Get(param.first));
+ }
+ }
+
+ scoped_assignments = std::move(variables_.Top());
+ }
+
+ // Propagate all variables assignments to the containing scope if the behavior is
+ // either 'Next' or 'Fallthrough'.
+ auto& behaviors = sem_.Get(b)->Behaviors();
+ if (behaviors.Contains(sem::Behavior::kNext) ||
+ behaviors.Contains(sem::Behavior::kFallthrough)) {
+ for (auto var : scoped_assignments) {
+ variables_.Set(var.first, var.second);
+ }
+ }
+
+ // Remove any variables declared in this scope from the set of in-scope variables.
+ for (auto* d : sem_.Get<sem::BlockStatement>(b)->Decls()) {
+ local_var_decls_.erase(sem_.Get<sem::LocalVariable>(d));
+ }
+
+ return cf;
+ },
+
+ [&](const ast::BreakStatement* b) {
+ // Find the loop or switch statement that we are in.
+ auto* parent = sem_.Get(b)
+ ->FindFirstParent<sem::SwitchStatement, sem::LoopStatement,
+ sem::ForLoopStatement>();
+ TINT_ASSERT(Resolver, loop_switch_infos_.count(parent));
+ auto& info = loop_switch_infos_.at(parent);
+
+ // Propagate variable values to the loop/switch exit nodes.
+ for (auto* var : local_var_decls_) {
+ // Skip variables that were declared inside this loop/switch.
+ if (auto* lv = var->As<sem::LocalVariable>();
+ lv &&
+ lv->Statement()->FindFirstParent([&](auto* s) { return s == parent; })) {
+ continue;
+ }
+
+ // Add an edge from the variable exit node to its value at this point.
+ auto* exit_node = utils::GetOrCreate(info.var_exit_nodes, var, [&]() {
+ auto name = builder_->Symbols().NameFor(var->Declaration()->symbol);
+ return CreateNode(name + "_value_" + info.type + "_exit");
+ });
+ exit_node->AddEdge(variables_.Get(var));
+ }
+
+ return cf;
+ },
+
+ [&](const ast::CallStatement* c) {
+ auto [cf1, _] = ProcessCall(cf, c->expr);
+ return cf1;
+ },
+
+ [&](const ast::CompoundAssignmentStatement* c) {
+ // The compound assignment statement `a += b` is equivalent to `a = a + b`.
+ auto [cf1, v1] = ProcessExpression(cf, c->lhs);
+ auto [cf2, v2] = ProcessExpression(cf1, c->rhs);
+ auto* result = CreateNode("binary_expr_result");
+ result->AddEdge(v1);
+ result->AddEdge(v2);
+
+ auto [cf3, l3] = ProcessLValueExpression(cf2, c->lhs);
+ l3->AddEdge(result);
+ return cf3;
+ },
+
+ [&](const ast::ContinueStatement* c) {
+ // Find the loop statement that we are in.
+ auto* parent =
+ sem_.Get(c)->FindFirstParent<sem::LoopStatement, sem::ForLoopStatement>();
+ TINT_ASSERT(Resolver, loop_switch_infos_.count(parent));
+ auto& info = loop_switch_infos_.at(parent);
+
+ // Propagate assignments to the loop input nodes.
+ for (auto* var : local_var_decls_) {
+ // Skip variables that were declared inside this loop.
+ if (auto* lv = var->As<sem::LocalVariable>();
+ lv &&
+ lv->Statement()->FindFirstParent([&](auto* s) { return s == parent; })) {
+ continue;
+ }
+
+ // Add an edge from the variable's loop input node to its value at this point.
+ TINT_ASSERT(Resolver, info.var_in_nodes.count(var));
+ auto* in_node = info.var_in_nodes.at(var);
+ auto* out_node = variables_.Get(var);
+ if (out_node != in_node) {
+ in_node->AddEdge(out_node);
+ }
+ }
+ return cf;
+ },
+
+ [&](const ast::DiscardStatement*) {
+ cf_return_->AddEdge(cf);
+ return cf;
+ },
+
+ [&](const ast::FallthroughStatement*) { return cf; },
+
+ [&](const ast::ForLoopStatement* f) {
+ auto* sem_loop = sem_.Get(f);
+ auto cfx = CreateNode("loop_start");
+
+ // Insert the initializer before the loop.
+ auto cf_init = cf;
+ if (f->initializer) {
+ cf_init = ProcessStatement(cf, f->initializer);
+ }
+ auto cf_start = cf_init;
+
+ auto& info = loop_switch_infos_[sem_loop];
+ info.type = "forloop";
+
+ // Create input nodes for any variables declared before this loop.
+ for (auto* v : local_var_decls_) {
+ auto name = builder_->Symbols().NameFor(v->Declaration()->symbol);
+ auto* in_node = CreateNode(name + "_value_forloop_in");
+ in_node->AddEdge(variables_.Get(v));
+ info.var_in_nodes[v] = in_node;
+ variables_.Set(v, in_node);
+ }
+
+ // Insert the condition at the start of the loop body.
+ if (f->condition) {
+ auto [cf_cond, v] = ProcessExpression(cfx, f->condition);
+ auto* cf_condition_end = CreateNode("for_condition_CFend");
+ cf_condition_end->AddEdge(v);
+ cf_start = cf_condition_end;
+
+ // Propagate assignments to the loop exit nodes.
+ for (auto* var : local_var_decls_) {
+ auto* exit_node = utils::GetOrCreate(info.var_exit_nodes, var, [&]() {
+ auto name = builder_->Symbols().NameFor(var->Declaration()->symbol);
+ return CreateNode(name + "_value_" + info.type + "_exit");
+ });
+ exit_node->AddEdge(variables_.Get(var));
+ }
+ }
+ auto cf1 = ProcessStatement(cf_start, f->body);
+
+ // Insert the continuing statement at the end of the loop body.
+ if (f->continuing) {
+ auto cf2 = ProcessStatement(cf1, f->continuing);
+ cfx->AddEdge(cf2);
+ } else {
+ cfx->AddEdge(cf1);
+ }
+ cfx->AddEdge(cf);
+
+ // Add edges from variable loop input nodes to their values at the end of the loop.
+ for (auto v : info.var_in_nodes) {
+ auto* in_node = v.second;
+ auto* out_node = variables_.Get(v.first);
+ if (out_node != in_node) {
+ in_node->AddEdge(out_node);
+ }
+ }
+
+ // Set each variable's exit node as its value in the outer scope.
+ for (auto v : info.var_exit_nodes) {
+ variables_.Set(v.first, v.second);
+ }
+
+ loop_switch_infos_.erase(sem_loop);
+
+ if (sem_loop->Behaviors() == sem::Behaviors{sem::Behavior::kNext}) {
+ return cf;
+ } else {
+ return cfx;
+ }
+ },
+
+ [&](const ast::IfStatement* i) {
+ auto* sem_if = sem_.Get(i);
+ auto [cfx, v] = ProcessExpression(cf, i->condition);
+
+ std::unordered_map<const sem::Variable*, Node*> true_vars;
+ std::unordered_map<const sem::Variable*, Node*> false_vars;
+
+ // Helper to process a statement with a new scope for variable assignments.
+ // Populates `assigned_vars` with new nodes for any variables that are assigned in
+ // this statement.
+ auto process_in_scope =
+ [&](Node* cf_in, const ast::Statement* s,
+ std::unordered_map<const sem::Variable*, Node*>& assigned_vars) {
+ // Push a new scope for variable assignments.
+ variables_.Push();
+
+ // Process the statement.
+ auto cf_out = ProcessStatement(cf_in, s);
+
+ assigned_vars = variables_.Top();
+
+ // Pop the scope and return.
+ variables_.Pop();
+ return cf_out;
+ };
+
+ auto cf1 = process_in_scope(v, i->body, true_vars);
+
+ bool true_has_next = sem_.Get(i->body)->Behaviors().Contains(sem::Behavior::kNext);
+ bool false_has_next = true;
+
+ Node* cf2 = nullptr;
+ if (i->else_statement) {
+ cf2 = process_in_scope(v, i->else_statement, false_vars);
+
+ false_has_next =
+ sem_.Get(i->else_statement)->Behaviors().Contains(sem::Behavior::kNext);
+ }
+
+ // Update values for any variables assigned in the if or else blocks.
+ for (auto var : local_var_decls_) {
+ // Skip variables not assigned in either block.
+ if (true_vars.count(var) == 0 && false_vars.count(var) == 0) {
+ continue;
+ }
+
+ // Create an exit node for the variable.
+ auto name = builder_->Symbols().NameFor(var->Declaration()->symbol);
+ auto* out_node = CreateNode(name + "_value_if_exit");
+
+ // Add edges to the assigned value or the initial value.
+ // Only add edges if the behavior for that block contains 'Next'.
+ if (true_has_next) {
+ if (true_vars.count(var)) {
+ out_node->AddEdge(true_vars.at(var));
+ } else {
+ out_node->AddEdge(variables_.Get(var));
+ }
+ }
+ if (false_has_next) {
+ if (false_vars.count(var)) {
+ out_node->AddEdge(false_vars.at(var));
+ } else {
+ out_node->AddEdge(variables_.Get(var));
+ }
+ }
+
+ variables_.Set(var, out_node);
+ }
+
+ if (sem_if->Behaviors() != sem::Behaviors{sem::Behavior::kNext}) {
+ auto* cf_end = CreateNode("if_CFend");
+ cf_end->AddEdge(cf1);
+ if (cf2) {
+ cf_end->AddEdge(cf2);
+ }
+ return cf_end;
+ }
+ return cf;
+ },
+
+ [&](const ast::IncrementDecrementStatement* i) {
+ // The increment/decrement statement `i++` is equivalent to `i = i + 1`.
+ auto [cf1, v1] = ProcessExpression(cf, i->lhs);
+ auto* result = CreateNode("incdec_result");
+ result->AddEdge(v1);
+ result->AddEdge(cf1);
+
+ auto [cf2, l2] = ProcessLValueExpression(cf1, i->lhs);
+ l2->AddEdge(result);
+ return cf2;
+ },
+
+ [&](const ast::LoopStatement* l) {
+ auto* sem_loop = sem_.Get(l);
+ auto cfx = CreateNode("loop_start");
+
+ auto& info = loop_switch_infos_[sem_loop];
+ info.type = "loop";
+
+ // Create input nodes for any variables declared before this loop.
+ for (auto* v : local_var_decls_) {
+ auto name = builder_->Symbols().NameFor(v->Declaration()->symbol);
+ auto* in_node = CreateNode(name + "_value_loop_in");
+ in_node->AddEdge(variables_.Get(v));
+ info.var_in_nodes[v] = in_node;
+ variables_.Set(v, in_node);
+ }
+
+ auto cf1 = ProcessStatement(cfx, l->body);
+ if (l->continuing) {
+ auto cf2 = ProcessStatement(cf1, l->continuing);
+ cfx->AddEdge(cf2);
+ } else {
+ cfx->AddEdge(cf1);
+ }
+ cfx->AddEdge(cf);
+
+ // Add edges from variable loop input nodes to their values at the end of the loop.
+ for (auto v : info.var_in_nodes) {
+ auto* in_node = v.second;
+ auto* out_node = variables_.Get(v.first);
+ if (out_node != in_node) {
+ in_node->AddEdge(out_node);
+ }
+ }
+
+ // Set each variable's exit node as its value in the outer scope.
+ for (auto v : info.var_exit_nodes) {
+ variables_.Set(v.first, v.second);
+ }
+
+ loop_switch_infos_.erase(sem_loop);
+
+ if (sem_loop->Behaviors() == sem::Behaviors{sem::Behavior::kNext}) {
+ return cf;
+ } else {
+ return cfx;
+ }
+ },
+ [&](const ast::ReturnStatement* r) {
+ Node* cf_ret;
+ if (r->value) {
+ auto [cf1, v] = ProcessExpression(cf, r->value);
+ cf_return_->AddEdge(cf1);
+ value_return_->AddEdge(v);
+ cf_ret = cf1;
+ } else {
+ TINT_ASSERT(Resolver, cf != nullptr);
+ cf_return_->AddEdge(cf);
+ cf_ret = cf;
+ }
+
+ // Add edges from each pointer parameter output to its current value.
+ for (auto param : pointer_param_returns_) {
+ param.second->AddEdge(variables_.Get(param.first));
+ }
+
+ return cf_ret;
+ },
+ [&](const ast::SwitchStatement* s) {
+ auto* sem_switch = sem_.Get(s);
+ auto [cfx, v] = ProcessExpression(cf, s->condition);
+
+ Node* cf_end = nullptr;
+ if (sem_switch->Behaviors() != sem::Behaviors{sem::Behavior::kNext}) {
+ cf_end = CreateNode("switch_CFend");
+ }
+
+ auto& info = loop_switch_infos_[sem_switch];
+ info.type = "switch";
+
+ auto cf_n = v;
+ bool previous_case_has_fallthrough = false;
+ for (auto* c : s->body) {
+ auto* sem_case = sem_.Get(c);
+
+ if (previous_case_has_fallthrough) {
+ cf_n = ProcessStatement(cf_n, c->body);
+ } else {
+ variables_.Push();
+ cf_n = ProcessStatement(v, c->body);
+ }
+
+ if (cf_end) {
+ cf_end->AddEdge(cf_n);
+ }
+
+ bool has_fallthrough =
+ sem_case->Behaviors().Contains(sem::Behavior::kFallthrough);
+ if (!has_fallthrough) {
+ if (sem_case->Behaviors().Contains(sem::Behavior::kNext)) {
+ // Propagate variable values to the switch exit nodes.
+ for (auto* var : local_var_decls_) {
+ // Skip variables that were declared inside the switch.
+ if (auto* lv = var->As<sem::LocalVariable>();
+ lv && lv->Statement()->FindFirstParent(
+ [&](auto* st) { return st == sem_switch; })) {
+ continue;
+ }
+
+ // Add an edge from the variable exit node to its new value.
+ auto* exit_node =
+ utils::GetOrCreate(info.var_exit_nodes, var, [&]() {
+ auto name =
+ builder_->Symbols().NameFor(var->Declaration()->symbol);
+ return CreateNode(name + "_value_" + info.type + "_exit");
+ });
+ exit_node->AddEdge(variables_.Get(var));
+ }
+ }
+ variables_.Pop();
+ }
+ previous_case_has_fallthrough = has_fallthrough;
+ }
+
+ // Update nodes for any variables assigned in the switch statement.
+ for (auto var : info.var_exit_nodes) {
+ variables_.Set(var.first, var.second);
+ }
+
+ return cf_end ? cf_end : cf;
+ },
+ [&](const ast::VariableDeclStatement* decl) {
+ Node* node;
+ if (decl->variable->constructor) {
+ auto [cf1, v] = ProcessExpression(cf, decl->variable->constructor);
+ cf = cf1;
+ node = v;
+ } else {
+ node = cf;
+ }
+ variables_.Set(sem_.Get(decl->variable), node);
+
+ if (!decl->variable->is_const) {
+ local_var_decls_.insert(sem_.Get<sem::LocalVariable>(decl->variable));
+ }
+
+ return cf;
+ },
+ [&](Default) {
+ TINT_ICE(Resolver, diagnostics_)
+ << "unknown statement type: " << std::string(stmt->TypeInfo().name);
+ return nullptr;
+ });
+ }
+
+ /// Process an identifier expression.
+ /// @param cf the input control flow node
+ /// @param ident the identifier expression to process
+ /// @returns a pair of (control flow node, value node)
+ std::pair<Node*, Node*> ProcessIdentExpression(Node* cf,
+ const ast::IdentifierExpression* ident) {
+ // Helper to check if the entry point attribute of `obj` indicates non-uniformity.
+ auto has_nonuniform_entry_point_attribute = [](auto* obj) {
+ // Only the num_workgroups and workgroup_id builtins are uniform.
+ if (auto* builtin = ast::GetAttribute<ast::BuiltinAttribute>(obj->attributes)) {
+ if (builtin->builtin == ast::Builtin::kNumWorkgroups ||
+ builtin->builtin == ast::Builtin::kWorkgroupId) {
+ return false;
+ }
+ }
+ return true;
+ };
+
+ auto name = builder_->Symbols().NameFor(ident->symbol);
+ auto* sem = sem_.Get<sem::VariableUser>(ident)->Variable();
+ return Switch(
+ sem,
+
+ [&](const sem::Parameter* param) {
+ auto* user_func = param->Owner()->As<sem::Function>();
+ if (user_func && user_func->Declaration()->IsEntryPoint()) {
+ if (auto* str = param->Type()->As<sem::Struct>()) {
+ // We consider the whole struct to be non-uniform if any one of its members
+ // is non-uniform.
+ bool uniform = true;
+ for (auto* member : str->Members()) {
+ if (has_nonuniform_entry_point_attribute(member->Declaration())) {
+ uniform = false;
+ }
+ }
+ return std::make_pair(cf, uniform ? cf : may_be_non_uniform_);
+ } else {
+ if (has_nonuniform_entry_point_attribute(param->Declaration())) {
+ return std::make_pair(cf, may_be_non_uniform_);
+ }
+ return std::make_pair(cf, cf);
+ }
+ } else {
+ auto* result = CreateNode(name + "_result");
+ auto* x = variables_.Get(param);
+ result->AddEdge(cf);
+ result->AddEdge(x);
+ return std::make_pair(cf, result);
+ }
+ },
+
+ [&](const sem::GlobalVariable* global) {
+ if (global->Declaration()->is_const || global->Access() == ast::Access::kRead) {
+ return std::make_pair(cf, cf);
+ } else {
+ return std::make_pair(cf, may_be_non_uniform_);
+ }
+ },
+
+ [&](const sem::LocalVariable* local) {
+ auto* result = CreateNode(name + "_result");
+ result->AddEdge(cf);
+ if (auto* x = variables_.Get(local)) {
+ result->AddEdge(x);
+ }
+ return std::make_pair(cf, result);
+ },
+
+ [&](Default) {
+ TINT_ICE(Resolver, diagnostics_)
+ << "unknown identifier expression type: " << std::string(sem->TypeInfo().name);
+ return std::pair<Node*, Node*>(nullptr, nullptr);
+ });
+ }
+
+ /// Process an expression.
+ /// @param cf the input control flow node
+ /// @param expr the expression to process
+ /// @returns a pair of (control flow node, value node)
+ std::pair<Node*, Node*> ProcessExpression(Node* cf, const ast::Expression* expr) {
+ return Switch(
+ expr,
+
+ [&](const ast::BinaryExpression* b) {
+ if (b->IsLogical()) {
+ // Short-circuiting binary operators are a special case.
+ auto [cf1, v1] = ProcessExpression(cf, b->lhs);
+ auto [cf2, v2] = ProcessExpression(v1, b->rhs);
+ return std::pair<Node*, Node*>(cf2, v2);
+ } else {
+ auto [cf1, v1] = ProcessExpression(cf, b->lhs);
+ auto [cf2, v2] = ProcessExpression(cf1, b->rhs);
+ auto* result = CreateNode("binary_expr_result");
+ result->AddEdge(v1);
+ result->AddEdge(v2);
+ return std::pair<Node*, Node*>(cf2, result);
+ }
+ },
+
+ [&](const ast::BitcastExpression* b) { return ProcessExpression(cf, b->expr); },
+
+ [&](const ast::CallExpression* c) { return ProcessCall(cf, c); },
+
+ [&](const ast::IdentifierExpression* i) { return ProcessIdentExpression(cf, i); },
+
+ [&](const ast::IndexAccessorExpression* i) {
+ auto [cf1, v1] = ProcessExpression(cf, i->object);
+ auto [cf2, v2] = ProcessExpression(cf1, i->index);
+ auto* result = CreateNode("index_accessor_result");
+ result->AddEdge(v1);
+ result->AddEdge(v2);
+ return std::pair<Node*, Node*>(cf2, result);
+ },
+
+ [&](const ast::LiteralExpression*) { return std::make_pair(cf, cf); },
+
+ [&](const ast::MemberAccessorExpression* m) {
+ return ProcessExpression(cf, m->structure);
+ },
+
+ [&](const ast::UnaryOpExpression* u) {
+ if (u->op == ast::UnaryOp::kIndirection) {
+ // Cut the analysis short, since we only need to know the originating variable
+ // which is being accessed.
+ auto* source_var = sem_.Get(u)->SourceVariable();
+ auto* value = variables_.Get(source_var);
+ if (!value) {
+ value = cf;
+ }
+ return std::pair<Node*, Node*>(cf, value);
+ }
+ return ProcessExpression(cf, u->expr);
+ },
+
+ [&](Default) {
+ TINT_ICE(Resolver, diagnostics_)
+ << "unknown expression type: " << std::string(expr->TypeInfo().name);
+ return std::pair<Node*, Node*>(nullptr, nullptr);
+ });
+ }
+
+ /// Process an LValue expression.
+ /// @param cf the input control flow node
+ /// @param expr the expression to process
+ /// @returns a pair of (control flow node, variable node)
+ std::pair<Node*, Node*> ProcessLValueExpression(Node* cf, const ast::Expression* expr) {
+ return Switch(
+ expr,
+
+ [&](const ast::IdentifierExpression* i) {
+ auto name = builder_->Symbols().NameFor(i->symbol);
+ auto* sem = sem_.Get<sem::VariableUser>(i);
+ if (sem->Variable()->Is<sem::GlobalVariable>()) {
+ return std::make_pair(cf, may_be_non_uniform_);
+ } else if (auto* local = sem->Variable()->As<sem::LocalVariable>()) {
+ // Create a new value node for this variable.
+ auto* value = CreateNode(name + "_lvalue");
+ auto* old_value = variables_.Set(local, value);
+
+ // Aggregate values link back to their previous value, as they can never become
+ // uniform again.
+ if (!local->Type()->UnwrapRef()->is_scalar() && old_value) {
+ value->AddEdge(old_value);
+ }
+
+ return std::make_pair(cf, value);
+ } else {
+ TINT_ICE(Resolver, diagnostics_)
+ << "unknown lvalue identifier expression type: "
+ << std::string(sem->Variable()->TypeInfo().name);
+ return std::pair<Node*, Node*>(nullptr, nullptr);
+ }
+ },
+
+ [&](const ast::IndexAccessorExpression* i) {
+ auto [cf1, l1] = ProcessLValueExpression(cf, i->object);
+ auto [cf2, v2] = ProcessExpression(cf1, i->index);
+ l1->AddEdge(v2);
+ return std::pair<Node*, Node*>(cf2, l1);
+ },
+
+ [&](const ast::MemberAccessorExpression* m) {
+ return ProcessLValueExpression(cf, m->structure);
+ },
+
+ [&](const ast::UnaryOpExpression* u) {
+ if (u->op == ast::UnaryOp::kIndirection) {
+ // Cut the analysis short, since we only need to know the originating variable
+ // that is being written to.
+ auto* source_var = sem_.Get(u)->SourceVariable();
+ auto name = builder_->Symbols().NameFor(source_var->Declaration()->symbol);
+ auto* deref = CreateNode(name + "_deref");
+ auto* old_value = variables_.Set(source_var, deref);
+
+ // Aggregate values link back to their previous value, as they can never become
+ // uniform again.
+ if (!source_var->Type()->UnwrapRef()->UnwrapPtr()->is_scalar() && old_value) {
+ deref->AddEdge(old_value);
+ }
+
+ return std::pair<Node*, Node*>(cf, deref);
+ }
+ return ProcessLValueExpression(cf, u->expr);
+ },
+
+ [&](Default) {
+ TINT_ICE(Resolver, diagnostics_)
+ << "unknown lvalue expression type: " << std::string(expr->TypeInfo().name);
+ return std::pair<Node*, Node*>(nullptr, nullptr);
+ });
+ }
+
+ /// Process a function call expression.
+ /// @param cf the input control flow node
+ /// @param call the function call to process
+ /// @returns a pair of (control flow node, value node)
+ std::pair<Node*, Node*> ProcessCall(Node* cf, const ast::CallExpression* call) {
+ std::string name;
+ if (call->target.name) {
+ name = builder_->Symbols().NameFor(call->target.name->symbol);
+ } else {
+ name = call->target.type->FriendlyName(builder_->Symbols());
+ }
+
+ // Process call arguments
+ Node* cf_last_arg = cf;
+ std::vector<Node*> args;
+ for (size_t i = 0; i < call->args.size(); i++) {
+ auto [cf_i, arg_i] = ProcessExpression(cf_last_arg, call->args[i]);
+
+ // Capture the index of this argument in a new node.
+ // Note: This is an additional node that isn't described in the specification, for the
+ // purpose of providing diagnostic information.
+ Node* arg_node = CreateNode(name + "_arg_" + std::to_string(i), call);
+ arg_node->arg_index = static_cast<uint32_t>(i);
+ arg_node->AddEdge(arg_i);
+
+ cf_last_arg = cf_i;
+ args.push_back(arg_node);
+ }
+
+ Node* result = CreateNode("Result_" + name);
+ Node* cf_after = CreateNode("CF_after_" + name, call);
+
+ // Get tags for the callee.
+ CallSiteTag callsite_tag = CallSiteNoRestriction;
+ FunctionTag function_tag = NoRestriction;
+ auto* sem = sem_.Get(call);
+ const FunctionInfo* func_info = nullptr;
+ Switch(
+ sem->Target(),
+ [&](const sem::Builtin* builtin) {
+ // Most builtins have no restrictions. The exceptions are barriers, derivatives, and
+ // some texture sampling builtins.
+ if (builtin->IsBarrier()) {
+ callsite_tag = CallSiteRequiredToBeUniform;
+ } else if (builtin->IsDerivative() ||
+ builtin->Type() == sem::BuiltinType::kTextureSample ||
+ builtin->Type() == sem::BuiltinType::kTextureSampleBias ||
+ builtin->Type() == sem::BuiltinType::kTextureSampleCompare) {
+ callsite_tag = CallSiteRequiredToBeUniform;
+ function_tag = ReturnValueMayBeNonUniform;
+ } else {
+ callsite_tag = CallSiteNoRestriction;
+ function_tag = NoRestriction;
+ }
+ },
+ [&](const sem::Function* func) {
+ // We must have already analyzed the user-defined function since we process
+ // functions in dependency order.
+ TINT_ASSERT(Resolver, functions_.count(func->Declaration()));
+ auto& info = functions_.at(func->Declaration());
+ callsite_tag = info.callsite_tag;
+ function_tag = info.function_tag;
+ func_info = &info;
+ },
+ [&](const sem::TypeConstructor*) {
+ callsite_tag = CallSiteNoRestriction;
+ function_tag = NoRestriction;
+ },
+ [&](const sem::TypeConversion*) {
+ callsite_tag = CallSiteNoRestriction;
+ function_tag = NoRestriction;
+ },
+ [&](Default) {
+ TINT_ICE(Resolver, diagnostics_) << "unhandled function call target: " << name;
+ });
+
+ if (callsite_tag == CallSiteRequiredToBeUniform) {
+ // Note: This deviates from the rules in the specification, which would add the edge
+ // directly to the incoming CF node. Going through CF_after instead makes it easier to
+ // produce diagnostics that can identify the function being called.
+ required_to_be_uniform_->AddEdge(cf_after);
+ }
+ cf_after->AddEdge(cf_last_arg);
+
+ if (function_tag == SubsequentControlFlowMayBeNonUniform) {
+ cf_after->AddEdge(may_be_non_uniform_);
+ } else if (function_tag == ReturnValueMayBeNonUniform) {
+ result->AddEdge(may_be_non_uniform_);
+ }
+
+ result->AddEdge(cf_after);
+
+ // For each argument, add edges based on parameter tags.
+ for (size_t i = 0; i < args.size(); i++) {
+ if (func_info) {
+ switch (func_info->parameters[i].tag) {
+ case ParameterRequiredToBeUniform:
+ required_to_be_uniform_->AddEdge(args[i]);
+ break;
+ case ParameterRequiredToBeUniformForSubsequentControlFlow:
+ cf_after->AddEdge(args[i]);
+ break;
+ case ParameterRequiredToBeUniformForReturnValue:
+ result->AddEdge(args[i]);
+ break;
+ case ParameterNoRestriction:
+ break;
+ }
+
+ auto* sem_arg = sem_.Get(call->args[i]);
+ if (sem_arg->Type()->Is<sem::Pointer>()) {
+ auto* ptr_result =
+ CreateNode(name + "_ptrarg_" + std::to_string(i) + "_result");
+ if (func_info->parameters[i].pointer_may_become_non_uniform) {
+ ptr_result->AddEdge(may_be_non_uniform_);
+ } else {
+ // Add edges from the resulting pointer value to any other arguments that
+ // feed it.
+ for (auto* source : func_info->parameters[i].pointer_param_output_sources) {
+ ptr_result->AddEdge(args[source->Index()]);
+ }
+ }
+
+ // Update the current stored value for this pointer argument.
+ auto* source_var = sem_arg->SourceVariable();
+ TINT_ASSERT(Resolver, source_var);
+ variables_.Set(source_var, ptr_result);
+ }
+ } else {
+ // All builtin function parameters are RequiredToBeUniformForReturnValue, as are
+ // parameters for type constructors and type conversions.
+ // The arrayLength() builtin is a special case, as there is currently no way for it
+ // to have a non-uniform return value.
+ auto* builtin = sem->Target()->As<sem::Builtin>();
+ if (!builtin || builtin->Type() != sem::BuiltinType::kArrayLength) {
+ result->AddEdge(args[i]);
+ }
+ }
+ }
+
+ return {cf_after, result};
+ }
+
+ /// Recursively traverse a graph starting at `node`, inserting all nodes that are reached into
+ /// `reachable`.
+ /// @param node the starting node
+ /// @param reachable the set of reachable nodes to populate
+ void Traverse(Node* node, utils::UniqueVector<Node*>& reachable) {
+ reachable.add(node);
+ for (auto* to : node->edges) {
+ if (to->visited_from == nullptr) {
+ to->visited_from = node;
+ Traverse(to, reachable);
+ }
+ }
+ }
+
+ /// Generate an error for a required_to_be_uniform->may_be_non_uniform path.
+ void MakeError() {
+ // Trace back to find a node that is required to be uniform that was reachable from a
+ // non-uniform value or control flow node.
+ Node* current = may_be_non_uniform_;
+ while (current) {
+ TINT_ASSERT(Resolver, current->visited_from);
+ if (current->visited_from == required_to_be_uniform_) {
+ break;
+ }
+ current = current->visited_from;
+ }
+
+ // The node will always have an corresponding call expression.
+ auto* call = current->ast->As<ast::CallExpression>();
+ TINT_ASSERT(Resolver, call);
+ auto* target = sem_.Get(call)->Target();
+
+ std::string name;
+ if (auto* builtin = target->As<sem::Builtin>()) {
+ name = builtin->str();
+ } else if (auto* user = target->As<sem::Function>()) {
+ name = builder_->Symbols().NameFor(user->Declaration()->symbol);
+ }
+
+ // TODO(jrprice): Switch to error instead of warning when feedback has settled.
+ if (current->arg_index != std::numeric_limits<uint32_t>::max()) {
+ // The requirement was on a function parameter.
+ auto param_name = builder_->Symbols().NameFor(
+ target->Parameters()[current->arg_index]->Declaration()->symbol);
+ diagnostics_.add_warning(
+ diag::System::Resolver,
+ "parameter '" + param_name + "' of '" + name + "' must be uniform",
+ call->args[current->arg_index]->source);
+ // TODO(jrprice): Show the reason why.
+ } else {
+ // The requirement was on a function callsite.
+ diagnostics_.add_warning(diag::System::Resolver,
+ "'" + name + "' must only be called from uniform control flow",
+ call->source);
+ // TODO(jrprice): Show full call stack to the problematic builtin.
+ }
+ }
+};
+
+} // namespace
+
+bool AnalyzeUniformity(ProgramBuilder* builder, const DependencyGraph& dependency_graph) {
+ if (builder->AST().Extensions().count(
+ ast::Enable::ExtensionKind::kChromiumDisableUniformityAnalysis)) {
+ return true;
+ }
+
+ UniformityGraph graph(builder);
+ return graph.Build(dependency_graph);
+}
+
+} // namespace tint::resolver
diff --git a/src/tint/resolver/uniformity.h b/src/tint/resolver/uniformity.h
new file mode 100644
index 0000000..39827cf
--- /dev/null
+++ b/src/tint/resolver/uniformity.h
@@ -0,0 +1,36 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_RESOLVER_UNIFORMITY_H_
+#define SRC_TINT_RESOLVER_UNIFORMITY_H_
+
+// Forward declarations.
+namespace tint {
+namespace resolver {
+struct DependencyGraph;
+} // namespace resolver
+class ProgramBuilder;
+} // namespace tint
+
+namespace tint::resolver {
+
+/// Analyze the uniformity of a program.
+/// @param builder the program to analyze
+/// @param dependency_graph the dependency-ordered module-scope declarations
+/// @returns true if there are no uniformity issues, false otherwise
+bool AnalyzeUniformity(ProgramBuilder* builder, const resolver::DependencyGraph& dependency_graph);
+
+} // namespace tint::resolver
+
+#endif // SRC_TINT_RESOLVER_UNIFORMITY_H_
diff --git a/src/tint/resolver/uniformity_test.cc b/src/tint/resolver/uniformity_test.cc
new file mode 100644
index 0000000..bb33f89
--- /dev/null
+++ b/src/tint/resolver/uniformity_test.cc
@@ -0,0 +1,4888 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <memory>
+#include <string>
+#include <tuple>
+#include <utility>
+
+#include "src/tint/program_builder.h"
+#include "src/tint/reader/wgsl/parser.h"
+#include "src/tint/resolver/uniformity.h"
+
+#include "gmock/gmock.h"
+#include "gtest/gtest.h"
+
+using namespace tint::number_suffixes; // NOLINT
+
+namespace tint::resolver {
+namespace {
+
+class UniformityAnalysisTestBase {
+ protected:
+ /// Parse and resolve a WGSL shader.
+ /// @param src the WGSL source code
+ /// @param should_pass true if `src` should pass the analysis, otherwise false
+ void RunTest(std::string src, bool should_pass) {
+ auto file = std::make_unique<Source::File>("test", src);
+ auto program = reader::wgsl::Parse(file.get());
+
+ diag::Formatter::Style style;
+ style.print_newline_at_end = false;
+ error_ = diag::Formatter(style).format(program.Diagnostics());
+
+ bool valid = program.IsValid();
+ if (should_pass) {
+ EXPECT_TRUE(valid) << error_;
+ if (program.Diagnostics().count() == 1u) {
+ EXPECT_THAT(program.Diagnostics().str(), ::testing::HasSubstr("unreachable"));
+ } else {
+ EXPECT_EQ(program.Diagnostics().count(), 0u) << error_;
+ }
+ } else {
+ // TODO(jrprice): expect false when uniformity issues become errors.
+ EXPECT_TRUE(valid) << error_;
+ }
+ }
+
+ /// Build and resolve a program from a ProgramBuilder object.
+ /// @param builder the program builder
+ /// @returns true on success, false on failure
+ bool RunTest(ProgramBuilder&& builder) {
+ auto program = Program(std::move(builder));
+
+ diag::Formatter::Style style;
+ style.print_newline_at_end = false;
+ error_ = diag::Formatter(style).format(program.Diagnostics());
+
+ return program.IsValid();
+ }
+
+ /// The error message from the parser or resolver, if any.
+ std::string error_;
+};
+
+class UniformityAnalysisTest : public UniformityAnalysisTestBase, public ::testing::Test {};
+
+class BasicTest : public UniformityAnalysisTestBase,
+ public ::testing::TestWithParam<std::tuple<int, int>> {
+ public:
+ /// Enum for the if-statement condition guarding a function call.
+ enum Condition {
+ // Uniform conditions:
+ kTrue,
+ kFalse,
+ kLiteral,
+ kModuleLet,
+ kPipelineOverridable,
+ kFuncLetUniformRhs,
+ kFuncVarUniform,
+ kFuncUniformRetVal,
+ kUniformBuffer,
+ kROStorageBuffer,
+ kLastUniformCondition = kROStorageBuffer,
+ // MayBeNonUniform conditions:
+ kFuncLetNonUniformRhs,
+ kFuncVarNonUniform,
+ kFuncNonUniformRetVal,
+ kRWStorageBuffer,
+ // End of range marker:
+ kEndOfConditionRange,
+ };
+
+ /// Enum for the function call statement.
+ enum Function {
+ // NoRestrictionFunctions:
+ kUserNoRestriction,
+ kMin,
+ kTextureSampleLevel,
+ kLastNoRestrictionFunction = kTextureSampleLevel,
+ // RequiredToBeUniform functions:
+ kUserRequiredToBeUniform,
+ kWorkgroupBarrier,
+ kStorageBarrier,
+ kTextureSample,
+ kTextureSampleBias,
+ kTextureSampleCompare,
+ kDpdx,
+ kDpdxCoarse,
+ kDpdxFine,
+ kDpdy,
+ kDpdyCoarse,
+ kDpdyFine,
+ kFwidth,
+ kFwidthCoarse,
+ kFwidthFine,
+ // End of range marker:
+ kEndOfFunctionRange,
+ };
+
+ /// Convert a condition to its string representation.
+ static std::string ConditionToStr(Condition c) {
+ switch (c) {
+ case kTrue:
+ return "true";
+ case kFalse:
+ return "false";
+ case kLiteral:
+ return "7 == 7";
+ case kModuleLet:
+ return "module_let == 0";
+ case kPipelineOverridable:
+ return "pipeline_overridable == 0";
+ case kFuncLetUniformRhs:
+ return "let_uniform_rhs == 0";
+ case kFuncVarUniform:
+ return "func_uniform == 0";
+ case kFuncUniformRetVal:
+ return "func_uniform_retval() == 0";
+ case kUniformBuffer:
+ return "u == 0";
+ case kROStorageBuffer:
+ return "ro == 0";
+ case kFuncLetNonUniformRhs:
+ return "let_nonuniform_rhs == 0";
+ case kFuncVarNonUniform:
+ return "func_non_uniform == 0";
+ case kFuncNonUniformRetVal:
+ return "func_nonuniform_retval() == 0";
+ case kRWStorageBuffer:
+ return "rw == 0";
+ case kEndOfConditionRange:
+ return "<invalid>";
+ }
+ return "<invalid>";
+ }
+
+ /// Convert a function call to its string representation.
+ static std::string FunctionToStr(Function f) {
+ switch (f) {
+ case kUserNoRestriction:
+ return "user_no_restriction()";
+ case kMin:
+ return "min(1, 1)";
+ case kTextureSampleLevel:
+ return "textureSampleLevel(t, s, vec2(0.5, 0.5), 0.0)";
+ case kUserRequiredToBeUniform:
+ return "user_required_to_be_uniform()";
+ case kWorkgroupBarrier:
+ return "workgroupBarrier()";
+ case kStorageBarrier:
+ return "storageBarrier()";
+ case kTextureSample:
+ return "textureSample(t, s, vec2(0.5, 0.5))";
+ case kTextureSampleBias:
+ return "textureSampleBias(t, s, vec2(0.5, 0.5), 2.0)";
+ case kTextureSampleCompare:
+ return "textureSampleCompare(td, sc, vec2(0.5, 0.5), 0.5)";
+ case kDpdx:
+ return "dpdx(1.0)";
+ case kDpdxCoarse:
+ return "dpdxCoarse(1.0)";
+ case kDpdxFine:
+ return "dpdxFine(1.0)";
+ case kDpdy:
+ return "dpdy(1.0)";
+ case kDpdyCoarse:
+ return "dpdyCoarse(1.0)";
+ case kDpdyFine:
+ return "dpdyFine(1.0)";
+ case kFwidth:
+ return "fwidth(1.0)";
+ case kFwidthCoarse:
+ return "fwidthCoarse(1.0)";
+ case kFwidthFine:
+ return "fwidthFine(1.0)";
+ case kEndOfFunctionRange:
+ return "<invalid>";
+ }
+ return "<invalid>";
+ }
+
+ /// @returns true if `c` is a condition that may be non-uniform.
+ static bool MayBeNonUniform(Condition c) { return c > kLastUniformCondition; }
+
+ /// @returns true if `f` is a function call that is required to be uniform.
+ static bool RequiredToBeUniform(Function f) { return f > kLastNoRestrictionFunction; }
+
+ /// Convert a test parameter pair of condition+function to a string that can be used as part of
+ /// a test name.
+ static std::string ParamsToName(::testing::TestParamInfo<ParamType> params) {
+ Condition c = static_cast<Condition>(std::get<0>(params.param));
+ Function f = static_cast<Function>(std::get<1>(params.param));
+ std::string name;
+#define CASE(c) \
+ case c: \
+ name += #c; \
+ break
+
+ switch (c) {
+ CASE(kTrue);
+ CASE(kFalse);
+ CASE(kLiteral);
+ CASE(kModuleLet);
+ CASE(kPipelineOverridable);
+ CASE(kFuncLetUniformRhs);
+ CASE(kFuncVarUniform);
+ CASE(kFuncUniformRetVal);
+ CASE(kUniformBuffer);
+ CASE(kROStorageBuffer);
+ CASE(kFuncLetNonUniformRhs);
+ CASE(kFuncVarNonUniform);
+ CASE(kFuncNonUniformRetVal);
+ CASE(kRWStorageBuffer);
+ case kEndOfConditionRange:
+ break;
+ }
+ name += "_";
+ switch (f) {
+ CASE(kUserNoRestriction);
+ CASE(kMin);
+ CASE(kTextureSampleLevel);
+ CASE(kUserRequiredToBeUniform);
+ CASE(kWorkgroupBarrier);
+ CASE(kStorageBarrier);
+ CASE(kTextureSample);
+ CASE(kTextureSampleBias);
+ CASE(kTextureSampleCompare);
+ CASE(kDpdx);
+ CASE(kDpdxCoarse);
+ CASE(kDpdxFine);
+ CASE(kDpdy);
+ CASE(kDpdyCoarse);
+ CASE(kDpdyFine);
+ CASE(kFwidth);
+ CASE(kFwidthCoarse);
+ CASE(kFwidthFine);
+ case kEndOfFunctionRange:
+ break;
+ }
+#undef CASE
+
+ return name;
+ }
+};
+
+// Test the uniformity constraints for a function call inside a conditional statement.
+TEST_P(BasicTest, ConditionalFunctionCall) {
+ auto condition = static_cast<Condition>(std::get<0>(GetParam()));
+ auto function = static_cast<Function>(std::get<1>(GetParam()));
+ auto src = R"(
+var<private> p : i32;
+var<workgroup> w : i32;
+@group(0) @binding(0) var<uniform> u : i32;
+@group(0) @binding(0) var<storage, read> ro : i32;
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+@group(1) @binding(0) var t : texture_2d<f32>;
+@group(1) @binding(1) var td : texture_depth_2d;
+@group(1) @binding(2) var s : sampler;
+@group(1) @binding(3) var sc : sampler_comparison;
+
+let module_let : i32 = 42;
+@id(42) override pipeline_overridable : i32;
+
+fn user_no_restriction() {}
+fn user_required_to_be_uniform() { workgroupBarrier(); }
+
+fn func_uniform_retval() -> i32 { return u; }
+fn func_nonuniform_retval() -> i32 { return rw; }
+
+fn foo() {
+ let let_uniform_rhs = 7;
+ let let_nonuniform_rhs = rw;
+
+ var func_uniform = 7;
+ var func_non_uniform = 7;
+ func_non_uniform = rw;
+
+ if ()" + ConditionToStr(condition) +
+ R"() {
+ )" + FunctionToStr(function) +
+ R"(;
+ }
+}
+)";
+
+ bool should_pass = !(MayBeNonUniform(condition) && RequiredToBeUniform(function));
+ RunTest(src, should_pass);
+ if (!should_pass) {
+ EXPECT_THAT(error_, ::testing::StartsWith("test:31:5 warning: "));
+ EXPECT_THAT(error_, ::testing::HasSubstr("must only be called from uniform control flow"));
+ }
+}
+
+INSTANTIATE_TEST_SUITE_P(
+ UniformityAnalysisTest,
+ BasicTest,
+ ::testing::Combine(::testing::Range<int>(0, BasicTest::kEndOfConditionRange),
+ ::testing::Range<int>(0, BasicTest::kEndOfFunctionRange)),
+ BasicTest::ParamsToName);
+
+////////////////////////////////////////////////////////////////////////////////
+/// Test specific function and parameter tags that are not tested above.
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(UniformityAnalysisTest, SubsequentControlFlowMayBeNonUniform_Pass) {
+ // Call a function that causes subsequent control flow to be non-uniform, and then call another
+ // function that doesn't require uniformity.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+var<private> p : i32;
+
+fn foo() {
+ if (rw == 0) {
+ p = 42;
+ return;
+ }
+ p = 5;
+ return;
+}
+
+fn bar() {
+ if (p == 42) {
+ p = 7;
+ }
+}
+
+fn main() {
+ foo();
+ bar();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, SubsequentControlFlowMayBeNonUniform_Fail) {
+ // Call a function that causes subsequent control flow to be non-uniform, and then call another
+ // function that requires uniformity.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+var<private> p : i32;
+
+fn foo() {
+ if (rw == 0) {
+ p = 42;
+ return;
+ }
+ p = 5;
+ return;
+}
+
+fn main() {
+ foo();
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:17:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ParameterNoRestriction_Pass) {
+ // Pass a non-uniform value as an argument, and then try to use the return value for
+ // control-flow guarding a barrier.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+var<private> p : i32;
+
+fn foo(i : i32) -> i32 {
+ if (i == 0) {
+ // This assignment is non-uniform, but shouldn't affect the return value.
+ p = 42;
+ }
+ return 7;
+}
+
+fn bar() {
+ let x = foo(rw);
+ if (x == 7) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ParameterRequiredToBeUniform_Pass) {
+ // Pass a uniform value as an argument to a function that uses that parameter for control-flow
+ // guarding a barrier.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> ro : i32;
+
+fn foo(i : i32) {
+ if (i == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn bar() {
+ foo(ro);
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ParameterRequiredToBeUniform_Fail) {
+ // Pass a non-uniform value as an argument to a function that uses that parameter for
+ // control-flow guarding a barrier.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo(i : i32) {
+ if (i == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn bar() {
+ foo(rw);
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:11:7 warning: parameter 'i' of 'foo' must be uniform
+ foo(rw);
+ ^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ParameterRequiredToBeUniformForReturnValue_Pass) {
+ // Pass a uniform value as an argument to a function that uses that parameter to produce the
+ // return value, and then use the return value for control-flow guarding a barrier.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> ro : i32;
+
+fn foo(i : i32) -> i32 {
+ return 1 + i;
+}
+
+fn bar() {
+ if (foo(ro) == 7) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ParameterRequiredToBeUniformForReturnValue_Fail) {
+ // Pass a non-uniform value as an argument to a function that uses that parameter to produce the
+ // return value, and then use the return value for control-flow guarding a barrier.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo(i : i32) -> i32 {
+ return 1 + i;
+}
+
+fn bar() {
+ if (foo(rw) == 7) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ParameterRequiredToBeUniformForSubsequentControlFlow_Pass) {
+ // Pass a uniform value as an argument to a function that uses that parameter return early, and
+ // then invoke a barrier after calling that function.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> ro : i32;
+
+var<private> p : i32;
+
+fn foo(i : i32) {
+ if (i == 0) {
+ p = 42;
+ return;
+ }
+ p = 5;
+ return;
+}
+
+fn bar() {
+ foo(ro);
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ParameterRequiredToBeUniformForSubsequentControlFlow_Fail) {
+ // Pass a non-uniform value as an argument to a function that uses that parameter return early,
+ // and then invoke a barrier after calling that function.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+var<private> p : i32;
+
+fn foo(i : i32) {
+ if (i == 0) {
+ p = 42;
+ return;
+ }
+ p = 5;
+ return;
+}
+
+fn bar() {
+ foo(rw);
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:17:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Test shader IO attributes.
+////////////////////////////////////////////////////////////////////////////////
+
+struct BuiltinEntry {
+ std::string name;
+ std::string type;
+ bool uniform;
+ BuiltinEntry(std::string n, std::string t, bool u) : name(n), type(t), uniform(u) {}
+};
+
+class ComputeBuiltin : public UniformityAnalysisTestBase,
+ public ::testing::TestWithParam<BuiltinEntry> {};
+TEST_P(ComputeBuiltin, AsParam) {
+ auto src = R"(
+@stage(compute) @workgroup_size(64)
+fn main(@builtin()" +
+ GetParam().name + R"() b : )" + GetParam().type + R"() {
+ if (all(vec3(b) == vec3(0u))) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ bool should_pass = GetParam().uniform;
+ RunTest(src, should_pass);
+ if (!should_pass) {
+ EXPECT_EQ(
+ error_,
+ R"(test:5:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+ }
+}
+
+TEST_P(ComputeBuiltin, InStruct) {
+ auto src = R"(
+struct S {
+ @builtin()" + GetParam().name +
+ R"() b : )" + GetParam().type + R"(
+}
+
+@stage(compute) @workgroup_size(64)
+fn main(s : S) {
+ if (all(vec3(s.b) == vec3(0u))) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ bool should_pass = GetParam().uniform;
+ RunTest(src, should_pass);
+ if (!should_pass) {
+ EXPECT_EQ(
+ error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+ }
+}
+
+INSTANTIATE_TEST_SUITE_P(UniformityAnalysisTest,
+ ComputeBuiltin,
+ ::testing::Values(BuiltinEntry{"local_invocation_id", "vec3<u32>", false},
+ BuiltinEntry{"local_invocation_index", "u32", false},
+ BuiltinEntry{"global_invocation_id", "vec3<u32>", false},
+ BuiltinEntry{"workgroup_id", "vec3<u32>", true},
+ BuiltinEntry{"num_workgroups", "vec3<u32>", true}),
+ [](const ::testing::TestParamInfo<ComputeBuiltin::ParamType>& p) {
+ return p.param.name;
+ });
+
+TEST_F(UniformityAnalysisTest, ComputeBuiltin_MixedAttributesInStruct) {
+ // Mix both non-uniform and uniform shader IO attributes in the same structure. Even accessing
+ // just uniform member causes non-uniformity in this case.
+ auto src = R"(
+struct S {
+ @builtin(num_workgroups) num_groups : vec3<u32>,
+ @builtin(local_invocation_index) idx : u32,
+}
+
+@stage(compute) @workgroup_size(64)
+fn main(s : S) {
+ if (s.num_groups.x == 0u) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+class FragmentBuiltin : public UniformityAnalysisTestBase,
+ public ::testing::TestWithParam<BuiltinEntry> {};
+TEST_P(FragmentBuiltin, AsParam) {
+ auto src = R"(
+@stage(fragment)
+fn main(@builtin()" +
+ GetParam().name + R"() b : )" + GetParam().type + R"() {
+ if (u32(vec4(b).x) == 0u) {
+ dpdx(0.5);
+ }
+}
+)";
+
+ bool should_pass = GetParam().uniform;
+ RunTest(src, should_pass);
+ if (!should_pass) {
+ EXPECT_EQ(error_,
+ R"(test:5:5 warning: 'dpdx' must only be called from uniform control flow
+ dpdx(0.5);
+ ^^^^
+)");
+ }
+}
+
+TEST_P(FragmentBuiltin, InStruct) {
+ auto src = R"(
+struct S {
+ @builtin()" + GetParam().name +
+ R"() b : )" + GetParam().type + R"(
+}
+
+@stage(fragment)
+fn main(s : S) {
+ if (u32(vec4(s.b).x) == 0u) {
+ dpdx(0.5);
+ }
+}
+)";
+
+ bool should_pass = GetParam().uniform;
+ RunTest(src, should_pass);
+ if (!should_pass) {
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'dpdx' must only be called from uniform control flow
+ dpdx(0.5);
+ ^^^^
+)");
+ }
+}
+
+INSTANTIATE_TEST_SUITE_P(UniformityAnalysisTest,
+ FragmentBuiltin,
+ ::testing::Values(BuiltinEntry{"position", "vec4<f32>", false},
+ BuiltinEntry{"front_facing", "bool", false},
+ BuiltinEntry{"sample_index", "u32", false},
+ BuiltinEntry{"sample_mask", "u32", false}),
+ [](const ::testing::TestParamInfo<FragmentBuiltin::ParamType>& p) {
+ return p.param.name;
+ });
+
+TEST_F(UniformityAnalysisTest, FragmentLocation) {
+ auto src = R"(
+@stage(fragment)
+fn main(@location(0) l : f32) {
+ if (l == 0.0) {
+ dpdx(0.5);
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:5:5 warning: 'dpdx' must only be called from uniform control flow
+ dpdx(0.5);
+ ^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, FragmentLocation_InStruct) {
+ auto src = R"(
+struct S {
+ @location(0) l : f32
+}
+
+@stage(fragment)
+fn main(s : S) {
+ if (s.l == 0.0) {
+ dpdx(0.5);
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'dpdx' must only be called from uniform control flow
+ dpdx(0.5);
+ ^^^^
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Test loop conditions and conditional break/continue statements.
+////////////////////////////////////////////////////////////////////////////////
+
+namespace LoopTest {
+
+enum ControlFlowInterrupt {
+ kBreak,
+ kContinue,
+ kReturn,
+ kDiscard,
+};
+enum Condition {
+ kNone,
+ kUniform,
+ kNonUniform,
+};
+
+using LoopTestParams = std::tuple<int, int>;
+
+static std::string ToStr(ControlFlowInterrupt interrupt) {
+ switch (interrupt) {
+ case kBreak:
+ return "break";
+ case kContinue:
+ return "continue";
+ case kReturn:
+ return "return";
+ case kDiscard:
+ return "discard";
+ }
+ return "";
+}
+
+static std::string ToStr(Condition condition) {
+ switch (condition) {
+ case kNone:
+ return "uncondtiional";
+ case kUniform:
+ return "uniform";
+ case kNonUniform:
+ return "nonuniform";
+ }
+ return "";
+}
+
+class LoopTest : public UniformityAnalysisTestBase,
+ public ::testing::TestWithParam<LoopTestParams> {
+ protected:
+ std::string MakeInterrupt(ControlFlowInterrupt interrupt, Condition condition) {
+ switch (condition) {
+ case kNone:
+ return ToStr(interrupt);
+ case kUniform:
+ return "if (uniform_var == 42) { " + ToStr(interrupt) + "; }";
+ case kNonUniform:
+ return "if (nonuniform_var == 42) { " + ToStr(interrupt) + "; }";
+ }
+ return "<invalid>";
+ }
+};
+
+INSTANTIATE_TEST_SUITE_P(UniformityAnalysisTest,
+ LoopTest,
+ ::testing::Combine(::testing::Range<int>(0, kDiscard + 1),
+ ::testing::Range<int>(0, kNonUniform + 1)),
+ [](const ::testing::TestParamInfo<LoopTestParams>& p) {
+ ControlFlowInterrupt interrupt =
+ static_cast<ControlFlowInterrupt>(std::get<0>(p.param));
+ auto condition = static_cast<Condition>(std::get<1>(p.param));
+ return ToStr(interrupt) + "_" + ToStr(condition);
+ });
+
+TEST_P(LoopTest, CallInBody_InterruptAfter) {
+ // Test control-flow interrupt in a loop after a function call that requires uniform control
+ // flow.
+ auto interrupt = static_cast<ControlFlowInterrupt>(std::get<0>(GetParam()));
+ auto condition = static_cast<Condition>(std::get<1>(GetParam()));
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> uniform_var : i32;
+@group(0) @binding(0) var<storage, read_write> nonuniform_var : i32;
+
+fn foo() {
+ loop {
+ // Pretend that this isn't an infinite loop, in case the interrupt is a
+ // continue statement.
+ if (false) {
+ break;
+ }
+
+ workgroupBarrier();
+ )" + MakeInterrupt(interrupt, condition) +
+ R"(;
+ }
+}
+)";
+
+ if (condition == kNonUniform) {
+ RunTest(src, false);
+ EXPECT_EQ(
+ error_,
+ R"(test:13:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+ } else {
+ RunTest(src, true);
+ }
+}
+
+TEST_P(LoopTest, CallInBody_InterruptBefore) {
+ // Test control-flow interrupt in a loop before a function call that requires uniform control
+ // flow.
+ auto interrupt = static_cast<ControlFlowInterrupt>(std::get<0>(GetParam()));
+ auto condition = static_cast<Condition>(std::get<1>(GetParam()));
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> uniform_var : i32;
+@group(0) @binding(0) var<storage, read_write> nonuniform_var : i32;
+
+fn foo() {
+ loop {
+ // Pretend that this isn't an infinite loop, in case the interrupt is a
+ // continue statement.
+ if (false) {
+ break;
+ }
+
+ )" + MakeInterrupt(interrupt, condition) +
+ R"(;
+ workgroupBarrier();
+ }
+}
+)";
+
+ if (condition == kNonUniform) {
+ RunTest(src, false);
+ EXPECT_EQ(
+ error_,
+ R"(test:14:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+ } else {
+ RunTest(src, true);
+ }
+}
+
+TEST_P(LoopTest, CallInContinuing_InterruptInBody) {
+ // Test control-flow interrupt in a loop with a function call that requires uniform control flow
+ // in the continuing statement.
+ auto interrupt = static_cast<ControlFlowInterrupt>(std::get<0>(GetParam()));
+ auto condition = static_cast<Condition>(std::get<1>(GetParam()));
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> uniform_var : i32;
+@group(0) @binding(0) var<storage, read_write> nonuniform_var : i32;
+
+fn foo() {
+ loop {
+ // Pretend that this isn't an infinite loop, in case the interrupt is a
+ // continue statement.
+ if (false) {
+ break;
+ }
+
+ )" + MakeInterrupt(interrupt, condition) +
+ R"(;
+ continuing {
+ workgroupBarrier();
+ }
+ }
+}
+)";
+
+ if (condition == kNonUniform) {
+ RunTest(src, false);
+ EXPECT_EQ(
+ error_,
+ R"(test:15:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+ } else {
+ RunTest(src, true);
+ }
+}
+
+TEST_F(UniformityAnalysisTest, Loop_CallInBody_UniformBreakInContinuing) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> n : i32;
+
+fn foo() {
+ var i = 0;
+ loop {
+ workgroupBarrier();
+ continuing {
+ i = i + 1;
+ if (i == n) {
+ break;
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_CallInBody_NonUniformBreakInContinuing) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn foo() {
+ var i = 0;
+ loop {
+ workgroupBarrier();
+ continuing {
+ i = i + 1;
+ if (i == n) {
+ break;
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:7:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_CallInContinuing_UniformBreakInContinuing) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> n : i32;
+
+fn foo() {
+ var i = 0;
+ loop {
+ continuing {
+ workgroupBarrier();
+ i = i + 1;
+ if (i == n) {
+ break;
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_CallInContinuing_NonUniformBreakInContinuing) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn foo() {
+ var i = 0;
+ loop {
+ continuing {
+ workgroupBarrier();
+ i = i + 1;
+ if (i == n) {
+ break;
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+class LoopDeadCodeTest : public UniformityAnalysisTestBase, public ::testing::TestWithParam<int> {};
+
+INSTANTIATE_TEST_SUITE_P(UniformityAnalysisTest,
+ LoopDeadCodeTest,
+ ::testing::Range<int>(0, kDiscard + 1),
+ [](const ::testing::TestParamInfo<LoopDeadCodeTest::ParamType>& p) {
+ return ToStr(static_cast<ControlFlowInterrupt>(p.param));
+ });
+
+TEST_P(LoopDeadCodeTest, AfterInterrupt) {
+ // Dead code after a control-flow interrupt in a loop shouldn't cause uniformity errors.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn foo() {
+ loop {
+ )" + ToStr(static_cast<ControlFlowInterrupt>(GetParam())) +
+ R"(;
+ if (n == 42) {
+ workgroupBarrier();
+ }
+ continuing {
+ // Pretend that this isn't an infinite loop, in case the interrupt is a
+ // continue statement.
+ if (false) {
+ break;
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_VarBecomesNonUniformInLoopAfterBarrier) {
+ // Use a variable for a conditional barrier in a loop, and then assign a non-uniform value to
+ // that variable later in that loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ v = non_uniform;
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_VarBecomesNonUniformInLoopAfterBarrier_BreakAtEnd) {
+ // Use a variable for a conditional barrier in a loop, and then assign a non-uniform value to
+ // that variable later in that loop. End the loop with a break statement to prevent the
+ // non-uniform value from causing an issue.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (v == 0) {
+ workgroupBarrier();
+ }
+
+ v = non_uniform;
+ break;
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_ConditionalAssignNonUniformWithBreak_BarrierInLoop) {
+ // In a conditional block, assign a non-uniform value and then break, then use a variable for a
+ // conditional barrier later in the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (true) {
+ v = non_uniform;
+ break;
+ }
+ if (v == 0) {
+ workgroupBarrier();
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_ConditionalAssignNonUniformWithConditionalBreak_BarrierInLoop) {
+ // In a conditional block, assign a non-uniform value and then conditionally break, then use a
+ // variable for a conditional barrier later in the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (true) {
+ v = non_uniform;
+ if (true) {
+ break;
+ }
+ }
+ if (v == 0) {
+ workgroupBarrier();
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_ConditionalAssignNonUniformWithBreak_BarrierAfterLoop) {
+ // In a conditional block, assign a non-uniform value and then break, then use a variable for a
+ // conditional barrier after the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (true) {
+ v = non_uniform;
+ break;
+ }
+ v = 5;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:15:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_VarBecomesUniformBeforeSomeExits_BarrierAfterLoop) {
+ // Assign a non-uniform value, have two exit points only one of which assigns a uniform value,
+ // then use a variable for a conditional barrier after the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (true) {
+ break;
+ }
+
+ v = non_uniform;
+
+ if (false) {
+ v = 6;
+ break;
+ }
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:20:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_VarBecomesUniformBeforeAllExits_BarrierAfterLoop) {
+ // Assign a non-uniform value, have two exit points both of which assigns a uniform value,
+ // then use a variable for a conditional barrier after the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (true) {
+ v = 5;
+ break;
+ }
+
+ v = non_uniform;
+
+ if (false) {
+ v = 6;
+ break;
+ }
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_AssignNonUniformBeforeConditionalBreak_BarrierAfterLoop) {
+ // Assign a non-uniform value and then break in a conditional block, then use a variable for a
+ // conditional barrier after the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ v = non_uniform;
+ if (true) {
+ if (false) {
+ v = 5;
+ } else {
+ break;
+ }
+ v = 5;
+ }
+ v = 5;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:20:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_VarBecomesNonUniformBeforeConditionalContinue_BarrierAtStart) {
+ // Use a variable for a conditional barrier in a loop, assign a non-uniform value to
+ // that variable later in that loop, then perform a conditional continue before assigning a
+ // uniform value to that variable.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ v = non_uniform;
+ if (true) {
+ continue;
+ }
+
+ v = 5;
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ Loop_VarBecomesUniformBeforeConditionalContinue_BarrierInContinuing) {
+ // Use a variable for a conditional barrier in the continuing statement of a loop, assign a
+ // non-uniform value to that variable later in that loop, then conditionally assign a uniform
+ // value before continuing.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ v = non_uniform;
+
+ if (false) {
+ v = 5;
+ continue;
+ }
+
+ continuing {
+ if (v == 0) {
+ workgroupBarrier();
+ }
+ if (true) {
+ break;
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:16:9 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_VarBecomesNonUniformBeforeConditionalContinue) {
+ // Use a variable for a conditional barrier in a loop, assign a non-uniform value to
+ // that variable later in that loop, then perform a conditional continue before assigning a
+ // uniform value to that variable.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ v = non_uniform;
+ if (true) {
+ continue;
+ }
+
+ v = 5;
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Loop_VarBecomesNonUniformInNestedLoopWithBreak_BarrierInLoop) {
+ // Use a variable for a conditional barrier in a loop, then conditionally assign a non-uniform
+ // value to that variable followed by a break in a nested loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ loop {
+ if (true) {
+ v = non_uniform;
+ break;
+ }
+ v = 5;
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ Loop_VarBecomesNonUniformInNestedLoopWithBreak_BecomesUniformAgain_BarrierAfterLoop) {
+ // Conditionally assign a non-uniform value followed by a break in a nested loop, assign a
+ // uniform value in the outer loop, and then use a variable for a conditional barrier after the
+ // loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ loop {
+ if (false) {
+ break;
+ }
+
+ loop {
+ if (true) {
+ v = non_uniform;
+ break;
+ }
+ }
+ v = 5;
+ }
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_NonUniformValueNeverReachesContinuing) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ loop {
+ var v = non_uniform;
+ return;
+
+ continuing {
+ if (v == 0) {
+ workgroupBarrier();
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_NonUniformBreakInBody_Reconverge) {
+ // Loops reconverge at exit, so test that we can call workgroupBarrier() after a loop that
+ // contains a non-uniform conditional break.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn foo() {
+ var i = 0;
+ loop {
+ if (i == n) {
+ break;
+ }
+ i = i + 1;
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_NonUniformFunctionInBody_Reconverge) {
+ // Loops reconverge at exit, so test that we can call workgroupBarrier() after a loop that
+ // contains a call to a function that causes non-uniform control flow.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn bar() {
+ if (n == 42) {
+ return;
+ } else {
+ return;
+ }
+}
+
+fn foo() {
+ loop {
+ bar();
+ break;
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Loop_NonUniformFunctionDiscard_NoReconvergence) {
+ // Loops should not reconverge after non-uniform discard statements.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn bar() {
+ if (n == 42) {
+ discard;
+ }
+}
+
+fn foo() {
+ loop {
+ bar();
+ break;
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:15:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_CallInside_UniformCondition) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> n : i32;
+
+fn foo() {
+ for (var i = 0; i < n; i = i + 1) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_CallInside_NonUniformCondition) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn foo() {
+ for (var i = 0; i < n; i = i + 1) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_CallInside_InitializerCausesNonUniformFlow) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn bar() -> i32 {
+ if (n == 42) {
+ return 1;
+ } else {
+ return 2;
+ }
+}
+
+fn foo() {
+ for (var i = bar(); i < 10; i = i + 1) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_CallInside_ContinuingCausesNonUniformFlow) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn bar() -> i32 {
+ if (n == 42) {
+ return 1;
+ } else {
+ return 2;
+ }
+}
+
+fn foo() {
+ for (var i = 0; i < 10; i = i + bar()) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_VarBecomesNonUniformInContinuing_BarrierInLoop) {
+ // Use a variable for a conditional barrier in a loop, and then assign a non-uniform value to
+ // that variable in the continuing statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; v = non_uniform) {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_VarBecomesUniformInContinuing_BarrierInLoop) {
+ // Use a variable for a conditional barrier in a loop, and then assign a uniform value to that
+ // variable in the continuing statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; v = 5) {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ v = non_uniform;
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_VarBecomesNonUniformInContinuing_BarrierAfterLoop) {
+ // Use a variable for a conditional barrier after a loop, and assign a non-uniform value to
+ // that variable in the continuing statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; v = non_uniform) {
+ v = 5;
+ }
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_VarBecomesUniformInContinuing_BarrierAfterLoop) {
+ // Use a variable for a conditional barrier after a loop, and assign a uniform value to that
+ // variable in the continuing statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; v = 5) {
+ v = non_uniform;
+ }
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_VarBecomesNonUniformInLoopAfterBarrier) {
+ // Use a variable for a conditional barrier in a loop, and then assign a non-uniform value to
+ // that variable later in that loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; i++) {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ v = non_uniform;
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_ConditionalAssignNonUniformWithBreak_BarrierInLoop) {
+ // In a conditional block, assign a non-uniform value and then break, then use a variable for a
+ // conditional barrier later in the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; i++) {
+ if (true) {
+ v = non_uniform;
+ break;
+ }
+ if (v == 0) {
+ workgroupBarrier();
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_ConditionalAssignNonUniformWithBreak_BarrierAfterLoop) {
+ // In a conditional block, assign a non-uniform value and then break, then use a variable for a
+ // conditional barrier after the loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; i++) {
+ if (true) {
+ v = non_uniform;
+ break;
+ }
+ v = 5;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:15:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_VarRemainsNonUniformAtLoopEnd_BarrierAfterLoop) {
+ // Assign a non-uniform value, assign a uniform value before all explicit break points but leave
+ // the value non-uniform at loop exit, then use a variable for a conditional barrier after the
+ // loop.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; i++) {
+ if (true) {
+ v = 5;
+ break;
+ }
+
+ v = non_uniform;
+
+ if (true) {
+ v = 6;
+ break;
+ }
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:21:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ ForLoop_VarBecomesNonUniformBeforeConditionalContinue_BarrierAtStart) {
+ // Use a variable for a conditional barrier in a loop, assign a non-uniform value to
+ // that variable later in that loop, then perform a conditional continue before assigning a
+ // uniform value to that variable.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; i++) {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ v = non_uniform;
+ if (true) {
+ continue;
+ }
+
+ v = 5;
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_VarBecomesNonUniformBeforeConditionalContinue) {
+ // Use a variable for a conditional barrier in a loop, assign a non-uniform value to
+ // that variable later in that loop, then perform a conditional continue before assigning a
+ // uniform value to that variable.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ for (var i = 0; i < 10; i++) {
+ if (v == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ v = non_uniform;
+ if (true) {
+ continue;
+ }
+
+ v = 5;
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ForLoop_NonUniformCondition_Reconverge) {
+ // Loops reconverge at exit, so test that we can call workgroupBarrier() after a loop that has a
+ // non-uniform condition.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn foo() {
+ for (var i = 0; i < n; i = i + 1) {
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+} // namespace LoopTest
+
+////////////////////////////////////////////////////////////////////////////////
+/// If-else statement tests.
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(UniformityAnalysisTest, IfElse_UniformCondition_BarrierInTrueBlock) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> uniform_global : i32;
+
+fn foo() {
+ if (uniform_global == 42) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_UniformCondition_BarrierInElseBlock) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> uniform_global : i32;
+
+fn foo() {
+ if (uniform_global == 42) {
+ } else {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_UniformCondition_BarrierInElseIfBlock) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> uniform_global : i32;
+
+fn foo() {
+ if (uniform_global == 42) {
+ } else if (true) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_NonUniformCondition_BarrierInTrueBlock) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (non_uniform == 42) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_NonUniformCondition_BarrierInElseBlock) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (non_uniform == 42) {
+ } else {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:7:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformLHS_And) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+var<private> p : i32;
+
+fn main() {
+ if ((non_uniform_global == 42) && false) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformRHS_And) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+var<private> p : i32;
+
+fn main() {
+ if (false && (non_uniform_global == 42)) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformLHS_Or) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+var<private> p : i32;
+
+fn main() {
+ if ((non_uniform_global == 42) || true) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformRHS_Or) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+var<private> p : i32;
+
+fn main() {
+ if (true || (non_uniform_global == 42)) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_NonUniformCondition_BarrierInElseIfBlock) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (non_uniform == 42) {
+ } else if (true) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:7:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_VarBecomesNonUniform_BeforeCondition) {
+ // Use a function-scope variable for control-flow guarding a barrier, and then assign to that
+ // variable before checking the condition.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v = 0;
+ v = rw;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_VarBecomesNonUniform_AfterCondition) {
+ // Use a function-scope variable for control-flow guarding a barrier, and then assign to that
+ // variable after checking the condition.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v = 0;
+ if (v == 0) {
+ v = rw;
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_VarBecomesNonUniformInIf_BarrierInElse) {
+ // Assign a non-uniform value to a variable in an if-block, and then use that variable for a
+ // conditional barrier in the else block.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ if (true) {
+ v = non_uniform;
+ } else {
+ if (v == 0) {
+ workgroupBarrier();
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_AssignNonUniformInIf_AssignUniformInElse) {
+ // Assign a non-uniform value to a variable in an if-block and a uniform value in the else
+ // block, and then use that variable for a conditional barrier after the if-else statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ if (true) {
+ if (true) {
+ v = non_uniform;
+ } else {
+ v = 5;
+ }
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:15:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_AssignNonUniformInIfWithReturn) {
+ // Assign a non-uniform value to a variable in an if-block followed by a return, and then use
+ // that variable for a conditional barrier after the if-else statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ if (true) {
+ v = non_uniform;
+ return;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_AssignNonUniformBeforeIf_BothBranchesAssignUniform) {
+ // Assign a non-uniform value to a variable before and if-else statement, assign uniform values
+ // in both branch of the if-else, and then use that variable for a conditional barrier after
+ // the if-else statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ v = non_uniform;
+ if (true) {
+ v = 5;
+ } else {
+ v = 6;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_AssignNonUniformBeforeIf_OnlyTrueBranchAssignsUniform) {
+ // Assign a non-uniform value to a variable before and if-else statement, assign a uniform value
+ // in the true branch of the if-else, and then use that variable for a conditional barrier after
+ // the if-else statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ v = non_uniform;
+ if (true) {
+ v = 5;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:12:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_AssignNonUniformBeforeIf_OnlyFalseBranchAssignsUniform) {
+ // Assign a non-uniform value to a variable before and if-else statement, assign a uniform value
+ // in the false branch of the if-else, and then use that variable for a conditional barrier
+ // after the if-else statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ v = non_uniform;
+ if (true) {
+ } else {
+ v = 5;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:13:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ IfElse_AssignNonUniformBeforeIf_OnlyTrueBranchAssignsUniform_FalseBranchReturns) {
+ // Assign a non-uniform value to a variable before and if-else statement, assign a uniform value
+ // in the true branch of the if-else, leave the variable untouched in the false branch and just
+ // return, and then use that variable for a conditional barrier after the if-else statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ v = non_uniform;
+ if (true) {
+ v = 5;
+ } else {
+ return;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest,
+ IfElse_AssignNonUniformBeforeIf_OnlyFalseBranchAssignsUniform_TrueBranchReturns) {
+ // Assign a non-uniform value to a variable before and if-else statement, assign a uniform value
+ // in the false branch of the if-else, leave the variable untouched in the true branch and just
+ // return, and then use that variable for a conditional barrier after the if-else statement.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ v = non_uniform;
+ if (true) {
+ return;
+ } else {
+ v = 5;
+ }
+
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_NonUniformCondition_Reconverge) {
+ // If statements reconverge at exit, so test that we can call workgroupBarrier() after an if
+ // statement with a non-uniform condition.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (non_uniform == 42) {
+ } else {
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingNonUniformConditionLHS_Reconverge) {
+ // If statements reconverge at exit, so test that we can call workgroupBarrier() after an if
+ // statement with a non-uniform condition that uses short-circuiting.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (non_uniform == 42 || true) {
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingNonUniformConditionRHS_Reconverge) {
+ // If statements reconverge at exit, so test that we can call workgroupBarrier() after an if
+ // statement with a non-uniform condition that uses short-circuiting.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (false && non_uniform == 42) {
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_NonUniformFunctionCall_Reconverge) {
+ // If statements reconverge at exit, so test that we can call workgroupBarrier() after an if
+ // statement with a non-uniform condition.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar() {
+ if (non_uniform == 42) {
+ return;
+ } else {
+ return;
+ }
+}
+
+fn foo() {
+ if (non_uniform == 42) {
+ bar();
+ } else {
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_NonUniformReturn_NoReconverge) {
+ // If statements should not reconverge after non-uniform returns.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (non_uniform == 42) {
+ return;
+ } else {
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, IfElse_NonUniformDiscard_NoReconverge) {
+ // If statements should not reconverge after non-uniform discards.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ if (non_uniform == 42) {
+ discard;
+ } else {
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Switch statement tests.
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformCondition_BarrierInCase) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ switch (non_uniform) {
+ case 42: {
+ workgroupBarrier();
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:7:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformCondition_BarrierInDefault) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ switch (non_uniform) {
+ default: {
+ workgroupBarrier();
+ break;
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:7:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformBreak) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ switch (condition) {
+ case 42: {
+ if (non_uniform == 42) {
+ break;
+ }
+ workgroupBarrier();
+ }
+ default: {
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:11:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformBreakInDifferentCase) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ switch (condition) {
+ case 0: {
+ if (non_uniform == 42) {
+ break;
+ }
+ }
+ case 42: {
+ workgroupBarrier();
+ }
+ default: {
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformBreakInDifferentCase_Fallthrough) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ switch (condition) {
+ case 0: {
+ if (non_uniform == 42) {
+ break;
+ }
+ fallthrough;
+ }
+ case 42: {
+ workgroupBarrier();
+ }
+ default: {
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesNonUniformInDifferentCase_WithBreak) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = 0;
+ switch (condition) {
+ case 0: {
+ x = non_uniform;
+ break;
+ }
+ case 42: {
+ if (x == 0) {
+ workgroupBarrier();
+ }
+ }
+ default: {
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesNonUniformInDifferentCase_WithFallthrough) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = 0;
+ switch (condition) {
+ case 0: {
+ x = non_uniform;
+ fallthrough;
+ }
+ case 42: {
+ if (x == 0) {
+ workgroupBarrier();
+ }
+ }
+ default: {
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:9 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInDifferentCase_WithBreak) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = non_uniform;
+ switch (condition) {
+ case 0: {
+ x = 5;
+ break;
+ }
+ case 42: {
+ if (x == 0) {
+ workgroupBarrier();
+ }
+ }
+ default: {
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:9 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInDifferentCase_WithFallthrough) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = non_uniform;
+ switch (condition) {
+ case 0: {
+ x = 5;
+ fallthrough;
+ }
+ case 42: {
+ if (x == 0) {
+ workgroupBarrier();
+ }
+ }
+ default: {
+ }
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesNonUniformInCase_BarrierAfter) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = 0;
+ switch (condition) {
+ case 0: {
+ x = non_uniform;
+ }
+ case 42: {
+ x = 5;
+ }
+ default: {
+ x = 6;
+ }
+ }
+ if (x == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:19:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInAllCases_BarrierAfter) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = non_uniform;
+ switch (condition) {
+ case 0: {
+ x = 4;
+ }
+ case 42: {
+ x = 5;
+ }
+ default: {
+ x = 6;
+ }
+ }
+ if (x == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInSomeCases_BarrierAfter) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = non_uniform;
+ switch (condition) {
+ case 0: {
+ x = 4;
+ }
+ case 42: {
+ }
+ default: {
+ x = 6;
+ }
+ }
+ if (x == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:18:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInCasesThatDontReturn_BarrierAfter) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = non_uniform;
+ switch (condition) {
+ case 0: {
+ x = 4;
+ }
+ case 42: {
+ return;
+ }
+ default: {
+ x = 6;
+ }
+ }
+ if (x == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformAfterConditionalBreak_BarrierAfter) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = non_uniform;
+ switch (condition) {
+ case 0: {
+ x = 4;
+ }
+ case 42: {
+ }
+ default: {
+ if (false) {
+ break;
+ }
+ x = 6;
+ }
+ }
+ if (x == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:21:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NestedInLoop_VarBecomesNonUniformWithBreak_BarrierInLoop) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = 0;
+ loop {
+ if (x == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ switch (condition) {
+ case 0: {
+ x = non_uniform;
+ break;
+ }
+ default: {
+ x = 6;
+ }
+ }
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NestedInLoop_VarBecomesNonUniformWithBreak_BarrierAfterLoop) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(0) var<uniform> condition : i32;
+
+fn foo() {
+ var x = 0;
+ loop {
+ if (false) {
+ break;
+ }
+ switch (condition) {
+ case 0: {
+ x = non_uniform;
+ break;
+ }
+ default: {
+ x = 6;
+ }
+ }
+ x = 5;
+ }
+ if (x == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformCondition_Reconverge) {
+ // Switch statements reconverge at exit, so test that we can call workgroupBarrier() after a
+ // switch statement that contains a non-uniform conditional break.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ switch (non_uniform) {
+ default: {
+ break;
+ }
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformBreak_Reconverge) {
+ // Switch statements reconverge at exit, so test that we can call workgroupBarrier() after a
+ // switch statement that contains a non-uniform conditional break.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ switch (42) {
+ default: {
+ if (non_uniform == 0) {
+ break;
+ }
+ break;
+ }
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformFunctionCall_Reconverge) {
+ // Switch statements reconverge at exit, so test that we can call workgroupBarrier() after a
+ // switch statement that contains a call to a function that causes non-uniform control flow.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn bar() {
+ if (n == 42) {
+ return;
+ } else {
+ return;
+ }
+}
+
+fn foo() {
+ switch (42) {
+ default: {
+ bar();
+ break;
+ }
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, Switch_NonUniformFunctionDiscard_NoReconvergence) {
+ // Switch statements should not reconverge after non-uniform discards.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> n : i32;
+
+fn bar() {
+ if (n == 42) {
+ discard;
+ }
+}
+
+fn foo() {
+ switch (42) {
+ default: {
+ bar();
+ break;
+ }
+ }
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:17:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Pointer tests.
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThroughPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ *&v = non_uniform;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThroughCapturedPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ let pv = &v;
+ *pv = non_uniform;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThroughPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = non_uniform;
+ *&v = 42;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThroughCapturedPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = non_uniform;
+ let pv = &v;
+ *pv = 42;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThroughCapturedPointer_InNonUniformControlFlow) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ let pv = &v;
+ if (non_uniform == 0) {
+ *pv = 42;
+ }
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:11:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, LoadNonUniformThroughPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = non_uniform;
+ if (*&v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:7:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, LoadNonUniformThroughCapturedPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = non_uniform;
+ let pv = &v;
+ if (*pv == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, LoadUniformThroughPointer) {
+ auto src = R"(
+fn foo() {
+ var v = 42;
+ if (*&v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, LoadUniformThroughCapturedPointer) {
+ auto src = R"(
+fn foo() {
+ var v = 42;
+ let pv = &v;
+ if (*pv == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, StoreNonUniformAfterCapturingPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ let pv = &v;
+ v = non_uniform;
+ if (*pv == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, StoreUniformAfterCapturingPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = non_uniform;
+ let pv = &v;
+ v = 42;
+ if (*pv == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThroughLongChainOfPointers) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ let pv1 = &*&v;
+ let pv2 = &*&*pv1;
+ *&*&*pv2 = non_uniform;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, LoadNonUniformThroughLongChainOfPointers) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = non_uniform;
+ let pv1 = &*&v;
+ let pv2 = &*&*pv1;
+ if (*&*&*pv2 == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThenNonUniformThroughDifferentPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ let pv1 = &v;
+ let pv2 = &v;
+ *pv1 = 42;
+ *pv2 = non_uniform;
+ if (*pv1 == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:11:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThenUniformThroughDifferentPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ var v = 0;
+ let pv1 = &v;
+ let pv2 = &v;
+ *pv1 = non_uniform;
+ *pv2 = 42;
+ if (*pv1 == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, UnmodifiedPointerParameterNonUniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+}
+
+fn foo() {
+ var v = non_uniform;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:11:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, UnmodifiedPointerParameterUniform) {
+ auto src = R"(
+fn bar(p : ptr<function, i32>) {
+}
+
+fn foo() {
+ var v = 42;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThroughPointerInFunctionCall) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+ *p = non_uniform;
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:12:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThroughPointerInFunctionCall) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+ *p = 42;
+}
+
+fn foo() {
+ var v = non_uniform;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThroughPointerInFunctionCallViaArg) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>, a : i32) {
+ *p = a;
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v, non_uniform);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:12:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThroughPointerInFunctionCallViaPointerArg) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>, a : ptr<function, i32>) {
+ *p = *a;
+}
+
+fn foo() {
+ var v = 0;
+ var a = non_uniform;
+ bar(&v, &a);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:13:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThroughPointerInFunctionCallViaArg) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>, a : i32) {
+ *p = a;
+}
+
+fn foo() {
+ var v = non_uniform;
+ bar(&v, 42);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThroughPointerInFunctionCallViaPointerArg) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>, a : ptr<function, i32>) {
+ *p = *a;
+}
+
+fn foo() {
+ var v = non_uniform;
+ var a = 42;
+ bar(&v, &a);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, AssignNonUniformThroughPointerInFunctionCallChain) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn f3(p : ptr<function, i32>, a : ptr<function, i32>) {
+ *p = *a;
+}
+
+fn f2(p : ptr<function, i32>, a : ptr<function, i32>) {
+ f3(p, a);
+}
+
+fn f1(p : ptr<function, i32>, a : ptr<function, i32>) {
+ f2(p, a);
+}
+
+fn foo() {
+ var v = 0;
+ var a = non_uniform;
+ f1(&v, &a);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:21:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, AssignUniformThroughPointerInFunctionCallChain) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn f3(p : ptr<function, i32>, a : ptr<function, i32>) {
+ *p = *a;
+}
+
+fn f2(p : ptr<function, i32>, a : ptr<function, i32>) {
+ f3(p, a);
+}
+
+fn f1(p : ptr<function, i32>, a : ptr<function, i32>) {
+ f2(p, a);
+}
+
+fn foo() {
+ var v = non_uniform;
+ var a = 42;
+ f1(&v, &a);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, MakePointerParamUniformInReturnExpression) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn zoo(p : ptr<function, i32>) -> i32 {
+ *p = 5;
+ return 6;
+}
+
+fn bar(p : ptr<function, i32>) -> i32 {
+ *p = non_uniform;
+ return zoo(p);
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, MakePointerParamNonUniformInReturnExpression) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn zoo(p : ptr<function, i32>) -> i32 {
+ *p = non_uniform;
+ return 6;
+}
+
+fn bar(p : ptr<function, i32>) -> i32 {
+ *p = 5;
+ return zoo(p);
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:18:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamAssignNonUniformInTrueAndUniformInFalse) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+ if (true) {
+ *p = non_uniform;
+ } else {
+ *p = 5;
+ }
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:16:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndReturn) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+ if (true) {
+ *p = non_uniform;
+ return;
+ }
+ *p = 5;
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:16:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndBreakFromSwitch) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+@group(0) @binding(1) var<uniform> condition : i32;
+
+fn bar(p : ptr<function, i32>) {
+ switch (condition) {
+ case 0 {
+ if (true) {
+ *p = non_uniform;
+ break;
+ }
+ *p = 5;
+ }
+ default {
+ *p = 6;
+ }
+ }
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:24:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndBreakFromLoop) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+ loop {
+ if (true) {
+ *p = non_uniform;
+ break;
+ }
+ *p = 5;
+ }
+}
+
+fn foo() {
+ var v = 0;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:18:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndContinue) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo(p : ptr<function, i32>) {
+ loop {
+ if (*p == 0) {
+ workgroupBarrier();
+ break;
+ }
+
+ if (true) {
+ *p = non_uniform;
+ continue;
+ }
+ *p = 5;
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:7:7 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamMaybeBecomesUniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+ if (true) {
+ *p = 5;
+ return;
+ }
+}
+
+fn foo() {
+ var v = non_uniform;
+ bar(&v);
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:15:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, NonUniformPointerParameterBecomesUniform_AfterUse) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(a : ptr<function, i32>, b : ptr<function, i32>) {
+ *b = *a;
+ *a = 0;
+}
+
+fn foo() {
+ var a = non_uniform;
+ var b = 0;
+ bar(&a, &b);
+ if (b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, NonUniformPointerParameterBecomesUniform_BeforeUse) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(a : ptr<function, i32>, b : ptr<function, i32>) {
+ *a = 0;
+ *b = *a;
+}
+
+fn foo() {
+ var a = non_uniform;
+ var b = 0;
+ bar(&a, &b);
+ if (b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, UniformPointerParameterBecomesNonUniform_BeforeUse) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(a : ptr<function, i32>, b : ptr<function, i32>) {
+ *a = non_uniform;
+ *b = *a;
+}
+
+fn foo() {
+ var a = 0;
+ var b = 0;
+ bar(&a, &b);
+ if (b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, UniformPointerParameterBecomesNonUniform_AfterUse) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(a : ptr<function, i32>, b : ptr<function, i32>) {
+ *b = *a;
+ *a = non_uniform;
+}
+
+fn foo() {
+ var a = 0;
+ var b = 0;
+ bar(&a, &b);
+ if (b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, NonUniformPointerParameterUpdatedInPlace) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(p : ptr<function, i32>) {
+ (*p)++;
+}
+
+fn foo() {
+ var v = non_uniform;
+ bar(&v);
+ if (v == 1) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:12:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, MultiplePointerParametersBecomeNonUniform) {
+ // The analysis traverses the tree for each pointer parameter, and we need to make sure that we
+ // reset the "visited" state of nodes in between these traversals to properly capture each of
+ // their uniformity states.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(a : ptr<function, i32>, b : ptr<function, i32>) {
+ *a = non_uniform;
+ *b = non_uniform;
+}
+
+fn foo() {
+ var a = 0;
+ var b = 0;
+ bar(&a, &b);
+ if (b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, MultiplePointerParametersWithEdgesToEachOther) {
+ // The analysis traverses the tree for each pointer parameter, and we need to make sure that we
+ // reset the "visited" state of nodes in between these traversals to properly capture each of
+ // their uniformity states.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn bar(a : ptr<function, i32>, b : ptr<function, i32>, c : ptr<function, i32>) {
+ *a = *a;
+ *b = *b;
+ *c = *a + *b;
+}
+
+fn foo() {
+ var a = non_uniform;
+ var b = 0;
+ var c = 0;
+ bar(&a, &b, &c);
+ if (c == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:16:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, MaximumNumberOfPointerParameters) {
+ // Create a function with the maximum number of parameters, all pointers, to stress the
+ // quadratic nature of the analysis.
+ ProgramBuilder b;
+ auto& ty = b.ty;
+
+ // fn foo(p0 : ptr<function, i32>, p1 : ptr<function, i32>, ...) {
+ // let rhs = *p0 + *p1 + ... + *p244;
+ // *p1 = rhs;
+ // *p2 = rhs;
+ // ...
+ // *p254 = rhs;
+ // }
+ ast::VariableList params;
+ ast::StatementList foo_body;
+ const ast::Expression* rhs_init = b.Deref("p0");
+ for (int i = 1; i < 255; i++) {
+ rhs_init = b.Add(rhs_init, b.Deref("p" + std::to_string(i)));
+ }
+ foo_body.push_back(b.Decl(b.Let("rhs", nullptr, rhs_init)));
+ for (int i = 0; i < 255; i++) {
+ params.push_back(
+ b.Param("p" + std::to_string(i), ty.pointer(ty.i32(), ast::StorageClass::kFunction)));
+ if (i > 0) {
+ foo_body.push_back(b.Assign(b.Deref("p" + std::to_string(i)), "rhs"));
+ }
+ }
+ b.Func("foo", std::move(params), ty.void_(), foo_body);
+
+ // var<private> non_uniform_global : i32;
+ // fn main() {
+ // var v0 : i32;
+ // var v1 : i32;
+ // ...
+ // var v254 : i32;
+ // v0 = non_uniform_global;
+ // foo(&v0, &v1, ..., &v254);
+ // if (v254 == 0) {
+ // workgroupBarrier();
+ // }
+ // }
+ b.Global("non_uniform_global", ty.i32(), ast::StorageClass::kPrivate);
+ ast::StatementList main_body;
+ ast::ExpressionList args;
+ for (int i = 0; i < 255; i++) {
+ auto name = "v" + std::to_string(i);
+ main_body.push_back(b.Decl(b.Var(name, ty.i32())));
+ args.push_back(b.AddressOf(name));
+ }
+ main_body.push_back(b.Assign("v0", "non_uniform_global"));
+ main_body.push_back(b.CallStmt(b.create<ast::CallExpression>(b.Expr("foo"), args)));
+ main_body.push_back(
+ b.If(b.Equal("v254", 0_i), b.Block(b.CallStmt(b.Call("workgroupBarrier")))));
+ b.Func("main", {}, ty.void_(), main_body);
+
+ // TODO(jrprice): Expect false when uniformity issues become errors.
+ EXPECT_TRUE(RunTest(std::move(b))) << error_;
+ EXPECT_EQ(error_,
+ R"(warning: 'workgroupBarrier' must only be called from uniform control flow)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Tests to cover access to aggregate types.
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(UniformityAnalysisTest, VectorElement_Uniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> v : vec4<i32>;
+
+fn foo() {
+ if (v[2] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, VectorElement_NonUniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> v : array<i32>;
+
+fn foo() {
+ if (v[2] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, VectorElement_BecomesNonUniform_BeforeCondition) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ v[2] = rw;
+ if (v[2] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, VectorElement_BecomesNonUniform_AfterCondition) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ if (v[2] == 0) {
+ v[2] = rw;
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, VectorElement_DifferentElementBecomesNonUniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ v[1] = rw;
+ if (v[2] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, VectorElement_ElementBecomesUniform) {
+ // For aggregate types, we conservatively consider them to be forever non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, that element is
+ // still considered to be non-uniform.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ v[1] = rw;
+ v[1] = 42;
+ if (v[1] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, VectorElement_DifferentElementBecomesUniform) {
+ // For aggregate types, we conservatively consider them to be forever non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, the whole vector
+ // is still considered to be non-uniform.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ v[1] = rw;
+ v[2] = 42;
+ if (v[1] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, VectorElement_NonUniform_AnyBuiltin) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ v[1] = non_uniform_global;
+ if (any(v == vec4(42))) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, StructMember_Uniform) {
+ auto src = R"(
+struct S {
+ a : i32,
+ b : i32,
+}
+@group(0) @binding(0) var<storage, read> s : S;
+
+fn foo() {
+ if (s.b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, StructMember_NonUniform) {
+ auto src = R"(
+struct S {
+ a : i32,
+ b : i32,
+}
+@group(0) @binding(0) var<storage, read_write> s : S;
+
+fn foo() {
+ if (s.b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, StructMember_BecomesNonUniform_BeforeCondition) {
+ auto src = R"(
+struct S {
+ a : i32,
+ b : i32,
+}
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var s : S;
+ s.b = rw;
+ if (s.b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:12:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, StructMember_BecomesNonUniform_AfterCondition) {
+ auto src = R"(
+struct S {
+ a : i32,
+ b : i32,
+}
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var s : S;
+ if (s.b == 0) {
+ s.b = rw;
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, StructMember_DifferentMemberBecomesNonUniform) {
+ auto src = R"(
+struct S {
+ a : i32,
+ b : i32,
+}
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var s : S;
+ s.a = rw;
+ if (s.b == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:12:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, StructMember_MemberBecomesUniform) {
+ // For aggregate types, we conservatively consider them to be forever non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to a member, that member is
+ // still considered to be non-uniform.
+ auto src = R"(
+struct S {
+ a : i32,
+ b : i32,
+}
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var s : S;
+ s.a = rw;
+ s.a = 0;
+ if (s.a == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:13:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, StructMember_DifferentMemberBecomesUniform) {
+ // For aggregate types, we conservatively consider them to be forever non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to a member, the whole struct
+ // is still considered to be non-uniform.
+ auto src = R"(
+struct S {
+ a : i32,
+ b : i32,
+}
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var s : S;
+ s.a = rw;
+ s.b = 0;
+ if (s.a == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:13:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_Uniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read> arr : array<i32>;
+
+fn foo() {
+ if (arr[7] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_NonUniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<i32>;
+
+fn foo() {
+ if (arr[7] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_BecomesNonUniform_BeforeCondition) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var arr : array<i32, 4>;
+ arr[2] = rw;
+ if (arr[2] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_BecomesNonUniform_AfterCondition) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var arr : array<i32, 4>;
+ if (arr[2] == 0) {
+ arr[2] = rw;
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_DifferentElementBecomesNonUniform) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var arr : array<i32, 4>;
+ arr[1] = rw;
+ if (arr[2] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_DifferentElementBecomesNonUniformThroughPointer) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var arr : array<i32, 4>;
+ let pa = &arr[1];
+ *pa = rw;
+ if (arr[2] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_ElementBecomesUniform) {
+ // For aggregate types, we conservatively consider them to be forever non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, that element is
+ // still considered to be non-uniform.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var arr : array<i32, 4>;
+ arr[1] = rw;
+ arr[1] = 42;
+ if (arr[1] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_DifferentElementBecomesUniform) {
+ // For aggregate types, we conservatively consider them to be forever non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, the whole array
+ // is still considered to be non-uniform.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var arr : array<i32, 4>;
+ arr[1] = rw;
+ arr[2] = 42;
+ if (arr[1] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:9:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ArrayElement_ElementBecomesUniformThroughPointer) {
+ // For aggregate types, we conservatively consider them to be forever non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element through a
+ // pointer, the whole array is still considered to be non-uniform.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var arr : array<i32, 4>;
+ let pa = &arr[2];
+ arr[1] = rw;
+ *pa = 42;
+ if (arr[1] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Miscellaneous statement and expression tests.
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(UniformityAnalysisTest, TypeConstructor) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo() {
+ if (i32(non_uniform_global) == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Conversion) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo() {
+ if (f32(non_uniform_global) == 0.0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, Bitcast) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo() {
+ if (bitcast<f32>(non_uniform_global) == 0.0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, CompoundAssignment_NonUniformRHS) {
+ // Use compound assignment with a non-uniform RHS on a variable.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v = 0;
+ v += rw;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, CompoundAssignment_UniformRHS_StillNonUniform) {
+ // Use compound assignment with a uniform RHS on a variable that is already non-uniform.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v = rw;
+ v += 1;
+ if (v == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PhonyAssignment_LhsCausesNonUniformControlFlow) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> nonuniform_var : i32;
+
+fn bar() -> i32 {
+ if (nonuniform_var == 42) {
+ return 1;
+ } else {
+ return 2;
+ }
+}
+
+fn foo() {
+ _ = bar();
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:14:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, ShortCircuiting_CausesNonUniformControlFlow) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
+
+var<private> p : i32;
+
+fn main() {
+ let b = (non_uniform_global == 42) && false;
+ workgroupBarrier();
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:8:3 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, DeadCode_AfterReturn) {
+ // Dead code after a return statement shouldn't cause uniformity errors.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ return;
+ if (non_uniform == 42) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, DeadCode_AfterDiscard) {
+ // Dead code after a discard statement shouldn't cause uniformity errors.
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
+
+fn foo() {
+ discard;
+ if (non_uniform == 42) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, ArrayLength) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<f32>;
+
+fn foo() {
+ for (var i = 0u; i < arrayLength(&arr); i++) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+TEST_F(UniformityAnalysisTest, WorkgroupAtomics) {
+ auto src = R"(
+var<workgroup> a : atomic<i32>;
+
+fn foo() {
+ if (atomicAdd(&a, 1) == 1) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, StorageAtomics) {
+ auto src = R"(
+@group(0) @binding(0) var<storage, read_write> a : atomic<i32>;
+
+fn foo() {
+ if (atomicAdd(&a, 1) == 1) {
+ storageBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 warning: 'storageBarrier' must only be called from uniform control flow
+ storageBarrier();
+ ^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, DisableAnalysisWithExtension) {
+ auto src = R"(
+enable chromium_disable_uniformity_analysis;
+
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ if (rw == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
+} // namespace
+} // namespace tint::resolver
diff --git a/src/tint/scope_stack.h b/src/tint/scope_stack.h
index 6892b40..6838f5b 100644
--- a/src/tint/scope_stack.h
+++ b/src/tint/scope_stack.h
@@ -71,6 +71,16 @@
return V{};
}
+ /// Return the top scope of the stack.
+ /// @returns the top scope of the stack
+ const std::unordered_map<K, V>& Top() const { return stack_.back(); }
+
+ /// Clear the scope stack.
+ void Clear() {
+ stack_.clear();
+ stack_.push_back({});
+ }
+
private:
std::vector<std::unordered_map<K, V>> stack_;
};
diff --git a/src/tint/scope_stack_test.cc b/src/tint/scope_stack_test.cc
index 8ed1f81..aeb7e73 100644
--- a/src/tint/scope_stack_test.cc
+++ b/src/tint/scope_stack_test.cc
@@ -67,5 +67,25 @@
EXPECT_EQ(s.Get(b), 25u);
}
+TEST_F(ScopeStackTest, Clear) {
+ ScopeStack<Symbol, uint32_t> s;
+ Symbol a(1, ID());
+ Symbol b(2, ID());
+
+ EXPECT_EQ(s.Set(a, 5u), 0u);
+ EXPECT_EQ(s.Get(a), 5u);
+
+ s.Push();
+
+ EXPECT_EQ(s.Set(b, 10u), 0u);
+ EXPECT_EQ(s.Get(b), 10u);
+
+ s.Push();
+
+ s.Clear();
+ EXPECT_EQ(s.Get(a), 0u);
+ EXPECT_EQ(s.Get(b), 0u);
+}
+
} // namespace
} // namespace tint
diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn
index 784fedb..e17c6c7 100644
--- a/test/tint/BUILD.gn
+++ b/test/tint/BUILD.gn
@@ -272,6 +272,7 @@
"../../src/tint/resolver/struct_storage_class_use_test.cc",
"../../src/tint/resolver/type_constructor_validation_test.cc",
"../../src/tint/resolver/type_validation_test.cc",
+ "../../src/tint/resolver/uniformity_test.cc",
"../../src/tint/resolver/validation_test.cc",
"../../src/tint/resolver/validator_is_storeable_test.cc",
"../../src/tint/resolver/var_let_test.cc",
diff --git a/test/tint/benchmark/uniformity-analysis-pointer-parameters.wgsl b/test/tint/benchmark/uniformity-analysis-pointer-parameters.wgsl
new file mode 100644
index 0000000..f73abef
--- /dev/null
+++ b/test/tint/benchmark/uniformity-analysis-pointer-parameters.wgsl
@@ -0,0 +1,1065 @@
+// Create a function with the maximum number of parameters, all pointers, to stress the
+// quadratic nature of the uniformity analysis.
+
+fn foo(
+ p0 : ptr<function, i32>,
+ p1 : ptr<function, i32>,
+ p2 : ptr<function, i32>,
+ p3 : ptr<function, i32>,
+ p4 : ptr<function, i32>,
+ p5 : ptr<function, i32>,
+ p6 : ptr<function, i32>,
+ p7 : ptr<function, i32>,
+ p8 : ptr<function, i32>,
+ p9 : ptr<function, i32>,
+ p10 : ptr<function, i32>,
+ p11 : ptr<function, i32>,
+ p12 : ptr<function, i32>,
+ p13 : ptr<function, i32>,
+ p14 : ptr<function, i32>,
+ p15 : ptr<function, i32>,
+ p16 : ptr<function, i32>,
+ p17 : ptr<function, i32>,
+ p18 : ptr<function, i32>,
+ p19 : ptr<function, i32>,
+ p20 : ptr<function, i32>,
+ p21 : ptr<function, i32>,
+ p22 : ptr<function, i32>,
+ p23 : ptr<function, i32>,
+ p24 : ptr<function, i32>,
+ p25 : ptr<function, i32>,
+ p26 : ptr<function, i32>,
+ p27 : ptr<function, i32>,
+ p28 : ptr<function, i32>,
+ p29 : ptr<function, i32>,
+ p30 : ptr<function, i32>,
+ p31 : ptr<function, i32>,
+ p32 : ptr<function, i32>,
+ p33 : ptr<function, i32>,
+ p34 : ptr<function, i32>,
+ p35 : ptr<function, i32>,
+ p36 : ptr<function, i32>,
+ p37 : ptr<function, i32>,
+ p38 : ptr<function, i32>,
+ p39 : ptr<function, i32>,
+ p40 : ptr<function, i32>,
+ p41 : ptr<function, i32>,
+ p42 : ptr<function, i32>,
+ p43 : ptr<function, i32>,
+ p44 : ptr<function, i32>,
+ p45 : ptr<function, i32>,
+ p46 : ptr<function, i32>,
+ p47 : ptr<function, i32>,
+ p48 : ptr<function, i32>,
+ p49 : ptr<function, i32>,
+ p50 : ptr<function, i32>,
+ p51 : ptr<function, i32>,
+ p52 : ptr<function, i32>,
+ p53 : ptr<function, i32>,
+ p54 : ptr<function, i32>,
+ p55 : ptr<function, i32>,
+ p56 : ptr<function, i32>,
+ p57 : ptr<function, i32>,
+ p58 : ptr<function, i32>,
+ p59 : ptr<function, i32>,
+ p60 : ptr<function, i32>,
+ p61 : ptr<function, i32>,
+ p62 : ptr<function, i32>,
+ p63 : ptr<function, i32>,
+ p64 : ptr<function, i32>,
+ p65 : ptr<function, i32>,
+ p66 : ptr<function, i32>,
+ p67 : ptr<function, i32>,
+ p68 : ptr<function, i32>,
+ p69 : ptr<function, i32>,
+ p70 : ptr<function, i32>,
+ p71 : ptr<function, i32>,
+ p72 : ptr<function, i32>,
+ p73 : ptr<function, i32>,
+ p74 : ptr<function, i32>,
+ p75 : ptr<function, i32>,
+ p76 : ptr<function, i32>,
+ p77 : ptr<function, i32>,
+ p78 : ptr<function, i32>,
+ p79 : ptr<function, i32>,
+ p80 : ptr<function, i32>,
+ p81 : ptr<function, i32>,
+ p82 : ptr<function, i32>,
+ p83 : ptr<function, i32>,
+ p84 : ptr<function, i32>,
+ p85 : ptr<function, i32>,
+ p86 : ptr<function, i32>,
+ p87 : ptr<function, i32>,
+ p88 : ptr<function, i32>,
+ p89 : ptr<function, i32>,
+ p90 : ptr<function, i32>,
+ p91 : ptr<function, i32>,
+ p92 : ptr<function, i32>,
+ p93 : ptr<function, i32>,
+ p94 : ptr<function, i32>,
+ p95 : ptr<function, i32>,
+ p96 : ptr<function, i32>,
+ p97 : ptr<function, i32>,
+ p98 : ptr<function, i32>,
+ p99 : ptr<function, i32>,
+ p100 : ptr<function, i32>,
+ p101 : ptr<function, i32>,
+ p102 : ptr<function, i32>,
+ p103 : ptr<function, i32>,
+ p104 : ptr<function, i32>,
+ p105 : ptr<function, i32>,
+ p106 : ptr<function, i32>,
+ p107 : ptr<function, i32>,
+ p108 : ptr<function, i32>,
+ p109 : ptr<function, i32>,
+ p110 : ptr<function, i32>,
+ p111 : ptr<function, i32>,
+ p112 : ptr<function, i32>,
+ p113 : ptr<function, i32>,
+ p114 : ptr<function, i32>,
+ p115 : ptr<function, i32>,
+ p116 : ptr<function, i32>,
+ p117 : ptr<function, i32>,
+ p118 : ptr<function, i32>,
+ p119 : ptr<function, i32>,
+ p120 : ptr<function, i32>,
+ p121 : ptr<function, i32>,
+ p122 : ptr<function, i32>,
+ p123 : ptr<function, i32>,
+ p124 : ptr<function, i32>,
+ p125 : ptr<function, i32>,
+ p126 : ptr<function, i32>,
+ p127 : ptr<function, i32>,
+ p128 : ptr<function, i32>,
+ p129 : ptr<function, i32>,
+ p130 : ptr<function, i32>,
+ p131 : ptr<function, i32>,
+ p132 : ptr<function, i32>,
+ p133 : ptr<function, i32>,
+ p134 : ptr<function, i32>,
+ p135 : ptr<function, i32>,
+ p136 : ptr<function, i32>,
+ p137 : ptr<function, i32>,
+ p138 : ptr<function, i32>,
+ p139 : ptr<function, i32>,
+ p140 : ptr<function, i32>,
+ p141 : ptr<function, i32>,
+ p142 : ptr<function, i32>,
+ p143 : ptr<function, i32>,
+ p144 : ptr<function, i32>,
+ p145 : ptr<function, i32>,
+ p146 : ptr<function, i32>,
+ p147 : ptr<function, i32>,
+ p148 : ptr<function, i32>,
+ p149 : ptr<function, i32>,
+ p150 : ptr<function, i32>,
+ p151 : ptr<function, i32>,
+ p152 : ptr<function, i32>,
+ p153 : ptr<function, i32>,
+ p154 : ptr<function, i32>,
+ p155 : ptr<function, i32>,
+ p156 : ptr<function, i32>,
+ p157 : ptr<function, i32>,
+ p158 : ptr<function, i32>,
+ p159 : ptr<function, i32>,
+ p160 : ptr<function, i32>,
+ p161 : ptr<function, i32>,
+ p162 : ptr<function, i32>,
+ p163 : ptr<function, i32>,
+ p164 : ptr<function, i32>,
+ p165 : ptr<function, i32>,
+ p166 : ptr<function, i32>,
+ p167 : ptr<function, i32>,
+ p168 : ptr<function, i32>,
+ p169 : ptr<function, i32>,
+ p170 : ptr<function, i32>,
+ p171 : ptr<function, i32>,
+ p172 : ptr<function, i32>,
+ p173 : ptr<function, i32>,
+ p174 : ptr<function, i32>,
+ p175 : ptr<function, i32>,
+ p176 : ptr<function, i32>,
+ p177 : ptr<function, i32>,
+ p178 : ptr<function, i32>,
+ p179 : ptr<function, i32>,
+ p180 : ptr<function, i32>,
+ p181 : ptr<function, i32>,
+ p182 : ptr<function, i32>,
+ p183 : ptr<function, i32>,
+ p184 : ptr<function, i32>,
+ p185 : ptr<function, i32>,
+ p186 : ptr<function, i32>,
+ p187 : ptr<function, i32>,
+ p188 : ptr<function, i32>,
+ p189 : ptr<function, i32>,
+ p190 : ptr<function, i32>,
+ p191 : ptr<function, i32>,
+ p192 : ptr<function, i32>,
+ p193 : ptr<function, i32>,
+ p194 : ptr<function, i32>,
+ p195 : ptr<function, i32>,
+ p196 : ptr<function, i32>,
+ p197 : ptr<function, i32>,
+ p198 : ptr<function, i32>,
+ p199 : ptr<function, i32>,
+ p200 : ptr<function, i32>,
+ p201 : ptr<function, i32>,
+ p202 : ptr<function, i32>,
+ p203 : ptr<function, i32>,
+ p204 : ptr<function, i32>,
+ p205 : ptr<function, i32>,
+ p206 : ptr<function, i32>,
+ p207 : ptr<function, i32>,
+ p208 : ptr<function, i32>,
+ p209 : ptr<function, i32>,
+ p210 : ptr<function, i32>,
+ p211 : ptr<function, i32>,
+ p212 : ptr<function, i32>,
+ p213 : ptr<function, i32>,
+ p214 : ptr<function, i32>,
+ p215 : ptr<function, i32>,
+ p216 : ptr<function, i32>,
+ p217 : ptr<function, i32>,
+ p218 : ptr<function, i32>,
+ p219 : ptr<function, i32>,
+ p220 : ptr<function, i32>,
+ p221 : ptr<function, i32>,
+ p222 : ptr<function, i32>,
+ p223 : ptr<function, i32>,
+ p224 : ptr<function, i32>,
+ p225 : ptr<function, i32>,
+ p226 : ptr<function, i32>,
+ p227 : ptr<function, i32>,
+ p228 : ptr<function, i32>,
+ p229 : ptr<function, i32>,
+ p230 : ptr<function, i32>,
+ p231 : ptr<function, i32>,
+ p232 : ptr<function, i32>,
+ p233 : ptr<function, i32>,
+ p234 : ptr<function, i32>,
+ p235 : ptr<function, i32>,
+ p236 : ptr<function, i32>,
+ p237 : ptr<function, i32>,
+ p238 : ptr<function, i32>,
+ p239 : ptr<function, i32>,
+ p240 : ptr<function, i32>,
+ p241 : ptr<function, i32>,
+ p242 : ptr<function, i32>,
+ p243 : ptr<function, i32>,
+ p244 : ptr<function, i32>,
+ p245 : ptr<function, i32>,
+ p246 : ptr<function, i32>,
+ p247 : ptr<function, i32>,
+ p248 : ptr<function, i32>,
+ p249 : ptr<function, i32>,
+ p250 : ptr<function, i32>,
+ p251 : ptr<function, i32>,
+ p252 : ptr<function, i32>,
+ p253 : ptr<function, i32>,
+ p254 : ptr<function, i32>,
+) {
+ var rhs = *p0 + *p1 + *p2 + *p3 + *p4 + *p5 + *p6 + *p7;
+ rhs += *p8 + *p9 + *p10 + *p11 + *p12 + *p13 + *p14 + *p15;
+ rhs += *p16 + *p17 + *p18 + *p19 + *p20 + *p21 + *p22 + *p23;
+ rhs += *p24 + *p25 + *p26 + *p27 + *p28 + *p29 + *p30 + *p31;
+ rhs += *p32 + *p33 + *p34 + *p35 + *p36 + *p37 + *p38 + *p39;
+ rhs += *p40 + *p41 + *p42 + *p43 + *p44 + *p45 + *p46 + *p47;
+ rhs += *p48 + *p49 + *p50 + *p51 + *p52 + *p53 + *p54 + *p55;
+ rhs += *p56 + *p57 + *p58 + *p59 + *p60 + *p61 + *p62 + *p63;
+ rhs += *p64 + *p65 + *p66 + *p67 + *p68 + *p69 + *p70 + *p71;
+ rhs += *p72 + *p73 + *p74 + *p75 + *p76 + *p77 + *p78 + *p79;
+ rhs += *p80 + *p81 + *p82 + *p83 + *p84 + *p85 + *p86 + *p87;
+ rhs += *p88 + *p89 + *p90 + *p91 + *p92 + *p93 + *p94 + *p95;
+ rhs += *p96 + *p97 + *p98 + *p99 + *p100 + *p101 + *p102 + *p103;
+ rhs += *p104 + *p105 + *p106 + *p107 + *p108 + *p109 + *p110 + *p111;
+ rhs += *p112 + *p113 + *p114 + *p115 + *p116 + *p117 + *p118 + *p119;
+ rhs += *p120 + *p121 + *p122 + *p123 + *p124 + *p125 + *p126 + *p127;
+ rhs += *p128 + *p129 + *p130 + *p131 + *p132 + *p133 + *p134 + *p135;
+ rhs += *p136 + *p137 + *p138 + *p139 + *p140 + *p141 + *p142 + *p143;
+ rhs += *p144 + *p145 + *p146 + *p147 + *p148 + *p149 + *p150 + *p151;
+ rhs += *p152 + *p153 + *p154 + *p155 + *p156 + *p157 + *p158 + *p159;
+ rhs += *p160 + *p161 + *p162 + *p163 + *p164 + *p165 + *p166 + *p167;
+ rhs += *p168 + *p169 + *p170 + *p171 + *p172 + *p173 + *p174 + *p175;
+ rhs += *p176 + *p177 + *p178 + *p179 + *p180 + *p181 + *p182 + *p183;
+ rhs += *p184 + *p185 + *p186 + *p187 + *p188 + *p189 + *p190 + *p191;
+ rhs += *p192 + *p193 + *p194 + *p195 + *p196 + *p197 + *p198 + *p199;
+ rhs += *p200 + *p201 + *p202 + *p203 + *p204 + *p205 + *p206 + *p207;
+ rhs += *p208 + *p209 + *p210 + *p211 + *p212 + *p213 + *p214 + *p215;
+ rhs += *p216 + *p217 + *p218 + *p219 + *p220 + *p221 + *p222 + *p223;
+ rhs += *p224 + *p225 + *p226 + *p227 + *p228 + *p229 + *p230 + *p231;
+ rhs += *p232 + *p233 + *p234 + *p235 + *p236 + *p237 + *p238 + *p239;
+ rhs += *p240 + *p241 + *p242 + *p243 + *p244 + *p245 + *p246 + *p247;
+ rhs += *p248 + *p249 + *p250 + *p251 + *p252 + *p253 + *p254;
+ *p1 = rhs;
+ *p2 = rhs;
+ *p3 = rhs;
+ *p4 = rhs;
+ *p5 = rhs;
+ *p6 = rhs;
+ *p7 = rhs;
+ *p8 = rhs;
+ *p9 = rhs;
+ *p10 = rhs;
+ *p11 = rhs;
+ *p12 = rhs;
+ *p13 = rhs;
+ *p14 = rhs;
+ *p15 = rhs;
+ *p16 = rhs;
+ *p17 = rhs;
+ *p18 = rhs;
+ *p19 = rhs;
+ *p20 = rhs;
+ *p21 = rhs;
+ *p22 = rhs;
+ *p23 = rhs;
+ *p24 = rhs;
+ *p25 = rhs;
+ *p26 = rhs;
+ *p27 = rhs;
+ *p28 = rhs;
+ *p29 = rhs;
+ *p30 = rhs;
+ *p31 = rhs;
+ *p32 = rhs;
+ *p33 = rhs;
+ *p34 = rhs;
+ *p35 = rhs;
+ *p36 = rhs;
+ *p37 = rhs;
+ *p38 = rhs;
+ *p39 = rhs;
+ *p40 = rhs;
+ *p41 = rhs;
+ *p42 = rhs;
+ *p43 = rhs;
+ *p44 = rhs;
+ *p45 = rhs;
+ *p46 = rhs;
+ *p47 = rhs;
+ *p48 = rhs;
+ *p49 = rhs;
+ *p50 = rhs;
+ *p51 = rhs;
+ *p52 = rhs;
+ *p53 = rhs;
+ *p54 = rhs;
+ *p55 = rhs;
+ *p56 = rhs;
+ *p57 = rhs;
+ *p58 = rhs;
+ *p59 = rhs;
+ *p60 = rhs;
+ *p61 = rhs;
+ *p62 = rhs;
+ *p63 = rhs;
+ *p64 = rhs;
+ *p65 = rhs;
+ *p66 = rhs;
+ *p67 = rhs;
+ *p68 = rhs;
+ *p69 = rhs;
+ *p70 = rhs;
+ *p71 = rhs;
+ *p72 = rhs;
+ *p73 = rhs;
+ *p74 = rhs;
+ *p75 = rhs;
+ *p76 = rhs;
+ *p77 = rhs;
+ *p78 = rhs;
+ *p79 = rhs;
+ *p80 = rhs;
+ *p81 = rhs;
+ *p82 = rhs;
+ *p83 = rhs;
+ *p84 = rhs;
+ *p85 = rhs;
+ *p86 = rhs;
+ *p87 = rhs;
+ *p88 = rhs;
+ *p89 = rhs;
+ *p90 = rhs;
+ *p91 = rhs;
+ *p92 = rhs;
+ *p93 = rhs;
+ *p94 = rhs;
+ *p95 = rhs;
+ *p96 = rhs;
+ *p97 = rhs;
+ *p98 = rhs;
+ *p99 = rhs;
+ *p100 = rhs;
+ *p101 = rhs;
+ *p102 = rhs;
+ *p103 = rhs;
+ *p104 = rhs;
+ *p105 = rhs;
+ *p106 = rhs;
+ *p107 = rhs;
+ *p108 = rhs;
+ *p109 = rhs;
+ *p110 = rhs;
+ *p111 = rhs;
+ *p112 = rhs;
+ *p113 = rhs;
+ *p114 = rhs;
+ *p115 = rhs;
+ *p116 = rhs;
+ *p117 = rhs;
+ *p118 = rhs;
+ *p119 = rhs;
+ *p120 = rhs;
+ *p121 = rhs;
+ *p122 = rhs;
+ *p123 = rhs;
+ *p124 = rhs;
+ *p125 = rhs;
+ *p126 = rhs;
+ *p127 = rhs;
+ *p128 = rhs;
+ *p129 = rhs;
+ *p130 = rhs;
+ *p131 = rhs;
+ *p132 = rhs;
+ *p133 = rhs;
+ *p134 = rhs;
+ *p135 = rhs;
+ *p136 = rhs;
+ *p137 = rhs;
+ *p138 = rhs;
+ *p139 = rhs;
+ *p140 = rhs;
+ *p141 = rhs;
+ *p142 = rhs;
+ *p143 = rhs;
+ *p144 = rhs;
+ *p145 = rhs;
+ *p146 = rhs;
+ *p147 = rhs;
+ *p148 = rhs;
+ *p149 = rhs;
+ *p150 = rhs;
+ *p151 = rhs;
+ *p152 = rhs;
+ *p153 = rhs;
+ *p154 = rhs;
+ *p155 = rhs;
+ *p156 = rhs;
+ *p157 = rhs;
+ *p158 = rhs;
+ *p159 = rhs;
+ *p160 = rhs;
+ *p161 = rhs;
+ *p162 = rhs;
+ *p163 = rhs;
+ *p164 = rhs;
+ *p165 = rhs;
+ *p166 = rhs;
+ *p167 = rhs;
+ *p168 = rhs;
+ *p169 = rhs;
+ *p170 = rhs;
+ *p171 = rhs;
+ *p172 = rhs;
+ *p173 = rhs;
+ *p174 = rhs;
+ *p175 = rhs;
+ *p176 = rhs;
+ *p177 = rhs;
+ *p178 = rhs;
+ *p179 = rhs;
+ *p180 = rhs;
+ *p181 = rhs;
+ *p182 = rhs;
+ *p183 = rhs;
+ *p184 = rhs;
+ *p185 = rhs;
+ *p186 = rhs;
+ *p187 = rhs;
+ *p188 = rhs;
+ *p189 = rhs;
+ *p190 = rhs;
+ *p191 = rhs;
+ *p192 = rhs;
+ *p193 = rhs;
+ *p194 = rhs;
+ *p195 = rhs;
+ *p196 = rhs;
+ *p197 = rhs;
+ *p198 = rhs;
+ *p199 = rhs;
+ *p200 = rhs;
+ *p201 = rhs;
+ *p202 = rhs;
+ *p203 = rhs;
+ *p204 = rhs;
+ *p205 = rhs;
+ *p206 = rhs;
+ *p207 = rhs;
+ *p208 = rhs;
+ *p209 = rhs;
+ *p210 = rhs;
+ *p211 = rhs;
+ *p212 = rhs;
+ *p213 = rhs;
+ *p214 = rhs;
+ *p215 = rhs;
+ *p216 = rhs;
+ *p217 = rhs;
+ *p218 = rhs;
+ *p219 = rhs;
+ *p220 = rhs;
+ *p221 = rhs;
+ *p222 = rhs;
+ *p223 = rhs;
+ *p224 = rhs;
+ *p225 = rhs;
+ *p226 = rhs;
+ *p227 = rhs;
+ *p228 = rhs;
+ *p229 = rhs;
+ *p230 = rhs;
+ *p231 = rhs;
+ *p232 = rhs;
+ *p233 = rhs;
+ *p234 = rhs;
+ *p235 = rhs;
+ *p236 = rhs;
+ *p237 = rhs;
+ *p238 = rhs;
+ *p239 = rhs;
+ *p240 = rhs;
+ *p241 = rhs;
+ *p242 = rhs;
+ *p243 = rhs;
+ *p244 = rhs;
+ *p245 = rhs;
+ *p246 = rhs;
+ *p247 = rhs;
+ *p248 = rhs;
+ *p249 = rhs;
+ *p250 = rhs;
+ *p251 = rhs;
+ *p252 = rhs;
+ *p253 = rhs;
+ *p254 = rhs;
+}
+
+fn main() {
+ var v0 : i32;
+ var v1 : i32;
+ var v2 : i32;
+ var v3 : i32;
+ var v4 : i32;
+ var v5 : i32;
+ var v6 : i32;
+ var v7 : i32;
+ var v8 : i32;
+ var v9 : i32;
+ var v10 : i32;
+ var v11 : i32;
+ var v12 : i32;
+ var v13 : i32;
+ var v14 : i32;
+ var v15 : i32;
+ var v16 : i32;
+ var v17 : i32;
+ var v18 : i32;
+ var v19 : i32;
+ var v20 : i32;
+ var v21 : i32;
+ var v22 : i32;
+ var v23 : i32;
+ var v24 : i32;
+ var v25 : i32;
+ var v26 : i32;
+ var v27 : i32;
+ var v28 : i32;
+ var v29 : i32;
+ var v30 : i32;
+ var v31 : i32;
+ var v32 : i32;
+ var v33 : i32;
+ var v34 : i32;
+ var v35 : i32;
+ var v36 : i32;
+ var v37 : i32;
+ var v38 : i32;
+ var v39 : i32;
+ var v40 : i32;
+ var v41 : i32;
+ var v42 : i32;
+ var v43 : i32;
+ var v44 : i32;
+ var v45 : i32;
+ var v46 : i32;
+ var v47 : i32;
+ var v48 : i32;
+ var v49 : i32;
+ var v50 : i32;
+ var v51 : i32;
+ var v52 : i32;
+ var v53 : i32;
+ var v54 : i32;
+ var v55 : i32;
+ var v56 : i32;
+ var v57 : i32;
+ var v58 : i32;
+ var v59 : i32;
+ var v60 : i32;
+ var v61 : i32;
+ var v62 : i32;
+ var v63 : i32;
+ var v64 : i32;
+ var v65 : i32;
+ var v66 : i32;
+ var v67 : i32;
+ var v68 : i32;
+ var v69 : i32;
+ var v70 : i32;
+ var v71 : i32;
+ var v72 : i32;
+ var v73 : i32;
+ var v74 : i32;
+ var v75 : i32;
+ var v76 : i32;
+ var v77 : i32;
+ var v78 : i32;
+ var v79 : i32;
+ var v80 : i32;
+ var v81 : i32;
+ var v82 : i32;
+ var v83 : i32;
+ var v84 : i32;
+ var v85 : i32;
+ var v86 : i32;
+ var v87 : i32;
+ var v88 : i32;
+ var v89 : i32;
+ var v90 : i32;
+ var v91 : i32;
+ var v92 : i32;
+ var v93 : i32;
+ var v94 : i32;
+ var v95 : i32;
+ var v96 : i32;
+ var v97 : i32;
+ var v98 : i32;
+ var v99 : i32;
+ var v100 : i32;
+ var v101 : i32;
+ var v102 : i32;
+ var v103 : i32;
+ var v104 : i32;
+ var v105 : i32;
+ var v106 : i32;
+ var v107 : i32;
+ var v108 : i32;
+ var v109 : i32;
+ var v110 : i32;
+ var v111 : i32;
+ var v112 : i32;
+ var v113 : i32;
+ var v114 : i32;
+ var v115 : i32;
+ var v116 : i32;
+ var v117 : i32;
+ var v118 : i32;
+ var v119 : i32;
+ var v120 : i32;
+ var v121 : i32;
+ var v122 : i32;
+ var v123 : i32;
+ var v124 : i32;
+ var v125 : i32;
+ var v126 : i32;
+ var v127 : i32;
+ var v128 : i32;
+ var v129 : i32;
+ var v130 : i32;
+ var v131 : i32;
+ var v132 : i32;
+ var v133 : i32;
+ var v134 : i32;
+ var v135 : i32;
+ var v136 : i32;
+ var v137 : i32;
+ var v138 : i32;
+ var v139 : i32;
+ var v140 : i32;
+ var v141 : i32;
+ var v142 : i32;
+ var v143 : i32;
+ var v144 : i32;
+ var v145 : i32;
+ var v146 : i32;
+ var v147 : i32;
+ var v148 : i32;
+ var v149 : i32;
+ var v150 : i32;
+ var v151 : i32;
+ var v152 : i32;
+ var v153 : i32;
+ var v154 : i32;
+ var v155 : i32;
+ var v156 : i32;
+ var v157 : i32;
+ var v158 : i32;
+ var v159 : i32;
+ var v160 : i32;
+ var v161 : i32;
+ var v162 : i32;
+ var v163 : i32;
+ var v164 : i32;
+ var v165 : i32;
+ var v166 : i32;
+ var v167 : i32;
+ var v168 : i32;
+ var v169 : i32;
+ var v170 : i32;
+ var v171 : i32;
+ var v172 : i32;
+ var v173 : i32;
+ var v174 : i32;
+ var v175 : i32;
+ var v176 : i32;
+ var v177 : i32;
+ var v178 : i32;
+ var v179 : i32;
+ var v180 : i32;
+ var v181 : i32;
+ var v182 : i32;
+ var v183 : i32;
+ var v184 : i32;
+ var v185 : i32;
+ var v186 : i32;
+ var v187 : i32;
+ var v188 : i32;
+ var v189 : i32;
+ var v190 : i32;
+ var v191 : i32;
+ var v192 : i32;
+ var v193 : i32;
+ var v194 : i32;
+ var v195 : i32;
+ var v196 : i32;
+ var v197 : i32;
+ var v198 : i32;
+ var v199 : i32;
+ var v200 : i32;
+ var v201 : i32;
+ var v202 : i32;
+ var v203 : i32;
+ var v204 : i32;
+ var v205 : i32;
+ var v206 : i32;
+ var v207 : i32;
+ var v208 : i32;
+ var v209 : i32;
+ var v210 : i32;
+ var v211 : i32;
+ var v212 : i32;
+ var v213 : i32;
+ var v214 : i32;
+ var v215 : i32;
+ var v216 : i32;
+ var v217 : i32;
+ var v218 : i32;
+ var v219 : i32;
+ var v220 : i32;
+ var v221 : i32;
+ var v222 : i32;
+ var v223 : i32;
+ var v224 : i32;
+ var v225 : i32;
+ var v226 : i32;
+ var v227 : i32;
+ var v228 : i32;
+ var v229 : i32;
+ var v230 : i32;
+ var v231 : i32;
+ var v232 : i32;
+ var v233 : i32;
+ var v234 : i32;
+ var v235 : i32;
+ var v236 : i32;
+ var v237 : i32;
+ var v238 : i32;
+ var v239 : i32;
+ var v240 : i32;
+ var v241 : i32;
+ var v242 : i32;
+ var v243 : i32;
+ var v244 : i32;
+ var v245 : i32;
+ var v246 : i32;
+ var v247 : i32;
+ var v248 : i32;
+ var v249 : i32;
+ var v250 : i32;
+ var v251 : i32;
+ var v252 : i32;
+ var v253 : i32;
+ var v254 : i32;
+ foo(
+ &v0,
+ &v1,
+ &v2,
+ &v3,
+ &v4,
+ &v5,
+ &v6,
+ &v7,
+ &v8,
+ &v9,
+ &v10,
+ &v11,
+ &v12,
+ &v13,
+ &v14,
+ &v15,
+ &v16,
+ &v17,
+ &v18,
+ &v19,
+ &v20,
+ &v21,
+ &v22,
+ &v23,
+ &v24,
+ &v25,
+ &v26,
+ &v27,
+ &v28,
+ &v29,
+ &v30,
+ &v31,
+ &v32,
+ &v33,
+ &v34,
+ &v35,
+ &v36,
+ &v37,
+ &v38,
+ &v39,
+ &v40,
+ &v41,
+ &v42,
+ &v43,
+ &v44,
+ &v45,
+ &v46,
+ &v47,
+ &v48,
+ &v49,
+ &v50,
+ &v51,
+ &v52,
+ &v53,
+ &v54,
+ &v55,
+ &v56,
+ &v57,
+ &v58,
+ &v59,
+ &v60,
+ &v61,
+ &v62,
+ &v63,
+ &v64,
+ &v65,
+ &v66,
+ &v67,
+ &v68,
+ &v69,
+ &v70,
+ &v71,
+ &v72,
+ &v73,
+ &v74,
+ &v75,
+ &v76,
+ &v77,
+ &v78,
+ &v79,
+ &v80,
+ &v81,
+ &v82,
+ &v83,
+ &v84,
+ &v85,
+ &v86,
+ &v87,
+ &v88,
+ &v89,
+ &v90,
+ &v91,
+ &v92,
+ &v93,
+ &v94,
+ &v95,
+ &v96,
+ &v97,
+ &v98,
+ &v99,
+ &v100,
+ &v101,
+ &v102,
+ &v103,
+ &v104,
+ &v105,
+ &v106,
+ &v107,
+ &v108,
+ &v109,
+ &v110,
+ &v111,
+ &v112,
+ &v113,
+ &v114,
+ &v115,
+ &v116,
+ &v117,
+ &v118,
+ &v119,
+ &v120,
+ &v121,
+ &v122,
+ &v123,
+ &v124,
+ &v125,
+ &v126,
+ &v127,
+ &v128,
+ &v129,
+ &v130,
+ &v131,
+ &v132,
+ &v133,
+ &v134,
+ &v135,
+ &v136,
+ &v137,
+ &v138,
+ &v139,
+ &v140,
+ &v141,
+ &v142,
+ &v143,
+ &v144,
+ &v145,
+ &v146,
+ &v147,
+ &v148,
+ &v149,
+ &v150,
+ &v151,
+ &v152,
+ &v153,
+ &v154,
+ &v155,
+ &v156,
+ &v157,
+ &v158,
+ &v159,
+ &v160,
+ &v161,
+ &v162,
+ &v163,
+ &v164,
+ &v165,
+ &v166,
+ &v167,
+ &v168,
+ &v169,
+ &v170,
+ &v171,
+ &v172,
+ &v173,
+ &v174,
+ &v175,
+ &v176,
+ &v177,
+ &v178,
+ &v179,
+ &v180,
+ &v181,
+ &v182,
+ &v183,
+ &v184,
+ &v185,
+ &v186,
+ &v187,
+ &v188,
+ &v189,
+ &v190,
+ &v191,
+ &v192,
+ &v193,
+ &v194,
+ &v195,
+ &v196,
+ &v197,
+ &v198,
+ &v199,
+ &v200,
+ &v201,
+ &v202,
+ &v203,
+ &v204,
+ &v205,
+ &v206,
+ &v207,
+ &v208,
+ &v209,
+ &v210,
+ &v211,
+ &v212,
+ &v213,
+ &v214,
+ &v215,
+ &v216,
+ &v217,
+ &v218,
+ &v219,
+ &v220,
+ &v221,
+ &v222,
+ &v223,
+ &v224,
+ &v225,
+ &v226,
+ &v227,
+ &v228,
+ &v229,
+ &v230,
+ &v231,
+ &v232,
+ &v233,
+ &v234,
+ &v235,
+ &v236,
+ &v237,
+ &v238,
+ &v239,
+ &v240,
+ &v241,
+ &v242,
+ &v243,
+ &v244,
+ &v245,
+ &v246,
+ &v247,
+ &v248,
+ &v249,
+ &v250,
+ &v251,
+ &v252,
+ &v253,
+ &v254,
+ );
+ if (v254 == 0) {
+ workgroupBarrier();
+ }
+}
diff --git a/test/tint/bug/dawn/947.wgsl.expected.glsl b/test/tint/bug/dawn/947.wgsl.expected.glsl
index 413ee0e..60c1105 100644
--- a/test/tint/bug/dawn/947.wgsl.expected.glsl
+++ b/test/tint/bug/dawn/947.wgsl.expected.glsl
@@ -1,3 +1,7 @@
+bug/dawn/947.wgsl:59:20 warning: 'textureSample' must only be called from uniform control flow
+ var srcColor = textureSample(myTexture, mySampler, texcoord);
+ ^^^^^^^^^^^^^
+
#version 310 es
layout(location = 0) out vec2 texcoords_1;
diff --git a/test/tint/bug/dawn/947.wgsl.expected.hlsl b/test/tint/bug/dawn/947.wgsl.expected.hlsl
index 7bcadcb..3b13d95 100644
--- a/test/tint/bug/dawn/947.wgsl.expected.hlsl
+++ b/test/tint/bug/dawn/947.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+bug/dawn/947.wgsl:59:20 warning: 'textureSample' must only be called from uniform control flow
+ var srcColor = textureSample(myTexture, mySampler, texcoord);
+ ^^^^^^^^^^^^^
+
cbuffer cbuffer_uniforms : register(b0, space0) {
uint4 uniforms[1];
};
diff --git a/test/tint/bug/dawn/947.wgsl.expected.msl b/test/tint/bug/dawn/947.wgsl.expected.msl
index 15efcd0..07a7f73 100644
--- a/test/tint/bug/dawn/947.wgsl.expected.msl
+++ b/test/tint/bug/dawn/947.wgsl.expected.msl
@@ -1,3 +1,7 @@
+bug/dawn/947.wgsl:59:20 warning: 'textureSample' must only be called from uniform control flow
+ var srcColor = textureSample(myTexture, mySampler, texcoord);
+ ^^^^^^^^^^^^^
+
#include <metal_stdlib>
using namespace metal;
diff --git a/test/tint/bug/dawn/947.wgsl.expected.spvasm b/test/tint/bug/dawn/947.wgsl.expected.spvasm
index df48a9a..a74165c 100644
--- a/test/tint/bug/dawn/947.wgsl.expected.spvasm
+++ b/test/tint/bug/dawn/947.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+bug/dawn/947.wgsl:59:20 warning: 'textureSample' must only be called from uniform control flow
+ var srcColor = textureSample(myTexture, mySampler, texcoord);
+ ^^^^^^^^^^^^^
+
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
diff --git a/test/tint/bug/dawn/947.wgsl.expected.wgsl b/test/tint/bug/dawn/947.wgsl.expected.wgsl
index 063f354..6f98645 100644
--- a/test/tint/bug/dawn/947.wgsl.expected.wgsl
+++ b/test/tint/bug/dawn/947.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+bug/dawn/947.wgsl:59:20 warning: 'textureSample' must only be called from uniform control flow
+ var srcColor = textureSample(myTexture, mySampler, texcoord);
+ ^^^^^^^^^^^^^
+
struct Uniforms {
u_scale : vec2<f32>,
u_offset : vec2<f32>,
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.glsl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.glsl
index b11b4b2..48926b7 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.glsl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.glsl
@@ -1,3 +1,7 @@
+bug/fxc/gradient_in_varying_loop/1112.wgsl:23:33 warning: 'textureSample' must only be called from uniform control flow
+ let sampleDepth : f32 = textureSample(depthTexture, Sampler, offset.xy).r;
+ ^^^^^^^^^^^^^
+
#version 310 es
precision mediump float;
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.hlsl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.hlsl
index 4c8362a..a6c6dd5 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.hlsl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+bug/fxc/gradient_in_varying_loop/1112.wgsl:23:33 warning: 'textureSample' must only be called from uniform control flow
+ let sampleDepth : f32 = textureSample(depthTexture, Sampler, offset.xy).r;
+ ^^^^^^^^^^^^^
+
SamplerState tint_symbol : register(s0, space0);
Texture2D<float4> randomTexture : register(t1, space0);
Texture2D<float4> depthTexture : register(t2, space0);
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
index 3b00074..62a41ed 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
@@ -1,3 +1,7 @@
+bug/fxc/gradient_in_varying_loop/1112.wgsl:23:33 warning: 'textureSample' must only be called from uniform control flow
+ let sampleDepth : f32 = textureSample(depthTexture, Sampler, offset.xy).r;
+ ^^^^^^^^^^^^^
+
#include <metal_stdlib>
using namespace metal;
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.spvasm b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.spvasm
index 49aff3b..31483ae 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.spvasm
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+bug/fxc/gradient_in_varying_loop/1112.wgsl:23:33 warning: 'textureSample' must only be called from uniform control flow
+ let sampleDepth : f32 = textureSample(depthTexture, Sampler, offset.xy).r;
+ ^^^^^^^^^^^^^
+
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.wgsl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.wgsl
index 0998c72..1761161 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.wgsl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+bug/fxc/gradient_in_varying_loop/1112.wgsl:23:33 warning: 'textureSample' must only be called from uniform control flow
+ let sampleDepth : f32 = textureSample(depthTexture, Sampler, offset.xy).r;
+ ^^^^^^^^^^^^^
+
@group(0) @binding(0) var Sampler : sampler;
@group(0) @binding(1) var randomTexture : texture_2d<f32>;
diff --git a/test/tint/bug/tint/1118.wgsl.expected.glsl b/test/tint/bug/tint/1118.wgsl.expected.glsl
index 1a8acc6..39bc459 100644
--- a/test/tint/bug/tint/1118.wgsl.expected.glsl
+++ b/test/tint/bug/tint/1118.wgsl.expected.glsl
@@ -1,3 +1,7 @@
+bug/tint/1118.wgsl:64:31 warning: 'dpdx' must only be called from uniform control flow
+ normalW = normalize(-(cross(dpdx(x_62), dpdy(x_64))));
+ ^^^^
+
#version 310 es
precision mediump float;
diff --git a/test/tint/bug/tint/1118.wgsl.expected.hlsl b/test/tint/bug/tint/1118.wgsl.expected.hlsl
index 117a578..467aea6 100644
--- a/test/tint/bug/tint/1118.wgsl.expected.hlsl
+++ b/test/tint/bug/tint/1118.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+bug/tint/1118.wgsl:64:31 warning: 'dpdx' must only be called from uniform control flow
+ normalW = normalize(-(cross(dpdx(x_62), dpdy(x_64))));
+ ^^^^
+
static float fClipDistance3 = 0.0f;
static float fClipDistance4 = 0.0f;
cbuffer cbuffer_x_29 : register(b0, space0) {
diff --git a/test/tint/bug/tint/1118.wgsl.expected.msl b/test/tint/bug/tint/1118.wgsl.expected.msl
index 06a0d7d..a8d14c5 100644
--- a/test/tint/bug/tint/1118.wgsl.expected.msl
+++ b/test/tint/bug/tint/1118.wgsl.expected.msl
@@ -1,3 +1,7 @@
+bug/tint/1118.wgsl:64:31 warning: 'dpdx' must only be called from uniform control flow
+ normalW = normalize(-(cross(dpdx(x_62), dpdy(x_64))));
+ ^^^^
+
#include <metal_stdlib>
using namespace metal;
diff --git a/test/tint/bug/tint/1118.wgsl.expected.spvasm b/test/tint/bug/tint/1118.wgsl.expected.spvasm
index 55d6f1d..26653b6 100644
--- a/test/tint/bug/tint/1118.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/1118.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+bug/tint/1118.wgsl:64:31 warning: 'dpdx' must only be called from uniform control flow
+ normalW = normalize(-(cross(dpdx(x_62), dpdy(x_64))));
+ ^^^^
+
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
diff --git a/test/tint/bug/tint/1118.wgsl.expected.wgsl b/test/tint/bug/tint/1118.wgsl.expected.wgsl
index 1274782..d2a3bc3 100644
--- a/test/tint/bug/tint/1118.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/1118.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+bug/tint/1118.wgsl:64:31 warning: 'dpdx' must only be called from uniform control flow
+ normalW = normalize(-(cross(dpdx(x_62), dpdy(x_64))));
+ ^^^^
+
struct Scene {
vEyePosition : vec4<f32>,
}
diff --git a/test/tint/bug/tint/943.spvasm.expected.glsl b/test/tint/bug/tint/943.spvasm.expected.glsl
index c488064..f581b5d 100644
--- a/test/tint/bug/tint/943.spvasm.expected.glsl
+++ b/test/tint/bug/tint/943.spvasm.expected.glsl
@@ -1,3 +1,4 @@
+warning: parameter 'dimInner' of 'mm_matMul_i1_i1_i1_' must be uniform
#version 310 es
struct Uniforms {
diff --git a/test/tint/bug/tint/943.spvasm.expected.hlsl b/test/tint/bug/tint/943.spvasm.expected.hlsl
index bdf9b12..89accf2 100644
--- a/test/tint/bug/tint/943.spvasm.expected.hlsl
+++ b/test/tint/bug/tint/943.spvasm.expected.hlsl
@@ -1,3 +1,4 @@
+warning: parameter 'dimInner' of 'mm_matMul_i1_i1_i1_' must be uniform
static int dimAOuter_1 = 0;
cbuffer cbuffer_x_48 : register(b3, space0) {
uint4 x_48[5];
diff --git a/test/tint/bug/tint/943.spvasm.expected.msl b/test/tint/bug/tint/943.spvasm.expected.msl
index 648486b..031bf63 100644
--- a/test/tint/bug/tint/943.spvasm.expected.msl
+++ b/test/tint/bug/tint/943.spvasm.expected.msl
@@ -1,3 +1,4 @@
+warning: parameter 'dimInner' of 'mm_matMul_i1_i1_i1_' must be uniform
#include <metal_stdlib>
using namespace metal;
diff --git a/test/tint/bug/tint/943.spvasm.expected.spvasm b/test/tint/bug/tint/943.spvasm.expected.spvasm
index 315f30e..228d7df 100644
--- a/test/tint/bug/tint/943.spvasm.expected.spvasm
+++ b/test/tint/bug/tint/943.spvasm.expected.spvasm
@@ -1,3 +1,4 @@
+warning: parameter 'dimInner' of 'mm_matMul_i1_i1_i1_' must be uniform
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
diff --git a/test/tint/bug/tint/943.spvasm.expected.wgsl b/test/tint/bug/tint/943.spvasm.expected.wgsl
index fe02e98..aa383da 100644
--- a/test/tint/bug/tint/943.spvasm.expected.wgsl
+++ b/test/tint/bug/tint/943.spvasm.expected.wgsl
@@ -1,3 +1,4 @@
+warning: parameter 'dimInner' of 'mm_matMul_i1_i1_i1_' must be uniform
struct Uniforms {
NAN : f32,
@size(12)
diff --git a/test/tint/bug/tint/948.wgsl.expected.glsl b/test/tint/bug/tint/948.wgsl.expected.glsl
index e5cd11e..a675fe9 100644
--- a/test/tint/bug/tint/948.wgsl.expected.glsl
+++ b/test/tint/bug/tint/948.wgsl.expected.glsl
@@ -1,3 +1,7 @@
+bug/tint/948.wgsl:146:33 warning: 'textureSampleBias' must only be called from uniform control flow
+ let x_217 : vec4<f32> = textureSampleBias(animationMapTexture, animationMapSampler, vec2<f32>(((x_208 + 0.5) / x_211), (0.125 * x_214)), 0.0);
+ ^^^^^^^^^^^^^^^^^
+
#version 310 es
precision mediump float;
diff --git a/test/tint/bug/tint/948.wgsl.expected.hlsl b/test/tint/bug/tint/948.wgsl.expected.hlsl
index 5a253ad..1869503 100644
--- a/test/tint/bug/tint/948.wgsl.expected.hlsl
+++ b/test/tint/bug/tint/948.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+bug/tint/948.wgsl:146:33 warning: 'textureSampleBias' must only be called from uniform control flow
+ let x_217 : vec4<f32> = textureSampleBias(animationMapTexture, animationMapSampler, vec2<f32>(((x_208 + 0.5) / x_211), (0.125 * x_214)), 0.0);
+ ^^^^^^^^^^^^^^^^^
+
cbuffer cbuffer_x_20 : register(b9, space2) {
uint4 x_20[8];
};
diff --git a/test/tint/bug/tint/948.wgsl.expected.msl b/test/tint/bug/tint/948.wgsl.expected.msl
index 25febaf..f799044 100644
--- a/test/tint/bug/tint/948.wgsl.expected.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.msl
@@ -1,3 +1,7 @@
+bug/tint/948.wgsl:146:33 warning: 'textureSampleBias' must only be called from uniform control flow
+ let x_217 : vec4<f32> = textureSampleBias(animationMapTexture, animationMapSampler, vec2<f32>(((x_208 + 0.5) / x_211), (0.125 * x_214)), 0.0);
+ ^^^^^^^^^^^^^^^^^
+
#include <metal_stdlib>
using namespace metal;
diff --git a/test/tint/bug/tint/948.wgsl.expected.spvasm b/test/tint/bug/tint/948.wgsl.expected.spvasm
index 51c0bb0..4f5102c 100644
--- a/test/tint/bug/tint/948.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/948.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+bug/tint/948.wgsl:146:33 warning: 'textureSampleBias' must only be called from uniform control flow
+ let x_217 : vec4<f32> = textureSampleBias(animationMapTexture, animationMapSampler, vec2<f32>(((x_208 + 0.5) / x_211), (0.125 * x_214)), 0.0);
+ ^^^^^^^^^^^^^^^^^
+
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
diff --git a/test/tint/bug/tint/948.wgsl.expected.wgsl b/test/tint/bug/tint/948.wgsl.expected.wgsl
index 3c11fec..791079b 100644
--- a/test/tint/bug/tint/948.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/948.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+bug/tint/948.wgsl:146:33 warning: 'textureSampleBias' must only be called from uniform control flow
+ let x_217 : vec4<f32> = textureSampleBias(animationMapTexture, animationMapSampler, vec2<f32>(((x_208 + 0.5) / x_211), (0.125 * x_214)), 0.0);
+ ^^^^^^^^^^^^^^^^^
+
struct LeftOver {
time : f32,
@size(12)
diff --git a/test/tint/bug/tint/949.wgsl.expected.glsl b/test/tint/bug/tint/949.wgsl.expected.glsl
index fbb9d32..e60090b 100644
--- a/test/tint/bug/tint/949.wgsl.expected.glsl
+++ b/test/tint/bug/tint/949.wgsl.expected.glsl
@@ -1,3 +1,7 @@
+bug/tint/949.wgsl:326:29 warning: 'textureSample' must only be called from uniform control flow
+ let x_397 : vec4<f32> = textureSample(TextureSamplerTexture, TextureSamplerSampler, (x_394 + x_395));
+ ^^^^^^^^^^^^^
+
#version 310 es
precision mediump float;
diff --git a/test/tint/bug/tint/949.wgsl.expected.hlsl b/test/tint/bug/tint/949.wgsl.expected.hlsl
index 4869166..2bdadf0 100644
--- a/test/tint/bug/tint/949.wgsl.expected.hlsl
+++ b/test/tint/bug/tint/949.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+bug/tint/949.wgsl:326:29 warning: 'textureSample' must only be called from uniform control flow
+ let x_397 : vec4<f32> = textureSample(TextureSamplerTexture, TextureSamplerSampler, (x_394 + x_395));
+ ^^^^^^^^^^^^^
+
struct lightingInfo {
float3 diffuse;
float3 specular;
diff --git a/test/tint/bug/tint/949.wgsl.expected.msl b/test/tint/bug/tint/949.wgsl.expected.msl
index fb89dec..0701a4a 100644
--- a/test/tint/bug/tint/949.wgsl.expected.msl
+++ b/test/tint/bug/tint/949.wgsl.expected.msl
@@ -1,3 +1,7 @@
+bug/tint/949.wgsl:326:29 warning: 'textureSample' must only be called from uniform control flow
+ let x_397 : vec4<f32> = textureSample(TextureSamplerTexture, TextureSamplerSampler, (x_394 + x_395));
+ ^^^^^^^^^^^^^
+
#include <metal_stdlib>
using namespace metal;
diff --git a/test/tint/bug/tint/949.wgsl.expected.spvasm b/test/tint/bug/tint/949.wgsl.expected.spvasm
index e2fce0a..287968a 100644
--- a/test/tint/bug/tint/949.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/949.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+bug/tint/949.wgsl:326:29 warning: 'textureSample' must only be called from uniform control flow
+ let x_397 : vec4<f32> = textureSample(TextureSamplerTexture, TextureSamplerSampler, (x_394 + x_395));
+ ^^^^^^^^^^^^^
+
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
diff --git a/test/tint/bug/tint/949.wgsl.expected.wgsl b/test/tint/bug/tint/949.wgsl.expected.wgsl
index c16bf96..48c21aa 100644
--- a/test/tint/bug/tint/949.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/949.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+bug/tint/949.wgsl:326:29 warning: 'textureSample' must only be called from uniform control flow
+ let x_397 : vec4<f32> = textureSample(TextureSamplerTexture, TextureSamplerSampler, (x_394 + x_395));
+ ^^^^^^^^^^^^^
+
struct lightingInfo {
diffuse : vec3<f32>,
specular : vec3<f32>,
