mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-09-01 16:28:39 +00:00
744d0eb4aa
Discard statements no longer affect the behavior or uniformity analysis. Update the resolver, validator, and several tests to reflect this. Some E2E tests were removed as they had loops that are now considered to be infinite. Use the DemoteToHelper transform to emulate the correct semantics on platforms where discard is (or may) terminate the invocation in a manner that would affect derivative operations. We no longer need the UnwindDiscardFunctions transform for HLSL, which already implements the correct semantics. However, we still run the DemoteToHelper transform for the HLSL backend due to issues with FXC's handling of discard statements (see crbug.com/tint/1118). Fixed: tint:1723 Change-Id: Ib49ff187919ae81c4af8675e1b66acd57e2ff7d2 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/109003 Reviewed-by: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: James Price <jrprice@google.com>
7282 lines
173 KiB
C++
7282 lines
173 KiB
C++
// 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:
|
|
/// Build and resolve a program from a ProgramBuilder object.
|
|
/// @param program the program
|
|
/// @param should_pass true if `builder` program should pass the analysis, otherwise false
|
|
void RunTest(Program&& program, bool should_pass) {
|
|
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 {
|
|
if (kUniformityFailuresAsError) {
|
|
EXPECT_FALSE(valid);
|
|
} else {
|
|
EXPECT_TRUE(valid) << error_;
|
|
}
|
|
}
|
|
}
|
|
|
|
/// 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());
|
|
return RunTest(std::move(program), should_pass);
|
|
}
|
|
|
|
/// Build and resolve a program from a ProgramBuilder object.
|
|
/// @param builder the program builder
|
|
/// @param should_pass true if `builder` program should pass the analysis, otherwise false
|
|
void RunTest(ProgramBuilder&& builder, bool should_pass) {
|
|
auto program = Program(std::move(builder));
|
|
return RunTest(std::move(program), should_pass);
|
|
}
|
|
|
|
/// 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, 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 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,
|
|
};
|
|
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";
|
|
}
|
|
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, kReturn + 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;
|
|
break if (i == n);
|
|
}
|
|
}
|
|
}
|
|
)";
|
|
|
|
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;
|
|
break if (i == n);
|
|
}
|
|
}
|
|
}
|
|
)";
|
|
|
|
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
|
|
break if (i == n);
|
|
^^^^^
|
|
|
|
test:10:22 note: reading from read_write storage buffer 'n' may result in a non-uniform value
|
|
break 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;
|
|
break if (i == n);
|
|
}
|
|
}
|
|
}
|
|
)";
|
|
|
|
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;
|
|
break if (i == n);
|
|
}
|
|
}
|
|
}
|
|
)";
|
|
|
|
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
|
|
break if (i == n);
|
|
^^^^^
|
|
|
|
test:10:22 note: reading from read_write storage buffer 'n' may result in a non-uniform value
|
|
break if (i == n);
|
|
^
|
|
)");
|
|
}
|
|
|
|
class LoopDeadCodeTest : public UniformityAnalysisTestBase, public ::testing::TestWithParam<int> {};
|
|
|
|
INSTANTIATE_TEST_SUITE_P(UniformityAnalysisTest,
|
|
LoopDeadCodeTest,
|
|
::testing::Range<int>(0, kReturn + 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.
|
|
break if (false);
|
|
}
|
|
}
|
|
}
|
|
)";
|
|
|
|
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();
|
|
}
|
|
break if (true);
|
|
}
|
|
}
|
|
}
|
|
)";
|
|
|
|
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, 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_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) {
|
|
^^^^^^^^^^^
|
|
)");
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////
|
|
/// 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:11:7 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
|
|
fallthrough;
|
|
^^^^^^^^^^^
|
|
|
|
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:10:7 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
|
|
fallthrough;
|
|
^^^^^^^^^^^
|
|
|
|
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_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);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////
|
|
/// 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;
|
|
// }
|
|
utils::Vector<const ast::Parameter*, 8> params;
|
|
utils::Vector<const ast::Statement*, 8> 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(b.Decl(b.Let("rhs", rhs_init)));
|
|
for (int i = 0; i < 255; i++) {
|
|
params.Push(
|
|
b.Param("p" + std::to_string(i), ty.pointer(ty.i32(), ast::AddressSpace::kFunction)));
|
|
if (i > 0) {
|
|
foo_body.Push(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.GlobalVar("non_uniform_global", ty.i32(), ast::AddressSpace::kPrivate);
|
|
utils::Vector<const ast::Statement*, 8> main_body;
|
|
utils::Vector<const ast::Expression*, 8> args;
|
|
for (int i = 0; i < 255; i++) {
|
|
auto name = "v" + std::to_string(i);
|
|
main_body.Push(b.Decl(b.Var(name, ty.i32())));
|
|
args.Push(b.AddressOf(name));
|
|
}
|
|
main_body.Push(b.Assign("v0", "non_uniform_global"));
|
|
main_body.Push(b.CallStmt(b.create<ast::CallExpression>(b.Expr("foo"), args)));
|
|
main_body.Push(b.If(b.Equal("v254", 0_i), b.Block(b.CallStmt(b.Call("workgroupBarrier")))));
|
|
b.Func("main", utils::Empty, ty.void_(), main_body);
|
|
|
|
RunTest(std::move(b), false);
|
|
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 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_VectorBecomesUniform_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var v : vec4<i32>;
|
|
v[1] = rw;
|
|
v = vec4(1, 2, 3, 4);
|
|
if (v[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, VectorElementViaMember_VectorBecomesUniform_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var v : vec4<i32>;
|
|
v.y = rw;
|
|
v = vec4(1, 2, 3, 4);
|
|
if (v.y == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, VectorElement_VectorBecomesUniform_ThroughPointer_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var v : vec4<i32>;
|
|
v[1] = rw;
|
|
*(&v) = vec4(1, 2, 3, 4);
|
|
if (v[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
VectorElement_VectorBecomesUniform_ThroughPointerChain_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var v : vec4<i32>;
|
|
v[1] = rw;
|
|
*(&(*(&(*(&v))))) = vec4(1, 2, 3, 4);
|
|
if (v[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
VectorElement_VectorBecomesUniform_ThroughCapturedPointer_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var v : vec4<i32>;
|
|
v[1] = rw;
|
|
let p = &v;
|
|
*p = vec4(1, 2, 3, 4);
|
|
if (v[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, VectorElement_VectorBecomesUniform_PartialAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var v : vec4<i32>;
|
|
v[1] = rw;
|
|
v = vec4(1, 2, 3, v[3]);
|
|
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, VectorElementViaMember_VectorBecomesUniform_PartialAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var v : vec4<i32>;
|
|
v.y = rw;
|
|
v = vec4(1, 2, 3, v.w);
|
|
if (v.y == 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.y == 0) {
|
|
^^
|
|
|
|
test:6:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
v.y = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, VectorElement_DifferentElementBecomesUniform) {
|
|
// For aggregate types, we conservatively consider them to be 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, MatrixElement_ElementBecomesUniform) {
|
|
// For aggregate types, we conservatively consider them to be 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 : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
m[1][1] = 42.0;
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:6:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_ElementBecomesUniform_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
m = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), vec3(7.0, 8.0, 9.0));
|
|
if (m[1][1] == 0.0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_ElementBecomesUniform_ThroughPointer_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
*(&m) = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), vec3(7.0, 8.0, 9.0));
|
|
if (m[1][1] == 0.0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
MatrixElement_ElementBecomesUniform_ThroughPointerChain_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
*(&(*(&(*(&m))))) = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), vec3(7.0, 8.0, 9.0));
|
|
if (m[1][1] == 0.0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
MatrixElement_ElementBecomesUniform_ThroughCapturedPointer_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
let p = &m;
|
|
*p = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), vec3(7.0, 8.0, 9.0));
|
|
if (m[1][1] == 0.0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_ColumnBecomesUniform) {
|
|
// For aggregate types, we conservatively consider them to be 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 : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
m[1] = vec3(0.0, 42.0, 0.0);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:6:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_ColumnBecomesUniform_ThroughPartialPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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 : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
*(&(m[1])) = vec3(0.0, 42.0, 0.0);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:6:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_ColumnBecomesUniform_ThroughPartialPointerChain) {
|
|
// For aggregate types, we conservatively consider them to be 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 : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
*(&(*(&(m[1])))) = vec3(0.0, 42.0, 0.0);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:6:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_ColumnBecomesUniform_ThroughCapturedPartialPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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 : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
let p = &m[1];
|
|
m[1][1] = rw;
|
|
*p = vec3(0.0, 42.0, 0.0);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
MatrixElement_ColumnBecomesUniform_ThroughCapturedPartialPointerChain) {
|
|
// For aggregate types, we conservatively consider them to be 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 : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
let p = &m[1];
|
|
m[1][1] = rw;
|
|
*(&(*p)) = vec3(0.0, 42.0, 0.0);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_ColumnBecomesUniform_ThroughCapturedPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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 : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
let p = &m;
|
|
m[1][1] = rw;
|
|
(*p)[1] = vec3(0.0, 42.0, 0.0);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_MatrixBecomesUniform_PartialAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
m = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), m[2]);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:6:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
MatrixElement_MatrixBecomesUniform_PartialAssignment_ThroughPointer) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
*(&m) = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), m[2]);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:6:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
MatrixElement_MatrixBecomesUniform_PartialAssignment_ThroughCapturedPointer) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
let p = &m;
|
|
m[1][1] = rw;
|
|
*p = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), (*p)[2]);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
MatrixElement_MatrixBecomesUniform_PartialAssignment_ThroughCapturedPointerChain) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
let p = &(*(&m));
|
|
m[1][1] = rw;
|
|
*p = mat3x3<f32>(vec3(1.0, 2.0, 3.0), vec3(4.0, 5.0, 6.0), (*p)[2]);
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, MatrixElement_DifferentElementBecomesUniform) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : f32;
|
|
|
|
fn foo() {
|
|
var m : mat3x3<f32>;
|
|
m[1][1] = rw;
|
|
m[2][2] = 42.0;
|
|
if (m[1][1] == 0.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 (m[1][1] == 0.0) {
|
|
^^
|
|
|
|
test:6:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
m[1][1] = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
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 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_MemberBecomesUniformThroughCapturedPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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;
|
|
let p = &s;
|
|
s.a = rw;
|
|
(*p).a = 0;
|
|
if (s.a == 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 (s.a == 0) {
|
|
^^
|
|
|
|
test:11:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
s.a = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, StructMember_MemberBecomesUniformThroughCapturedPartialPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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;
|
|
let p = &s.a;
|
|
s.a = rw;
|
|
(*p) = 0;
|
|
if (s.a == 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 (s.a == 0) {
|
|
^^
|
|
|
|
test:11:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
s.a = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, StructMember_StructBecomesUniform_FullAssignment) {
|
|
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 = S(1, 2);
|
|
if (s.a == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, StructMember_StructBecomesUniform_PartialAssignment) {
|
|
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 = S(1, s.b);
|
|
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_StructBecomesUniform_FullAssignment_ThroughPointer) {
|
|
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) = S(1, 2);
|
|
if (s.a == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
StructMember_StructBecomesUniform_FullAssignment_ThroughCapturedPointer) {
|
|
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;
|
|
let p = &s;
|
|
s.a = rw;
|
|
*p = S(1, 2);
|
|
if (s.a == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
StructMember_StructBecomesUniform_FullAssignment_ThroughCapturedPointerChain) {
|
|
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;
|
|
let p = &(*(&s));
|
|
s.a = rw;
|
|
*p = S(1, 2);
|
|
if (s.a == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, StructMember_StructBecomesUniform_PartialAssignment_ThroughPointer) {
|
|
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) = S(1, (*(&s)).b);
|
|
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_StructBecomesUniform_PartialAssignment_ThroughCapturedPointer) {
|
|
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;
|
|
let p = &s;
|
|
s.a = rw;
|
|
*p = S(1, (*p).b);
|
|
if (s.a == 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 (s.a == 0) {
|
|
^^
|
|
|
|
test:11:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
|
|
s.a = rw;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
StructMember_StructBecomesUniform_PartialAssignment_ThroughCapturedPointerChain) {
|
|
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;
|
|
let p = &(*(&s));
|
|
s.a = rw;
|
|
*p = S(1, (*p).b);
|
|
if (s.a == 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 (s.a == 0) {
|
|
^^
|
|
|
|
test:11: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 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_DifferentElementBecomesNonUniformThroughPartialPointer) {
|
|
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_ElementBecomesUniform_FullAssignment) {
|
|
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 = array<i32, 4>(1, 2, 3, 4);
|
|
if (arr[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, ArrayElement_ElementBecomesUniform_PartialAssignment) {
|
|
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 = array<i32, 4>(1, 2, 3, arr[3]);
|
|
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 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_ElementBecomesUniform_ThroughPartialPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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>;
|
|
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_ElementBecomesUniform_ThroughPartialPointerChain) {
|
|
// For aggregate types, we conservatively consider them to be 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>;
|
|
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_ElementBecomesUniform_ThroughCapturedPartialPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
ArrayElement_ElementBecomesUniform_ThroughCapturedPartialPointerChain) {
|
|
// For aggregate types, we conservatively consider them to be 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;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, ArrayElement_ElementBecomesUniform_ThroughCapturedPointer) {
|
|
// For aggregate types, we conservatively consider them to be 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;
|
|
arr[1] = rw;
|
|
(*pa)[2] = 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;
|
|
^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, ArrayElement_ArrayBecomesUniform_ThroughPointer_FullAssignment) {
|
|
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) = array<i32, 4>();
|
|
if (arr[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
ArrayElement_ArrayBecomesUniform_ThroughPointerChain_FullAssignment) {
|
|
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))))) = array<i32, 4>();
|
|
if (arr[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
ArrayElement_ArrayBecomesUniform_ThroughCapturedPointer_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var arr : array<i32, 4>;
|
|
let pa = &arr;
|
|
arr[1] = rw;
|
|
*pa = array<i32, 4>();
|
|
if (arr[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest,
|
|
ArrayElement_ArrayBecomesUniform_ThroughCapturedPointerChain_FullAssignment) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> rw : i32;
|
|
|
|
fn foo() {
|
|
var arr : array<i32, 4>;
|
|
let pa = &(*(&arr));
|
|
arr[1] = rw;
|
|
*pa = array<i32, 4>();
|
|
if (arr[1] == 0) {
|
|
workgroupBarrier();
|
|
}
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
////////////////////////////////////////////////////////////////////////////////
|
|
/// Miscellaneous statement and expression tests.
|
|
////////////////////////////////////////////////////////////////////////////////
|
|
|
|
TEST_F(UniformityAnalysisTest, NonUniformDiscard) {
|
|
// Non-uniform discard statements should not cause uniformity issues.
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
|
|
|
|
fn foo() {
|
|
if (non_uniform == 42) {
|
|
discard;
|
|
}
|
|
_ = dpdx(1.0);
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, FunctionReconvergesOnExit) {
|
|
// Call a function that has returns during non-uniform control flow, and test that the analysis
|
|
// reconverges when returning to the caller.
|
|
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, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, TypeInitializer) {
|
|
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, ShortCircuiting_UniformLHS) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read> uniform_global : i32;
|
|
|
|
fn main() {
|
|
let b = (uniform_global == 0) && (dpdx(1.0) == 0.0);
|
|
}
|
|
)";
|
|
|
|
RunTest(src, true);
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, ShortCircuiting_NonUniformLHS) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
|
|
|
|
fn main() {
|
|
let b = (non_uniform_global == 0) && (dpdx(1.0) == 0.0);
|
|
}
|
|
)";
|
|
|
|
RunTest(src, false);
|
|
EXPECT_EQ(error_,
|
|
R"(test:5:41 warning: 'dpdx' must only be called from uniform control flow
|
|
let b = (non_uniform_global == 0) && (dpdx(1.0) == 0.0);
|
|
^^^^
|
|
|
|
test:5:37 note: control flow depends on non-uniform value
|
|
let b = (non_uniform_global == 0) && (dpdx(1.0) == 0.0);
|
|
^^
|
|
|
|
test:5:12 note: reading from read_write storage buffer 'non_uniform_global' may result in a non-uniform value
|
|
let b = (non_uniform_global == 0) && (dpdx(1.0) == 0.0);
|
|
^^^^^^^^^^^^^^^^^^
|
|
)");
|
|
}
|
|
|
|
TEST_F(UniformityAnalysisTest, ShortCircuiting_ReconvergeLHS) {
|
|
std::string src = R"(
|
|
@group(0) @binding(0) var<storage, read_write> non_uniform_global : i32;
|
|
|
|
fn main() {
|
|
let b = (non_uniform_global == 0) && 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;
|
|
|
|
fn main() {
|
|
let b = false && (non_uniform_global == 0);
|
|
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;
|
|
|
|
fn main() {
|
|
let b = (non_uniform_global != 0) && (non_uniform_global != 42);
|
|
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, 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.GlobalVar("v0", ty.i32(), ast::AddressSpace::kPrivate, b.Expr(0_i));
|
|
utils::Vector<const ast::Statement*, 8> foo_body;
|
|
std::string v_last = "v0";
|
|
for (int i = 1; i < 100000; i++) {
|
|
auto v = "v" + std::to_string(i);
|
|
foo_body.Push(b.Decl(b.Var(v, b.Expr(v_last))));
|
|
v_last = v;
|
|
}
|
|
foo_body.Push(b.If(b.Equal(v_last, 0_i), b.Block(b.CallStmt(b.Call("workgroupBarrier")))));
|
|
b.Func("foo", utils::Empty, ty.void_(), foo_body);
|
|
|
|
RunTest(std::move(b), false);
|
|
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) {
|
|
^^^
|
|
)");
|
|
}
|
|
|
|
} // namespace
|
|
} // namespace tint::resolver
|