| // 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, |
| kModuleConst, |
| 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 kModuleConst: |
| return "module_const == 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(kModuleConst); |
| 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; |
| |
| const module_const : 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"( |
| |