| // 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())); |
| std::string 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. |
| std::string 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. |
| std::string 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:16:3 note: calling 'foo' may cause subsequent control flow to be non-uniform |
| foo(); |
| ^^^ |
| |
| test:7:3 note: control flow depends on non-uniform value |
| if (rw == 0) { |
| ^^ |
| |
| test:7:7 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| if (rw == 0) { |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, SubsequentControlFlowMayBeNonUniform_Nested_Fail) { |
| // Indirectly call a function that causes subsequent control flow to be non-uniform, and then |
| // call another function that requires uniformity. |
| // The lack of return statement in `foo()` requires that we implicitly add an edge from |
| // CF_return to that last control flow node of the function. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> rw : i32; |
| |
| var<private> p : i32; |
| |
| fn bar() { |
| if (rw == 0) { |
| p = 42; |
| return; |
| } |
| p = 5; |
| return; |
| } |
| |
| fn foo() { |
| bar(); |
| } |
| |
| fn main() { |
| foo(); |
| workgroupBarrier(); |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:21:3 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:20:3 note: calling 'foo' may cause subsequent control flow to be non-uniform |
| foo(); |
| ^^^ |
| |
| test:16:3 note: calling 'bar' may cause subsequent control flow to be non-uniform |
| bar(); |
| ^^^ |
| |
| test:7:3 note: control flow depends on non-uniform value |
| if (rw == 0) { |
| ^^ |
| |
| test:7:7 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| if (rw == 0) { |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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. |
| std::string 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:6:5 note: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:11:7 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| 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. |
| std::string 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. |
| std::string 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:9:3 note: control flow depends on non-uniform value |
| if (foo(rw) == 7) { |
| ^^ |
| |
| test:9:11 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| if (foo(rw) == 7) { |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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:16:7 note: non-uniform function call argument causes subsequent control flow to be non-uniform |
| foo(rw); |
| ^^ |
| |
| test:7:3 note: control flow depends on non-uniform value |
| if (i == 0) { |
| ^^ |
| |
| test:7:7 note: reading from 'i' may result in a non-uniform value |
| if (i == 0) { |
| ^ |
| |
| test:16:7 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| foo(rw); |
| ^^ |
| )"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// 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) { |
| std::string src = R"( |
| @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:4:3 note: control flow depends on non-uniform value |
| if (all(vec3(b) == vec3(0u))) { |
| ^^ |
| |
| test:4:16 note: reading from builtin 'b' may result in a non-uniform value |
| if (all(vec3(b) == vec3(0u))) { |
| ^ |
| )"); |
| } |
| } |
| |
| TEST_P(ComputeBuiltin, InStruct) { |
| std::string src = R"( |
| struct S { |
| @builtin()" + GetParam().name + |
| R"() b : )" + GetParam().type + R"( |
| } |
| |
| @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(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:8:3 note: control flow depends on non-uniform value |
| if (all(vec3(s.b) == vec3(0u))) { |
| ^^ |
| |
| test:8:16 note: reading from 's' may result in a non-uniform value |
| if (all(vec3(s.b) == vec3(0u))) { |
| ^ |
| )"); |
| } |
| } |
| |
| 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. |
| std::string src = R"( |
| struct S { |
| @builtin(num_workgroups) num_groups : vec3<u32>, |
| @builtin(local_invocation_index) idx : u32, |
| } |
| |
| @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(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:9:3 note: control flow depends on non-uniform value |
| if (s.num_groups.x == 0u) { |
| ^^ |
| |
| test:9:7 note: reading from 's' may result in a non-uniform value |
| if (s.num_groups.x == 0u) { |
| ^ |
| )"); |
| } |
| |
| class FragmentBuiltin : public UniformityAnalysisTestBase, |
| public ::testing::TestWithParam<BuiltinEntry> {}; |
| TEST_P(FragmentBuiltin, AsParam) { |
| std::string src = R"( |
| @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:4:3 note: control flow depends on non-uniform value |
| if (u32(vec4(b).x) == 0u) { |
| ^^ |
| |
| test:4:16 note: reading from builtin 'b' may result in a non-uniform value |
| if (u32(vec4(b).x) == 0u) { |
| ^ |
| )"); |
| } |
| } |
| |
| TEST_P(FragmentBuiltin, InStruct) { |
| std::string src = R"( |
| struct S { |
| @builtin()" + GetParam().name + |
| R"() b : )" + GetParam().type + R"( |
| } |
| |
| @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); |
| ^^^^ |
| |
| test:8:3 note: control flow depends on non-uniform value |
| if (u32(vec4(s.b).x) == 0u) { |
| ^^ |
| |
| test:8:16 note: reading from 's' may result in a non-uniform value |
| if (u32(vec4(s.b).x) == 0u) { |
| ^ |
| )"); |
| } |
| } |
| |
| 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) { |
| std::string src = R"( |
| @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:4:3 note: control flow depends on non-uniform value |
| if (l == 0.0) { |
| ^^ |
| |
| test:4:7 note: reading from user-defined input 'l' may result in a non-uniform value |
| if (l == 0.0) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, FragmentLocation_InStruct) { |
| std::string src = R"( |
| struct S { |
| @location(0) l : f32 |
| } |
| |
| @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:8:3 note: control flow depends on non-uniform value |
| if (s.l == 0.0) { |
| ^^ |
| |
| test:8:7 note: reading from 's' may result in a non-uniform value |
| if (s.l == 0.0) { |
| ^ |
| )"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// 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())); |
| std::string 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_THAT( |
| error_, |
| ::testing::StartsWith( |
| R"(test:13:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier();)")); |
| EXPECT_THAT(error_, |
| ::testing::HasSubstr("test:14:9 note: reading from read_write storage buffer " |
| "'nonuniform_var' may result in a non-uniform value")); |
| } 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())); |
| std::string 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_THAT( |
| error_, |
| ::testing::StartsWith( |
| R"(test:14:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier();)")); |
| EXPECT_THAT(error_, |
| ::testing::HasSubstr("test:13:9 note: reading from read_write storage buffer " |
| "'nonuniform_var' may result in a non-uniform value")); |
| } 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())); |
| std::string 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_THAT( |
| error_, |
| ::testing::StartsWith( |
| R"(test:15:7 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier();)")); |
| EXPECT_THAT(error_, |
| ::testing::HasSubstr("test:13:9 note: reading from read_write storage buffer " |
| "'nonuniform_var' may result in a non-uniform value")); |
| } else { |
| RunTest(src, true); |
| } |
| } |
| |
| TEST_F(UniformityAnalysisTest, Loop_CallInBody_UniformBreakInContinuing) { |
| std::string 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) { |
| std::string 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:10:7 note: control flow depends on non-uniform value |
| if (i == n) { |
| ^^ |
| |
| test:10:16 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| if (i == n) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Loop_CallInContinuing_UniformBreakInContinuing) { |
| std::string 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) { |
| std::string 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(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:10:7 note: control flow depends on non-uniform value |
| if (i == n) { |
| ^^ |
| |
| test:10:16 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| if (i == n) { |
| ^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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. |
| std::string 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:13:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:8:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:14:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:8:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:19:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:11:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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:19:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:7:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:15:7 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:7:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:14:13 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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) { |
| std::string 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. |
| std::string 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. |
| std::string 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. |
| std::string 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:12:5 note: calling 'bar' may cause subsequent control flow to be non-uniform |
| bar(); |
| ^^^ |
| |
| test:5:3 note: control flow depends on non-uniform value |
| if (n == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| if (n == 42) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ForLoop_CallInside_UniformCondition) { |
| std::string 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) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| for (var i = 0; i < n; i = i + 1) { |
| ^^^ |
| |
| test:5:23 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| for (var i = 0; i < n; i = i + 1) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ForLoop_CallInside_InitializerCausesNonUniformFlow) { |
| std::string 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:13:16 note: calling 'bar' may cause subsequent control flow to be non-uniform |
| for (var i = bar(); i < 10; i = i + 1) { |
| ^^^ |
| |
| test:5:3 note: control flow depends on non-uniform value |
| if (n == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| if (n == 42) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ForLoop_CallInside_ContinuingCausesNonUniformFlow) { |
| std::string 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:13:35 note: calling 'bar' may cause subsequent control flow to be non-uniform |
| for (var i = 0; i < 10; i = i + bar()) { |
| ^^^ |
| |
| test:5:3 note: control flow depends on non-uniform value |
| if (n == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| if (n == 42) { |
| ^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:6:31 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| for (var i = 0; i < 10; v = non_uniform) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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:9:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:6:31 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| for (var i = 0; i < 10; v = non_uniform) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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:14:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:8:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:20:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:7:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_CallInside_UniformCondition) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read> n : i32; |
| |
| fn foo() { |
| var i = 0; |
| while (i < n) { |
| workgroupBarrier(); |
| i = i + 1; |
| } |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_CallInside_NonUniformCondition) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> n : i32; |
| |
| fn foo() { |
| var i = 0; |
| while (i < n) { |
| workgroupBarrier(); |
| i = i + 1; |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:7:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:6:3 note: control flow depends on non-uniform value |
| while (i < n) { |
| ^^^^^ |
| |
| test:6:14 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| while (i < n) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_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. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn foo() { |
| var v = 0; |
| var i = 0; |
| while (i < 10) { |
| if (v == 0) { |
| workgroupBarrier(); |
| break; |
| } |
| |
| v = non_uniform; |
| i++; |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:9:7 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:8:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:13:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_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. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn foo() { |
| var v = 0; |
| var i = 0; |
| while (i < 10) { |
| if (true) { |
| v = non_uniform; |
| break; |
| } |
| if (v == 0) { |
| workgroupBarrier(); |
| } |
| i++; |
| } |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_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. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn foo() { |
| var v = 0; |
| var i = 0; |
| while (i < 10) { |
| if (true) { |
| v = non_uniform; |
| break; |
| } |
| v = 5; |
| i++; |
| } |
| |
| if (v == 0) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:17:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:16:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:9:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_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. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn foo() { |
| var v = 0; |
| var i = 0; |
| while (i < 10) { |
| if (true) { |
| v = 5; |
| break; |
| } |
| |
| v = non_uniform; |
| |
| if (true) { |
| v = 6; |
| break; |
| } |
| i++; |
| } |
| |
| if (v == 0) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:23:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:22:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:13:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_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. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn foo() { |
| var v = 0; |
| var i = 0; |
| while (i < 10) { |
| if (v == 0) { |
| workgroupBarrier(); |
| break; |
| } |
| |
| v = non_uniform; |
| if (true) { |
| continue; |
| } |
| |
| v = 5; |
| i++; |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:9:7 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:8:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:13:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_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. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn foo() { |
| var v = 0; |
| var i = 0; |
| while (i < 10) { |
| if (v == 0) { |
| workgroupBarrier(); |
| break; |
| } |
| |
| v = non_uniform; |
| if (true) { |
| continue; |
| } |
| |
| v = 5; |
| i++; |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:9:7 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:8:5 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:13:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, While_NonUniformCondition_Reconverge) { |
| // Loops reconverge at exit, so test that we can call workgroupBarrier() after a loop that has a |
| // non-uniform condition. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> n : i32; |
| |
| fn foo() { |
| var i = 0; |
| while (i < n) { |
| } |
| workgroupBarrier(); |
| i = i + 1; |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| } // namespace LoopTest |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// If-else statement tests. |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_F(UniformityAnalysisTest, IfElse_UniformCondition_BarrierInTrueBlock) { |
| std::string 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) { |
| std::string 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) { |
| std::string 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) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, IfElse_NonUniformCondition_BarrierInElseBlock) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformLHS_And) { |
| std::string 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:7:34 note: control flow depends on non-uniform value |
| if ((non_uniform_global == 42) && false) { |
| ^^ |
| |
| test:7:8 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if ((non_uniform_global == 42) && false) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformRHS_And) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (false && (non_uniform_global == 42)) { |
| ^^ |
| |
| test:7:17 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (false && (non_uniform_global == 42)) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformLHS_Or) { |
| std::string 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:7:34 note: control flow depends on non-uniform value |
| if ((non_uniform_global == 42) || true) { |
| ^^ |
| |
| test:7:8 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if ((non_uniform_global == 42) || true) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, IfElse_ShortCircuitingCondition_NonUniformRHS_Or) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (true || (non_uniform_global == 42)) { |
| ^^ |
| |
| test:7:16 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (true || (non_uniform_global == 42)) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, IfElse_NonUniformCondition_BarrierInElseIfBlock) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:6:7 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| v = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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. |
| std::string 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:14:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:8:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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. |
| std::string 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:11:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:6:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:12:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:6:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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. |
| std::string 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. |
| std::string 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. |
| std::string 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. |
| std::string 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. |
| std::string 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. |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, IfElse_NonUniformDiscard_NoReconverge) { |
| // If statements should not reconverge after non-uniform discards. |
| std::string 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(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:5:3 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// Switch statement tests. |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_F(UniformityAnalysisTest, Switch_NonUniformCondition_BarrierInCase) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| switch (non_uniform) { |
| ^^^^^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| switch (non_uniform) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_NonUniformCondition_BarrierInDefault) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| switch (non_uniform) { |
| ^^^^^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| switch (non_uniform) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_NonUniformBreak) { |
| std::string 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:8:7 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:8:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_NonUniformBreakInDifferentCase) { |
| std::string 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) { |
| std::string 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:8:7 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:8:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_VarBecomesNonUniformInDifferentCase_WithBreak) { |
| std::string 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) { |
| std::string 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:13:7 note: control flow depends on non-uniform value |
| if (x == 0) { |
| ^^ |
| |
| test:9:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| x = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInDifferentCase_WithBreak) { |
| std::string 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:13:7 note: control flow depends on non-uniform value |
| if (x == 0) { |
| ^^ |
| |
| test:6:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var x = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInDifferentCase_WithFallthrough) { |
| std::string 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) { |
| std::string 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:18:3 note: control flow depends on non-uniform value |
| if (x == 0) { |
| ^^ |
| |
| test:9:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| x = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInAllCases_BarrierAfter) { |
| std::string 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) { |
| std::string 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:17:3 note: control flow depends on non-uniform value |
| if (x == 0) { |
| ^^ |
| |
| test:6:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var x = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_VarBecomesUniformInCasesThatDontReturn_BarrierAfter) { |
| std::string 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) { |
| std::string 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:20:3 note: control flow depends on non-uniform value |
| if (x == 0) { |
| ^^ |
| |
| test:6:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var x = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_NestedInLoop_VarBecomesNonUniformWithBreak_BarrierInLoop) { |
| std::string 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:8:5 note: control flow depends on non-uniform value |
| if (x == 0) { |
| ^^ |
| |
| test:15:13 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| x = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Switch_NestedInLoop_VarBecomesNonUniformWithBreak_BarrierAfterLoop) { |
| std::string 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. |
| std::string 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. |
| std::string 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. |
| std::string 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. |
| std::string 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(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:13:7 note: calling 'bar' may cause subsequent control flow to be non-uniform |
| bar(); |
| ^^^ |
| |
| test:5:3 note: control flow depends on non-uniform value |
| if (n == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'n' may result in a non-uniform value |
| if (n == 42) { |
| ^ |
| )"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// Pointer tests. |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_F(UniformityAnalysisTest, AssignNonUniformThroughPointer) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:6:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| *&v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignNonUniformThroughCapturedPointer) { |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:7:9 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| *pv = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignUniformThroughPointer) { |
| std::string 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) { |
| std::string 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) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (non_uniform == 0) { |
| ^^ |
| |
| test:7:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 0) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, LoadNonUniformThroughPointer) { |
| std::string 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:6:3 note: control flow depends on non-uniform value |
| if (*&v == 0) { |
| ^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, LoadNonUniformThroughCapturedPointer) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (*pv == 0) { |
| ^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, LoadNonUniformThroughPointerParameter) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn bar(p : ptr<function, i32>) { |
| if (*p == 0) { |
| workgroupBarrier(); |
| } |
| } |
| |
| fn foo() { |
| var v = non_uniform; |
| bar(&v); |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:12:7 warning: parameter 'p' of 'bar' must be uniform |
| bar(&v); |
| ^ |
| |
| test:6:5 note: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:11:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, LoadUniformThroughPointer) { |
| std::string src = R"( |
| fn foo() { |
| var v = 42; |
| if (*&v == 0) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, LoadUniformThroughCapturedPointer) { |
| std::string src = R"( |
| fn foo() { |
| var v = 42; |
| let pv = &v; |
| if (*pv == 0) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, LoadUniformThroughPointerParameter) { |
| std::string src = R"( |
| fn bar(p : ptr<function, i32>) { |
| if (*p == 0) { |
| workgroupBarrier(); |
| } |
| } |
| |
| fn foo() { |
| var v = 42; |
| bar(&v); |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, StoreNonUniformAfterCapturingPointer) { |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (*pv == 0) { |
| ^^ |
| |
| test:7:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, StoreUniformAfterCapturingPointer) { |
| std::string 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) { |
| std::string 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:9:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:8:14 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| *&*&*pv2 = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, LoadNonUniformThroughLongChainOfPointers) { |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (*&*&*pv2 == 0) { |
| ^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignUniformThenNonUniformThroughDifferentPointer) { |
| std::string 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:10:3 note: control flow depends on non-uniform value |
| if (*pv1 == 0) { |
| ^^ |
| |
| test:9:10 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| *pv2 = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignNonUniformThenUniformThroughDifferentPointer) { |
| std::string 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) { |
| std::string 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:10:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:8:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, UnmodifiedPointerParameterUniform) { |
| std::string 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) { |
| std::string 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:11:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:10:7 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&v); |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignUniformThroughPointerInFunctionCall) { |
| std::string 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) { |
| std::string 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:11:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:10:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| bar(&v, non_uniform); |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignNonUniformThroughPointerInFunctionCallViaPointerArg) { |
| std::string 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:12:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:10:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var a = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignUniformThroughPointerInFunctionCallViaArg) { |
| std::string 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) { |
| std::string 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) { |
| std::string 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:20:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:18:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var a = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, AssignUniformThroughPointerInFunctionCallChain) { |
| std::string 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) { |
| std::string 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) { |
| std::string 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:17:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:16:7 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&v); |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, PointerParamAssignNonUniformInTrueAndUniformInFalse) { |
| std::string 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:15:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:14:7 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&v); |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndReturn) { |
| std::string 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:15:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:14:7 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&v); |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndBreakFromSwitch) { |
| std::string 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:23:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:22:7 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&v); |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndBreakFromLoop) { |
| std::string 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:17:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:16:7 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&v); |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ConditionalAssignNonUniformToPointerParamAndContinue) { |
| std::string 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:6:5 note: control flow depends on non-uniform value |
| if (*p == 0) { |
| ^^ |
| |
| test:12:12 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| *p = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, PointerParamMaybeBecomesUniform) { |
| std::string 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:14:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:12:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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; |
| |
| 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:13:3 note: control flow depends on non-uniform value |
| if (b == 0) { |
| ^^ |
| |
| test:10:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var a = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, NonUniformPointerParameterBecomesUniform_BeforeUse) { |
| std::string 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) { |
| std::string 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:13:3 note: control flow depends on non-uniform value |
| if (b == 0) { |
| ^^ |
| |
| test:12:11 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&a, &b); |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, UniformPointerParameterBecomesNonUniform_AfterUse) { |
| std::string 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) { |
| std::string 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:11:3 note: control flow depends on non-uniform value |
| if (v == 1) { |
| ^^ |
| |
| test:9:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var v = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:13:3 note: control flow depends on non-uniform value |
| if (b == 0) { |
| ^^ |
| |
| test:12:11 note: pointer contents may become non-uniform after calling 'bar' |
| bar(&a, &b); |
| ^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:15:3 note: control flow depends on non-uniform value |
| if (c == 0) { |
| ^^ |
| |
| test:11:11 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| var a = non_uniform; |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| 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::ParameterList 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 |
| note: control flow depends on non-uniform value |
| note: reading from module-scope private variable 'non_uniform_global' may result in a non-uniform value)"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// Tests to cover access to aggregate types. |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_F(UniformityAnalysisTest, VectorElement_Uniform) { |
| std::string 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) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (v[2] == 0) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'v' may result in a non-uniform value |
| if (v[2] == 0) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, VectorElement_BecomesNonUniform_BeforeCondition) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (v[2] == 0) { |
| ^^ |
| |
| test:6:10 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| v[2] = rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, VectorElement_BecomesNonUniform_AfterCondition) { |
| std::string 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) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (v[2] == 0) { |
| ^^ |
| |
| test:6:10 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| v[1] = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (v[1] == 0) { |
| ^^ |
| |
| test:6:10 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| v[1] = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (v[1] == 0) { |
| ^^ |
| |
| test:6:10 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| v[1] = rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, VectorElement_NonUniform_AnyBuiltin) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (any(v == vec4(42))) { |
| ^^ |
| |
| test:6:10 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| v[1] = non_uniform_global; |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, StructMember_Uniform) { |
| std::string 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) { |
| std::string 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:9:3 note: control flow depends on non-uniform value |
| if (s.b == 0) { |
| ^^ |
| |
| test:9:7 note: reading from read_write storage buffer 's' may result in a non-uniform value |
| if (s.b == 0) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, StructMember_BecomesNonUniform_BeforeCondition) { |
| std::string 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:11:3 note: control flow depends on non-uniform value |
| if (s.b == 0) { |
| ^^ |
| |
| test:10:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| s.b = rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, StructMember_BecomesNonUniform_AfterCondition) { |
| std::string 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) { |
| std::string 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:11:3 note: control flow depends on non-uniform value |
| if (s.b == 0) { |
| ^^ |
| |
| test:10:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| s.a = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:12:3 note: control flow depends on non-uniform value |
| if (s.a == 0) { |
| ^^ |
| |
| test:10:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| s.a = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:12:3 note: control flow depends on non-uniform value |
| if (s.a == 0) { |
| ^^ |
| |
| test:10:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| s.a = rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ArrayElement_Uniform) { |
| std::string 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) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (arr[7] == 0) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'arr' may result in a non-uniform value |
| if (arr[7] == 0) { |
| ^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ArrayElement_BecomesNonUniform_BeforeCondition) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (arr[2] == 0) { |
| ^^ |
| |
| test:6:12 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| arr[2] = rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ArrayElement_BecomesNonUniform_AfterCondition) { |
| std::string 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) { |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (arr[2] == 0) { |
| ^^ |
| |
| test:6:12 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| arr[1] = rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ArrayElement_DifferentElementBecomesNonUniformThroughPointer) { |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (arr[2] == 0) { |
| ^^ |
| |
| test:7:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| *pa = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (arr[1] == 0) { |
| ^^ |
| |
| test:6:12 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| arr[1] = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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:8:3 note: control flow depends on non-uniform value |
| if (arr[1] == 0) { |
| ^^ |
| |
| test:6:12 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| arr[1] = rw; |
| ^^ |
| )"); |
| } |
| |
| 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. |
| std::string 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(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:9:3 note: control flow depends on non-uniform value |
| if (arr[1] == 0) { |
| ^^ |
| |
| test:7:12 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| arr[1] = rw; |
| ^^ |
| )"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// Miscellaneous statement and expression tests. |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_F(UniformityAnalysisTest, FunctionRequiresUniformFlowAndCausesNonUniformFlow) { |
| // Test that a function that requires uniform flow and then causes non-uniform flow can be |
| // called without error. |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform_global : i32; |
| |
| fn foo() { |
| _ = dpdx(0.5); |
| |
| if (non_uniform_global == 0) { |
| discard; |
| } |
| } |
| |
| @fragment |
| fn main() { |
| foo(); |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, TypeConstructor) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (i32(non_uniform_global) == 0) { |
| ^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (i32(non_uniform_global) == 0) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Conversion) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (f32(non_uniform_global) == 0.0) { |
| ^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (f32(non_uniform_global) == 0.0) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Bitcast) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (bitcast<f32>(non_uniform_global) == 0.0) { |
| ^^ |
| |
| test:5:20 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (bitcast<f32>(non_uniform_global) == 0.0) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, CompoundAssignment_NonUniformRHS) { |
| // Use compound assignment with a non-uniform RHS on a variable. |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:6:8 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| v += rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, CompoundAssignment_UniformRHS_StillNonUniform) { |
| // Use compound assignment with a uniform RHS on a variable that is already non-uniform. |
| std::string 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:7:3 note: control flow depends on non-uniform value |
| if (v == 0) { |
| ^^ |
| |
| test:5:11 note: reading from read_write storage buffer 'rw' may result in a non-uniform value |
| var v = rw; |
| ^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, PhonyAssignment_LhsCausesNonUniformControlFlow) { |
| std::string 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:13:7 note: calling 'bar' may cause subsequent control flow to be non-uniform |
| _ = bar(); |
| ^^^ |
| |
| test:5:3 note: control flow depends on non-uniform value |
| if (nonuniform_var == 42) { |
| ^^ |
| |
| test:5:7 note: reading from read_write storage buffer 'nonuniform_var' may result in a non-uniform value |
| if (nonuniform_var == 42) { |
| ^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ShortCircuiting_NoReconvergeLHS) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform_global : i32; |
| |
| var<private> p : i32; |
| |
| fn non_uniform_discard_func() -> bool { |
| if (non_uniform_global == 42) { |
| discard; |
| } |
| return false; |
| } |
| |
| fn main() { |
| let b = non_uniform_discard_func() && false; |
| workgroupBarrier(); |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:15:3 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:14:11 note: calling 'non_uniform_discard_func' may cause subsequent control flow to be non-uniform |
| let b = non_uniform_discard_func() && false; |
| ^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| test:7:3 note: control flow depends on non-uniform value |
| if (non_uniform_global == 42) { |
| ^^ |
| |
| test:7:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (non_uniform_global == 42) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ShortCircuiting_NoReconvergeRHS) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform_global : i32; |
| |
| var<private> p : i32; |
| |
| fn non_uniform_discard_func() -> bool { |
| if (non_uniform_global == 42) { |
| discard; |
| } |
| return false; |
| } |
| |
| fn main() { |
| let b = false && non_uniform_discard_func(); |
| workgroupBarrier(); |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:15:3 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:14:20 note: calling 'non_uniform_discard_func' may cause subsequent control flow to be non-uniform |
| let b = false && non_uniform_discard_func(); |
| ^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| test:7:3 note: control flow depends on non-uniform value |
| if (non_uniform_global == 42) { |
| ^^ |
| |
| test:7:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (non_uniform_global == 42) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ShortCircuiting_NoReconvergeBoth) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform_global : i32; |
| |
| var<private> p : i32; |
| |
| fn non_uniform_discard_func() -> bool { |
| if (non_uniform_global == 42) { |
| discard; |
| } |
| return false; |
| } |
| |
| fn main() { |
| let b = non_uniform_discard_func() && non_uniform_discard_func(); |
| workgroupBarrier(); |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:15:3 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:14:41 note: calling 'non_uniform_discard_func' may cause subsequent control flow to be non-uniform |
| let b = non_uniform_discard_func() && non_uniform_discard_func(); |
| ^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| test:7:3 note: control flow depends on non-uniform value |
| if (non_uniform_global == 42) { |
| ^^ |
| |
| test:7:7 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value |
| if (non_uniform_global == 42) { |
| ^^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ShortCircuiting_ReconvergeLHS) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform_global : i32; |
| |
| var<private> p : i32; |
| |
| fn uniform_discard_func() -> bool { |
| if (true) { |
| discard; |
| } |
| return false; |
| } |
| |
| fn main() { |
| let b = uniform_discard_func() && false; |
| workgroupBarrier(); |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ShortCircuiting_ReconvergeRHS) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform_global : i32; |
| |
| var<private> p : i32; |
| |
| fn uniform_discard_func() -> bool { |
| if (true) { |
| discard; |
| } |
| return false; |
| } |
| |
| fn main() { |
| let b = false && uniform_discard_func(); |
| workgroupBarrier(); |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, ShortCircuiting_ReconvergeBoth) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform_global : i32; |
| |
| var<private> p : i32; |
| |
| fn uniform_discard_func() -> bool { |
| if (true) { |
| discard; |
| } |
| return false; |
| } |
| |
| fn main() { |
| let b = uniform_discard_func() && uniform_discard_func(); |
| workgroupBarrier(); |
| } |
| )"; |
| |
| RunTest(src, true); |
| } |
| |
| TEST_F(UniformityAnalysisTest, DeadCode_AfterReturn) { |
| // Dead code after a return statement shouldn't cause uniformity errors. |
| std::string 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. |
| std::string 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) { |
| std::string 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) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (atomicAdd(&a, 1) == 1) { |
| ^^ |
| |
| test:5:18 note: reading from workgroup storage variable 'a' may result in a non-uniform value |
| if (atomicAdd(&a, 1) == 1) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, StorageAtomics) { |
| std::string 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:5:3 note: control flow depends on non-uniform value |
| if (atomicAdd(&a, 1) == 1) { |
| ^^ |
| |
| test:5:18 note: reading from read_write storage buffer 'a' may result in a non-uniform value |
| if (atomicAdd(&a, 1) == 1) { |
| ^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, DisableAnalysisWithExtension) { |
| std::string 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); |
| } |
| |
| TEST_F(UniformityAnalysisTest, StressGraphTraversalDepth) { |
| // Create a function with a very long sequence of variable declarations and assignments to |
| // test traversals of very deep graphs. This requires a non-recursive traversal algorithm. |
| ProgramBuilder b; |
| auto& ty = b.ty; |
| |
| // var<private> v0 : i32 = 0i; |
| // fn foo() { |
| // let v1 = v0; |
| // let v2 = v1; |
| // ... |
| // let v{N} = v{N-1}; |
| // if (v{N} == 0) { |
| // workgroupBarrier(); |
| // } |
| // } |
| b.Global("v0", ty.i32(), ast::StorageClass::kPrivate, b.Expr(0_i)); |
| ast::StatementList foo_body; |
| std::string v_last = "v0"; |
| for (int i = 1; i < 100000; i++) { |
| auto v = "v" + std::to_string(i); |
| foo_body.push_back(b.Decl(b.Var(v, nullptr, b.Expr(v_last)))); |
| v_last = v; |
| } |
| foo_body.push_back(b.If(b.Equal(v_last, 0_i), b.Block(b.CallStmt(b.Call("workgroupBarrier"))))); |
| b.Func("foo", {}, ty.void_(), foo_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 |
| note: control flow depends on non-uniform value |
| note: reading from module-scope private variable 'v0' may result in a non-uniform value)"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| /// Tests for the quality of the error messages produced by the analysis. |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_F(UniformityAnalysisTest, Error_CallUserThatCallsBuiltinDirectly) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn foo() { |
| workgroupBarrier(); |
| } |
| |
| fn main() { |
| if (non_uniform == 42) { |
| foo(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:10:5 warning: 'foo' must only be called from uniform control flow |
| foo(); |
| ^^^ |
| |
| test:5:3 note: 'foo' requires uniformity because it calls workgroupBarrier |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:9:3 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:9:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Error_CallUserThatCallsBuiltinIndirectly) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn zoo() { |
| workgroupBarrier(); |
| } |
| |
| fn bar() { |
| zoo(); |
| } |
| |
| fn foo() { |
| bar(); |
| } |
| |
| fn main() { |
| if (non_uniform == 42) { |
| foo(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:18:5 warning: 'foo' must only be called from uniform control flow |
| foo(); |
| ^^^ |
| |
| test:5:3 note: 'foo' requires uniformity because it indirectly calls workgroupBarrier |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:17:3 note: control flow depends on non-uniform value |
| if (non_uniform == 42) { |
| ^^ |
| |
| test:17:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| if (non_uniform == 42) { |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Error_ParametersRequireUniformityInChain) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn zoo(a : i32) { |
| if (a == 42) { |
| workgroupBarrier(); |
| } |
| } |
| |
| fn bar(b : i32) { |
| zoo(b); |
| } |
| |
| fn foo(c : i32) { |
| bar(c); |
| } |
| |
| fn main() { |
| foo(non_uniform); |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:19:7 warning: parameter 'c' of 'foo' must be uniform |
| foo(non_uniform); |
| ^^^^^^^^^^^ |
| |
| test:15:7 note: parameter 'b' of 'bar' must be uniform |
| bar(c); |
| ^ |
| |
| test:11:7 note: parameter 'a' of 'zoo' must be uniform |
| zoo(b); |
| ^ |
| |
| test:6:5 note: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:19:7 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value |
| foo(non_uniform); |
| ^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Error_ReturnValueMayBeNonUniformChain) { |
| std::string src = R"( |
| @group(0) @binding(0) var<storage, read_write> non_uniform : i32; |
| |
| fn zoo() -> i32 { |
| return non_uniform; |
| } |
| |
| fn bar() -> i32 { |
| return zoo(); |
| } |
| |
| fn foo() -> i32 { |
| return bar(); |
| } |
| |
| fn main() { |
| if (foo() == 42) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:18:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:17:3 note: control flow depends on non-uniform value |
| if (foo() == 42) { |
| ^^ |
| |
| test:17:7 note: return value of 'foo' may be non-uniform |
| if (foo() == 42) { |
| ^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Error_SubsequentControlFlowMayBeNonUniform) { |
| // Make sure we correctly identify the function call as the source of non-uniform control flow |
| // and not the if statement with the uniform condition. |
| std::string src = R"( |
| @group(0) @binding(0) var<uniform> uniform_value : i32; |
| @group(0) @binding(1) var<storage, read_write> non_uniform_value : i32; |
| |
| fn foo() -> i32 { |
| if (non_uniform_value == 0) { |
| return 5; |
| } |
| return 6; |
| } |
| |
| fn main() { |
| foo(); |
| if (uniform_value == 42) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:15:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:13:3 note: calling 'foo' may cause subsequent control flow to be non-uniform |
| foo(); |
| ^^^ |
| |
| test:6:3 note: control flow depends on non-uniform value |
| if (non_uniform_value == 0) { |
| ^^ |
| |
| test:6:7 note: reading from read_write storage buffer 'non_uniform_value' may result in a non-uniform value |
| if (non_uniform_value == 0) { |
| ^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Error_ParameterRequiredToBeUniformForSubsequentControlFlow) { |
| // Make sure we correctly identify the function call as the source of non-uniform control flow |
| // and not the if statement with the uniform condition. |
| std::string src = R"( |
| @group(0) @binding(0) var<uniform> uniform_value : i32; |
| @group(0) @binding(1) var<storage, read_write> non_uniform_value : i32; |
| |
| fn foo(x : i32) -> i32 { |
| if (x == 0) { |
| return 5; |
| } |
| return 6; |
| } |
| |
| fn main() { |
| foo(non_uniform_value); |
| if (uniform_value == 42) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:15:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:13:7 note: non-uniform function call argument causes subsequent control flow to be non-uniform |
| foo(non_uniform_value); |
| ^^^^^^^^^^^^^^^^^ |
| |
| test:6:3 note: control flow depends on non-uniform value |
| if (x == 0) { |
| ^^ |
| |
| test:6:7 note: reading from 'x' may result in a non-uniform value |
| if (x == 0) { |
| ^ |
| |
| test:13:7 note: reading from read_write storage buffer 'non_uniform_value' may result in a non-uniform value |
| foo(non_uniform_value); |
| ^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| TEST_F(UniformityAnalysisTest, Error_ShortCircuitingExprCausesNonUniformControlFlow) { |
| // Make sure we correctly identify the short-circuit as the source of non-uniform control flow |
| // and not the if statement with the uniform condition. |
| std::string src = R"( |
| @group(0) @binding(0) var<uniform> uniform_value : i32; |
| @group(0) @binding(1) var<storage, read_write> non_uniform_value : i32; |
| |
| fn non_uniform_discard_func() -> bool { |
| if (non_uniform_value == 42) { |
| discard; |
| } |
| return false; |
| } |
| |
| fn main() { |
| let b = non_uniform_discard_func() && true; |
| if (uniform_value == 42) { |
| workgroupBarrier(); |
| } |
| } |
| )"; |
| |
| RunTest(src, false); |
| EXPECT_EQ(error_, |
| R"(test:15:5 warning: 'workgroupBarrier' must only be called from uniform control flow |
| workgroupBarrier(); |
| ^^^^^^^^^^^^^^^^ |
| |
| test:13:11 note: calling 'non_uniform_discard_func' may cause subsequent control flow to be non-uniform |
| let b = non_uniform_discard_func() && true; |
| ^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| test:6:3 note: control flow depends on non-uniform value |
| if (non_uniform_value == 42) { |
| ^^ |
| |
| test:6:7 note: reading from read_write storage buffer 'non_uniform_value' may result in a non-uniform value |
| if (non_uniform_value == 42) { |
| ^^^^^^^^^^^^^^^^^ |
| )"); |
| } |
| |
| } // namespace |
| } // namespace tint::resolver |