Import Tint changes from Dawn
Changes:
- 142bddf7dabe2cbbbb2adfaaf0a9708aad0ba1a8 Fix #ifdefs for spvtools by Austin Eng <enga@chromium.org>
- f244bffdd51d8c4aeddd4181025559a41de4154f Fix build of tint_cmd_helper by Austin Eng <enga@chromium.org>
- 8525ff29da6b005381abe94023a303edc0f39c44 tint/transform/robustness: Implement predicated mode by Ben Clayton <bclayton@google.com>
- 5e5655155110e121763e3dca13d20c3eda09167e Revert "Add writer to emit the AST." by Dan Sinclair <dsinclair@chromium.org>
- 5ab2928f6ab2c2c9d899bf1220d4be7740f09664 Convert `strtoll` to use `std::from_chars`. by dan sinclair <dsinclair@chromium.org>
- a753ad47a48fc0fba6e617b7ffcddc0e1fdc4b2e Add writer to emit the AST. by dan sinclair <dsinclair@chromium.org>
- 63d0fabeb1fc11e4829256804bf776a23e8f2849 tint: Fix HLSL emission for out-of-order storage / unifor... by Ben Clayton <bclayton@google.com>
- fd387a37c3b936a547288347bf854405ca2ec647 tint: Fix WGSL emission of const_assert by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: 142bddf7dabe2cbbbb2adfaaf0a9708aad0ba1a8
Change-Id: Ic0f52f73c9a9fa70f04d272f616db8e5abd82136
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/122880
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/ast/disable_validation_attribute.cc b/src/tint/ast/disable_validation_attribute.cc
index eff1c1f..27eaef0 100644
--- a/src/tint/ast/disable_validation_attribute.cc
+++ b/src/tint/ast/disable_validation_attribute.cc
@@ -23,7 +23,7 @@
DisableValidationAttribute::DisableValidationAttribute(ProgramID pid,
NodeID nid,
DisabledValidation val)
- : Base(pid, nid), validation(val) {}
+ : Base(pid, nid, utils::Empty), validation(val) {}
DisableValidationAttribute::~DisableValidationAttribute() = default;
diff --git a/src/tint/ast/internal_attribute.cc b/src/tint/ast/internal_attribute.cc
index 1b4ca9e..c5c5f2e 100644
--- a/src/tint/ast/internal_attribute.cc
+++ b/src/tint/ast/internal_attribute.cc
@@ -14,11 +14,16 @@
#include "src/tint/ast/internal_attribute.h"
+#include <utility>
+
TINT_INSTANTIATE_TYPEINFO(tint::ast::InternalAttribute);
namespace tint::ast {
-InternalAttribute::InternalAttribute(ProgramID pid, NodeID nid) : Base(pid, nid, Source{}) {}
+InternalAttribute::InternalAttribute(ProgramID pid,
+ NodeID nid,
+ utils::VectorRef<const IdentifierExpression*> deps)
+ : Base(pid, nid, Source{}), dependencies(std::move(deps)) {}
InternalAttribute::~InternalAttribute() = default;
diff --git a/src/tint/ast/internal_attribute.h b/src/tint/ast/internal_attribute.h
index 9904af8..36f2a98 100644
--- a/src/tint/ast/internal_attribute.h
+++ b/src/tint/ast/internal_attribute.h
@@ -18,6 +18,12 @@
#include <string>
#include "src/tint/ast/attribute.h"
+#include "src/tint/utils/vector.h"
+
+// Forward declarations
+namespace tint::ast {
+class IdentifierExpression;
+} // namespace tint::ast
namespace tint::ast {
@@ -29,7 +35,10 @@
/// Constructor
/// @param program_id the identifier of the program that owns this node
/// @param nid the unique node identifier
- explicit InternalAttribute(ProgramID program_id, NodeID nid);
+ /// @param deps a list of identifiers that this attribute is dependent on
+ InternalAttribute(ProgramID program_id,
+ NodeID nid,
+ utils::VectorRef<const IdentifierExpression*> deps);
/// Destructor
~InternalAttribute() override;
@@ -40,6 +49,9 @@
/// @returns the WGSL name for the attribute
std::string Name() const override;
+
+ /// A list of identifiers that this attribute is dependent on
+ const utils::Vector<const IdentifierExpression*, 1> dependencies;
};
} // namespace tint::ast
diff --git a/src/tint/cmd/BUILD.gn b/src/tint/cmd/BUILD.gn
index ab0c964..b4ff6d5 100644
--- a/src/tint/cmd/BUILD.gn
+++ b/src/tint/cmd/BUILD.gn
@@ -21,10 +21,10 @@
"helper.h",
]
- deps = [
- "${tint_root_dir}/src/tint:libtint",
- "${tint_spirv_tools_dir}/:spvtools_headers",
- ]
+ deps = [ "${tint_root_dir}/src/tint:libtint" ]
+ if (tint_build_spv_reader || tint_build_spv_writer) {
+ deps += [ "${tint_spirv_tools_dir}/:spvtools_headers" ]
+ }
}
executable("tint") {
diff --git a/src/tint/fuzzers/tint_common_fuzzer.cc b/src/tint/fuzzers/tint_common_fuzzer.cc
index c4f1fd1..7c390af 100644
--- a/src/tint/fuzzers/tint_common_fuzzer.cc
+++ b/src/tint/fuzzers/tint_common_fuzzer.cc
@@ -24,9 +24,9 @@
#include <utility>
#include <vector>
-#if TINT_BUILD_SPV_READER
+#if TINT_BUILD_SPV_READER || TINT_BUILD_SPV_WRITER
#include "spirv-tools/libspirv.hpp"
-#endif // TINT_BUILD_SPV_READER
+#endif // TINT_BUILD_SPV_READER || TINT_BUILD_SPV_WRITER
#include "src/tint/ast/module.h"
#include "src/tint/diagnostic/formatter.h"
diff --git a/src/tint/reader/wgsl/lexer.cc b/src/tint/reader/wgsl/lexer.cc
index 5171e4c..28fb135 100644
--- a/src/tint/reader/wgsl/lexer.cc
+++ b/src/tint/reader/wgsl/lexer.cc
@@ -15,6 +15,7 @@
#include "src/tint/reader/wgsl/lexer.h"
#include <cctype>
+#include <charconv>
#include <cmath>
#include <cstring>
#include <functional>
@@ -804,41 +805,48 @@
return {Token::Type::kFloatLiteral, source, result_f64};
}
-Token Lexer::build_token_from_int_if_possible(Source source, size_t start, int32_t base) {
+Token Lexer::build_token_from_int_if_possible(Source source,
+ size_t start,
+ size_t prefix_count,
+ int32_t base) {
const char* start_ptr = &at(start);
- char* end_ptr = nullptr;
+ // The call to `from_chars` will return the pointer to just after the last parsed character.
+ // We also need to tell it the maximum end character to parse. So, instead of walking all the
+ // characters to find the last possible and using that, we just provide the end of the string.
+ // We then calculate the count based off the provided end pointer and the start pointer. The
+ // extra `prefix_count` is to handle a `0x` which is not included in the `start` value.
+ const char* end_ptr = &at(length() - 1) + 1;
- errno = 0;
- int64_t res = strtoll(start_ptr, &end_ptr, base);
- const bool overflow = errno == ERANGE;
-
- if (end_ptr) {
- advance(static_cast<size_t>(end_ptr - start_ptr));
- }
+ int64_t value = 0;
+ auto res = std::from_chars(start_ptr, end_ptr, value, base);
+ const bool overflow = res.ec != std::errc();
+ advance(static_cast<size_t>(res.ptr - start_ptr) + prefix_count);
if (matches(pos(), 'u')) {
- if (!overflow && CheckedConvert<u32>(AInt(res))) {
+ if (!overflow && CheckedConvert<u32>(AInt(value))) {
advance(1);
end_source(source);
- return {Token::Type::kIntLiteral_U, source, res};
+ return {Token::Type::kIntLiteral_U, source, value};
}
return {Token::Type::kError, source, "value cannot be represented as 'u32'"};
}
if (matches(pos(), 'i')) {
- if (!overflow && CheckedConvert<i32>(AInt(res))) {
+ if (!overflow && CheckedConvert<i32>(AInt(value))) {
advance(1);
end_source(source);
- return {Token::Type::kIntLiteral_I, source, res};
+ return {Token::Type::kIntLiteral_I, source, value};
}
return {Token::Type::kError, source, "value cannot be represented as 'i32'"};
}
- end_source(source);
+ // Check this last in order to get the more specific sized error messages
if (overflow) {
return {Token::Type::kError, source, "value cannot be represented as 'abstract-int'"};
}
- return {Token::Type::kIntLiteral, source, res};
+
+ end_source(source);
+ return {Token::Type::kIntLiteral, source, value};
}
Token Lexer::try_hex_integer() {
@@ -858,7 +866,7 @@
"integer or float hex literal has no significant digits"};
}
- return build_token_from_int_if_possible(source, start, 16);
+ return build_token_from_int_if_possible(source, curr, curr - start, 16);
}
Token Lexer::try_integer() {
@@ -879,7 +887,7 @@
}
}
- return build_token_from_int_if_possible(source, start, 10);
+ return build_token_from_int_if_possible(source, start, 0, 10);
}
Token Lexer::try_ident() {
diff --git a/src/tint/reader/wgsl/lexer.h b/src/tint/reader/wgsl/lexer.h
index 8e0306b..ac41dae 100644
--- a/src/tint/reader/wgsl/lexer.h
+++ b/src/tint/reader/wgsl/lexer.h
@@ -47,7 +47,10 @@
/// @returns uninitialized token on success, or error
Token skip_comment();
- Token build_token_from_int_if_possible(Source source, size_t start, int32_t base);
+ Token build_token_from_int_if_possible(Source source,
+ size_t start,
+ size_t prefix_count,
+ int32_t base);
Token check_keyword(const Source&, std::string_view);
diff --git a/src/tint/resolver/attribute_validation_test.cc b/src/tint/resolver/attribute_validation_test.cc
index b0a863d..5c09c2a 100644
--- a/src/tint/resolver/attribute_validation_test.cc
+++ b/src/tint/resolver/attribute_validation_test.cc
@@ -1977,4 +1977,40 @@
} // namespace
} // namespace MustUseTests
+namespace InternalAttributeDeps {
+namespace {
+
+class TestAttribute : public Castable<TestAttribute, ast::InternalAttribute> {
+ public:
+ TestAttribute(ProgramID pid, ast::NodeID nid, const ast::IdentifierExpression* dep)
+ : Base(pid, nid, utils::Vector{dep}) {}
+ std::string InternalName() const override { return "test_attribute"; }
+ const Cloneable* Clone(CloneContext*) const override { return nullptr; }
+};
+
+using InternalAttributeDepsTest = ResolverTest;
+TEST_F(InternalAttributeDepsTest, Dependency) {
+ auto* ident = Expr("v");
+ auto* attr = ASTNodes().Create<TestAttribute>(ID(), AllocateNodeID(), ident);
+ auto* f = Func("f", utils::Empty, ty.void_(), utils::Empty, utils::Vector{attr});
+ auto* v = GlobalVar("v", ty.i32(), builtin::AddressSpace::kPrivate);
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* user = As<sem::VariableUser>(Sem().Get(ident));
+ ASSERT_NE(user, nullptr);
+
+ auto* var = Sem().Get(v);
+ EXPECT_EQ(user->Variable(), var);
+
+ auto* fn = Sem().Get(f);
+ EXPECT_THAT(fn->DirectlyReferencedGlobals(), testing::ElementsAre(var));
+ EXPECT_THAT(fn->TransitivelyReferencedGlobals(), testing::ElementsAre(var));
+}
+
+} // namespace
+} // namespace InternalAttributeDeps
+
} // namespace tint::resolver
+
+TINT_INSTANTIATE_TYPEINFO(tint::resolver::InternalAttributeDeps::TestAttribute);
diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc
index 8292adb..f718a46 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -414,12 +414,18 @@
TraverseExpression(wg->y);
TraverseExpression(wg->z);
return true;
+ },
+ [&](const ast::InternalAttribute* i) {
+ for (auto* dep : i->dependencies) {
+ TraverseExpression(dep);
+ }
+ return true;
});
if (handled) {
return;
}
- if (attr->IsAnyOf<ast::BuiltinAttribute, ast::DiagnosticAttribute, ast::InternalAttribute,
+ if (attr->IsAnyOf<ast::BuiltinAttribute, ast::DiagnosticAttribute,
ast::InterpolateAttribute, ast::InvariantAttribute, ast::MustUseAttribute,
ast::StageAttribute, ast::StrideAttribute,
ast::StructMemberOffsetAttribute>()) {
diff --git a/src/tint/resolver/intrinsic_table.cc b/src/tint/resolver/intrinsic_table.cc
index 0d2cc10..e8fe271 100644
--- a/src/tint/resolver/intrinsic_table.cc
+++ b/src/tint/resolver/intrinsic_table.cc
@@ -1495,7 +1495,7 @@
// Was this overload a constructor or conversion?
if (match.overload->flags.Contains(OverloadFlag::kIsConstructor)) {
- utils::Vector<const sem::Parameter*, 8> params;
+ utils::Vector<sem::Parameter*, 8> params;
params.Reserve(match.parameters.Length());
for (auto& p : match.parameters) {
params.Push(builder.create<sem::Parameter>(
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index cd426f7..f51e5cc 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -856,9 +856,9 @@
sem::Function* Resolver::Function(const ast::Function* decl) {
Mark(decl->name);
- uint32_t parameter_index = 0;
- utils::Hashmap<Symbol, Source, 8> parameter_names;
- utils::Vector<sem::Parameter*, 8> parameters;
+ auto* func = builder_->create<sem::Function>(decl);
+ builder_->Sem().Add(decl, func);
+ TINT_SCOPED_ASSIGNMENT(current_function_, func);
validator_.DiagnosticFilters().Push();
TINT_DEFER(validator_.DiagnosticFilters().Pop());
@@ -872,6 +872,8 @@
}
// Resolve all the parameters
+ uint32_t parameter_index = 0;
+ utils::Hashmap<Symbol, Source, 8> parameter_names;
for (auto* param : decl->params) {
Mark(param);
@@ -893,7 +895,7 @@
return nullptr;
}
- parameters.Push(p);
+ func->AddParameter(p);
auto* p_ty = const_cast<type::Type*>(p->Type());
if (auto* str = p_ty->As<sem::Struct>()) {
@@ -923,9 +925,9 @@
} else {
return_type = builder_->create<type::Void>();
}
+ func->SetReturnType(return_type);
// Determine if the return type has a location
- std::optional<uint32_t> return_location;
for (auto* attr : decl->return_type_attributes) {
if (!Attribute(attr)) {
return nullptr;
@@ -936,7 +938,7 @@
if (!value) {
return nullptr;
}
- return_location = value.Get();
+ func->SetReturnLocation(value.Get());
}
}
@@ -963,12 +965,7 @@
}
}
- auto* func =
- builder_->create<sem::Function>(decl, return_type, return_location, std::move(parameters));
ApplyDiagnosticSeverities(func);
- builder_->Sem().Add(decl, func);
-
- TINT_SCOPED_ASSIGNMENT(current_function_, func);
if (!WorkgroupSize(decl)) {
return nullptr;
@@ -2089,7 +2086,7 @@
auto* call_target = struct_ctors_.GetOrCreate(
StructConstructorSig{{str, args.Length(), args_stage}},
[&]() -> sem::ValueConstructor* {
- utils::Vector<const sem::Parameter*, 8> params;
+ utils::Vector<sem::Parameter*, 8> params;
params.Resize(std::min(args.Length(), str->Members().Length()));
for (size_t i = 0, n = params.Length(); i < n; i++) {
params[i] = builder_->create<sem::Parameter>(
@@ -3436,6 +3433,7 @@
[&](const ast::BuiltinAttribute* b) { return BuiltinAttribute(b); },
[&](const ast::DiagnosticAttribute* d) { return DiagnosticControl(d->control); },
[&](const ast::InterpolateAttribute* i) { return InterpolateAttribute(i); },
+ [&](const ast::InternalAttribute* i) { return InternalAttribute(i); },
[&](Default) { return true; });
}
@@ -3460,6 +3458,15 @@
return true;
}
+bool Resolver::InternalAttribute(const ast::InternalAttribute* attr) {
+ for (auto* dep : attr->dependencies) {
+ if (!Expression(dep)) {
+ return false;
+ }
+ }
+ return true;
+}
+
bool Resolver::DiagnosticControl(const ast::DiagnosticControl& control) {
Mark(control.rule_name);
diff --git a/src/tint/resolver/resolver.h b/src/tint/resolver/resolver.h
index fff7506..8acf602 100644
--- a/src/tint/resolver/resolver.h
+++ b/src/tint/resolver/resolver.h
@@ -321,6 +321,10 @@
/// @returns true on success, false on failure
bool InterpolateAttribute(const ast::InterpolateAttribute* attr);
+ /// Resolves the internal attribute @p attr
+ /// @returns true on success, false on failure
+ bool InternalAttribute(const ast::InternalAttribute* attr);
+
/// @param control the diagnostic control
/// @returns true on success, false on failure
bool DiagnosticControl(const ast::DiagnosticControl& control);
diff --git a/src/tint/sem/builtin.cc b/src/tint/sem/builtin.cc
index 8bcd8af..4fd2c34 100644
--- a/src/tint/sem/builtin.cc
+++ b/src/tint/sem/builtin.cc
@@ -25,17 +25,6 @@
TINT_INSTANTIATE_TYPEINFO(tint::sem::Builtin);
namespace tint::sem {
-namespace {
-
-utils::VectorRef<const Parameter*> SetOwner(utils::VectorRef<Parameter*> parameters,
- const tint::sem::CallTarget* owner) {
- for (auto* parameter : parameters) {
- parameter->SetOwner(owner);
- }
- return parameters;
-}
-
-} // namespace
const char* Builtin::str() const {
return sem::str(type_);
@@ -112,7 +101,7 @@
PipelineStageSet supported_stages,
bool is_deprecated,
bool must_use)
- : Base(return_type, SetOwner(std::move(parameters), this), eval_stage, must_use),
+ : Base(return_type, std::move(parameters), eval_stage, must_use),
type_(type),
supported_stages_(supported_stages),
is_deprecated_(is_deprecated) {}
diff --git a/src/tint/sem/call_target.cc b/src/tint/sem/call_target.cc
index 76c1ab1..9ef0059 100644
--- a/src/tint/sem/call_target.cc
+++ b/src/tint/sem/call_target.cc
@@ -23,17 +23,25 @@
namespace tint::sem {
+CallTarget::CallTarget(EvaluationStage stage, bool must_use) : stage_(stage), must_use_(must_use) {}
+
CallTarget::CallTarget(const type::Type* return_type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage,
bool must_use)
- : signature_{return_type, std::move(parameters)}, stage_(stage), must_use_(must_use) {
+ : stage_(stage), must_use_(must_use) {
+ SetReturnType(return_type);
+ for (auto* param : parameters) {
+ AddParameter(param);
+ }
TINT_ASSERT(Semantic, return_type);
}
CallTarget::CallTarget(const CallTarget&) = default;
CallTarget::~CallTarget() = default;
+CallTargetSignature::CallTargetSignature() = default;
+
CallTargetSignature::CallTargetSignature(const type::Type* ret_ty,
utils::VectorRef<const sem::Parameter*> params)
: return_type(ret_ty), parameters(std::move(params)) {}
diff --git a/src/tint/sem/call_target.h b/src/tint/sem/call_target.h
index 4f0ab33..096aa8f 100644
--- a/src/tint/sem/call_target.h
+++ b/src/tint/sem/call_target.h
@@ -28,6 +28,9 @@
/// CallTargetSignature holds the return type and parameters for a call target
struct CallTargetSignature {
/// Constructor
+ CallTargetSignature();
+
+ /// Constructor
/// @param ret_ty the call target return type
/// @param params the call target parameters
CallTargetSignature(const type::Type* ret_ty, utils::VectorRef<const Parameter*> params);
@@ -39,9 +42,9 @@
~CallTargetSignature();
/// The type of the call target return value
- const type::Type* const return_type = nullptr;
+ const type::Type* return_type = nullptr;
/// The parameters of the call target
- const utils::Vector<const sem::Parameter*, 8> parameters;
+ utils::Vector<const sem::Parameter*, 8> parameters;
/// Equality operator
/// @param other the signature to compare this to
@@ -67,13 +70,19 @@
class CallTarget : public Castable<CallTarget, Node> {
public:
/// Constructor
+ /// @param stage the earliest evaluation stage for a call to this target
+ /// @param must_use the result of the call target must be used, i.e. it cannot be used as a call
+ /// statement.
+ CallTarget(EvaluationStage stage, bool must_use);
+
+ /// Constructor
/// @param return_type the return type of the call target
/// @param parameters the parameters for the call target
/// @param stage the earliest evaluation stage for a call to this target
/// @param must_use the result of the call target must be used, i.e. it cannot be used as a call
/// statement.
CallTarget(const type::Type* return_type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage,
bool must_use);
@@ -83,9 +92,20 @@
/// Destructor
~CallTarget() override;
+ /// Sets the call target's return type
+ /// @param ty the parameter
+ void SetReturnType(const type::Type* ty) { signature_.return_type = ty; }
+
/// @return the return type of the call target
const type::Type* ReturnType() const { return signature_.return_type; }
+ /// Adds a parameter to the call target
+ /// @param parameter the parameter
+ void AddParameter(Parameter* parameter) {
+ parameter->SetOwner(this);
+ signature_.parameters.Push(parameter);
+ }
+
/// @return the parameters of the call target
auto& Parameters() const { return signature_.parameters; }
diff --git a/src/tint/sem/function.cc b/src/tint/sem/function.cc
index 0859dd8..ec51b7b 100644
--- a/src/tint/sem/function.cc
+++ b/src/tint/sem/function.cc
@@ -28,29 +28,12 @@
TINT_INSTANTIATE_TYPEINFO(tint::sem::Function);
namespace tint::sem {
-namespace {
-utils::VectorRef<const Parameter*> SetOwner(utils::VectorRef<Parameter*> parameters,
- const tint::sem::CallTarget* owner) {
- for (auto* parameter : parameters) {
- parameter->SetOwner(owner);
- }
- return parameters;
-}
-
-} // namespace
-
-Function::Function(const ast::Function* declaration,
- type::Type* return_type,
- std::optional<uint32_t> return_location,
- utils::VectorRef<Parameter*> parameters)
- : Base(return_type,
- SetOwner(std::move(parameters), this),
- EvaluationStage::kRuntime,
+Function::Function(const ast::Function* declaration)
+ : Base(EvaluationStage::kRuntime,
ast::HasAttribute<ast::MustUseAttribute>(declaration->attributes)),
declaration_(declaration),
- workgroup_size_{1, 1, 1},
- return_location_(return_location) {}
+ workgroup_size_{1, 1, 1} {}
Function::~Function() = default;
diff --git a/src/tint/sem/function.h b/src/tint/sem/function.h
index 528ebef..b157e7d 100644
--- a/src/tint/sem/function.h
+++ b/src/tint/sem/function.h
@@ -54,17 +54,15 @@
/// Constructor
/// @param declaration the ast::Function
- /// @param return_type the return type of the function
- /// @param return_location the location value for the return, if provided
- /// @param parameters the parameters to the function
- Function(const ast::Function* declaration,
- type::Type* return_type,
- std::optional<uint32_t> return_location,
- utils::VectorRef<Parameter*> parameters);
+ explicit Function(const ast::Function* declaration);
/// Destructor
~Function() override;
+ /// Sets the function's return location
+ /// @param return_location the location value
+ void SetReturnLocation(uint32_t return_location) { return_location_ = return_location; }
+
/// @returns the ast::Function declaration
const ast::Function* Declaration() const { return declaration_; }
diff --git a/src/tint/sem/value_constructor.cc b/src/tint/sem/value_constructor.cc
index 62be478..b270414 100644
--- a/src/tint/sem/value_constructor.cc
+++ b/src/tint/sem/value_constructor.cc
@@ -21,7 +21,7 @@
namespace tint::sem {
ValueConstructor::ValueConstructor(const type::Type* type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage)
: Base(type, std::move(parameters), stage, /* must_use */ true) {}
diff --git a/src/tint/sem/value_constructor.h b/src/tint/sem/value_constructor.h
index aef827b..34c3b43 100644
--- a/src/tint/sem/value_constructor.h
+++ b/src/tint/sem/value_constructor.h
@@ -28,7 +28,7 @@
/// @param parameters the constructor parameters
/// @param stage the earliest evaluation stage for the expression
ValueConstructor(const type::Type* type,
- utils::VectorRef<const Parameter*> parameters,
+ utils::VectorRef<Parameter*> parameters,
EvaluationStage stage);
/// Destructor
diff --git a/src/tint/sem/value_conversion.cc b/src/tint/sem/value_conversion.cc
index 37331af..95587fe 100644
--- a/src/tint/sem/value_conversion.cc
+++ b/src/tint/sem/value_conversion.cc
@@ -19,9 +19,9 @@
namespace tint::sem {
ValueConversion::ValueConversion(const type::Type* type,
- const sem::Parameter* parameter,
+ sem::Parameter* parameter,
EvaluationStage stage)
- : Base(type, utils::Vector<const sem::Parameter*, 1>{parameter}, stage, /* must_use */ true) {}
+ : Base(type, utils::Vector<sem::Parameter*, 1>{parameter}, stage, /* must_use */ true) {}
ValueConversion::~ValueConversion() = default;
diff --git a/src/tint/sem/value_conversion.h b/src/tint/sem/value_conversion.h
index b79caa3..2d2ab38 100644
--- a/src/tint/sem/value_conversion.h
+++ b/src/tint/sem/value_conversion.h
@@ -26,7 +26,7 @@
/// @param type the target type of the cast
/// @param parameter the type cast parameter
/// @param stage the earliest evaluation stage for the expression
- ValueConversion(const type::Type* type, const sem::Parameter* parameter, EvaluationStage stage);
+ ValueConversion(const type::Type* type, sem::Parameter* parameter, EvaluationStage stage);
/// Destructor
~ValueConversion() override;
diff --git a/src/tint/transform/add_block_attribute.cc b/src/tint/transform/add_block_attribute.cc
index c63bd04..c486936 100644
--- a/src/tint/transform/add_block_attribute.cc
+++ b/src/tint/transform/add_block_attribute.cc
@@ -102,7 +102,7 @@
}
AddBlockAttribute::BlockAttribute::BlockAttribute(ProgramID pid, ast::NodeID nid)
- : Base(pid, nid) {}
+ : Base(pid, nid, utils::Empty) {}
AddBlockAttribute::BlockAttribute::~BlockAttribute() = default;
std::string AddBlockAttribute::BlockAttribute::InternalName() const {
return "block";
diff --git a/src/tint/transform/calculate_array_length.cc b/src/tint/transform/calculate_array_length.cc
index a25c89e..9f5b659 100644
--- a/src/tint/transform/calculate_array_length.cc
+++ b/src/tint/transform/calculate_array_length.cc
@@ -71,7 +71,7 @@
} // namespace
CalculateArrayLength::BufferSizeIntrinsic::BufferSizeIntrinsic(ProgramID pid, ast::NodeID nid)
- : Base(pid, nid) {}
+ : Base(pid, nid, utils::Empty) {}
CalculateArrayLength::BufferSizeIntrinsic::~BufferSizeIntrinsic() = default;
std::string CalculateArrayLength::BufferSizeIntrinsic::InternalName() const {
return "intrinsic_buffer_size";
diff --git a/src/tint/transform/decompose_memory_access.cc b/src/tint/transform/decompose_memory_access.cc
index 0030ca1..f3d92c5 100644
--- a/src/tint/transform/decompose_memory_access.cc
+++ b/src/tint/transform/decompose_memory_access.cc
@@ -228,7 +228,7 @@
}
return builder->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
builder->ID(), builder->AllocateNodeID(), DecomposeMemoryAccess::Intrinsic::Op::kLoad, type,
- address_space, buffer);
+ address_space, builder->Expr(buffer));
}
/// @returns a DecomposeMemoryAccess::Intrinsic attribute that can be applied to a stub function to
@@ -242,7 +242,7 @@
}
return builder->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
builder->ID(), builder->AllocateNodeID(), DecomposeMemoryAccess::Intrinsic::Op::kStore,
- type, builtin::AddressSpace::kStorage, buffer);
+ type, builtin::AddressSpace::kStorage, builder->Expr(buffer));
}
/// @returns a DecomposeMemoryAccess::Intrinsic attribute that can be applied to a stub function for
@@ -299,7 +299,7 @@
}
return builder->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
builder->ID(), builder->AllocateNodeID(), op, type, builtin::AddressSpace::kStorage,
- buffer);
+ builder->Expr(buffer));
}
/// BufferAccess describes a single storage or uniform buffer access
@@ -692,8 +692,8 @@
Op o,
DataType ty,
builtin::AddressSpace as,
- const Symbol& buf)
- : Base(pid, nid), op(o), type(ty), address_space(as), buffer(buf) {}
+ const ast::IdentifierExpression* buf)
+ : Base(pid, nid, utils::Vector{buf}), op(o), type(ty), address_space(as) {}
DecomposeMemoryAccess::Intrinsic::~Intrinsic() = default;
std::string DecomposeMemoryAccess::Intrinsic::InternalName() const {
utils::StringStream ss;
@@ -794,7 +794,7 @@
const DecomposeMemoryAccess::Intrinsic* DecomposeMemoryAccess::Intrinsic::Clone(
CloneContext* ctx) const {
- auto buf = ctx->Clone(buffer);
+ auto buf = ctx->Clone(Buffer());
return ctx->dst->ASTNodes().Create<DecomposeMemoryAccess::Intrinsic>(
ctx->dst->ID(), ctx->dst->AllocateNodeID(), op, type, address_space, buf);
}
@@ -803,6 +803,10 @@
return op != Op::kLoad && op != Op::kStore;
}
+const ast::IdentifierExpression* DecomposeMemoryAccess::Intrinsic::Buffer() const {
+ return dependencies[0];
+}
+
DecomposeMemoryAccess::DecomposeMemoryAccess() = default;
DecomposeMemoryAccess::~DecomposeMemoryAccess() = default;
diff --git a/src/tint/transform/decompose_memory_access.h b/src/tint/transform/decompose_memory_access.h
index e34cd63..f85ad6d 100644
--- a/src/tint/transform/decompose_memory_access.h
+++ b/src/tint/transform/decompose_memory_access.h
@@ -80,13 +80,13 @@
/// @param o the op of the intrinsic
/// @param type the data type of the intrinsic
/// @param address_space the address space of the buffer
- /// @param buffer the storage or uniform buffer name
+ /// @param buffer the storage or uniform buffer identifier
Intrinsic(ProgramID pid,
ast::NodeID nid,
Op o,
DataType type,
builtin::AddressSpace address_space,
- const Symbol& buffer);
+ const ast::IdentifierExpression* buffer);
/// Destructor
~Intrinsic() override;
@@ -102,6 +102,9 @@
/// @return true if op is atomic
bool IsAtomic() const;
+ /// @return the buffer that this intrinsic operates on
+ const ast::IdentifierExpression* Buffer() const;
+
/// The op of the intrinsic
const Op op;
@@ -110,9 +113,6 @@
/// The address space of the buffer this intrinsic operates on
const builtin::AddressSpace address_space;
-
- /// The buffer name
- const Symbol buffer;
};
/// Constructor
diff --git a/src/tint/transform/robustness.cc b/src/tint/transform/robustness.cc
index c6bf218..b79a029 100644
--- a/src/tint/transform/robustness.cc
+++ b/src/tint/transform/robustness.cc
@@ -20,10 +20,15 @@
#include "src/tint/program_builder.h"
#include "src/tint/sem/block_statement.h"
+#include "src/tint/sem/builtin.h"
#include "src/tint/sem/call.h"
+#include "src/tint/sem/function.h"
#include "src/tint/sem/index_accessor_expression.h"
+#include "src/tint/sem/load.h"
+#include "src/tint/sem/member_accessor_expression.h"
#include "src/tint/sem/statement.h"
#include "src/tint/sem/value_expression.h"
+#include "src/tint/transform/utils/hoist_to_decl_before.h"
#include "src/tint/type/reference.h"
TINT_INSTANTIATE_TYPEINFO(tint::transform::Robustness);
@@ -36,16 +41,147 @@
/// PIMPL state for the transform
struct Robustness::State {
/// Constructor
- /// @param program the source program
- /// @param omitted the omitted address spaces
- State(const Program* program, std::unordered_set<builtin::AddressSpace>&& omitted)
- : src(program), omitted_address_spaces(std::move(omitted)) {}
+ /// @param p the source program
+ /// @param c the transform config
+ State(const Program* p, Config&& c) : src(p), cfg(std::move(c)) {}
/// Runs the transform
/// @returns the new program or SkipTransform if the transform is not required
ApplyResult Run() {
- ctx.ReplaceAll([&](const ast::IndexAccessorExpression* expr) { return Transform(expr); });
- ctx.ReplaceAll([&](const ast::CallExpression* expr) { return Transform(expr); });
+ if (HasAction(Action::kPredicate)) {
+ AddPredicateParameters();
+ }
+
+ // Walk all the AST nodes in the module, starting with the leaf nodes.
+ // The most deeply nested expressions will come first.
+ for (auto* node : ctx.src->ASTNodes().Objects()) {
+ Switch(
+ node, //
+ [&](const ast::IndexAccessorExpression* e) {
+ // obj[idx]
+ // Array, matrix and vector indexing may require robustness transformation.
+ auto* expr = sem.Get(e)->Unwrap()->As<sem::IndexAccessorExpression>();
+ switch (ActionFor(expr)) {
+ case Action::kPredicate:
+ PredicateIndexAccessor(expr);
+ break;
+ case Action::kClamp:
+ ClampIndexAccessor(expr);
+ break;
+ case Action::kIgnore:
+ break;
+ }
+ },
+ [&](const ast::IdentifierExpression* e) {
+ // Identifiers may resolve to pointer lets, which may be predicated.
+ // Inspect.
+ if (auto* user = sem.Get<sem::VariableUser>(e)) {
+ auto* v = user->Variable();
+ if (v->Type()->Is<type::Pointer>()) {
+ // Propagate predicate from pointer
+ if (auto pred = predicates.Get(v->Declaration()->initializer)) {
+ predicates.Add(e, *pred);
+ }
+ }
+ }
+ },
+ [&](const ast::AccessorExpression* e) {
+ // obj.member
+ // Propagate the predication from the object to this expression.
+ if (auto pred = predicates.Get(e->object)) {
+ predicates.Add(e, *pred);
+ }
+ },
+ [&](const ast::UnaryOpExpression* e) {
+ // Includes address-of, or indirection
+ // Propagate the predication from the inner expression to this expression.
+ if (auto pred = predicates.Get(e->expr)) {
+ predicates.Add(e, *pred);
+ }
+ },
+ [&](const ast::AssignmentStatement* s) {
+ if (auto pred = predicates.Get(s->lhs)) {
+ // Assignment target is predicated
+ // Replace statement with condition on the predicate
+ ctx.Replace(s, b.If(*pred, b.Block(ctx.Clone(s))));
+ }
+ },
+ [&](const ast::CompoundAssignmentStatement* s) {
+ if (auto pred = predicates.Get(s->lhs)) {
+ // Assignment expression is predicated
+ // Replace statement with condition on the predicate
+ ctx.Replace(s, b.If(*pred, b.Block(ctx.Clone(s))));
+ }
+ },
+ [&](const ast::IncrementDecrementStatement* s) {
+ if (auto pred = predicates.Get(s->lhs)) {
+ // Assignment expression is predicated
+ // Replace statement with condition on the predicate
+ ctx.Replace(s, b.If(*pred, b.Block(ctx.Clone(s))));
+ }
+ },
+ [&](const ast::CallExpression* e) {
+ if (auto* call = sem.Get<sem::Call>(e)) {
+ Switch(
+ call->Target(), //
+ [&](const sem::Builtin* builtin) {
+ // Calls to builtins may require robustness transformation.
+ // Inspect.
+ if (builtin->IsTexture()) {
+ switch (cfg.texture_action) {
+ case Action::kPredicate:
+ PredicateTextureBuiltin(call, builtin);
+ break;
+ case Action::kClamp:
+ ClampTextureBuiltin(call, builtin);
+ break;
+ case Action::kIgnore:
+ break;
+ }
+ } else {
+ MaybePredicateNonTextureBuiltin(call, builtin);
+ }
+ },
+ [&](const sem::Function* fn) {
+ // Calls to user function may require passing additional predicate
+ // arguments.
+ InsertPredicateArguments(call, fn);
+ });
+ }
+ });
+
+ // Check whether the node is an expression that:
+ // * Has a predicate
+ // * Is of a non-pointer or non-reference type
+ // If the above is true, then we need to predicate evaluation of this expression by
+ // replacing `expr` with `predicated_expr` and injecting the following above the
+ // expression's statement:
+ //
+ // var predicated_expr : expr_ty;
+ // if (predicate) {
+ // predicated_expr = expr;
+ // }
+ //
+ if (auto* expr = node->As<ast::Expression>()) {
+ if (auto pred = predicates.Get(expr)) {
+ // Expression is predicated
+ auto* sem_expr = sem.GetVal(expr);
+ if (!sem_expr->Type()->IsAnyOf<type::Reference, type::Pointer>()) {
+ auto pred_load = b.Symbols().New("predicated_expr");
+ auto ty = CreateASTTypeFor(ctx, sem_expr->Type());
+ hoist.InsertBefore(sem_expr->Stmt(), b.Decl(b.Var(pred_load, ty)));
+ hoist.InsertBefore(
+ sem_expr->Stmt(),
+ b.If(*pred, b.Block(b.Assign(pred_load, ctx.Clone(expr)))));
+ ctx.Replace(expr, b.Expr(pred_load));
+
+ // The predication has been consumed for this expression.
+ // Don't predicate expressions that use this expression.
+ predicates.Remove(expr);
+ }
+ }
+ }
+ }
ctx.Clone();
return Program(std::move(b));
@@ -54,227 +190,486 @@
private:
/// The source program
const Program* const src;
+ /// The transform's config
+ Config cfg;
/// The target program builder
- ProgramBuilder b;
+ ProgramBuilder b{};
/// The clone context
CloneContext ctx = {&b, src, /* auto_clone_symbols */ true};
+ /// Helper for hoisting declarations
+ HoistToDeclBefore hoist{ctx};
+ /// Alias to the source program's semantic info
+ const sem::Info& sem = ctx.src->Sem();
+ /// Map of expression to predicate condition
+ utils::Hashmap<const ast::Expression*, Symbol, 32> predicates{};
- /// Set of address spaces to not apply the transform to
- std::unordered_set<builtin::AddressSpace> omitted_address_spaces;
-
- /// Apply bounds clamping to array, vector and matrix indexing
- /// @param expr the array, vector or matrix index expression
- /// @return the clamped replacement expression, or nullptr if `expr` should be cloned without
- /// changes.
- const ast::IndexAccessorExpression* Transform(const ast::IndexAccessorExpression* expr) {
- auto* sem = src->Sem().Get(expr)->Unwrap()->As<sem::IndexAccessorExpression>();
- auto* ret_type = sem->Type();
-
- auto* ref = ret_type->As<type::Reference>();
- if (ref && omitted_address_spaces.count(ref->AddressSpace()) != 0) {
- return nullptr;
- }
-
- // idx return the cloned index expression, as a u32.
- auto idx = [&]() -> const ast::Expression* {
- auto* i = ctx.Clone(expr->index);
- if (sem->Index()->Type()->is_signed_integer_scalar()) {
- return b.Call<u32>(i); // u32(idx)
- }
- return i;
- };
-
- auto* clamped_idx = Switch(
- sem->Object()->Type()->UnwrapRef(), //
+ /// @return the `u32` typed expression that represents the maximum indexable value for the index
+ /// accessor @p expr, or nullptr if there is no robustness limit for this expression.
+ const ast::Expression* DynamicLimitFor(const sem::IndexAccessorExpression* expr) {
+ auto* obj_type = expr->Object()->Type();
+ return Switch(
+ obj_type->UnwrapRef(), //
[&](const type::Vector* vec) -> const ast::Expression* {
- if (sem->Index()->ConstantValue()) {
+ if (expr->Index()->ConstantValue() || expr->Index()->Is<sem::Swizzle>()) {
// Index and size is constant.
// Validation will have rejected any OOB accesses.
return nullptr;
}
-
- return b.Call("min", idx(), u32(vec->Width() - 1u));
+ return b.Expr(u32(vec->Width() - 1u));
},
[&](const type::Matrix* mat) -> const ast::Expression* {
- if (sem->Index()->ConstantValue()) {
+ if (expr->Index()->ConstantValue()) {
// Index and size is constant.
// Validation will have rejected any OOB accesses.
return nullptr;
}
-
- return b.Call("min", idx(), u32(mat->columns() - 1u));
+ return b.Expr(u32(mat->columns() - 1u));
},
[&](const type::Array* arr) -> const ast::Expression* {
- const ast::Expression* max = nullptr;
if (arr->Count()->Is<type::RuntimeArrayCount>()) {
// Size is unknown until runtime.
// Must clamp, even if the index is constant.
- auto* arr_ptr = b.AddressOf(ctx.Clone(expr->object));
- max = b.Sub(b.Call("arrayLength", arr_ptr), 1_u);
- } else if (auto count = arr->ConstantCount()) {
- if (sem->Index()->ConstantValue()) {
+
+ auto* arr_ptr = b.AddressOf(ctx.Clone(expr->Object()->Declaration()));
+ return b.Sub(b.Call(sem::BuiltinType::kArrayLength, arr_ptr), 1_u);
+ }
+ if (auto count = arr->ConstantCount()) {
+ if (expr->Index()->ConstantValue()) {
// Index and size is constant.
// Validation will have rejected any OOB accesses.
return nullptr;
}
- max = b.Expr(u32(count.value() - 1u));
- } else {
- // Note: Don't be tempted to use the array override variable as an expression
- // here, the name might be shadowed!
- b.Diagnostics().add_error(diag::System::Transform,
- type::Array::kErrExpectedConstantCount);
- return nullptr;
+ return b.Expr(u32(count.value() - 1u));
}
-
- return b.Call("min", idx(), max);
+ // Note: Don't be tempted to use the array override variable as an expression here,
+ // the name might be shadowed!
+ b.Diagnostics().add_error(diag::System::Transform,
+ type::Array::kErrExpectedConstantCount);
+ return nullptr;
},
- [&](Default) {
+ [&](Default) -> const ast::Expression* {
TINT_ICE(Transform, b.Diagnostics())
<< "unhandled object type in robustness of array index: "
- << src->FriendlyName(ret_type->UnwrapRef());
+ << src->FriendlyName(obj_type->UnwrapRef());
return nullptr;
});
-
- if (!clamped_idx) {
- return nullptr; // Clamping not needed
- }
-
- auto idx_src = ctx.Clone(expr->source);
- auto* idx_obj = ctx.Clone(expr->object);
- return b.IndexAccessor(idx_src, idx_obj, clamped_idx);
}
- /// @param type builtin type
- /// @returns true if the given builtin is a texture function that requires
- /// argument clamping,
- bool TextureBuiltinNeedsClamping(sem::BuiltinType type) {
- return type == sem::BuiltinType::kTextureLoad || type == sem::BuiltinType::kTextureStore;
+ /// Transform the program to insert additional predicate parameters to all user functions that
+ /// have a pointer parameter type in an address space that has predicate action.
+ void AddPredicateParameters() {
+ for (auto* fn : src->AST().Functions()) {
+ for (auto* param : fn->params) {
+ auto* sem_param = sem.Get(param);
+ if (auto* ptr = sem_param->Type()->As<type::Pointer>()) {
+ if (ActionFor(ptr->AddressSpace()) == Action::kPredicate) {
+ auto name = b.Symbols().New(src->Symbols().NameFor(param->name->symbol) +
+ "_predicate");
+ ctx.InsertAfter(fn->params, param, b.Param(name, b.ty.bool_()));
+
+ // Associate the pointer parameter expressions with the predicate.
+ for (auto* user : sem_param->Users()) {
+ predicates.Add(user->Declaration(), name);
+ }
+ }
+ }
+ }
+ }
}
- /// Apply bounds clamping to the coordinates, array index and level arguments
- /// of the `textureLoad()` and `textureStore()` builtins.
- /// @param expr the builtin call expression
- /// @return the clamped replacement call expression, or nullptr if `expr`
- /// should be cloned without changes.
- const ast::CallExpression* Transform(const ast::CallExpression* expr) {
- auto* call = src->Sem().Get(expr)->UnwrapMaterialize()->As<sem::Call>();
- auto* call_target = call->Target();
- auto* builtin = call_target->As<sem::Builtin>();
- if (!builtin || !TextureBuiltinNeedsClamping(builtin->Type())) {
- return nullptr; // No transform, just clone.
+ /// Transforms call expressions to user functions, inserting additional predicate arguments
+ /// after all pointer parameters with a type in an address space that has predicate action.
+ void InsertPredicateArguments(const sem::Call* call, const sem::Function* fn) {
+ auto* expr = call->Declaration();
+ for (size_t i = 0; i < fn->Parameters().Length(); i++) {
+ auto* param = fn->Parameters()[i];
+ if (auto* ptr = param->Type()->As<type::Pointer>()) {
+ if (ActionFor(ptr->AddressSpace()) == Action::kPredicate) {
+ auto* arg = expr->args[i];
+ if (auto predicate = predicates.Get(arg)) {
+ ctx.InsertAfter(expr->args, arg, b.Expr(*predicate));
+ } else {
+ ctx.InsertAfter(expr->args, arg, b.Expr(true));
+ }
+ }
+ }
}
+ }
+
+ /// Applies predication to the index on an array, vector or matrix.
+ /// @param expr the index accessor expression.
+ void PredicateIndexAccessor(const sem::IndexAccessorExpression* expr) {
+ auto* obj = expr->Object()->Declaration();
+ auto* idx = expr->Index()->Declaration();
+ auto* max = DynamicLimitFor(expr);
+ if (!max) {
+ // robustness is not required
+ // Just propagate predicate from object
+ if (auto pred = predicates.Get(obj)) {
+ predicates.Add(expr->Declaration(), *pred);
+ }
+ return;
+ }
+
+ auto* stmt = expr->Stmt();
+ auto obj_pred = *predicates.GetOrZero(obj);
+
+ auto idx_let = b.Symbols().New("index");
+ auto pred = b.Symbols().New("predicate");
+
+ hoist.InsertBefore(stmt, b.Decl(b.Let(idx_let, ctx.Clone(idx))));
+ ctx.Replace(idx, b.Expr(idx_let));
+
+ auto* cond = b.LessThanEqual(b.Call<u32>(b.Expr(idx_let)), max);
+ if (obj_pred.IsValid()) {
+ cond = b.And(b.Expr(obj_pred), cond);
+ }
+ hoist.InsertBefore(stmt, b.Decl(b.Let(pred, cond)));
+
+ predicates.Add(expr->Declaration(), pred);
+ }
+
+ /// Applies bounds clamping to the index on an array, vector or matrix.
+ /// @param expr the index accessor expression.
+ void ClampIndexAccessor(const sem::IndexAccessorExpression* expr) {
+ auto* max = DynamicLimitFor(expr);
+ if (!max) {
+ return; // robustness is not required
+ }
+
+ auto* expr_sem = expr->Unwrap()->As<sem::IndexAccessorExpression>();
+
+ auto idx = ctx.Clone(expr->Declaration()->index);
+ if (expr_sem->Index()->Type()->is_signed_integer_scalar()) {
+ idx = b.Call<u32>(idx); // u32(idx)
+ }
+ auto* clamped_idx = b.Call(sem::BuiltinType::kMin, idx, max);
+ ctx.Replace(expr->Declaration()->index, clamped_idx);
+ }
+
+ /// Applies predication to the non-texture builtin call, if required.
+ void MaybePredicateNonTextureBuiltin(const sem::Call* call, const sem::Builtin* builtin) {
+ // Gather the predications for the builtin arguments
+ const ast::Expression* predicate = nullptr;
+ for (auto* arg : call->Declaration()->args) {
+ if (auto pred = predicates.Get(arg)) {
+ predicate = And(predicate, b.Expr(*pred));
+ }
+ }
+
+ if (predicate) {
+ if (builtin->Type() == sem::BuiltinType::kWorkgroupUniformLoad) {
+ // https://www.w3.org/TR/WGSL/#workgroupUniformLoad-builtin:
+ // "Executes a control barrier synchronization function that affects memory and
+ // atomic operations in the workgroup address space."
+ // Because the call acts like a control barrier, we need to make sure that we still
+ // trigger a workgroup barrier if the predicate fails.
+ PredicateCall(call, predicate,
+ b.Block(b.CallStmt(b.Call(sem::BuiltinType::kWorkgroupBarrier))));
+ } else {
+ PredicateCall(call, predicate);
+ }
+ }
+ }
+
+ /// Applies predication to texture builtins, based on whether the coordinates, array index and
+ /// level arguments are all in bounds.
+ void PredicateTextureBuiltin(const sem::Call* call, const sem::Builtin* builtin) {
+ if (!TextureBuiltinNeedsRobustness(builtin->Type())) {
+ return;
+ }
+
+ auto* expr = call->Declaration();
+ auto* stmt = call->Stmt();
// Indices of the mandatory texture and coords parameters, and the optional
// array and level parameters.
auto& signature = builtin->Signature();
- auto texture_idx = signature.IndexOf(sem::ParameterUsage::kTexture);
- auto coords_idx = signature.IndexOf(sem::ParameterUsage::kCoords);
- auto array_idx = signature.IndexOf(sem::ParameterUsage::kArrayIndex);
- auto level_idx = signature.IndexOf(sem::ParameterUsage::kLevel);
+ auto texture_arg_idx = signature.IndexOf(sem::ParameterUsage::kTexture);
+ auto coords_arg_idx = signature.IndexOf(sem::ParameterUsage::kCoords);
+ auto array_arg_idx = signature.IndexOf(sem::ParameterUsage::kArrayIndex);
+ auto level_arg_idx = signature.IndexOf(sem::ParameterUsage::kLevel);
- auto* texture_arg = expr->args[static_cast<size_t>(texture_idx)];
- auto* coords_arg = expr->args[static_cast<size_t>(coords_idx)];
- auto* coords_ty = builtin->Parameters()[static_cast<size_t>(coords_idx)]->Type();
+ auto* texture_arg = expr->args[static_cast<size_t>(texture_arg_idx)];
- auto width_of = [&](const type::Type* ty) {
- if (auto* vec = ty->As<type::Vector>()) {
- return vec->Width();
+ // Build the builtin predicate from the arguments
+ const ast::Expression* predicate = nullptr;
+
+ Symbol level_idx, num_levels;
+ if (level_arg_idx >= 0) {
+ auto* param = builtin->Parameters()[static_cast<size_t>(level_arg_idx)];
+ if (param->Type()->is_integer_scalar()) {
+ // let level_idx = u32(level-arg);
+ level_idx = b.Symbols().New("level_idx");
+ auto* arg = expr->args[static_cast<size_t>(level_arg_idx)];
+ hoist.InsertBefore(stmt,
+ b.Decl(b.Let(level_idx, CastToUnsigned(ctx.Clone(arg), 1u))));
+
+ // let num_levels = textureNumLevels(texture-arg);
+ num_levels = b.Symbols().New("num_levels");
+ hoist.InsertBefore(
+ stmt, b.Decl(b.Let(num_levels, b.Call(sem::BuiltinType::kTextureNumLevels,
+ ctx.Clone(texture_arg)))));
+
+ // predicate: level_idx < num_levels
+ predicate = And(predicate, b.LessThan(level_idx, num_levels));
+
+ // Replace the level argument with `level_idx`
+ ctx.Replace(arg, b.Expr(level_idx));
}
- return 1u;
- };
- auto scalar_or_vec_ty = [&](ast::Type scalar, uint32_t width) {
- if (width > 1) {
- return b.ty.vec(scalar, width);
+ }
+
+ Symbol coords;
+ if (coords_arg_idx >= 0) {
+ auto* param = builtin->Parameters()[static_cast<size_t>(coords_arg_idx)];
+ if (param->Type()->is_integer_scalar_or_vector()) {
+ // let coords = u32(coords-arg)
+ coords = b.Symbols().New("coords");
+ auto* arg = expr->args[static_cast<size_t>(coords_arg_idx)];
+ hoist.InsertBefore(stmt,
+ b.Decl(b.Let(coords, CastToUnsigned(b.Expr(ctx.Clone(arg)),
+ WidthOf(param->Type())))));
+
+ // predicate: all(coords < textureDimensions(texture))
+ auto* dimensions =
+ level_idx.IsValid()
+ ? b.Call(sem::BuiltinType::kTextureDimensions, ctx.Clone(texture_arg),
+ b.Call(sem::BuiltinType::kMin, b.Expr(level_idx),
+ b.Sub(num_levels, 1_a)))
+ : b.Call(sem::BuiltinType::kTextureDimensions, ctx.Clone(texture_arg));
+ predicate =
+ And(predicate, b.Call(sem::BuiltinType::kAll, b.LessThan(coords, dimensions)));
+
+ // Replace the level argument with `coord`
+ ctx.Replace(arg, b.Expr(coords));
}
- return scalar;
- };
- auto scalar_or_vec = [&](const ast::Expression* scalar,
- uint32_t width) -> const ast::Expression* {
- if (width > 1) {
- return b.Call(b.ty.vec<Infer>(width), scalar);
+ }
+
+ if (array_arg_idx >= 0) {
+ // let array_idx = u32(array-arg)
+ auto* arg = expr->args[static_cast<size_t>(array_arg_idx)];
+ auto* num_layers = b.Call(sem::BuiltinType::kTextureNumLayers, ctx.Clone(texture_arg));
+ auto array_idx = b.Symbols().New("array_idx");
+ hoist.InsertBefore(stmt, b.Decl(b.Let(array_idx, CastToUnsigned(ctx.Clone(arg), 1u))));
+
+ // predicate: array_idx < textureNumLayers(texture)
+ predicate = And(predicate, b.LessThan(array_idx, num_layers));
+
+ // Replace the array index argument with `array_idx`
+ ctx.Replace(arg, b.Expr(array_idx));
+ }
+
+ if (predicate) {
+ PredicateCall(call, predicate);
+ }
+ }
+
+ /// Applies bounds clamping to the coordinates, array index and level arguments of the texture
+ /// builtin.
+ void ClampTextureBuiltin(const sem::Call* call, const sem::Builtin* builtin) {
+ if (!TextureBuiltinNeedsRobustness(builtin->Type())) {
+ return;
+ }
+
+ auto* expr = call->Declaration();
+ auto* stmt = call->Stmt();
+
+ // Indices of the mandatory texture and coords parameters, and the optional
+ // array and level parameters.
+ auto& signature = builtin->Signature();
+ auto texture_arg_idx = signature.IndexOf(sem::ParameterUsage::kTexture);
+ auto coords_arg_idx = signature.IndexOf(sem::ParameterUsage::kCoords);
+ auto array_arg_idx = signature.IndexOf(sem::ParameterUsage::kArrayIndex);
+ auto level_arg_idx = signature.IndexOf(sem::ParameterUsage::kLevel);
+
+ auto* texture_arg = expr->args[static_cast<size_t>(texture_arg_idx)];
+
+ // If the level is provided, then we need to clamp this. As the level is used by
+ // textureDimensions() and the texture[Load|Store]() calls, we need to clamp both usages.
+ Symbol level_idx;
+ if (level_arg_idx >= 0) {
+ const auto* param = builtin->Parameters()[static_cast<size_t>(level_arg_idx)];
+ if (param->Type()->is_integer_scalar()) {
+ const auto* arg = expr->args[static_cast<size_t>(level_arg_idx)];
+ level_idx = b.Symbols().New("level_idx");
+ const auto* num_levels =
+ b.Call(sem::BuiltinType::kTextureNumLevels, ctx.Clone(texture_arg));
+ const auto* max = b.Sub(num_levels, 1_a);
+ hoist.InsertBefore(
+ stmt, b.Decl(b.Let(level_idx, b.Call(sem::BuiltinType::kMin,
+ b.Call<u32>(ctx.Clone(arg)), max))));
+ ctx.Replace(arg, b.Expr(level_idx));
}
- return scalar;
- };
- auto cast_to_signed = [&](const ast::Expression* val, uint32_t width) {
- return b.Call(scalar_or_vec_ty(b.ty.i32(), width), val);
- };
- auto cast_to_unsigned = [&](const ast::Expression* val, uint32_t width) {
- return b.Call(scalar_or_vec_ty(b.ty.u32(), width), val);
- };
-
- // If the level is provided, then we need to clamp this. As the level is
- // used by textureDimensions() and the texture[Load|Store]() calls, we need
- // to clamp both usages.
- // TODO(bclayton): We probably want to place this into a let so that the
- // calculation can be reused. This is fiddly to get right.
- std::function<const ast::Expression*()> level_arg;
- if (level_idx >= 0) {
- level_arg = [&] {
- const auto* arg = expr->args[static_cast<size_t>(level_idx)];
- const auto* target_ty =
- builtin->Parameters()[static_cast<size_t>(level_idx)]->Type();
- const auto* num_levels = b.Call("textureNumLevels", ctx.Clone(texture_arg));
-
- // TODO(crbug.com/tint/1526) remove when num_levels returns u32
- num_levels = cast_to_unsigned(num_levels, 1u);
-
- const auto* unsigned_max = b.Sub(num_levels, 1_a);
- if (target_ty->is_signed_integer_scalar()) {
- const auto* signed_max = cast_to_signed(unsigned_max, 1u);
- return b.Call("clamp", ctx.Clone(arg), 0_a, signed_max);
- } else {
- return b.Call("min", ctx.Clone(arg), unsigned_max);
- }
- };
}
// Clamp the coordinates argument
- {
- const auto* target_ty = coords_ty;
- const auto width = width_of(target_ty);
- const auto* texture_dims =
- level_arg ? b.Call("textureDimensions", ctx.Clone(texture_arg), level_arg())
- : b.Call("textureDimensions", ctx.Clone(texture_arg));
+ if (coords_arg_idx >= 0) {
+ const auto* param = builtin->Parameters()[static_cast<size_t>(coords_arg_idx)];
+ if (param->Type()->is_integer_scalar_or_vector()) {
+ auto* arg = expr->args[static_cast<size_t>(coords_arg_idx)];
+ const auto width = WidthOf(param->Type());
+ const auto* dimensions =
+ level_idx.IsValid()
+ ? b.Call(sem::BuiltinType::kTextureDimensions, ctx.Clone(texture_arg),
+ level_idx)
+ : b.Call(sem::BuiltinType::kTextureDimensions, ctx.Clone(texture_arg));
- // TODO(crbug.com/tint/1526) remove when texture_dims returns u32 or vecN<u32>
- texture_dims = cast_to_unsigned(texture_dims, width);
-
- // texture_dims is u32 or vecN<u32>
- const auto* unsigned_max = b.Sub(texture_dims, scalar_or_vec(b.Expr(1_a), width));
- if (target_ty->is_signed_integer_scalar_or_vector()) {
- const auto* zero = scalar_or_vec(b.Expr(0_a), width);
- const auto* signed_max = cast_to_signed(unsigned_max, width);
- ctx.Replace(coords_arg, b.Call("clamp", ctx.Clone(coords_arg), zero, signed_max));
- } else {
- ctx.Replace(coords_arg, b.Call("min", ctx.Clone(coords_arg), unsigned_max));
+ // dimensions is u32 or vecN<u32>
+ const auto* unsigned_max = b.Sub(dimensions, ScalarOrVec(b.Expr(1_a), width));
+ if (param->Type()->is_signed_integer_scalar_or_vector()) {
+ const auto* zero = ScalarOrVec(b.Expr(0_a), width);
+ const auto* signed_max = CastToSigned(unsigned_max, width);
+ ctx.Replace(arg,
+ b.Call(sem::BuiltinType::kClamp, ctx.Clone(arg), zero, signed_max));
+ } else {
+ ctx.Replace(arg, b.Call(sem::BuiltinType::kMin, ctx.Clone(arg), unsigned_max));
+ }
}
}
// Clamp the array_index argument, if provided
- if (array_idx >= 0) {
- auto* target_ty = builtin->Parameters()[static_cast<size_t>(array_idx)]->Type();
- auto* arg = expr->args[static_cast<size_t>(array_idx)];
- auto* num_layers = b.Call("textureNumLayers", ctx.Clone(texture_arg));
-
- // TODO(crbug.com/tint/1526) remove when num_layers returns u32
- num_layers = cast_to_unsigned(num_layers, 1u);
+ if (array_arg_idx >= 0) {
+ auto* param = builtin->Parameters()[static_cast<size_t>(array_arg_idx)];
+ auto* arg = expr->args[static_cast<size_t>(array_arg_idx)];
+ auto* num_layers = b.Call(sem::BuiltinType::kTextureNumLayers, ctx.Clone(texture_arg));
const auto* unsigned_max = b.Sub(num_layers, 1_a);
- if (target_ty->is_signed_integer_scalar()) {
- const auto* signed_max = cast_to_signed(unsigned_max, 1u);
- ctx.Replace(arg, b.Call("clamp", ctx.Clone(arg), 0_a, signed_max));
+ if (param->Type()->is_signed_integer_scalar()) {
+ const auto* signed_max = CastToSigned(unsigned_max, 1u);
+ ctx.Replace(arg, b.Call(sem::BuiltinType::kClamp, ctx.Clone(arg), 0_a, signed_max));
} else {
- ctx.Replace(arg, b.Call("min", ctx.Clone(arg), unsigned_max));
+ ctx.Replace(arg, b.Call(sem::BuiltinType::kMin, ctx.Clone(arg), unsigned_max));
}
}
+ }
- // Clamp the level argument, if provided
- if (level_idx >= 0) {
- auto* arg = expr->args[static_cast<size_t>(level_idx)];
- ctx.Replace(arg, level_arg ? level_arg() : b.Expr(0_a));
+ /// @param type builtin type
+ /// @returns true if the given builtin is a texture function that requires predication or
+ /// clamping of arguments.
+ bool TextureBuiltinNeedsRobustness(sem::BuiltinType type) {
+ return type == sem::BuiltinType::kTextureLoad || type == sem::BuiltinType::kTextureStore ||
+ type == sem::BuiltinType::kTextureDimensions;
+ }
+
+ /// @returns a bitwise and of the two expressions, or the other expression if one is null.
+ const ast::Expression* And(const ast::Expression* lhs, const ast::Expression* rhs) {
+ if (lhs && rhs) {
+ return b.And(lhs, rhs);
}
+ if (lhs) {
+ return lhs;
+ }
+ return rhs;
+ }
- return nullptr; // Clone, which will use the argument replacements above.
+ /// Transforms a call statement or expression so that the expression is predicated by @p
+ /// predicate.
+ /// @param else_stmt - the statement to execute for the predication failure
+ void PredicateCall(const sem::Call* call,
+ const ast::Expression* predicate,
+ const ast::BlockStatement* else_stmt = nullptr) {
+ auto* expr = call->Declaration();
+ auto* stmt = call->Stmt();
+ auto* call_stmt = stmt->Declaration()->As<ast::CallStatement>();
+ if (call_stmt && call_stmt->expr == expr) {
+ // Wrap the statement in an if-statement with the predicate condition.
+ hoist.Replace(stmt, b.If(predicate, b.Block(ctx.Clone(stmt->Declaration())),
+ ProgramBuilder::ElseStmt(else_stmt)));
+ } else {
+ // Emit the following before the expression's statement:
+ // var predicated_value : return-type;
+ // if (predicate) {
+ // predicated_value = call(...);
+ // }
+ auto value = b.Symbols().New("predicated_value");
+ hoist.InsertBefore(stmt, b.Decl(b.Var(value, CreateASTTypeFor(ctx, call->Type()))));
+ hoist.InsertBefore(stmt, b.If(predicate, b.Block(b.Assign(value, ctx.Clone(expr))),
+ ProgramBuilder::ElseStmt(else_stmt)));
+
+ // Replace the call expression with `predicated_value`
+ ctx.Replace(expr, b.Expr(value));
+ }
+ }
+
+ /// @returns true if @p action is enabled for any address space
+ bool HasAction(Action action) const {
+ return action == cfg.function_action || //
+ action == cfg.texture_action || //
+ action == cfg.private_action || //
+ action == cfg.push_constant_action || //
+ action == cfg.storage_action || //
+ action == cfg.uniform_action || //
+ action == cfg.workgroup_action;
+ }
+
+ /// @returns the robustness action to perform for an OOB access with the expression @p expr
+ Action ActionFor(const sem::ValueExpression* expr) {
+ return Switch(
+ expr->Type(), //
+ [&](const type::Reference* t) { return ActionFor(t->AddressSpace()); },
+ [&](Default) { return cfg.value_action; });
+ }
+
+ /// @returns the robustness action to perform for an OOB access in the address space @p
+ /// address_space
+ Action ActionFor(builtin::AddressSpace address_space) {
+ switch (address_space) {
+ case builtin::AddressSpace::kFunction:
+ return cfg.function_action;
+ case builtin::AddressSpace::kHandle:
+ return cfg.texture_action;
+ case builtin::AddressSpace::kPrivate:
+ return cfg.private_action;
+ case builtin::AddressSpace::kPushConstant:
+ return cfg.push_constant_action;
+ case builtin::AddressSpace::kStorage:
+ return cfg.storage_action;
+ case builtin::AddressSpace::kUniform:
+ return cfg.uniform_action;
+ case builtin::AddressSpace::kWorkgroup:
+ return cfg.workgroup_action;
+ default:
+ break;
+ }
+ TINT_UNREACHABLE(Transform, b.Diagnostics()) << "unhandled address space" << address_space;
+ return Action::kDefault;
+ }
+
+ /// @returns the vector width of @p ty, or 1 if @p ty is not a vector
+ static uint32_t WidthOf(const type::Type* ty) {
+ if (auto* vec = ty->As<type::Vector>()) {
+ return vec->Width();
+ }
+ return 1u;
+ }
+
+ /// @returns a scalar or vector type with the element type @p scalar and width @p width
+ ast::Type ScalarOrVecTy(ast::Type scalar, uint32_t width) const {
+ if (width > 1) {
+ return b.ty.vec(scalar, width);
+ }
+ return scalar;
+ }
+
+ /// @returns a vector constructed with the scalar expression @p scalar if @p width > 1,
+ /// otherwise returns @p scalar.
+ const ast::Expression* ScalarOrVec(const ast::Expression* scalar, uint32_t width) {
+ if (width > 1) {
+ return b.Call(b.ty.vec<Infer>(width), scalar);
+ }
+ return scalar;
+ }
+
+ /// @returns @p val cast to a `vecN<i32>`, where `N` is @p width, or cast to i32 if @p width
+ /// is 1.
+ const ast::CallExpression* CastToSigned(const ast::Expression* val, uint32_t width) {
+ return b.Call(ScalarOrVecTy(b.ty.i32(), width), val);
+ }
+
+ /// @returns @p val cast to a `vecN<u32>`, where `N` is @p width, or cast to u32 if @p width
+ /// is 1.
+ const ast::CallExpression* CastToUnsigned(const ast::Expression* val, uint32_t width) {
+ return b.Call(ScalarOrVecTy(b.ty.u32(), width), val);
}
};
@@ -294,19 +689,7 @@
cfg = *cfg_data;
}
- std::unordered_set<builtin::AddressSpace> omitted_address_spaces;
- for (auto sc : cfg.omitted_address_spaces) {
- switch (sc) {
- case AddressSpace::kUniform:
- omitted_address_spaces.insert(builtin::AddressSpace::kUniform);
- break;
- case AddressSpace::kStorage:
- omitted_address_spaces.insert(builtin::AddressSpace::kStorage);
- break;
- }
- }
-
- return State{src, std::move(omitted_address_spaces)}.Run();
+ return State{src, std::move(cfg)}.Run();
}
} // namespace tint::transform
diff --git a/src/tint/transform/robustness.h b/src/tint/transform/robustness.h
index 780eab9..596db4d 100644
--- a/src/tint/transform/robustness.h
+++ b/src/tint/transform/robustness.h
@@ -15,29 +15,33 @@
#ifndef SRC_TINT_TRANSFORM_ROBUSTNESS_H_
#define SRC_TINT_TRANSFORM_ROBUSTNESS_H_
-#include <unordered_set>
-
#include "src/tint/transform/transform.h"
-// Forward declarations
-namespace tint::ast {
-class IndexAccessorExpression;
-class CallExpression;
-} // namespace tint::ast
-
namespace tint::transform {
-/// This transform is responsible for clamping all array accesses to be within
-/// the bounds of the array. Any access before the start of the array will clamp
-/// to zero and any access past the end of the array will clamp to
-/// (array length - 1).
-/// @note This transform must come before the BuiltinPolyfill transform
+/// This transform is responsible for ensuring that all out of bounds accesses are prevented,
+/// either by conditioning the access (predication) or through clamping of the index to keep the
+/// access in bounds.
+/// @note Robustness must come after:
+/// * PromoteSideEffectsToDecl as Robustness requires side-effecting expressions to be hoisted
+/// to their own statements.
+/// Robustness must come before:
+/// * BuiltinPolyfill as 'clamp' and binary operators may need to be polyfilled.
+/// * CanonicalizeEntryPointIO as the transform does not support the 'in' and 'out' address
+/// spaces.
class Robustness final : public Castable<Robustness, Transform> {
public:
- /// Address space to be skipped in the transform
- enum class AddressSpace {
- kUniform,
- kStorage,
+ /// Robustness action for out-of-bounds indexing.
+ enum class Action {
+ /// Do nothing to prevent the out-of-bounds action.
+ kIgnore,
+ /// Clamp the index to be within bounds.
+ kClamp,
+ /// Do not execute the read or write if the index is out-of-bounds.
+ kPredicate,
+
+ /// The default action
+ kDefault = kClamp,
};
/// Configuration options for the transform
@@ -55,9 +59,24 @@
/// @returns this Config
Config& operator=(const Config&);
- /// Address spaces to omit from apply the transform to.
- /// This allows for optimizing on hardware that provide safe accesses.
- std::unordered_set<AddressSpace> omitted_address_spaces;
+ /// Robustness action for values
+ Action value_action = Action::kDefault;
+
+ /// Robustness action for non-sampling texture operations
+ Action texture_action = Action::kDefault;
+
+ /// Robustness action for variables in the 'function' address space
+ Action function_action = Action::kDefault;
+ /// Robustness action for variables in the 'private' address space
+ Action private_action = Action::kDefault;
+ /// Robustness action for variables in the 'push_constant' address space
+ Action push_constant_action = Action::kDefault;
+ /// Robustness action for variables in the 'storage' address space
+ Action storage_action = Action::kDefault;
+ /// Robustness action for variables in the 'uniform' address space
+ Action uniform_action = Action::kDefault;
+ /// Robustness action for variables in the 'workgroup' address space
+ Action workgroup_action = Action::kDefault;
};
/// Constructor
diff --git a/src/tint/transform/robustness_test.cc b/src/tint/transform/robustness_test.cc
index 9bb9527..ee583cf 100644
--- a/src/tint/transform/robustness_test.cc
+++ b/src/tint/transform/robustness_test.cc
@@ -17,11 +17,345 @@
#include "src/tint/transform/test_helper.h"
namespace tint::transform {
+
+static std::ostream& operator<<(std::ostream& out, Robustness::Action action) {
+ switch (action) {
+ case Robustness::Action::kIgnore:
+ return out << "ignore";
+ case Robustness::Action::kClamp:
+ return out << "clamp";
+ case Robustness::Action::kPredicate:
+ return out << "predicate";
+ }
+ return out << "unknown";
+}
+
namespace {
-using RobustnessTest = TransformTest;
+DataMap Config(Robustness::Action action) {
+ Robustness::Config cfg;
+ cfg.value_action = action;
+ cfg.texture_action = action;
+ cfg.function_action = action;
+ cfg.private_action = action;
+ cfg.push_constant_action = action;
+ cfg.storage_action = action;
+ cfg.uniform_action = action;
+ cfg.workgroup_action = action;
+ DataMap data;
+ data.Add<Robustness::Config>(cfg);
+ return data;
+}
-TEST_F(RobustnessTest, Array_Let_Idx_Clamp) {
+const char* Expect(Robustness::Action action,
+ const char* expect_ignore,
+ const char* expect_clamp,
+ const char* expect_predicate) {
+ switch (action) {
+ case Robustness::Action::kIgnore:
+ return expect_ignore;
+ case Robustness::Action::kClamp:
+ return expect_clamp;
+ case Robustness::Action::kPredicate:
+ return expect_predicate;
+ }
+ return "<invalid action>";
+}
+
+using RobustnessTest = TransformTestWithParam<Robustness::Action>;
+
+////////////////////////////////////////////////////////////////////////////////
+// Constant sized array
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithLiteral) {
+ auto* src = R"(
+fn f() {
+ var b : f32 = array<f32, 3>()[1i];
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithConst) {
+ auto* src = R"(
+const c : u32 = 1u;
+
+fn f() {
+ let b : f32 = array<f32, 3>()[c];
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithLet) {
+ auto* src = R"(
+fn f() {
+ let l : u32 = 1u;
+ let b : f32 = array<f32, 3>()[l];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+fn f() {
+ let l : u32 = 1u;
+ let b : f32 = array<f32, 3>()[min(l, 2u)];
+}
+)",
+ /* predicate */ R"(
+fn f() {
+ let l : u32 = 1u;
+ let index = l;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = array<f32, 3>()[index];
+ }
+ let b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithRuntimeArrayIndex) {
+ auto* src = R"(
+var<private> i : u32;
+
+fn f() {
+ let a = array<f32, 3>();
+ let b = array<i32, 5>();
+ var c : f32 = a[b[i]];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> i : u32;
+
+fn f() {
+ let a = array<f32, 3>();
+ let b = array<i32, 5>();
+ var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)];
+}
+)",
+ /* predicate */ R"(
+var<private> i : u32;
+
+fn f() {
+ let a = array<f32, 3>();
+ let b = array<i32, 5>();
+ let index = i;
+ let predicate = (u32(index) <= 4u);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = b[index];
+ }
+ let index_1 = predicated_expr;
+ let predicate_1 = (u32(index_1) <= 2u);
+ var predicated_expr_1 : f32;
+ if (predicate_1) {
+ predicated_expr_1 = a[index_1];
+ }
+ var c : f32 = predicated_expr_1;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithRuntimeExpression) {
+ auto* src = R"(
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = array<f32, 3>()[((c + 2) - 3)];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = array<f32, 3>()[min(u32(((c + 2) - 3)), 2u)];
+}
+)",
+ /* predicate */ R"(
+var<private> c : i32;
+
+fn f() {
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = array<f32, 3>()[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_NestedConstantSizedArraysVal_IndexWithRuntimeExpressions) {
+ auto* src = R"(
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let a = array<array<array<f32, 1>, 2>, 3>();
+ var r = a[x][y][z];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let a = array<array<array<f32, 1>, 2>, 3>();
+ var r = a[min(u32(x), 2u)][min(u32(y), 1u)][min(u32(z), 0u)];
+}
+)",
+ /* predicate */ R"(
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let a = array<array<array<f32, 1>, 2>, 3>();
+ let index = x;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : array<array<f32, 1u>, 2u>;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ let index_1 = y;
+ let predicate_1 = (u32(index_1) <= 1u);
+ var predicated_expr_1 : array<f32, 1u>;
+ if (predicate_1) {
+ predicated_expr_1 = predicated_expr[index_1];
+ }
+ let index_2 = z;
+ let predicate_2 = (u32(index_2) <= 0u);
+ var predicated_expr_2 : f32;
+ if (predicate_2) {
+ predicated_expr_2 = predicated_expr_1[index_2];
+ }
+ var r = predicated_expr_2;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithOverride) {
+ auto* src = R"(
+@id(1300) override idx : i32;
+
+fn f() {
+ let a = array<f32, 4>();
+ var b : f32 = a[idx];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@id(1300) override idx : i32;
+
+fn f() {
+ let a = array<f32, 4>();
+ var b : f32 = a[min(u32(idx), 3u)];
+}
+)",
+ /* predicate */ R"(
+@id(1300) override idx : i32;
+
+fn f() {
+ let a = array<f32, 4>();
+ let index = idx;
+ let predicate = (u32(index) <= 3u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithLiteral) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ var b : f32 = a[1i];
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithConst) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+const c : u32 = 1u;
+
+fn f() {
+ let b : f32 = a[c];
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithLet) {
auto* src = R"(
var<private> a : array<f32, 3>;
@@ -31,97 +365,37 @@
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
var<private> a : array<f32, 3>;
fn f() {
let l : u32 = 1u;
let b : f32 = a[min(l, 2u)];
}
-)";
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
- auto got = Run<Robustness>(src);
+fn f() {
+ let l : u32 = 1u;
+ let index = l;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ let b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, Array_Let_Idx_Clamp_OutOfOrder) {
- auto* src = R"(
-fn f() {
- let c : u32 = 1u;
- let b : f32 = a[c];
-}
-
-var<private> a : array<f32, 3>;
-)";
-
- auto* expect = R"(
-fn f() {
- let c : u32 = 1u;
- let b : f32 = a[min(c, 2u)];
-}
-
-var<private> a : array<f32, 3>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Const_Idx_Clamp) {
- auto* src = R"(
-var<private> a : array<f32, 3>;
-
-const c : u32 = 1u;
-
-fn f() {
- let b : f32 = a[c];
-}
-)";
-
- auto* expect = R"(
-var<private> a : array<f32, 3>;
-
-const c : u32 = 1u;
-
-fn f() {
- let b : f32 = a[c];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Const_Idx_Clamp_OutOfOrder) {
- auto* src = R"(
-fn f() {
- let b : f32 = a[c];
-}
-
-const c : u32 = 1u;
-
-var<private> a : array<f32, 3>;
-)";
-
- auto* expect = R"(
-fn f() {
- let b : f32 = a[c];
-}
-
-const c : u32 = 1u;
-
-var<private> a : array<f32, 3>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Idx_Nested_Scalar) {
+TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeArrayIndex) {
auto* src = R"(
var<private> a : array<f32, 3>;
@@ -130,11 +404,13 @@
var<private> i : u32;
fn f() {
- var c : f32 = a[ b[i] ];
+ var c : f32 = a[b[i]];
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
var<private> a : array<f32, 3>;
var<private> b : array<i32, 5>;
@@ -144,99 +420,50 @@
fn f() {
var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)];
}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Idx_Nested_Scalar_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var c : f32 = a[ b[i] ];
-}
-
-var<private> i : u32;
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
var<private> b : array<i32, 5>;
-var<private> a : array<f32, 3>;
-)";
-
- auto* expect = R"(
-fn f() {
- var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)];
-}
-
var<private> i : u32;
-var<private> b : array<i32, 5>;
+fn f() {
+ let index = i;
+ let predicate = (u32(index) <= 4u);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = b[index];
+ }
+ let index_1 = predicated_expr;
+ let predicate_1 = (u32(index_1) <= 2u);
+ var predicated_expr_1 : f32;
+ if (predicate_1) {
+ predicated_expr_1 = a[index_1];
+ }
+ var c : f32 = predicated_expr_1;
+}
+)");
-var<private> a : array<f32, 3>;
-)";
-
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, Array_Idx_Scalar) {
- auto* src = R"(
-var<private> a : array<f32, 3>;
-
-fn f() {
- var b : f32 = a[1i];
-}
-)";
-
- auto* expect = R"(
-var<private> a : array<f32, 3>;
-
-fn f() {
- var b : f32 = a[1i];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Idx_Scalar_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a[1i];
-}
-
-var<private> a : array<f32, 3>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a[1i];
-}
-
-var<private> a : array<f32, 3>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Idx_Expr) {
+TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeExpression) {
auto* src = R"(
var<private> a : array<f32, 3>;
var<private> c : i32;
fn f() {
- var b : f32 = a[c + 2 - 3];
+ var b : f32 = a[((c + 2) - 3)];
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
var<private> a : array<f32, 3>;
var<private> c : i32;
@@ -244,852 +471,3040 @@
fn f() {
var b : f32 = a[min(u32(((c + 2) - 3)), 2u)];
}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Idx_Expr_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a[c + 2 - 3];
-}
-
-var<private> c : i32;
-
+)",
+ /* predicate */ R"(
var<private> a : array<f32, 3>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a[min(u32(((c + 2) - 3)), 2u)];
-}
var<private> c : i32;
-var<private> a : array<f32, 3>;
-)";
+fn f() {
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, Vector_Idx_Scalar) {
+TEST_P(RobustnessTest, Read_NestedConstantSizedArraysRef_IndexWithRuntimeExpressions) {
auto* src = R"(
-var<private> a : vec3<f32>;
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
fn f() {
- var b : f32 = a[1i];
+ var r = a[x][y][z];
}
)";
- auto* expect = R"(
-var<private> a : vec3<f32>;
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
fn f() {
- var b : f32 = a[1i];
+ var r = a[min(u32(x), 2u)][min(u32(y), 1u)][min(u32(z), 0u)];
}
-)";
+)",
+ /* predicate */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
- auto got = Run<Robustness>(src);
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let index = x;
+ let predicate = (u32(index) <= 2u);
+ let index_1 = y;
+ let predicate_1 = (predicate & (u32(index_1) <= 1u));
+ let index_2 = z;
+ let predicate_2 = (predicate_1 & (u32(index_2) <= 0u));
+ var predicated_expr : f32;
+ if (predicate_2) {
+ predicated_expr = a[index][index_1][index_2];
+ }
+ var r = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, Vector_Idx_Scalar_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a[1i];
-}
-
-var<private> a : vec3<f32>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a[1i];
-}
-
-var<private> a : vec3<f32>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Vector_Idx_Expr) {
- auto* src = R"(
-var<private> a : vec3<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a[c + 2 - 3];
-}
-)";
-
- auto* expect = R"(
-var<private> a : vec3<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a[min(u32(((c + 2) - 3)), 2u)];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Vector_Idx_Expr_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a[c + 2 - 3];
-}
-
-var<private> c : i32;
-
-var<private> a : vec3<f32>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a[min(u32(((c + 2) - 3)), 2u)];
-}
-
-var<private> c : i32;
-
-var<private> a : vec3<f32>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Vector_Swizzle_Idx_Var) {
- auto* src = R"(
-var<private> a : vec3<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a.xy[c];
-}
-)";
-
- auto* expect = R"(
-var<private> a : vec3<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a.xy[min(u32(c), 1u)];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Vector_Swizzle_Idx_Var_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a.xy[c];
-}
-
-var<private> c : i32;
-
-var<private> a : vec3<f32>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a.xy[min(u32(c), 1u)];
-}
-
-var<private> c : i32;
-
-var<private> a : vec3<f32>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Vector_Swizzle_Idx_Expr) {
- auto* src = R"(
-var<private> a : vec3<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a.xy[c + 2 - 3];
-}
-)";
-
- auto* expect = R"(
-var<private> a : vec3<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a.xy[min(u32(((c + 2) - 3)), 1u)];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Vector_Swizzle_Idx_Expr_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a.xy[c + 2 - 3];
-}
-
-var<private> c : i32;
-
-var<private> a : vec3<f32>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a.xy[min(u32(((c + 2) - 3)), 1u)];
-}
-
-var<private> c : i32;
-
-var<private> a : vec3<f32>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Matrix_Idx_Scalar) {
- auto* src = R"(
-var<private> a : mat3x2<f32>;
-
-fn f() {
- var b : f32 = a[2i][1i];
-}
-)";
-
- auto* expect = R"(
-var<private> a : mat3x2<f32>;
-
-fn f() {
- var b : f32 = a[2i][1i];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Matrix_Idx_Scalar_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a[2i][1i];
-}
-
-var<private> a : mat3x2<f32>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a[2i][1i];
-}
-
-var<private> a : mat3x2<f32>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Matrix_Idx_Expr_Column) {
- auto* src = R"(
-var<private> a : mat3x2<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a[c + 2 - 3][1];
-}
-)";
-
- auto* expect = R"(
-var<private> a : mat3x2<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a[min(u32(((c + 2) - 3)), 2u)][1];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Matrix_Idx_Expr_Column_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a[c + 2 - 3][1];
-}
-
-var<private> c : i32;
-
-var<private> a : mat3x2<f32>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a[min(u32(((c + 2) - 3)), 2u)][1];
-}
-
-var<private> c : i32;
-
-var<private> a : mat3x2<f32>;
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Matrix_Idx_Expr_Row) {
- auto* src = R"(
-var<private> a : mat3x2<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a[1][c + 2 - 3];
-}
-)";
-
- auto* expect = R"(
-var<private> a : mat3x2<f32>;
-
-var<private> c : i32;
-
-fn f() {
- var b : f32 = a[1][min(u32(((c + 2) - 3)), 1u)];
-}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Matrix_Idx_Expr_Row_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var b : f32 = a[1][c + 2 - 3];
-}
-
-var<private> c : i32;
-
-var<private> a : mat3x2<f32>;
-)";
-
- auto* expect = R"(
-fn f() {
- var b : f32 = a[1][min(u32(((c + 2) - 3)), 1u)];
-}
-
-var<private> c : i32;
-
-var<private> a : mat3x2<f32>;
-)";
-
- auto got = Run<Robustness>(src);
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Vector_Constant_Id_Clamps) {
+TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithOverride) {
auto* src = R"(
@id(1300) override idx : i32;
-fn f() {
- var a : vec3<f32>;
- var b : f32 = a[idx];
-}
-)";
- auto* expect = R"(
-@id(1300) override idx : i32;
-
-fn f() {
- var a : vec3<f32>;
- var b : f32 = a[min(u32(idx), 2u)];
-}
-)";
-
- auto got = Run<Robustness>(src);
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, Array_Constant_Id_Clamps) {
- auto* src = R"(
-@id(1300) override idx : i32;
fn f() {
var a : array<f32, 4>;
var b : f32 = a[idx];
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
@id(1300) override idx : i32;
fn f() {
var a : array<f32, 4>;
var b : f32 = a[min(u32(idx), 3u)];
}
-)";
+)",
+ /* predicate */ R"(
+@id(1300) override idx : i32;
- auto got = Run<Robustness>(src);
+fn f() {
+ var a : array<f32, 4>;
+ let index = idx;
+ let predicate = (u32(index) <= 3u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, Matrix_Column_Constant_Id_Clamps) {
+TEST_P(RobustnessTest, Read_ConstantSizedArrayPtr_IndexWithLet) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let p = &(a[l]);
+ let f : f32 = *(p);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let p = &(a[min(l, 2u)]);
+ let f : f32 = *(p);
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let index = l;
+ let predicate = (u32(index) <= 2u);
+ let p = &(a[index]);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = *(p);
+ }
+ let f : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_ConstantSizedArrayPtr_IndexWithRuntimeArrayIndex) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let pa = &(a);
+ let pb = &(b);
+ let p0 = &((*(pb))[i]);
+ let p1 = &(a[*(p0)]);
+ var x : f32 = *(p1);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let pa = &(a);
+ let pb = &(b);
+ let p0 = &((*(pb))[min(i, 4u)]);
+ let p1 = &(a[min(u32(*(p0)), 2u)]);
+ var x : f32 = *(p1);
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let pa = &(a);
+ let pb = &(b);
+ let index = i;
+ let predicate = (u32(index) <= 4u);
+ let p0 = &((*(pb))[index]);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = *(p0);
+ }
+ let index_1 = predicated_expr;
+ let predicate_1 = (u32(index_1) <= 2u);
+ let p1 = &(a[index_1]);
+ var predicated_expr_1 : f32;
+ if (predicate_1) {
+ predicated_expr_1 = *(p1);
+ }
+ var x : f32 = predicated_expr_1;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_NestedConstantSizedArraysPtr_IndexWithRuntimeExpressions) {
+ auto* src = R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[x]);
+ let p2 = &((*(p1))[y]);
+ let p3 = &((*(p2))[z]);
+ var r = *(p3);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[min(u32(x), 2u)]);
+ let p2 = &((*(p1))[min(u32(y), 1u)]);
+ let p3 = &((*(p2))[min(u32(z), 0u)]);
+ var r = *(p3);
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let index = x;
+ let predicate = (u32(index) <= 2u);
+ let p1 = &((*(p0))[index]);
+ let index_1 = y;
+ let predicate_1 = (predicate & (u32(index_1) <= 1u));
+ let p2 = &((*(p1))[index_1]);
+ let index_2 = z;
+ let predicate_2 = (predicate_1 & (u32(index_2) <= 0u));
+ let p3 = &((*(p2))[index_2]);
+ var predicated_expr : f32;
+ if (predicate_2) {
+ predicated_expr = *(p3);
+ }
+ var r = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_NestedConstantSizedArrays_MixedAccess) {
+ auto* src = R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+const y = 1;
+
+override z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[x]);
+ let p2 = &((*(p1))[y]);
+ let p3 = &((*(p2))[z]);
+ var r = *(p3);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+const y = 1;
+
+override z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[min(u32(x), 2u)]);
+ let p2 = &((*(p1))[y]);
+ let p3 = &((*(p2))[min(u32(z), 0u)]);
+ var r = *(p3);
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+const y = 1;
+
+override z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let index = x;
+ let predicate = (u32(index) <= 2u);
+ let p1 = &((*(p0))[index]);
+ let p2 = &((*(p1))[y]);
+ let index_1 = z;
+ let predicate_1 = (predicate & (u32(index_1) <= 0u));
+ let p3 = &((*(p2))[index_1]);
+ var predicated_expr : f32;
+ if (predicate_1) {
+ predicated_expr = *(p3);
+ }
+ var r = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Assign_ConstantSizedArray_IndexWithLet) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ a[l] = 42.0f;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ a[min(l, 2u)] = 42.0f;
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let index = l;
+ let predicate = (u32(index) <= 2u);
+ if (predicate) {
+ a[index] = 42.0f;
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Assign_ConstantSizedArrayPtr_IndexWithLet) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let p = &(a[l]);
+ *(p) = 42.0f;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let p = &(a[min(l, 2u)]);
+ *(p) = 42.0f;
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let index = l;
+ let predicate = (u32(index) <= 2u);
+ let p = &(a[index]);
+ if (predicate) {
+ *(p) = 42.0f;
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Assign_ConstantSizedArrayPtr_IndexWithRuntimeArrayIndex) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let pa = &(a);
+ let pb = &(b);
+ let p0 = &((*(pb))[i]);
+ let p1 = &(a[*(p0)]);
+ *(p1) = 42.0f;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let pa = &(a);
+ let pb = &(b);
+ let p0 = &((*(pb))[min(i, 4u)]);
+ let p1 = &(a[min(u32(*(p0)), 2u)]);
+ *(p1) = 42.0f;
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let pa = &(a);
+ let pb = &(b);
+ let index = i;
+ let predicate = (u32(index) <= 4u);
+ let p0 = &((*(pb))[index]);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = *(p0);
+ }
+ let index_1 = predicated_expr;
+ let predicate_1 = (u32(index_1) <= 2u);
+ let p1 = &(a[index_1]);
+ if (predicate_1) {
+ *(p1) = 42.0f;
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Assign_NestedConstantSizedArraysPtr_IndexWithRuntimeExpressions) {
+ auto* src = R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[x]);
+ let p2 = &((*(p1))[y]);
+ let p3 = &((*(p2))[z]);
+ *(p3) = 42.0f;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[min(u32(x), 2u)]);
+ let p2 = &((*(p1))[min(u32(y), 1u)]);
+ let p3 = &((*(p2))[min(u32(z), 0u)]);
+ *(p3) = 42.0f;
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+var<private> y : i32;
+
+var<private> z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let index = x;
+ let predicate = (u32(index) <= 2u);
+ let p1 = &((*(p0))[index]);
+ let index_1 = y;
+ let predicate_1 = (predicate & (u32(index_1) <= 1u));
+ let p2 = &((*(p1))[index_1]);
+ let index_2 = z;
+ let predicate_2 = (predicate_1 & (u32(index_2) <= 0u));
+ let p3 = &((*(p2))[index_2]);
+ if (predicate_2) {
+ *(p3) = 42.0f;
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Assign_NestedConstantSizedArrays_MixedAccess) {
+ auto* src = R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+const y = 1;
+
+override z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[x]);
+ let p2 = &((*(p1))[y]);
+ let p3 = &((*(p2))[z]);
+ *(p3) = 42.0f;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+const y = 1;
+
+override z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let p1 = &((*(p0))[min(u32(x), 2u)]);
+ let p2 = &((*(p1))[y]);
+ let p3 = &((*(p2))[min(u32(z), 0u)]);
+ *(p3) = 42.0f;
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<array<array<f32, 1>, 2>, 3>;
+
+var<private> x : i32;
+
+const y = 1;
+
+override z : i32;
+
+fn f() {
+ let p0 = &(a);
+ let index = x;
+ let predicate = (u32(index) <= 2u);
+ let p1 = &((*(p0))[index]);
+ let p2 = &((*(p1))[y]);
+ let index_1 = z;
+ let predicate_1 = (predicate & (u32(index_1) <= 0u));
+ let p3 = &((*(p2))[index_1]);
+ if (predicate_1) {
+ *(p3) = 42.0f;
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, CompoundAssign_ConstantSizedArray_IndexWithLet) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ a[l] += 42.0f;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ a[min(l, 2u)] += 42.0f;
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let index = l;
+ let predicate = (u32(index) <= 2u);
+ if (predicate) {
+ a[index] += 42.0f;
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Increment_ConstantSizedArray_IndexWithLet) {
+ auto* src = R"(
+var<private> a : array<i32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ a[l]++;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<i32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ a[min(l, 2u)]++;
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<i32, 3>;
+
+fn f() {
+ let l : u32 = 1u;
+ let index = l;
+ let predicate = (u32(index) <= 2u);
+ if (predicate) {
+ a[index]++;
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Runtime sized array
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, Read_RuntimeArray_IndexWithLiteral) {
+ auto* src = R"(
+struct S {
+ a : f32,
+ b : array<f32>,
+}
+
+@group(0) @binding(0) var<storage, read> s : S;
+
+fn f() {
+ var d : f32 = s.b[25];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+struct S {
+ a : f32,
+ b : array<f32>,
+}
+
+@group(0) @binding(0) var<storage, read> s : S;
+
+fn f() {
+ var d : f32 = s.b[min(u32(25), (arrayLength(&(s.b)) - 1u))];
+}
+)",
+ /* predicate */ R"(
+struct S {
+ a : f32,
+ b : array<f32>,
+}
+
+@group(0) @binding(0) var<storage, read> s : S;
+
+fn f() {
+ let index = 25;
+ let predicate = (u32(index) <= (arrayLength(&(s.b)) - 1u));
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = s.b[index];
+ }
+ var d : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Vector
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, Read_Vector_IndexWithLiteral) {
+ auto* src = R"(
+var<private> a : vec3<f32>;
+
+fn f() {
+ var b : f32 = a[1i];
+}
+)";
+
+ auto* expect = src;
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_Vector_IndexWithConst) {
+ auto* src = R"(
+var<private> a : vec3<f32>;
+
+fn f() {
+ const i = 1;
+ var b : f32 = a[i];
+}
+)";
+
+ auto* expect = src;
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_Vector_IndexWithLet) {
+ auto* src = R"(
+fn f() {
+ let i = 99;
+ let v = vec4<f32>()[i];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+fn f() {
+ let i = 99;
+ let v = vec4<f32>()[min(u32(i), 3u)];
+}
+)",
+ /* predicate */ R"(
+fn f() {
+ let i = 99;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = vec4<f32>()[index];
+ }
+ let v = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_Vector_IndexWithRuntimeExpression) {
+ auto* src = R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a[((c + 2) - 3)];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a[min(u32(((c + 2) - 3)), 2u)];
+}
+)",
+ /* predicate */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithGlobalVar) {
+ auto* src = R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a.xy[c];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a.xy[min(u32(c), 1u)];
+}
+)",
+ /* predicate */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let index = c;
+ let predicate = (u32(index) <= 1u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a.xy[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithRuntimeExpression) {
+ auto* src = R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a.xy[((c + 2) - 3)];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a.xy[min(u32(((c + 2) - 3)), 1u)];
+}
+)",
+ /* predicate */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 1u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a.xy[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_Vector_IndexWithOverride) {
auto* src = R"(
@id(1300) override idx : i32;
+
+fn f() {
+ var a : vec3<f32>;
+ var b : f32 = a[idx];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@id(1300) override idx : i32;
+
+fn f() {
+ var a : vec3<f32>;
+ var b : f32 = a[min(u32(idx), 2u)];
+}
+)",
+ /* predicate */ R"(
+@id(1300) override idx : i32;
+
+fn f() {
+ var a : vec3<f32>;
+ let index = idx;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Matrix
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, Read_MatrixRef_IndexingWithLiterals) {
+ auto* src = R"(
+var<private> a : mat3x2<f32>;
+
+fn f() {
+ var b : f32 = a[2i][1i];
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_MatrixRef_IndexWithRuntimeExpressionThenLiteral) {
+ auto* src = R"(
+var<private> a : mat3x2<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a[((c + 2) - 3)][1];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : mat3x2<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a[min(u32(((c + 2) - 3)), 2u)][1];
+}
+)",
+ /* predicate */ R"(
+var<private> a : mat3x2<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index][1];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_MatrixRef_IndexWithLiteralThenRuntimeExpression) {
+ auto* src = R"(
+var<private> a : mat3x2<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a[1][((c + 2) - 3)];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : mat3x2<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ var b : f32 = a[1][min(u32(((c + 2) - 3)), 1u)];
+}
+)",
+ /* predicate */ R"(
+var<private> a : mat3x2<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 1u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[1][index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_MatrixRef_IndexWithOverrideThenLiteral) {
+ auto* src = R"(
+@id(1300) override idx : i32;
+
fn f() {
var a : mat3x2<f32>;
var b : f32 = a[idx][1];
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
@id(1300) override idx : i32;
fn f() {
var a : mat3x2<f32>;
var b : f32 = a[min(u32(idx), 2u)][1];
}
-)";
+)",
+ /* predicate */ R"(
+@id(1300) override idx : i32;
- auto got = Run<Robustness>(src);
+fn f() {
+ var a : mat3x2<f32>;
+ let index = idx;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index][1];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, Matrix_Row_Constant_Id_Clamps) {
+TEST_P(RobustnessTest, Read_MatrixRef_IndexWithLetThenSwizzle) {
+ auto* src = R"(
+fn f() {
+ let i = 1;
+ var m = mat3x2<f32>();
+ var v = m[i].yx;
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+fn f() {
+ let i = 1;
+ var m = mat3x2<f32>();
+ var v = m[min(u32(i), 2u)].yx;
+}
+)",
+ /* predicate */ R"(
+fn f() {
+ let i = 1;
+ var m = mat3x2<f32>();
+ let index = i;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : vec2<f32>;
+ if (predicate) {
+ predicated_expr = m[index];
+ }
+ var v = predicated_expr.yx;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_MatrixRef_IndexWithLiteralThenOverride) {
auto* src = R"(
@id(1300) override idx : i32;
+
fn f() {
var a : mat3x2<f32>;
var b : f32 = a[1][idx];
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
@id(1300) override idx : i32;
fn f() {
var a : mat3x2<f32>;
var b : f32 = a[1][min(u32(idx), 1u)];
}
-)";
-
- auto got = Run<Robustness>(src);
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, RuntimeArray_Clamps) {
- auto* src = R"(
-struct S {
- a : f32,
- b : array<f32>,
-};
-@group(0) @binding(0) var<storage, read> s : S;
+)",
+ /* predicate */ R"(
+@id(1300) override idx : i32;
fn f() {
- var d : f32 = s.b[25];
+ var a : mat3x2<f32>;
+ let index = idx;
+ let predicate = (u32(index) <= 1u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[1][index];
+ }
+ var b : f32 = predicated_expr;
}
-)";
+)");
- auto* expect = R"(
-struct S {
- a : f32,
- b : array<f32>,
+ auto got = Run<Robustness>(src, Config(GetParam()));
+ EXPECT_EQ(expect, str(got));
}
-@group(0) @binding(0) var<storage, read> s : S;
+TEST_P(RobustnessTest, Assign_Matrix_IndexWithLet) {
+ auto* src = R"(
+var<private> m : mat3x4f;
fn f() {
- var d : f32 = s.b[min(u32(25), (arrayLength(&(s.b)) - 1u))];
+ let c = 1;
+ m[c] = vec4f(1);
}
)";
- auto got = Run<Robustness>(src);
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> m : mat3x4f;
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, RuntimeArray_Clamps_OutOfOrder) {
- auto* src = R"(
fn f() {
- var d : f32 = s.b[25];
+ let c = 1;
+ m[min(u32(c), 2u)] = vec4f(1);
}
+)",
+ /* predicate */ R"(
+var<private> m : mat3x4f;
-@group(0) @binding(0) var<storage, read> s : S;
-
-struct S {
- a : f32,
- b : array<f32>,
-};
-)";
-
- auto* expect = R"(
fn f() {
- var d : f32 = s.b[min(u32(25), (arrayLength(&(s.b)) - 1u))];
+ let c = 1;
+ let index = c;
+ let predicate = (u32(index) <= 2u);
+ if (predicate) {
+ m[index] = vec4f(1);
+ }
}
+)");
-@group(0) @binding(0) var<storage, read> s : S;
-
-struct S {
- a : f32,
- b : array<f32>,
-}
-)";
-
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-// Clamp textureLoad() coord, array_index and level values
-TEST_F(RobustnessTest, TextureLoad_Clamp) {
+TEST_P(RobustnessTest, CompoundAssign_Matrix_IndexWithLet) {
auto* src = R"(
-@group(0) @binding(0) var tex_1d : texture_1d<f32>;
-@group(0) @binding(0) var tex_2d : texture_2d<f32>;
-@group(0) @binding(0) var tex_2d_arr : texture_2d_array<f32>;
-@group(0) @binding(0) var tex_3d : texture_3d<f32>;
-@group(0) @binding(0) var tex_ms_2d : texture_multisampled_2d<f32>;
-@group(0) @binding(0) var tex_depth_2d : texture_depth_2d;
-@group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
-@group(0) @binding(0) var tex_external : texture_external;
+var<private> m : mat3x4f;
-fn idx_signed() {
- var array_idx : i32;
- var level_idx : i32;
- var sample_idx : i32;
-
- _ = textureLoad(tex_1d, 1i, level_idx);
- _ = textureLoad(tex_2d, vec2<i32>(1, 2), level_idx);
- _ = textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx);
- _ = textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx);
- _ = textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx);
- _ = textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_external, vec2<i32>(1, 2));
-}
-
-fn idx_unsigned() {
- var array_idx : u32;
- var level_idx : u32;
- var sample_idx : u32;
-
- _ = textureLoad(tex_1d, 1u, level_idx);
- _ = textureLoad(tex_2d, vec2<u32>(1, 2), level_idx);
- _ = textureLoad(tex_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_3d, vec3<u32>(1, 2, 3), level_idx);
- _ = textureLoad(tex_ms_2d, vec2<u32>(1, 2), sample_idx);
- _ = textureLoad(tex_depth_2d, vec2<u32>(1, 2), level_idx);
- _ = textureLoad(tex_depth_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_external, vec2<u32>(1, 2));
+fn f() {
+ let c = 1;
+ m[c] += vec4f(1);
}
)";
- auto* expect =
- R"(
-@group(0) @binding(0) var tex_1d : texture_1d<f32>;
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> m : mat3x4f;
-@group(0) @binding(0) var tex_2d : texture_2d<f32>;
-
-@group(0) @binding(0) var tex_2d_arr : texture_2d_array<f32>;
-
-@group(0) @binding(0) var tex_3d : texture_3d<f32>;
-
-@group(0) @binding(0) var tex_ms_2d : texture_multisampled_2d<f32>;
-
-@group(0) @binding(0) var tex_depth_2d : texture_depth_2d;
-
-@group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
-
-@group(0) @binding(0) var tex_external : texture_external;
-
-fn idx_signed() {
- var array_idx : i32;
- var level_idx : i32;
- var sample_idx : i32;
- _ = textureLoad(tex_1d, clamp(1i, 0, i32((u32(textureDimensions(tex_1d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))))) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))));
- _ = textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))));
- _ = textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))));
- _ = textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex_3d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))))) - vec3(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))));
- _ = textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1)))), sample_idx);
- _ = textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))));
- _ = textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_depth_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))));
- _ = textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
+fn f() {
+ let c = 1;
+ m[min(u32(c), 2u)] += vec4f(1);
}
+)",
+ /* predicate */ R"(
+var<private> m : mat3x4f;
-fn idx_unsigned() {
- var array_idx : u32;
- var level_idx : u32;
- var sample_idx : u32;
- _ = textureLoad(tex_1d, min(1u, (u32(textureDimensions(tex_1d, min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)))) - 1)), min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)));
- _ = textureLoad(tex_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d, min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)));
- _ = textureLoad(tex_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d_arr, min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)));
- _ = textureLoad(tex_3d, min(vec3<u32>(1, 2, 3), (vec3<u32>(textureDimensions(tex_3d, min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)))) - vec3(1))), min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)));
- _ = textureLoad(tex_ms_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1))), sample_idx);
- _ = textureLoad(tex_depth_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d, min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)));
- _ = textureLoad(tex_depth_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d_arr, min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_depth_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)));
- _ = textureLoad(tex_external, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_external)) - vec2(1))));
+fn f() {
+ let c = 1;
+ let index = c;
+ let predicate = (u32(index) <= 2u);
+ if (predicate) {
+ m[index] += vec4f(1);
+ }
}
-)";
+)");
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-// Clamp textureLoad() coord, array_index and level values
-TEST_F(RobustnessTest, TextureLoad_Clamp_OutOfOrder) {
+////////////////////////////////////////////////////////////////////////////////
+// Texture
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, TextureDimensions) {
auto* src = R"(
-fn idx_signed() {
- var array_idx : i32;
- var level_idx : i32;
- var sample_idx : i32;
+@group(0) @binding(0) var t : texture_2d<f32>;
- _ = textureLoad(tex_1d, 1i, level_idx);
- _ = textureLoad(tex_2d, vec2<i32>(1, 2), level_idx);
- _ = textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx);
- _ = textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx);
- _ = textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx);
- _ = textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_external, vec2<i32>(1, 2));
+fn dimensions() {
+ let l = textureDimensions(t);
}
-
-fn idx_unsigned() {
- var array_idx : u32;
- var level_idx : u32;
- var sample_idx : u32;
-
- _ = textureLoad(tex_1d, 1u, level_idx);
- _ = textureLoad(tex_2d, vec2<u32>(1, 2), level_idx);
- _ = textureLoad(tex_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_3d, vec3<u32>(1, 2, 3), level_idx);
- _ = textureLoad(tex_ms_2d, vec2<u32>(1, 2), sample_idx);
- _ = textureLoad(tex_depth_2d, vec2<u32>(1, 2), level_idx);
- _ = textureLoad(tex_depth_2d_arr, vec2<u32>(1, 2), array_idx, level_idx);
- _ = textureLoad(tex_external, vec2<u32>(1, 2));
-}
-
-@group(0) @binding(0) var tex_1d : texture_1d<f32>;
-@group(0) @binding(0) var tex_2d : texture_2d<f32>;
-@group(0) @binding(0) var tex_2d_arr : texture_2d_array<f32>;
-@group(0) @binding(0) var tex_3d : texture_3d<f32>;
-@group(0) @binding(0) var tex_ms_2d : texture_multisampled_2d<f32>;
-@group(0) @binding(0) var tex_depth_2d : texture_depth_2d;
-@group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
-@group(0) @binding(0) var tex_external : texture_external;
)";
- auto* expect =
- R"(
-fn idx_signed() {
- var array_idx : i32;
- var level_idx : i32;
- var sample_idx : i32;
- _ = textureLoad(tex_1d, clamp(1i, 0, i32((u32(textureDimensions(tex_1d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))))) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_1d)) - 1))));
- _ = textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d)) - 1))));
- _ = textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_2d_arr)) - 1))));
- _ = textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex_3d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))))) - vec3(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_3d)) - 1))));
- _ = textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1)))), sample_idx);
- _ = textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))))) - vec2(1)))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d)) - 1))));
- _ = textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))))) - vec2(1)))), clamp(array_idx, 0, i32((u32(textureNumLayers(tex_depth_2d_arr)) - 1))), clamp(level_idx, 0, i32((u32(textureNumLevels(tex_depth_2d_arr)) - 1))));
- _ = textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex_external)) - vec2(1)))));
-}
+ auto* expect = src;
-fn idx_unsigned() {
- var array_idx : u32;
- var level_idx : u32;
- var sample_idx : u32;
- _ = textureLoad(tex_1d, min(1u, (u32(textureDimensions(tex_1d, min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)))) - 1)), min(level_idx, (u32(textureNumLevels(tex_1d)) - 1)));
- _ = textureLoad(tex_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d, min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_2d)) - 1)));
- _ = textureLoad(tex_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_2d_arr, min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_2d_arr)) - 1)));
- _ = textureLoad(tex_3d, min(vec3<u32>(1, 2, 3), (vec3<u32>(textureDimensions(tex_3d, min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)))) - vec3(1))), min(level_idx, (u32(textureNumLevels(tex_3d)) - 1)));
- _ = textureLoad(tex_ms_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_ms_2d)) - vec2(1))), sample_idx);
- _ = textureLoad(tex_depth_2d, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d, min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)))) - vec2(1))), min(level_idx, (u32(textureNumLevels(tex_depth_2d)) - 1)));
- _ = textureLoad(tex_depth_2d_arr, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_depth_2d_arr, min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)))) - vec2(1))), min(array_idx, (u32(textureNumLayers(tex_depth_2d_arr)) - 1)), min(level_idx, (u32(textureNumLevels(tex_depth_2d_arr)) - 1)));
- _ = textureLoad(tex_external, min(vec2<u32>(1, 2), (vec2<u32>(textureDimensions(tex_external)) - vec2(1))));
-}
-
-@group(0) @binding(0) var tex_1d : texture_1d<f32>;
-
-@group(0) @binding(0) var tex_2d : texture_2d<f32>;
-
-@group(0) @binding(0) var tex_2d_arr : texture_2d_array<f32>;
-
-@group(0) @binding(0) var tex_3d : texture_3d<f32>;
-
-@group(0) @binding(0) var tex_ms_2d : texture_multisampled_2d<f32>;
-
-@group(0) @binding(0) var tex_depth_2d : texture_depth_2d;
-
-@group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
-
-@group(0) @binding(0) var tex_external : texture_external;
-)";
-
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-// Clamp textureStore() coord, array_index and level values
-TEST_F(RobustnessTest, TextureStore_Clamp) {
+TEST_P(RobustnessTest, TextureDimensions_Level) {
auto* src = R"(
-@group(0) @binding(0) var tex1d : texture_storage_1d<rgba8sint, write>;
+@group(0) @binding(0) var t : texture_2d<f32>;
-@group(0) @binding(1) var tex2d : texture_storage_2d<rgba8sint, write>;
-
-@group(0) @binding(2) var tex2d_arr : texture_storage_2d_array<rgba8sint, write>;
-
-@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
-
-fn idx_signed() {
- textureStore(tex1d, 10i, vec4<i32>());
- textureStore(tex2d, vec2<i32>(10, 20), vec4<i32>());
- textureStore(tex2d_arr, vec2<i32>(10, 20), 50i, vec4<i32>());
- textureStore(tex3d, vec3<i32>(10, 20, 30), vec4<i32>());
+fn dimensions_signed(level : i32) {
+ let l = textureDimensions(t, level);
}
-fn idx_unsigned() {
- textureStore(tex1d, 10u, vec4<i32>());
- textureStore(tex2d, vec2<u32>(10, 20), vec4<i32>());
- textureStore(tex2d_arr, vec2<u32>(10, 20), 50u, vec4<i32>());
- textureStore(tex3d, vec3<u32>(10, 20, 30), vec4<i32>());
+fn dimensions_unsigned(level : u32) {
+ let l = textureDimensions(t, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ dimensions_signed(i32(non_uniform.x));
+ dimensions_unsigned(u32(non_uniform.x));
}
)";
- auto* expect = R"(
-@group(0) @binding(0) var tex1d : texture_storage_1d<rgba8sint, write>;
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
-@group(0) @binding(1) var tex2d : texture_storage_2d<rgba8sint, write>;
-
-@group(0) @binding(2) var tex2d_arr : texture_storage_2d_array<rgba8sint, write>;
-
-@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
-
-fn idx_signed() {
- textureStore(tex1d, clamp(10i, 0, i32((u32(textureDimensions(tex1d)) - 1))), vec4<i32>());
- textureStore(tex2d, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d)) - vec2(1)))), vec4<i32>());
- textureStore(tex2d_arr, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1)))), clamp(50i, 0, i32((u32(textureNumLayers(tex2d_arr)) - 1))), vec4<i32>());
- textureStore(tex3d, clamp(vec3<i32>(10, 20, 30), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex3d)) - vec3(1)))), vec4<i32>());
+fn dimensions_signed(level : i32) {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureDimensions(t, level_idx);
}
-fn idx_unsigned() {
- textureStore(tex1d, min(10u, (u32(textureDimensions(tex1d)) - 1)), vec4<i32>());
- textureStore(tex2d, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d)) - vec2(1))), vec4<i32>());
- textureStore(tex2d_arr, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1))), min(50u, (u32(textureNumLayers(tex2d_arr)) - 1)), vec4<i32>());
- textureStore(tex3d, min(vec3<u32>(10, 20, 30), (vec3<u32>(textureDimensions(tex3d)) - vec3(1))), vec4<i32>());
+fn dimensions_unsigned(level : u32) {
+ let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureDimensions(t, level_idx_1);
}
-)";
- auto got = Run<Robustness>(src);
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ dimensions_signed(i32(non_uniform.x));
+ dimensions_unsigned(u32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+fn dimensions_signed(level : i32) {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ var predicated_value : vec2<u32>;
+ if ((level_idx < num_levels)) {
+ predicated_value = textureDimensions(t, level_idx);
+ }
+ let l = predicated_value;
+}
+
+fn dimensions_unsigned(level : u32) {
+ let level_idx_1 = u32(level);
+ let num_levels_1 = textureNumLevels(t);
+ var predicated_value_1 : vec2<u32>;
+ if ((level_idx_1 < num_levels_1)) {
+ predicated_value_1 = textureDimensions(t, level_idx_1);
+ }
+ let l = predicated_value_1;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ dimensions_signed(i32(non_uniform.x));
+ dimensions_unsigned(u32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-// Clamp textureStore() coord, array_index and level values
-TEST_F(RobustnessTest, TextureStore_Clamp_OutOfOrder) {
+TEST_P(RobustnessTest, TextureGather) {
auto* src = R"(
-fn idx_signed() {
- textureStore(tex1d, 10i, vec4<i32>());
- textureStore(tex2d, vec2<i32>(10, 20), vec4<i32>());
- textureStore(tex2d_arr, vec2<i32>(10, 20), 50i, vec4<i32>());
- textureStore(tex3d, vec3<i32>(10, 20, 30), vec4<i32>());
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn gather(coords : vec2f) {
+ let l = textureGather(0, t, s, coords);
}
-
-fn idx_unsigned() {
- textureStore(tex1d, 10u, vec4<i32>());
- textureStore(tex2d, vec2<u32>(10, 20), vec4<i32>());
- textureStore(tex2d_arr, vec2<u32>(10, 20), 50u, vec4<i32>());
- textureStore(tex3d, vec3<u32>(10, 20, 30), vec4<i32>());
-}
-
-@group(0) @binding(0) var tex1d : texture_storage_1d<rgba8sint, write>;
-
-@group(0) @binding(1) var tex2d : texture_storage_2d<rgba8sint, write>;
-
-@group(0) @binding(2) var tex2d_arr : texture_storage_2d_array<rgba8sint, write>;
-
-@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
-
)";
- auto* expect = R"(
-fn idx_signed() {
- textureStore(tex1d, clamp(10i, 0, i32((u32(textureDimensions(tex1d)) - 1))), vec4<i32>());
- textureStore(tex2d, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d)) - vec2(1)))), vec4<i32>());
- textureStore(tex2d_arr, clamp(vec2<i32>(10, 20), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1)))), clamp(50i, 0, i32((u32(textureNumLayers(tex2d_arr)) - 1))), vec4<i32>());
- textureStore(tex3d, clamp(vec3<i32>(10, 20, 30), vec3(0), vec3<i32>((vec3<u32>(textureDimensions(tex3d)) - vec3(1)))), vec4<i32>());
-}
+ auto* expect = src;
-fn idx_unsigned() {
- textureStore(tex1d, min(10u, (u32(textureDimensions(tex1d)) - 1)), vec4<i32>());
- textureStore(tex2d, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d)) - vec2(1))), vec4<i32>());
- textureStore(tex2d_arr, min(vec2<u32>(10, 20), (vec2<u32>(textureDimensions(tex2d_arr)) - vec2(1))), min(50u, (u32(textureNumLayers(tex2d_arr)) - 1)), vec4<i32>());
- textureStore(tex3d, min(vec3<u32>(10, 20, 30), (vec3<u32>(textureDimensions(tex3d)) - vec3(1))), vec4<i32>());
-}
-
-@group(0) @binding(0) var tex1d : texture_storage_1d<rgba8sint, write>;
-
-@group(0) @binding(1) var tex2d : texture_storage_2d<rgba8sint, write>;
-
-@group(0) @binding(2) var tex2d_arr : texture_storage_2d_array<rgba8sint, write>;
-
-@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
-)";
-
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, Shadowed_Variable) {
+TEST_P(RobustnessTest, TextureGather_Array) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn gather_signed(coords : vec2f, array : i32) {
+ let l = textureGather(1, t, s, coords, array);
+}
+
+fn gather_unsigned(coords : vec2f, array : u32) {
+ let l = textureGather(1, t, s, coords, array);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ gather_signed(non_uniform.xy, i32(non_uniform.x));
+ gather_unsigned(non_uniform.xy, u32(non_uniform.x));
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureGatherCompare) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn gather(coords : vec2f, depth_ref : f32) {
+ let l = textureGatherCompare(t, s, coords, depth_ref);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureGatherCompare_Array) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d_array;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn gather_signed(coords : vec2f, array : i32, depth_ref : f32) {
+ let l = textureGatherCompare(t, s, coords, array, depth_ref);
+}
+
+fn gather_unsigned(coords : vec2f, array : u32, depth_ref : f32) {
+ let l = textureGatherCompare(t, s, coords, array, depth_ref);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ gather_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x);
+ gather_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_1D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_1d<f32>;
+
+fn load_signed(coords : i32, level : i32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_unsigned(coords : u32, level : u32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_mixed(coords : i32, level : u32) {
+ let l = textureLoad(t, coords, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(i32(non_uniform.x), u32(non_uniform.x));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_1d<f32>;
+
+fn load_signed(coords : i32, level : i32) {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, 0, i32((textureDimensions(t, level_idx) - 1))), level_idx);
+}
+
+fn load_unsigned(coords : u32, level : u32) {
+ let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - 1)), level_idx_1);
+}
+
+fn load_mixed(coords : i32, level : u32) {
+ let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, 0, i32((textureDimensions(t, level_idx_2) - 1))), level_idx_2);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(i32(non_uniform.x), u32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_1d<f32>;
+
+fn load_signed(coords : i32, level : i32) {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = u32(coords);
+ var predicated_value : vec4<f32>;
+ if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) {
+ predicated_value = textureLoad(t, coords_1, level_idx);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : u32, level : u32) {
+ let level_idx_1 = u32(level);
+ let num_levels_1 = textureNumLevels(t);
+ let coords_2 = u32(coords);
+ var predicated_value_1 : vec4<f32>;
+ if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) {
+ predicated_value_1 = textureLoad(t, coords_2, level_idx_1);
+ }
+ let l = predicated_value_1;
+}
+
+fn load_mixed(coords : i32, level : u32) {
+ let level_idx_2 = u32(level);
+ let num_levels_2 = textureNumLevels(t);
+ let coords_3 = u32(coords);
+ var predicated_value_2 : vec4<f32>;
+ if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) {
+ predicated_value_2 = textureLoad(t, coords_3, level_idx_2);
+ }
+ let l = predicated_value_2;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(i32(non_uniform.x), u32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_2D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+fn load_signed(coords : vec2i, level : i32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_unsigned(coords : vec2u, level : u32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_mixed(coords : vec2u, level : i32) {
+ let l = textureLoad(t, coords, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+fn load_signed(coords : vec2i, level : i32) {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t, level_idx) - vec2(1)))), level_idx);
+}
+
+fn load_unsigned(coords : vec2u, level : u32) {
+ let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), level_idx_1);
+}
+
+fn load_mixed(coords : vec2u, level : i32) {
+ let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec2(1))), level_idx_2);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+fn load_signed(coords : vec2i, level : i32) {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = vec2<u32>(coords);
+ var predicated_value : vec4<f32>;
+ if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) {
+ predicated_value = textureLoad(t, coords_1, level_idx);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : vec2u, level : u32) {
+ let level_idx_1 = u32(level);
+ let num_levels_1 = textureNumLevels(t);
+ let coords_2 = vec2<u32>(coords);
+ var predicated_value_1 : vec4<f32>;
+ if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) {
+ predicated_value_1 = textureLoad(t, coords_2, level_idx_1);
+ }
+ let l = predicated_value_1;
+}
+
+fn load_mixed(coords : vec2u, level : i32) {
+ let level_idx_2 = u32(level);
+ let num_levels_2 = textureNumLevels(t);
+ let coords_3 = vec2<u32>(coords);
+ var predicated_value_2 : vec4<f32>;
+ if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) {
+ predicated_value_2 = textureLoad(t, coords_3, level_idx_2);
+ }
+ let l = predicated_value_2;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_2DArray) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+fn load_signed(coords : vec2i, array : i32, level : i32) {
+ let l = textureLoad(t, coords, array, level);
+}
+
+fn load_unsigned(coords : vec2u, array : u32, level : u32) {
+ let l = textureLoad(t, coords, array, level);
+}
+
+fn load_mixed(coords : vec2u, array : i32, level : u32) {
+ let l = textureLoad(t, coords, array, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+fn load_signed(coords : vec2i, array : i32, level : i32) {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t, level_idx) - vec2(1)))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx);
+}
+
+fn load_unsigned(coords : vec2u, array : u32, level : u32) {
+ let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), min(array, (textureNumLayers(t) - 1)), level_idx_1);
+}
+
+fn load_mixed(coords : vec2u, array : i32, level : u32) {
+ let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec2(1))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx_2);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+fn load_signed(coords : vec2i, array : i32, level : i32) {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = vec2<u32>(coords);
+ let array_idx = u32(array);
+ var predicated_value : vec4<f32>;
+ if ((((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1)))))) & (array_idx < textureNumLayers(t)))) {
+ predicated_value = textureLoad(t, coords_1, array_idx, level_idx);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : vec2u, array : u32, level : u32) {
+ let level_idx_1 = u32(level);
+ let num_levels_1 = textureNumLevels(t);
+ let coords_2 = vec2<u32>(coords);
+ let array_idx_1 = u32(array);
+ var predicated_value_1 : vec4<f32>;
+ if ((((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1)))))) & (array_idx_1 < textureNumLayers(t)))) {
+ predicated_value_1 = textureLoad(t, coords_2, array_idx_1, level_idx_1);
+ }
+ let l = predicated_value_1;
+}
+
+fn load_mixed(coords : vec2u, array : i32, level : u32) {
+ let level_idx_2 = u32(level);
+ let num_levels_2 = textureNumLevels(t);
+ let coords_3 = vec2<u32>(coords);
+ let array_idx_2 = u32(array);
+ var predicated_value_2 : vec4<f32>;
+ if ((((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1)))))) & (array_idx_2 < textureNumLayers(t)))) {
+ predicated_value_2 = textureLoad(t, coords_3, array_idx_2, level_idx_2);
+ }
+ let l = predicated_value_2;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_3D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_3d<f32>;
+
+fn load_signed(coords : vec3i, level : i32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_unsigned(coords : vec3u, level : u32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_mixed(coords : vec3u, level : i32) {
+ let l = textureLoad(t, coords, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec3i(non_uniform.xyz), i32(non_uniform.x));
+ load_unsigned(vec3u(non_uniform.xyz), u32(non_uniform.x));
+ load_mixed(vec3u(non_uniform.xyz), i32(non_uniform.x));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_3d<f32>;
+
+fn load_signed(coords : vec3i, level : i32) {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, vec3(0), vec3<i32>((textureDimensions(t, level_idx) - vec3(1)))), level_idx);
+}
+
+fn load_unsigned(coords : vec3u, level : u32) {
+ let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec3(1))), level_idx_1);
+}
+
+fn load_mixed(coords : vec3u, level : i32) {
+ let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec3(1))), level_idx_2);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec3i(non_uniform.xyz), i32(non_uniform.x));
+ load_unsigned(vec3u(non_uniform.xyz), u32(non_uniform.x));
+ load_mixed(vec3u(non_uniform.xyz), i32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_3d<f32>;
+
+fn load_signed(coords : vec3i, level : i32) {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = vec3<u32>(coords);
+ var predicated_value : vec4<f32>;
+ if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) {
+ predicated_value = textureLoad(t, coords_1, level_idx);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : vec3u, level : u32) {
+ let level_idx_1 = u32(level);
+ let num_levels_1 = textureNumLevels(t);
+ let coords_2 = vec3<u32>(coords);
+ var predicated_value_1 : vec4<f32>;
+ if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) {
+ predicated_value_1 = textureLoad(t, coords_2, level_idx_1);
+ }
+ let l = predicated_value_1;
+}
+
+fn load_mixed(coords : vec3u, level : i32) {
+ let level_idx_2 = u32(level);
+ let num_levels_2 = textureNumLevels(t);
+ let coords_3 = vec3<u32>(coords);
+ var predicated_value_2 : vec4<f32>;
+ if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) {
+ predicated_value_2 = textureLoad(t, coords_3, level_idx_2);
+ }
+ let l = predicated_value_2;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec3i(non_uniform.xyz), i32(non_uniform.x));
+ load_unsigned(vec3u(non_uniform.xyz), u32(non_uniform.x));
+ load_mixed(vec3u(non_uniform.xyz), i32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_Multisampled2D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_multisampled_2d<f32>;
+
+fn load_signed(coords : vec2i, sample : i32) {
+ let l = textureLoad(t, coords, sample);
+}
+
+fn load_unsigned(coords : vec2u, sample : u32) {
+ let l = textureLoad(t, coords, sample);
+}
+
+fn load_mixed(coords : vec2i, sample : u32) {
+ let l = textureLoad(t, coords, sample);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_multisampled_2d<f32>;
+
+fn load_signed(coords : vec2i, sample : i32) {
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t) - vec2(1)))), sample);
+}
+
+fn load_unsigned(coords : vec2u, sample : u32) {
+ let l = textureLoad(t, min(coords, (textureDimensions(t) - vec2(1))), sample);
+}
+
+fn load_mixed(coords : vec2i, sample : u32) {
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t) - vec2(1)))), sample);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_multisampled_2d<f32>;
+
+fn load_signed(coords : vec2i, sample : i32) {
+ let coords_1 = vec2<u32>(coords);
+ var predicated_value : vec4<f32>;
+ if (all((coords_1 < textureDimensions(t)))) {
+ predicated_value = textureLoad(t, coords_1, sample);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : vec2u, sample : u32) {
+ let coords_2 = vec2<u32>(coords);
+ var predicated_value_1 : vec4<f32>;
+ if (all((coords_2 < textureDimensions(t)))) {
+ predicated_value_1 = textureLoad(t, coords_2, sample);
+ }
+ let l = predicated_value_1;
+}
+
+fn load_mixed(coords : vec2i, sample : u32) {
+ let coords_3 = vec2<u32>(coords);
+ var predicated_value_2 : vec4<f32>;
+ if (all((coords_3 < textureDimensions(t)))) {
+ predicated_value_2 = textureLoad(t, coords_3, sample);
+ }
+ let l = predicated_value_2;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_Depth2D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+fn load_signed(coords : vec2i, level : i32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_unsigned(coords : vec2u, level : u32) {
+ let l = textureLoad(t, coords, level);
+}
+
+fn load_mixed(coords : vec2i, level : u32) {
+ let l = textureLoad(t, coords, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+fn load_signed(coords : vec2i, level : i32) {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t, level_idx) - vec2(1)))), level_idx);
+}
+
+fn load_unsigned(coords : vec2u, level : u32) {
+ let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), level_idx_1);
+}
+
+fn load_mixed(coords : vec2i, level : u32) {
+ let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t, level_idx_2) - vec2(1)))), level_idx_2);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+fn load_signed(coords : vec2i, level : i32) {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = vec2<u32>(coords);
+ var predicated_value : f32;
+ if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) {
+ predicated_value = textureLoad(t, coords_1, level_idx);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : vec2u, level : u32) {
+ let level_idx_1 = u32(level);
+ let num_levels_1 = textureNumLevels(t);
+ let coords_2 = vec2<u32>(coords);
+ var predicated_value_1 : f32;
+ if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) {
+ predicated_value_1 = textureLoad(t, coords_2, level_idx_1);
+ }
+ let l = predicated_value_1;
+}
+
+fn load_mixed(coords : vec2i, level : u32) {
+ let level_idx_2 = u32(level);
+ let num_levels_2 = textureNumLevels(t);
+ let coords_3 = vec2<u32>(coords);
+ var predicated_value_2 : f32;
+ if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) {
+ predicated_value_2 = textureLoad(t, coords_3, level_idx_2);
+ }
+ let l = predicated_value_2;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x));
+ load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_Depth2DArray) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d_array;
+
+fn load_signed(coords : vec2i, array : i32, level : i32) {
+ let l = textureLoad(t, coords, array, level);
+}
+
+fn load_unsigned(coords : vec2u, array : u32, level : u32) {
+ let l = textureLoad(t, coords, array, level);
+}
+
+fn load_mixed(coords : vec2u, array : i32, level : u32) {
+ let l = textureLoad(t, coords, array, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_depth_2d_array;
+
+fn load_signed(coords : vec2i, array : i32, level : i32) {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t, level_idx) - vec2(1)))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx);
+}
+
+fn load_unsigned(coords : vec2u, array : u32, level : u32) {
+ let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), min(array, (textureNumLayers(t) - 1)), level_idx_1);
+}
+
+fn load_mixed(coords : vec2u, array : i32, level : u32) {
+ let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1));
+ let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec2(1))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx_2);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_depth_2d_array;
+
+fn load_signed(coords : vec2i, array : i32, level : i32) {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = vec2<u32>(coords);
+ let array_idx = u32(array);
+ var predicated_value : f32;
+ if ((((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1)))))) & (array_idx < textureNumLayers(t)))) {
+ predicated_value = textureLoad(t, coords_1, array_idx, level_idx);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : vec2u, array : u32, level : u32) {
+ let level_idx_1 = u32(level);
+ let num_levels_1 = textureNumLevels(t);
+ let coords_2 = vec2<u32>(coords);
+ let array_idx_1 = u32(array);
+ var predicated_value_1 : f32;
+ if ((((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1)))))) & (array_idx_1 < textureNumLayers(t)))) {
+ predicated_value_1 = textureLoad(t, coords_2, array_idx_1, level_idx_1);
+ }
+ let l = predicated_value_1;
+}
+
+fn load_mixed(coords : vec2u, array : i32, level : u32) {
+ let level_idx_2 = u32(level);
+ let num_levels_2 = textureNumLevels(t);
+ let coords_3 = vec2<u32>(coords);
+ let array_idx_2 = u32(array);
+ var predicated_value_2 : f32;
+ if ((((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1)))))) & (array_idx_2 < textureNumLayers(t)))) {
+ predicated_value_2 = textureLoad(t, coords_3, array_idx_2, level_idx_2);
+ }
+ let l = predicated_value_2;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x));
+ load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x));
+ load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_External) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_external;
+
+fn load_signed(coords : vec2i) {
+ let l = textureLoad(t, coords);
+}
+
+fn load_unsigned(coords : vec2u) {
+ let l = textureLoad(t, coords);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy));
+ load_unsigned(vec2u(non_uniform.xy));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_external;
+
+fn load_signed(coords : vec2i) {
+ let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t) - vec2(1)))));
+}
+
+fn load_unsigned(coords : vec2u) {
+ let l = textureLoad(t, min(coords, (textureDimensions(t) - vec2(1))));
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy));
+ load_unsigned(vec2u(non_uniform.xy));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_external;
+
+fn load_signed(coords : vec2i) {
+ let coords_1 = vec2<u32>(coords);
+ var predicated_value : vec4<f32>;
+ if (all((coords_1 < textureDimensions(t)))) {
+ predicated_value = textureLoad(t, coords_1);
+ }
+ let l = predicated_value;
+}
+
+fn load_unsigned(coords : vec2u) {
+ let coords_2 = vec2<u32>(coords);
+ var predicated_value_1 : vec4<f32>;
+ if (all((coords_2 < textureDimensions(t)))) {
+ predicated_value_1 = textureLoad(t, coords_2);
+ }
+ let l = predicated_value_1;
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ load_signed(vec2i(non_uniform.xy));
+ load_unsigned(vec2u(non_uniform.xy));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureNumLayers) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d_array;
+
+fn num_layers(coords : vec2f, depth_ref : f32) {
+ let l = textureNumLayers(t);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureNumLevels) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+fn num_levels(coords : vec2f, depth_ref : f32) {
+ let l = textureNumLevels(t);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureNumSamples) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_multisampled_2d;
+
+fn num_levels(coords : vec2f, depth_ref : f32) {
+ let l = textureNumSamples(t);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSample) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample(coords : vec2f) {
+ let l = textureSample(t, s, coords);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSample_Offset) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample(coords : vec2f) {
+ const offset = vec2i(1);
+ let l = textureSample(t, s, coords, offset);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSample_ArrayIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_signed(coords : vec2f, array : i32) {
+ let l = textureSample(t, s, coords, array);
+}
+
+fn sample_unsigned(coords : vec2f, array : u32) {
+ let l = textureSample(t, s, coords, array);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ sample_signed(non_uniform.xy, i32(non_uniform.x));
+ sample_unsigned(non_uniform.xy, u32(non_uniform.x));
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleBias) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_bias(coords : vec2f, bias : f32) {
+ let l = textureSampleBias(t, s, coords, bias);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleBias_Offset) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_bias(coords : vec2f, bias : f32) {
+ const offset = vec2i(1);
+ let l = textureSampleBias(t, s, coords, bias, offset);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleBias_ArrayIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_bias_signed(coords : vec2f, array : i32, bias : f32) {
+ let l = textureSampleBias(t, s, coords, array, bias);
+}
+
+fn sample_bias_unsigned(coords : vec2f, array : u32, bias : f32) {
+ let l = textureSampleBias(t, s, coords, array, bias);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ sample_bias_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x);
+ sample_bias_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleCompare) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn sample_compare(coords : vec2f, depth_ref : f32) {
+ let l = textureSampleCompare(t, s, coords, depth_ref);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleCompare_Offset) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn sample_compare(coords : vec2f, depth_ref : f32) {
+ const offset = vec2i(1);
+ let l = textureSampleCompare(t, s, coords, depth_ref, offset);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleCompare_ArrayIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d_array;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn sample_compare_signed(coords : vec2f, array : i32, depth_ref : f32) {
+ let l = textureSampleCompare(t, s, coords, array, depth_ref);
+}
+
+fn sample_compare_unsigned(coords : vec2f, array : u32, depth_ref : f32) {
+ let l = textureSampleCompare(t, s, coords, array, depth_ref);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ sample_compare_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x);
+ sample_compare_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleCompareLevel) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn sample_compare_level(coords : vec2f, depth_ref : f32) {
+ let l = textureSampleCompareLevel(t, s, coords, depth_ref);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleCompareLevel_Offset) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn sample_compare_level(coords : vec2f, depth_ref : f32) {
+ const offset = vec2i(1);
+ let l = textureSampleCompareLevel(t, s, coords, depth_ref, offset);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleCompareLevel_ArrayIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_depth_2d_array;
+
+@group(0) @binding(1) var s : sampler_comparison;
+
+fn sample_compare_level_signed(coords : vec2f, array : i32, depth_ref : f32) {
+ let l = textureSampleCompareLevel(t, s, coords, array, depth_ref);
+}
+
+fn sample_compare_level_unsigned(coords : vec2f, array : u32, depth_ref : f32) {
+ let l = textureSampleCompareLevel(t, s, coords, array, depth_ref);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ sample_compare_level_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x);
+ sample_compare_level_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleGrad) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_compare_level(coords : vec2f, ddx : vec2f, ddy : vec2f) {
+ let l = textureSampleGrad(t, s, coords, ddx, ddy);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleGrad_Offset) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_compare_level(coords : vec2f, ddx : vec2f, ddy : vec2f) {
+ const offset = vec2i(1);
+ let l = textureSampleGrad(t, s, coords, ddx, ddy, offset);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleGrad_ArrayIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_grad_signed(coords : vec2f, array : i32, ddx : vec2f, ddy : vec2f) {
+ let l = textureSampleGrad(t, s, coords, array, ddx, ddy);
+}
+
+fn sample_grad_unsigned(coords : vec2f, array : u32, ddx : vec2f, ddy : vec2f) {
+ let l = textureSampleGrad(t, s, coords, array, ddx, ddy);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ sample_grad_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.xy, non_uniform.xy);
+ sample_grad_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.xy, non_uniform.xy);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleLevel) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_compare_level(coords : vec2f, level : f32) {
+ let l = textureSampleLevel(t, s, coords, level);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleLevel_Offset) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_compare_level(coords : vec2f, level : f32) {
+ const offset = vec2i(1);
+ let l = textureSampleLevel(t, s, coords, level, offset);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleLevel_ArrayIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d_array<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_compare_level_signed(coords : vec2f, array : i32, level : f32) {
+ let l = textureSampleLevel(t, s, coords, array, level);
+}
+
+fn sample_compare_level_unsigned(coords : vec2f, array : u32, level : f32) {
+ let l = textureSampleLevel(t, s, coords, array, level);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ sample_compare_level_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x);
+ sample_compare_level_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureSampleBaseClampToEdge) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_external;
+
+@group(0) @binding(1) var s : sampler;
+
+fn sample_base_clamp_to_edge(coords : vec2f) {
+ let l = textureSampleBaseClampToEdge(t, s, coords);
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureStore_1D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn store_signed(coords : i32, value : vec4i) {
+ textureStore(t, coords, value);
+}
+
+fn store_unsigned(coords : u32, value : vec4i) {
+ textureStore(t, coords, value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(i32(non_uniform.x), vec4i(non_uniform));
+ store_unsigned(u32(non_uniform.x), vec4i(non_uniform));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn store_signed(coords : i32, value : vec4i) {
+ textureStore(t, clamp(coords, 0, i32((textureDimensions(t) - 1))), value);
+}
+
+fn store_unsigned(coords : u32, value : vec4i) {
+ textureStore(t, min(coords, (textureDimensions(t) - 1)), value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(i32(non_uniform.x), vec4i(non_uniform));
+ store_unsigned(u32(non_uniform.x), vec4i(non_uniform));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn store_signed(coords : i32, value : vec4i) {
+ let coords_1 = u32(coords);
+ if (all((coords_1 < textureDimensions(t)))) {
+ textureStore(t, coords_1, value);
+ }
+}
+
+fn store_unsigned(coords : u32, value : vec4i) {
+ let coords_2 = u32(coords);
+ if (all((coords_2 < textureDimensions(t)))) {
+ textureStore(t, coords_2, value);
+ }
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(i32(non_uniform.x), vec4i(non_uniform));
+ store_unsigned(u32(non_uniform.x), vec4i(non_uniform));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureStore_2D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_storage_2d<rgba8sint, write>;
+
+fn store_signed(coords : vec2i, value : vec4i) {
+ textureStore(t, coords, value);
+}
+
+fn store_unsigned(coords : vec2u, value : vec4i) {
+ textureStore(t, coords, value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec2i(non_uniform.xy), vec4i(non_uniform));
+ store_unsigned(vec2u(non_uniform.xy), vec4i(non_uniform));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_storage_2d<rgba8sint, write>;
+
+fn store_signed(coords : vec2i, value : vec4i) {
+ textureStore(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t) - vec2(1)))), value);
+}
+
+fn store_unsigned(coords : vec2u, value : vec4i) {
+ textureStore(t, min(coords, (textureDimensions(t) - vec2(1))), value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec2i(non_uniform.xy), vec4i(non_uniform));
+ store_unsigned(vec2u(non_uniform.xy), vec4i(non_uniform));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_storage_2d<rgba8sint, write>;
+
+fn store_signed(coords : vec2i, value : vec4i) {
+ let coords_1 = vec2<u32>(coords);
+ if (all((coords_1 < textureDimensions(t)))) {
+ textureStore(t, coords_1, value);
+ }
+}
+
+fn store_unsigned(coords : vec2u, value : vec4i) {
+ let coords_2 = vec2<u32>(coords);
+ if (all((coords_2 < textureDimensions(t)))) {
+ textureStore(t, coords_2, value);
+ }
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec2i(non_uniform.xy), vec4i(non_uniform));
+ store_unsigned(vec2u(non_uniform.xy), vec4i(non_uniform));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureStore_2DArray) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_storage_2d_array<rgba8sint, write>;
+
+fn store_signed(coords : vec2i, array : i32, value : vec4i) {
+ textureStore(t, coords, array, value);
+}
+
+fn store_unsigned(coords : vec2u, array : u32, value : vec4i) {
+ textureStore(t, coords, array, value);
+}
+
+fn store_mixed(coords : vec2i, array : u32, value : vec4i) {
+ textureStore(t, coords, array, value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec2i(non_uniform.xy), i32(non_uniform.x), vec4i(non_uniform));
+ store_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform));
+ store_mixed(vec2i(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_storage_2d_array<rgba8sint, write>;
+
+fn store_signed(coords : vec2i, array : i32, value : vec4i) {
+ textureStore(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t) - vec2(1)))), clamp(array, 0, i32((textureNumLayers(t) - 1))), value);
+}
+
+fn store_unsigned(coords : vec2u, array : u32, value : vec4i) {
+ textureStore(t, min(coords, (textureDimensions(t) - vec2(1))), min(array, (textureNumLayers(t) - 1)), value);
+}
+
+fn store_mixed(coords : vec2i, array : u32, value : vec4i) {
+ textureStore(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t) - vec2(1)))), min(array, (textureNumLayers(t) - 1)), value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec2i(non_uniform.xy), i32(non_uniform.x), vec4i(non_uniform));
+ store_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform));
+ store_mixed(vec2i(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_storage_2d_array<rgba8sint, write>;
+
+fn store_signed(coords : vec2i, array : i32, value : vec4i) {
+ let coords_1 = vec2<u32>(coords);
+ let array_idx = u32(array);
+ if ((all((coords_1 < textureDimensions(t))) & (array_idx < textureNumLayers(t)))) {
+ textureStore(t, coords_1, array_idx, value);
+ }
+}
+
+fn store_unsigned(coords : vec2u, array : u32, value : vec4i) {
+ let coords_2 = vec2<u32>(coords);
+ let array_idx_1 = u32(array);
+ if ((all((coords_2 < textureDimensions(t))) & (array_idx_1 < textureNumLayers(t)))) {
+ textureStore(t, coords_2, array_idx_1, value);
+ }
+}
+
+fn store_mixed(coords : vec2i, array : u32, value : vec4i) {
+ let coords_3 = vec2<u32>(coords);
+ let array_idx_2 = u32(array);
+ if ((all((coords_3 < textureDimensions(t))) & (array_idx_2 < textureNumLayers(t)))) {
+ textureStore(t, coords_3, array_idx_2, value);
+ }
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec2i(non_uniform.xy), i32(non_uniform.x), vec4i(non_uniform));
+ store_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform));
+ store_mixed(vec2i(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureStore_3D) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_storage_3d<rgba8sint, write>;
+
+fn store_signed(coords : vec3i, value : vec4i) {
+ textureStore(t, coords, value);
+}
+
+fn store_unsigned(coords : vec3u, value : vec4i) {
+ textureStore(t, coords, value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec3i(non_uniform.xyz), vec4i(non_uniform));
+ store_unsigned(vec3u(non_uniform.xyz), vec4i(non_uniform));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_storage_3d<rgba8sint, write>;
+
+fn store_signed(coords : vec3i, value : vec4i) {
+ textureStore(t, clamp(coords, vec3(0), vec3<i32>((textureDimensions(t) - vec3(1)))), value);
+}
+
+fn store_unsigned(coords : vec3u, value : vec4i) {
+ textureStore(t, min(coords, (textureDimensions(t) - vec3(1))), value);
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec3i(non_uniform.xyz), vec4i(non_uniform));
+ store_unsigned(vec3u(non_uniform.xyz), vec4i(non_uniform));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_storage_3d<rgba8sint, write>;
+
+fn store_signed(coords : vec3i, value : vec4i) {
+ let coords_1 = vec3<u32>(coords);
+ if (all((coords_1 < textureDimensions(t)))) {
+ textureStore(t, coords_1, value);
+ }
+}
+
+fn store_unsigned(coords : vec3u, value : vec4i) {
+ let coords_2 = vec3<u32>(coords);
+ if (all((coords_2 < textureDimensions(t)))) {
+ textureStore(t, coords_2, value);
+ }
+}
+
+@fragment
+fn main(@builtin(position) non_uniform : vec4f) {
+ store_signed(vec3i(non_uniform.xyz), vec4i(non_uniform));
+ store_unsigned(vec3u(non_uniform.xyz), vec4i(non_uniform));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Other
+////////////////////////////////////////////////////////////////////////////////
+TEST_P(RobustnessTest, ShadowedVariable) {
auto* src = R"(
fn f() {
var a : array<f32, 3>;
var i : u32;
{
- var a : array<f32, 5>;
- var b : f32 = a[i];
+ var a : array<f32, 5>;
+ var b : f32 = a[i];
}
var c : f32 = a[i];
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
fn f() {
var a : array<f32, 3>;
var i : u32;
@@ -1099,19 +3514,42 @@
}
var c : f32 = a[min(i, 2u)];
}
-)";
+)",
+ /* predicate */ R"(
+fn f() {
+ var a : array<f32, 3>;
+ var i : u32;
+ {
+ var a : array<f32, 5>;
+ let index = i;
+ let predicate = (u32(index) <= 4u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ var b : f32 = predicated_expr;
+ }
+ let index_1 = i;
+ let predicate_1 = (u32(index_1) <= 2u);
+ var predicated_expr_1 : f32;
+ if (predicate_1) {
+ predicated_expr_1 = a[index_1];
+ }
+ var c : f32 = predicated_expr_1;
+}
+)");
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
// Check that existing use of min() and arrayLength() do not get renamed.
-TEST_F(RobustnessTest, DontRenameSymbols) {
+TEST_P(RobustnessTest, DontRenameSymbols) {
auto* src = R"(
struct S {
a : f32,
b : array<f32>,
-};
+}
@group(0) @binding(0) var<storage, read> s : S;
@@ -1120,11 +3558,13 @@
fn f() {
let b : f32 = s.b[c];
let x : i32 = min(1, 2);
- let y : u32 = arrayLength(&s.b);
+ let y : u32 = arrayLength(&(s.b));
}
)";
- auto* expect = R"(
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
struct S {
a : f32,
b : array<f32>,
@@ -1139,257 +3579,39 @@
let x : i32 = min(1, 2);
let y : u32 = arrayLength(&(s.b));
}
-)";
-
- auto got = Run<Robustness>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-const char* kOmitSourceShader = R"(
+)",
+ /* predicate */ R"(
struct S {
- vector : vec3<f32>,
- fixed_arr : array<f32, 4>,
- runtime_arr : array<f32>,
-};
-@group(0) @binding(0) var<storage, read> s : S;
-
-struct U {
- vector : vec4<f32>,
- fixed_arr : array<vec4<f32>, 4>,
-};
-@group(1) @binding(0) var<uniform> u : U;
-
-fn f() {
- // i32
- {
- let i = 0i;
- var storage_vector : f32 = s.vector[i];
- var storage_fixed_arr : f32 = s.fixed_arr[i];
- var storage_runtime_arr : f32 = s.runtime_arr[i];
- var uniform_vector : f32 = u.vector[i];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[i];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][i];
- }
- // u32
- {
- let i = 0u;
- var storage_vector : f32 = s.vector[i];
- var storage_fixed_arr : f32 = s.fixed_arr[i];
- var storage_runtime_arr : f32 = s.runtime_arr[i];
- var uniform_vector : f32 = u.vector[i];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[i];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][i];
- }
-}
-)";
-
-TEST_F(RobustnessTest, OmitNone) {
- auto* expect =
- R"(
-struct S {
- vector : vec3<f32>,
- fixed_arr : array<f32, 4>,
- runtime_arr : array<f32>,
+ a : f32,
+ b : array<f32>,
}
@group(0) @binding(0) var<storage, read> s : S;
-struct U {
- vector : vec4<f32>,
- fixed_arr : array<vec4<f32>, 4>,
-}
-
-@group(1) @binding(0) var<uniform> u : U;
+const c : u32 = 1u;
fn f() {
- {
- let i = 0i;
- var storage_vector : f32 = s.vector[min(u32(i), 2u)];
- var storage_fixed_arr : f32 = s.fixed_arr[min(u32(i), 3u)];
- var storage_runtime_arr : f32 = s.runtime_arr[min(u32(i), (arrayLength(&(s.runtime_arr)) - 1u))];
- var uniform_vector : f32 = u.vector[min(u32(i), 3u)];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[min(u32(i), 3u)];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][min(u32(i), 3u)];
+ let index = c;
+ let predicate = (u32(index) <= (arrayLength(&(s.b)) - 1u));
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = s.b[index];
}
- {
- let i = 0u;
- var storage_vector : f32 = s.vector[min(i, 2u)];
- var storage_fixed_arr : f32 = s.fixed_arr[min(i, 3u)];
- var storage_runtime_arr : f32 = s.runtime_arr[min(i, (arrayLength(&(s.runtime_arr)) - 1u))];
- var uniform_vector : f32 = u.vector[min(i, 3u)];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[min(i, 3u)];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][min(i, 3u)];
- }
+ let b : f32 = predicated_expr;
+ let x : i32 = min(1, 2);
+ let y : u32 = arrayLength(&(s.b));
}
-)";
+)");
- Robustness::Config cfg;
- DataMap data;
- data.Add<Robustness::Config>(cfg);
-
- auto got = Run<Robustness>(kOmitSourceShader, data);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
-TEST_F(RobustnessTest, OmitStorage) {
- auto* expect =
- R"(
-struct S {
- vector : vec3<f32>,
- fixed_arr : array<f32, 4>,
- runtime_arr : array<f32>,
-}
-
-@group(0) @binding(0) var<storage, read> s : S;
-
-struct U {
- vector : vec4<f32>,
- fixed_arr : array<vec4<f32>, 4>,
-}
-
-@group(1) @binding(0) var<uniform> u : U;
-
-fn f() {
- {
- let i = 0i;
- var storage_vector : f32 = s.vector[i];
- var storage_fixed_arr : f32 = s.fixed_arr[i];
- var storage_runtime_arr : f32 = s.runtime_arr[i];
- var uniform_vector : f32 = u.vector[min(u32(i), 3u)];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[min(u32(i), 3u)];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][min(u32(i), 3u)];
- }
- {
- let i = 0u;
- var storage_vector : f32 = s.vector[i];
- var storage_fixed_arr : f32 = s.fixed_arr[i];
- var storage_runtime_arr : f32 = s.runtime_arr[i];
- var uniform_vector : f32 = u.vector[min(i, 3u)];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[min(i, 3u)];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][min(i, 3u)];
- }
-}
-)";
-
- Robustness::Config cfg;
- cfg.omitted_address_spaces.insert(Robustness::AddressSpace::kStorage);
-
- DataMap data;
- data.Add<Robustness::Config>(cfg);
-
- auto got = Run<Robustness>(kOmitSourceShader, data);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, OmitUniform) {
- auto* expect =
- R"(
-struct S {
- vector : vec3<f32>,
- fixed_arr : array<f32, 4>,
- runtime_arr : array<f32>,
-}
-
-@group(0) @binding(0) var<storage, read> s : S;
-
-struct U {
- vector : vec4<f32>,
- fixed_arr : array<vec4<f32>, 4>,
-}
-
-@group(1) @binding(0) var<uniform> u : U;
-
-fn f() {
- {
- let i = 0i;
- var storage_vector : f32 = s.vector[min(u32(i), 2u)];
- var storage_fixed_arr : f32 = s.fixed_arr[min(u32(i), 3u)];
- var storage_runtime_arr : f32 = s.runtime_arr[min(u32(i), (arrayLength(&(s.runtime_arr)) - 1u))];
- var uniform_vector : f32 = u.vector[i];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[i];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][i];
- }
- {
- let i = 0u;
- var storage_vector : f32 = s.vector[min(i, 2u)];
- var storage_fixed_arr : f32 = s.fixed_arr[min(i, 3u)];
- var storage_runtime_arr : f32 = s.runtime_arr[min(i, (arrayLength(&(s.runtime_arr)) - 1u))];
- var uniform_vector : f32 = u.vector[i];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[i];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][i];
- }
-}
-)";
-
- Robustness::Config cfg;
- cfg.omitted_address_spaces.insert(Robustness::AddressSpace::kUniform);
-
- DataMap data;
- data.Add<Robustness::Config>(cfg);
-
- auto got = Run<Robustness>(kOmitSourceShader, data);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, OmitBoth) {
- auto* expect =
- R"(
-struct S {
- vector : vec3<f32>,
- fixed_arr : array<f32, 4>,
- runtime_arr : array<f32>,
-}
-
-@group(0) @binding(0) var<storage, read> s : S;
-
-struct U {
- vector : vec4<f32>,
- fixed_arr : array<vec4<f32>, 4>,
-}
-
-@group(1) @binding(0) var<uniform> u : U;
-
-fn f() {
- {
- let i = 0i;
- var storage_vector : f32 = s.vector[i];
- var storage_fixed_arr : f32 = s.fixed_arr[i];
- var storage_runtime_arr : f32 = s.runtime_arr[i];
- var uniform_vector : f32 = u.vector[i];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[i];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][i];
- }
- {
- let i = 0u;
- var storage_vector : f32 = s.vector[i];
- var storage_fixed_arr : f32 = s.fixed_arr[i];
- var storage_runtime_arr : f32 = s.runtime_arr[i];
- var uniform_vector : f32 = u.vector[i];
- var uniform_fixed_arr : vec4<f32> = u.fixed_arr[i];
- var uniform_fixed_arr_vector : f32 = u.fixed_arr[0][i];
- }
-}
-)";
-
- Robustness::Config cfg;
- cfg.omitted_address_spaces.insert(Robustness::AddressSpace::kStorage);
- cfg.omitted_address_spaces.insert(Robustness::AddressSpace::kUniform);
-
- DataMap data;
- data.Add<Robustness::Config>(cfg);
-
- auto got = Run<Robustness>(kOmitSourceShader, data);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(RobustnessTest, WorkgroupOverrideCount) {
+TEST_P(RobustnessTest, WorkgroupOverrideCount) {
auto* src = R"(
override N = 123;
+
var<workgroup> w : array<f32, N>;
fn f() {
@@ -1398,13 +3620,1880 @@
)";
auto* expect =
- R"(error: array size is an override-expression, when expected a constant-expression.
-Was the SubstituteOverride transform run?)";
+ Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */
+ R"(error: array size is an override-expression, when expected a constant-expression.
+Was the SubstituteOverride transform run?)",
+ /* predicate */
+ R"(error: array size is an override-expression, when expected a constant-expression.
+Was the SubstituteOverride transform run?)");
- auto got = Run<Robustness>(src);
+ auto got = Run<Robustness>(src, Config(GetParam()));
EXPECT_EQ(expect, str(got));
}
+////////////////////////////////////////////////////////////////////////////////
+// atomic predication
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, AtomicLoad) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ let i = 0;
+ let r = atomicLoad(&(a[i]));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ let i = 0;
+ let r = atomicLoad(&(a[min(u32(i), 3u)]));
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ var predicated_value : i32;
+ if (predicate) {
+ predicated_value = atomicLoad(&(a[index]));
+ }
+ let r = predicated_value;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, AtomicAnd) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ let i = 0;
+ atomicAnd(&(a[i]), 10);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ let i = 0;
+ atomicAnd(&(a[min(u32(i), 3u)]), 10);
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ if (predicate) {
+ atomicAnd(&(a[index]), 10);
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// workgroupUniformLoad
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, WorkgroupUniformLoad) {
+ auto* src = R"(
+var<workgroup> a : array<u32, 32>;
+
+@compute @workgroup_size(1)
+fn f() {
+ let i = 1;
+ let p = &(a[i]);
+ let v = workgroupUniformLoad(p);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<workgroup> a : array<u32, 32>;
+
+@compute @workgroup_size(1)
+fn f() {
+ let i = 1;
+ let p = &(a[min(u32(i), 31u)]);
+ let v = workgroupUniformLoad(p);
+}
+)",
+ /* predicate */ R"(
+var<workgroup> a : array<u32, 32>;
+
+@compute @workgroup_size(1)
+fn f() {
+ let i = 1;
+ let index = i;
+ let predicate = (u32(index) <= 31u);
+ let p = &(a[index]);
+ var predicated_value : u32;
+ if (predicate) {
+ predicated_value = workgroupUniformLoad(p);
+ } else {
+ workgroupBarrier();
+ }
+ let v = predicated_value;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Usage in loops
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, ArrayVal_ForLoopInit) {
+ auto* src = R"(
+fn f() {
+ let a = array<i32, 3>();
+ var v = 1;
+ for(var i = a[v]; (i < 3); i++) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+fn f() {
+ let a = array<i32, 3>();
+ var v = 1;
+ for(var i = a[min(u32(v), 2u)]; (i < 3); i++) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+fn f() {
+ let a = array<i32, 3>();
+ var v = 1;
+ {
+ let index = v;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ var i = predicated_expr;
+ loop {
+ if (!((i < 3))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, ArrayVal_ForLoopCond) {
+ auto* src = R"(
+fn f() {
+ let a = array<i32, 3>();
+ var v = 1;
+ for(var i = 0; (i < a[v]); i++) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+fn f() {
+ let a = array<i32, 3>();
+ var v = 1;
+ for(var i = 0; (i < a[min(u32(v), 2u)]); i++) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+fn f() {
+ let a = array<i32, 3>();
+ var v = 1;
+ {
+ var i = 0;
+ loop {
+ let index = v;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ if (!((i < predicated_expr))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, ArrayVal_ForLoopCont) {
+ auto* src = R"(
+fn f() {
+ let a = array<i32, 3>();
+ for(var i = 0; (i < 5); i = a[i]) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+fn f() {
+ let a = array<i32, 3>();
+ for(var i = 0; (i < 5); i = a[min(u32(i), 2u)]) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+fn f() {
+ let a = array<i32, 3>();
+ {
+ var i = 0;
+ loop {
+ if (!((i < 5))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ let index = i;
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = a[index];
+ }
+ i = predicated_expr;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_ForLoopInit) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var coords = vec2(1);
+ var level = 1;
+ for(var i = textureLoad(t, coords, level).x; (i < 3); i++) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var coords = vec2(1);
+ var level = 1;
+ {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ var i = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t, level_idx) - vec2(1)))), level_idx).x;
+ loop {
+ if (!((i < 3))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var coords = vec2(1);
+ var level = 1;
+ {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = vec2<u32>(coords);
+ var predicated_value : vec4<i32>;
+ if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) {
+ predicated_value = textureLoad(t, coords_1, level_idx);
+ }
+ var i = predicated_value.x;
+ loop {
+ if (!((i < 3))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_ForLoopCond) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var coords = vec2(1);
+ var level = 1;
+ for(var i = 0; (i < textureLoad(t, coords, level).x); i++) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var coords = vec2(1);
+ var level = 1;
+ {
+ var i = 0;
+ loop {
+ let level_idx = min(u32(level), (textureNumLevels(t) - 1));
+ if (!((i < textureLoad(t, clamp(coords, vec2(0), vec2<i32>((textureDimensions(t, level_idx) - vec2(1)))), level_idx).x))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var coords = vec2(1);
+ var level = 1;
+ {
+ var i = 0;
+ loop {
+ let level_idx = u32(level);
+ let num_levels = textureNumLevels(t);
+ let coords_1 = vec2<u32>(coords);
+ var predicated_value : vec4<i32>;
+ if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) {
+ predicated_value = textureLoad(t, coords_1, level_idx);
+ }
+ if (!((i < predicated_value.x))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureLoad_ForLoopCont) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var level = 1;
+ for(var i = 0; (i < 5); i = textureLoad(t, vec2(i), i).x) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var level = 1;
+ {
+ var i = 0;
+ loop {
+ if (!((i < 5))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ let level_idx = min(u32(i), (textureNumLevels(t) - 1));
+ i = textureLoad(t, clamp(vec2(i), vec2(0), vec2<i32>((textureDimensions(t, level_idx) - vec2(1)))), level_idx).x;
+ }
+ }
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_2d<i32>;
+
+fn f() {
+ var level = 1;
+ {
+ var i = 0;
+ loop {
+ if (!((i < 5))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ let level_idx = u32(i);
+ let num_levels = textureNumLevels(t);
+ let coords = vec2<u32>(vec2(i));
+ var predicated_value : vec4<i32>;
+ if (((level_idx < num_levels) & all((coords < textureDimensions(t, min(level_idx, (num_levels - 1))))))) {
+ predicated_value = textureLoad(t, coords, level_idx);
+ }
+ i = predicated_value.x;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureStore_ForLoopInit) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn f() {
+ var coords = 1;
+ var value = vec4(1);
+ var i = 0;
+ for(textureStore(t, coords, value); (i < 3); i++) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn f() {
+ var coords = 1;
+ var value = vec4(1);
+ var i = 0;
+ for(textureStore(t, clamp(coords, 0, i32((textureDimensions(t) - 1))), value); (i < 3); i++) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn f() {
+ var coords = 1;
+ var value = vec4(1);
+ var i = 0;
+ {
+ let coords_1 = u32(coords);
+ if (all((coords_1 < textureDimensions(t)))) {
+ textureStore(t, coords_1, value);
+ }
+ loop {
+ if (!((i < 3))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, TextureStore_ForLoopCont) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn f() {
+ var level = 1;
+ var value = vec4(1);
+ for(var i = 0; (i < 3); textureStore(t, i, value)) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn f() {
+ var level = 1;
+ var value = vec4(1);
+ for(var i = 0; (i < 3); textureStore(t, clamp(i, 0, i32((textureDimensions(t) - 1))), value)) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var t : texture_storage_1d<rgba8sint, write>;
+
+fn f() {
+ var level = 1;
+ var value = vec4(1);
+ {
+ var i = 0;
+ loop {
+ if (!((i < 3))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ let coords = u32(i);
+ if (all((coords < textureDimensions(t)))) {
+ textureStore(t, coords, value);
+ }
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, AtomicXor_ForLoopInit) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ var i = 0;
+ for(atomicXor(&(a[i]), 4); (i < 3); i++) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ var i = 0;
+ for(atomicXor(&(a[min(u32(i), 3u)]), 4); (i < 3); i++) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ var i = 0;
+ {
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ if (predicate) {
+ atomicXor(&(a[index]), 4);
+ }
+ loop {
+ if (!((i < 3))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, AtomicXor_ForLoopCond) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ for(var i = 0; (i < atomicXor(&(a[i]), 4)); i++) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ for(var i = 0; (i < atomicXor(&(a[min(u32(i), 3u)]), 4)); i++) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ {
+ var i = 0;
+ loop {
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ var predicated_value : i32;
+ if (predicate) {
+ predicated_value = atomicXor(&(a[index]), 4);
+ }
+ if (!((i < predicated_value))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ i++;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, AtomicXor_ForLoopCont) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ for(var i = 0; (i < 5); i = atomicXor(&(a[i]), 4)) {
+ var in_loop = 42;
+ }
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ for(var i = 0; (i < 5); i = atomicXor(&(a[min(u32(i), 3u)]), 4)) {
+ var in_loop = 42;
+ }
+}
+)",
+ /* predicate */ R"(
+@group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>;
+
+fn f() {
+ {
+ var i = 0;
+ loop {
+ if (!((i < 5))) {
+ break;
+ }
+ {
+ var in_loop = 42;
+ }
+
+ continuing {
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ var predicated_value : i32;
+ if (predicate) {
+ predicated_value = atomicXor(&(a[index]), 4);
+ }
+ i = predicated_value;
+ }
+ }
+ }
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// User pointer parameters
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_P(RobustnessTest, Read_PrivatePointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ y(1, &(a[1]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) -> i32 {
+ var predicated_expr : i32;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, p_predicate_1 : bool, post : i32) -> i32 {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ y(1, &(a[1]), true, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ y(1, &(a[1]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) -> i32 {
+ var predicated_expr : i32;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, p_predicate_1 : bool, post : i32) -> i32 {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ y(1, &(a[1]), true, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_UniformPointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
+ return x(pre, p, post);
+}
+
+fn z() {
+ y(vec4(1), &(a[1]), vec4(2));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
+ var predicated_expr : vec4<i32>;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate_1 : bool, post : vec4i) -> vec4i {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ y(vec4(1), &(a[1]), true, vec4(2));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_StoragePointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
+ return x(pre, p, post);
+}
+
+fn z() {
+ y(vec4(1), &(a[1]), vec4(2));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<storage, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
+ var predicated_expr : vec4<i32>;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : vec4i, p : ptr<storage, vec4i>, p_predicate_1 : bool, post : vec4i) -> vec4i {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ y(vec4(1), &(a[1]), true, vec4(2));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_FunctionPointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ y(1, &(a[1]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) -> i32 {
+ var predicated_expr : i32;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, p_predicate_1 : bool, post : i32) -> i32 {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ y(1, &(a[1]), true, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_PrivatePointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[i]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[min(u32(i), 3u)]), 2);
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) -> i32 {
+ var predicated_expr : i32;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, p_predicate_1 : bool, post : i32) -> i32 {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(1, &(a[index]), predicate, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[i]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[min(u32(i), 3u)]), 2);
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) -> i32 {
+ var predicated_expr : i32;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, p_predicate_1 : bool, post : i32) -> i32 {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(1, &(a[index]), predicate, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_UniformPointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(vec4(1), &(a[i]), vec4(2));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(vec4(1), &(a[min(u32(i), 3u)]), vec4(2));
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
+ var predicated_expr : vec4<i32>;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate_1 : bool, post : vec4i) -> vec4i {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(vec4(1), &(a[index]), predicate, vec4(2));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_StoragePointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(vec4(1), &(a[i]), vec4(2));
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
+ return x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(vec4(1), &(a[min(u32(i), 3u)]), vec4(2));
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
+
+fn x(pre : vec4i, p : ptr<storage, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
+ var predicated_expr : vec4<i32>;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : vec4i, p : ptr<storage, vec4i>, p_predicate_1 : bool, post : vec4i) -> vec4i {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(vec4(1), &(a[index]), predicate, vec4(2));
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Read_FunctionPointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ let i = 0;
+ y(1, &(a[i]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
+ return ((pre + *(p)) + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
+ return x(pre, p, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ let i = 0;
+ y(1, &(a[min(u32(i), 3u)]), 2);
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) -> i32 {
+ var predicated_expr : i32;
+ if (p_predicate) {
+ predicated_expr = *(p);
+ }
+ return ((pre + predicated_expr) + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, p_predicate_1 : bool, post : i32) -> i32 {
+ return x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(1, &(a[index]), predicate, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ y(1, &(a[1]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<private, i32>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ y(1, &(a[1]), true, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ y(1, &(a[1]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ y(1, &(a[1]), true, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_StoragePointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ y(1, &(a[1]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<storage, i32, read_write>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<storage, i32, read_write>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ y(1, &(a[1]), true, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_FunctionPointerParameter_IndexWithConstant) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ y(1, &(a[1]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ src,
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<function, i32>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ y(1, &(a[1]), true, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[i]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<private, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[min(u32(i), 3u)]), 2);
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<private> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<private, i32>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(1, &(a[index]), predicate, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[i]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[min(u32(i), 3u)]), 2);
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+var<workgroup> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<workgroup, i32>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(1, &(a[index]), predicate, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_StoragePointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[i]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ let i = 0;
+ y(1, &(a[min(u32(i), 3u)]), 2);
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
+
+fn x(pre : i32, p : ptr<storage, i32, read_write>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<storage, i32, read_write>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(1, &(a[index]), predicate, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_P(RobustnessTest, Write_FunctionPointerParameter_IndexWithLet) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ let i = 0;
+ y(1, &(a[i]), 2);
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, post : i32) {
+ *(p) = (pre + post);
+}
+
+fn y(pre : i32, p : ptr<function, i32>, post : i32) {
+ x(pre, p, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ let i = 0;
+ y(1, &(a[min(u32(i), 3u)]), 2);
+}
+)",
+ /* predicate */ R"(
+enable chromium_experimental_full_ptr_parameters;
+
+fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) {
+ if (p_predicate) {
+ *(p) = (pre + post);
+ }
+}
+
+fn y(pre : i32, p : ptr<function, i32>, p_predicate_1 : bool, post : i32) {
+ x(pre, p, p_predicate_1, post);
+}
+
+fn z() {
+ var a : array<i32, 4>;
+ let i = 0;
+ let index = i;
+ let predicate = (u32(index) <= 3u);
+ y(1, &(a[index]), predicate, 2);
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+INSTANTIATE_TEST_SUITE_P(,
+ RobustnessTest,
+ testing::Values(Robustness::Action::kIgnore,
+ Robustness::Action::kClamp,
+ Robustness::Action::kPredicate));
} // namespace
} // namespace tint::transform
diff --git a/src/tint/transform/spirv_atomic.cc b/src/tint/transform/spirv_atomic.cc
index b3924cc..043861b 100644
--- a/src/tint/transform/spirv_atomic.cc
+++ b/src/tint/transform/spirv_atomic.cc
@@ -294,7 +294,7 @@
SpirvAtomic::~SpirvAtomic() = default;
SpirvAtomic::Stub::Stub(ProgramID pid, ast::NodeID nid, sem::BuiltinType b)
- : Base(pid, nid), builtin(b) {}
+ : Base(pid, nid, utils::Empty), builtin(b) {}
SpirvAtomic::Stub::~Stub() = default;
std::string SpirvAtomic::Stub::InternalName() const {
return "@internal(spirv-atomic " + std::string(sem::str(builtin)) + ")";
diff --git a/src/tint/utils/unique_vector.h b/src/tint/utils/unique_vector.h
index 6cb8f88..eca0d4c 100644
--- a/src/tint/utils/unique_vector.h
+++ b/src/tint/utils/unique_vector.h
@@ -30,6 +30,9 @@
/// Attempting to add a duplicate is a no-op.
template <typename T, size_t N, typename HASH = std::hash<T>, typename EQUAL = std::equal_to<T>>
struct UniqueVector {
+ /// STL-friendly alias to T. Used by gmock.
+ using value_type = T;
+
/// Constructor
UniqueVector() = default;
diff --git a/src/tint/writer/append_vector.cc b/src/tint/writer/append_vector.cc
index c31667d..98d9798 100644
--- a/src/tint/writer/append_vector.cc
+++ b/src/tint/writer/append_vector.cc
@@ -156,13 +156,12 @@
}));
auto* ctor_target = b->create<sem::ValueConstructor>(
packed_sem_ty,
- utils::Transform(
- packed,
- [&](const tint::sem::ValueExpression* arg, size_t i) -> const sem::Parameter* {
- return b->create<sem::Parameter>(
- nullptr, static_cast<uint32_t>(i), arg->Type()->UnwrapRef(),
- builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
- }),
+ utils::Transform(packed,
+ [&](const tint::sem::ValueExpression* arg, size_t i) {
+ return b->create<sem::Parameter>(
+ nullptr, static_cast<uint32_t>(i), arg->Type()->UnwrapRef(),
+ builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ }),
sem::EvaluationStage::kRuntime);
auto* ctor_sem = b->create<sem::Call>(ctor_ast, ctor_target, sem::EvaluationStage::kRuntime,
std::move(packed), statement,
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 24c6386..02c1bac 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -166,7 +166,8 @@
manager.Add<transform::PromoteSideEffectsToDecl>();
if (!options.disable_robustness) {
- // Robustness must come before BuiltinPolyfill
+ // Robustness must come after PromoteSideEffectsToDecl
+ // Robustness must come before BuiltinPolyfill and CanonicalizeEntryPointIO
manager.Add<transform::Robustness>();
}
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 823e6bd..f0247c4 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -182,7 +182,8 @@
manager.Add<transform::PromoteSideEffectsToDecl>();
if (!options.disable_robustness) {
- // Robustness must come before BuiltinPolyfill
+ // Robustness must come after PromoteSideEffectsToDecl
+ // Robustness must come before BuiltinPolyfill and CanonicalizeEntryPointIO
manager.Add<transform::Robustness>();
}
@@ -1125,7 +1126,7 @@
utils::StringStream& out,
const ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
- auto const buffer = program_->Symbols().NameFor(intrinsic->buffer);
+ auto const buffer = program_->Symbols().NameFor(intrinsic->Buffer()->identifier->symbol);
auto* const offset = expr->args[0];
// offset in bytes
@@ -1413,7 +1414,7 @@
utils::StringStream& out,
const ast::CallExpression* expr,
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
- auto const buffer = program_->Symbols().NameFor(intrinsic->buffer);
+ auto const buffer = program_->Symbols().NameFor(intrinsic->Buffer()->identifier->symbol);
auto* const offset = expr->args[0];
auto* const value = expr->args[1];
@@ -1581,7 +1582,7 @@
const auto name = builder_.Symbols().NameFor(func->name->symbol);
auto& buf = *current_buffer_;
- auto const buffer = program_->Symbols().NameFor(intrinsic->buffer);
+ auto const buffer = program_->Symbols().NameFor(intrinsic->Buffer()->identifier->symbol);
auto rmw = [&](const char* hlsl) -> bool {
{
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index b2c9ffa..5c9e11b 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -207,7 +207,8 @@
manager.Add<transform::PromoteSideEffectsToDecl>();
if (!options.disable_robustness) {
- // Robustness must come before BuiltinPolyfill
+ // Robustness must come after PromoteSideEffectsToDecl
+ // Robustness must come before BuiltinPolyfill and CanonicalizeEntryPointIO
manager.Add<transform::Robustness>();
}
diff --git a/src/tint/writer/spirv/generator_impl.cc b/src/tint/writer/spirv/generator_impl.cc
index c46bf9a..72a3629 100644
--- a/src/tint/writer/spirv/generator_impl.cc
+++ b/src/tint/writer/spirv/generator_impl.cc
@@ -68,7 +68,8 @@
manager.Add<transform::MergeReturn>();
if (!options.disable_robustness) {
- // Robustness must come before BuiltinPolyfill
+ // Robustness must come after PromoteSideEffectsToDecl
+ // Robustness must come before BuiltinPolyfill and CanonicalizeEntryPointIO
manager.Add<transform::Robustness>();
}
diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc
index c052075..69aaabd 100644
--- a/src/tint/writer/wgsl/generator_impl.cc
+++ b/src/tint/writer/wgsl/generator_impl.cc
@@ -1149,7 +1149,7 @@
bool GeneratorImpl::EmitConstAssert(const ast::ConstAssert* stmt) {
auto out = line();
- out << "static_assert ";
+ out << "const_assert ";
if (!EmitExpression(out, stmt->condition)) {
return false;
}
diff --git a/src/tint/writer/wgsl/generator_impl_const_assert_test.cc b/src/tint/writer/wgsl/generator_impl_const_assert_test.cc
index 9f76d65..b785357 100644
--- a/src/tint/writer/wgsl/generator_impl_const_assert_test.cc
+++ b/src/tint/writer/wgsl/generator_impl_const_assert_test.cc
@@ -27,7 +27,7 @@
GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error();
- EXPECT_EQ(gen.result(), R"(static_assert true;
+ EXPECT_EQ(gen.result(), R"(const_assert true;
)");
}
@@ -38,7 +38,7 @@
ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_EQ(gen.result(), R"(fn f() {
- static_assert true;
+ const_assert true;
}
)");
}
