tint: uniformity: detect pointers assigned to in non-uniform control flow
Bug: tint:1558
Change-Id: Ia92258f1fb40b008a6052ce2ea5a20ec29351ce5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/93264
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/resolver/uniformity.cc b/src/tint/resolver/uniformity.cc
index 97612a4..2bbb487 100644
--- a/src/tint/resolver/uniformity.cc
+++ b/src/tint/resolver/uniformity.cc
@@ -305,7 +305,7 @@
/// @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) {
- return current_function_->CreateNode(tag, ast);
+ return current_function_->CreateNode(std::move(tag), ast);
}
/// Process a function.
@@ -1248,6 +1248,10 @@
if (func_info->parameters[i].pointer_may_become_non_uniform) {
ptr_result->AddEdge(current_function_->may_be_non_uniform);
} else {
+ // Add edge to the call to catch when it's called in non-uniform control
+ // flow.
+ ptr_result->AddEdge(call_node);
+
// 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) {
diff --git a/src/tint/resolver/uniformity_test.cc b/src/tint/resolver/uniformity_test.cc
index 7ed4a6b..d661099 100644
--- a/src/tint/resolver/uniformity_test.cc
+++ b/src/tint/resolver/uniformity_test.cc
@@ -4624,6 +4624,145 @@
)");
}
+TEST_F(UniformityAnalysisTest, PointerParamModifiedInNonUniformControlFlow) {
+ std::string src = R"(
+@binding(0) @group(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo(p : ptr<function, i32>) {
+ *p = 42;
+}
+
+@compute @workgroup_size(64)
+fn main() {
+ var a : i32;
+ if (non_uniform_global == 0) {
+ foo(&a);
+ }
+
+ if (a == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:16:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:11:3 note: control flow depends on non-uniform value
+ if (non_uniform_global == 0) {
+ ^^
+
+test:11:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
+ if (non_uniform_global == 0) {
+ ^^^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamAssumedModifiedInNonUniformControlFlow) {
+ std::string src = R"(
+@binding(0) @group(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo(p : ptr<function, i32>) {
+ // Do not modify 'p', uniformity analysis presently assumes it will be.
+}
+
+@compute @workgroup_size(64)
+fn main() {
+ var a : i32;
+ if (non_uniform_global == 0) {
+ foo(&a);
+ }
+
+ if (a == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:16:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:11:3 note: control flow depends on non-uniform value
+ if (non_uniform_global == 0) {
+ ^^
+
+test:11:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
+ if (non_uniform_global == 0) {
+ ^^^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamModifiedInNonUniformControlFlow_NestedCall) {
+ std::string src = R"(
+@binding(0) @group(0) var<storage, read_write> non_uniform_global : i32;
+
+fn foo2(p : ptr<function, i32>) {
+ *p = 42;
+}
+
+fn foo(p : ptr<function, i32>) {
+ foo2(p);
+}
+
+@compute @workgroup_size(64)
+fn main() {
+ var a : i32;
+ if (non_uniform_global == 0) {
+ foo(&a);
+ }
+
+ if (a == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:20:5 warning: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:15:3 note: control flow depends on non-uniform value
+ if (non_uniform_global == 0) {
+ ^^
+
+test:15:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
+ if (non_uniform_global == 0) {
+ ^^^^^^^^^^^^^^^^^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, PointerParamModifiedInUniformControlFlow) {
+ std::string src = R"(
+@binding(0) @group(0) var<uniform> uniform_global : i32;
+
+fn foo(p : ptr<function, i32>) {
+ *p = 42;
+}
+
+@compute @workgroup_size(64)
+fn main() {
+ var a : i32;
+ if (uniform_global == 0) {
+ foo(&a);
+ }
+
+ if (a == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, true);
+}
+
TEST_F(UniformityAnalysisTest, NonUniformPointerParameterBecomesUniform_AfterUse) {
std::string src = R"(
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
