Import Tint changes from Dawn
Changes:
- 3b8b9699d69f1eeb6fadddcbe56ffc68ae57715e tint: Implement const eval of unary minus by Antonio Maiorano <amaiorano@google.com>
- eb02cd33012e1a461109b13c38335db9e711355e Regex fuzzer: replace function calls with builtins by Alastair F. Donaldson <alastair.donaldson@imperial.ac.uk>
- ca21fa019fa6e706320661a35e78ddcc6132384c Update dictionary for fuzzing by Alastair F. Donaldson <alastair.donaldson@imperial.ac.uk>
- 46c32d882d6d7a83e6c11e0b986a3f9d91200f1b Early out Is checks. by dan sinclair <dsinclair@chromium.org>
- 633cdf41cc40ff97619093c0784a87cd65a2b187 tint/resolver: Make the F16 ban error message show a sing... by Corentin Wallez <cwallez@chromium.org>
- e4df87fd017cd1d3eacbfbed57e521f6f7733936 Regex fuzzer: Add break and continue statements by Alastair F. Donaldson <alastair.donaldson@imperial.ac.uk>
- 84f7830874fc28c173448f8291ad8c2571fda573 Regex fuzzer: replace operators by Alastair F. Donaldson <alastair.donaldson@imperial.ac.uk>
- 056f97a9e53215398dfbd9145714e678c872a070 tint: Add a benchmark for atan2 const eval by Ben Clayton <bclayton@google.com>
- 853cbadc8fa6c5eaa2c6d2d55a5f10b6a9ab94d4 Clean regex fuzzer API by Alastair F. Donaldson <alastair.donaldson@imperial.ac.uk>
GitOrigin-RevId: 3b8b9699d69f1eeb6fadddcbe56ffc68ae57715e
Change-Id: I6a489442f5fff4d39396cc9ced675564730d2aa7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/96660
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/ast/f16.h b/src/tint/ast/f16.h
index c444a20..bae6291 100644
--- a/src/tint/ast/f16.h
+++ b/src/tint/ast/f16.h
@@ -22,7 +22,7 @@
namespace tint::ast {
/// A float 16 type
-class F16 : public Castable<F16, Type> {
+class F16 final : public Castable<F16, Type> {
public:
/// Constructor
/// @param pid the identifier of the program that owns this node
diff --git a/src/tint/ast/int_literal_expression.h b/src/tint/ast/int_literal_expression.h
index b4e184a..10cbbee 100644
--- a/src/tint/ast/int_literal_expression.h
+++ b/src/tint/ast/int_literal_expression.h
@@ -20,7 +20,7 @@
namespace tint::ast {
/// An integer literal. The literal may have an 'i', 'u' or no suffix.
-class IntLiteralExpression : public Castable<IntLiteralExpression, LiteralExpression> {
+class IntLiteralExpression final : public Castable<IntLiteralExpression, LiteralExpression> {
public:
/// Literal suffix
enum class Suffix {
diff --git a/src/tint/bench/benchmark.h b/src/tint/bench/benchmark.h
index d52a0d6..6801585 100644
--- a/src/tint/bench/benchmark.h
+++ b/src/tint/bench/benchmark.h
@@ -58,6 +58,7 @@
/// files in `<tint>/test/benchmark`.
#define TINT_BENCHMARK_WGSL_PROGRAMS(FUNC) \
TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "animometer.wgsl"); \
+ TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "atan2-const-eval.wgsl"); \
TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "bloom-vertical-blur.wgsl"); \
TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "cluster-lights.wgsl"); \
TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "empty.wgsl"); \
diff --git a/src/tint/castable.h b/src/tint/castable.h
index c7b8608..806d0f0 100644
--- a/src/tint/castable.h
+++ b/src/tint/castable.h
@@ -101,6 +101,36 @@
/// The type hash code bitwise-or'd with all ancestor's hashcodes.
const HashCode full_hashcode;
+ /// @returns true if `type` derives from the class `TO`
+ /// @param object the object type to test from, which must be, or derive from
+ /// type `FROM`.
+ /// @see CastFlags
+ template <typename TO, typename FROM, int FLAGS = 0>
+ static inline bool Is(const tint::TypeInfo* object) {
+ constexpr const bool downcast = std::is_base_of<FROM, TO>::value;
+ constexpr const bool upcast = std::is_base_of<TO, FROM>::value;
+ constexpr const bool nocast = std::is_same<FROM, TO>::value;
+ constexpr const bool assert_is_castable = (FLAGS & kDontErrorOnImpossibleCast) == 0;
+
+ static_assert(upcast || downcast || nocast || !assert_is_castable, "impossible cast");
+
+ return upcast || nocast || object->Is<TO>();
+ }
+
+ /// @returns true if this type derives from the class `T`
+ template <typename T>
+ inline bool Is() const {
+ auto* type = &Of<std::remove_cv_t<T>>();
+
+ if constexpr (std::is_final_v<T>) {
+ // T is final, so nothing can derive from T.
+ // We do not need to check ancestors, only whether this type is equal to the type T.
+ return type == this;
+ } else {
+ return Is(type);
+ }
+ }
+
/// @param type the test type info
/// @returns true if the class with this TypeInfo is of, or derives from the
/// class with the given TypeInfo.
@@ -112,8 +142,8 @@
return false;
}
- // Walk the base types, starting with this TypeInfo, to see if any of the
- // pointers match `type`.
+ // Walk the base types, starting with this TypeInfo, to see if any of the pointers match
+ // `type`.
for (auto* ti = this; ti != nullptr; ti = ti->base) {
if (ti == type) {
return true;
@@ -122,26 +152,6 @@
return false;
}
- /// @returns true if `type` derives from the class `TO`
- /// @param type the object type to test from, which must be, or derive from
- /// type `FROM`.
- /// @see CastFlags
- template <typename TO, typename FROM, int FLAGS = 0>
- static inline bool Is(const tint::TypeInfo* type) {
- constexpr const bool downcast = std::is_base_of<FROM, TO>::value;
- constexpr const bool upcast = std::is_base_of<TO, FROM>::value;
- constexpr const bool nocast = std::is_same<FROM, TO>::value;
- constexpr const bool assert_is_castable = (FLAGS & kDontErrorOnImpossibleCast) == 0;
-
- static_assert(upcast || downcast || nocast || !assert_is_castable, "impossible cast");
-
- if (upcast || nocast) {
- return true;
- }
-
- return type->Is(&Of<std::remove_cv_t<TO>>());
- }
-
/// @returns the static TypeInfo for the type T
template <typename T>
static const TypeInfo& Of() {
@@ -211,14 +221,12 @@
if constexpr (kCount == 0) {
return false;
} else if constexpr (kCount == 1) {
- return Is(&Of<std::tuple_element_t<0, TUPLE>>());
+ return Is<std::tuple_element_t<0, TUPLE>>();
} else if constexpr (kCount == 2) {
- return Is(&Of<std::tuple_element_t<0, TUPLE>>()) ||
- Is(&Of<std::tuple_element_t<1, TUPLE>>());
+ return Is<std::tuple_element_t<0, TUPLE>>() || Is<std::tuple_element_t<1, TUPLE>>();
} else if constexpr (kCount == 3) {
- return Is(&Of<std::tuple_element_t<0, TUPLE>>()) ||
- Is(&Of<std::tuple_element_t<1, TUPLE>>()) ||
- Is(&Of<std::tuple_element_t<2, TUPLE>>());
+ return Is<std::tuple_element_t<0, TUPLE>>() || Is<std::tuple_element_t<1, TUPLE>>() ||
+ Is<std::tuple_element_t<2, TUPLE>>();
} else {
// Optimization: Compare the object's hashcode to the bitwise-or of all
// the tested type's hashcodes. If there's no intersection of bits in
@@ -587,7 +595,7 @@
// Attempt to dynamically cast the object to the handler type. If that
// succeeds, call the case handler with the cast object.
using CaseType = SwitchCaseType<CaseFunc>;
- if (type->Is(&TypeInfo::Of<CaseType>())) {
+ if (type->Is<CaseType>()) {
auto* ptr = static_cast<CaseType*>(object);
if constexpr (kHasReturnType) {
new (result) RETURN_TYPE(static_cast<RETURN_TYPE>(std::get<0>(cases)(ptr)));
diff --git a/src/tint/fuzzers/dictionary.txt b/src/tint/fuzzers/dictionary.txt
index 8f15ff6..005735f 100644
--- a/src/tint/fuzzers/dictionary.txt
+++ b/src/tint/fuzzers/dictionary.txt
@@ -1,63 +1,165 @@
+"!"
+"!="
+"%"
+"%="
"&"
"&&"
-"->"
-"[["
-"]]"
-"/"
-"!"
-"["
-"]"
-"{"
-"}"
-":"
+"&="
+"("
+")"
+"*"
+"*="
+"+"
+"++"
+"+="
","
+"-"
+"--"
+"-="
+"->"
+"."
+"/"
+"/="
+":"
+";"
+"<"
+"<<"
+"<<="
+"<="
"="
"=="
">"
">="
-"<"
-"<="
-"%"
-"-"
-"::"
-"!="
-"."
-"+"
-"|"
-"||"
-"("
-")"
-";"
-"*"
+">>"
+">>="
+"@"
+"["
+"]"
"^"
+"^="
+"_"
+"{"
+"|"
+"|="
+"||"
+"}"
+"~"
+"a"
+"abs"
+"acos"
+"acosh"
+"@align"
+"all"
+"any"
"array"
-"binding"
+"arrayLength"
+"asin"
+"asinh"
+"atan"
+"atan2"
+"atanh"
+"atomic"
+"atomicAdd"
+"atomicAnd"
+"atomicLoad"
+"atomicMax"
+"atomicMin"
+"atomicOr"
+"atomicStore"
+"atomicSub"
+"atomicXor"
+"b"
+"@binding"
"bitcast"
"bool"
-"block"
"break"
-"builtin"
+"@builtin"
+"@builtin(frag_depth)"
+"@builtin(front_facing)"
+"@builtin(global_invocation_id)"
+"@builtin(instance_index)"
+"@builtin(local_invocation_id)"
+"@builtin(local_invocation_index)"
+"@builtin(num_workgroups)"
+"@builtin(position)"
+"@builtin(sample_index)"
+"@builtin(sample_mask)"
+"@builtin(vertex_index)"
+"@builtin(workgroup_id)"
"case"
-"compute"
+"ceil"
+"center"
+"centroid"
+"clamp"
+"@compute"
+"@const"
"const"
"continue"
"continuing"
-"discard"
+"cos"
+"cosh"
+"countLeadingZeros"
+"countOneBits"
+"countTrailingZeros"
+"cross"
"default"
+"degrees"
+"determinant"
+"discard"
+"distance"
+"dot"
+"dpdx"
+"dpdxCoarse"
+"dpdxFine"
+"dpdy"
+"dpdyCoarse"
+"dpdyFine"
"else"
-"elseif"
+"enable"
+"exp"
+"exp2"
+"extractBits"
+"f16"
"f32"
+"faceForward"
"fallthrough"
"false"
+"firstLeadingBit"
+"firstTrailingBit"
+"flat"
+"floor"
+"fma"
"fn"
-"fragment"
+"for"
+"fract"
+"frag_depth"
+"@fragment"
+"frexp"
+"front_facing"
"function"
+"fwidth"
+"fwidthCoarse"
+"fwidthFine"
+"g"
+"global_invocation_id"
+"@group"
"i32"
+"@id"
"if"
-"image"
-"import"
-"in"
-"location"
+"insertBits"
+"instance_index"
+"@interpolate"
+"@invariant"
+"inverseSqrt"
+"ldexp"
+"length"
+"let"
+"linear"
+"local_invocation_id"
+"local_invocation_index"
+"@location"
+"log"
+"log2"
"loop"
"mat2x2"
"mat2x3"
@@ -68,45 +170,124 @@
"mat4x2"
"mat4x3"
"mat4x4"
-"offset"
-"out"
+"max"
+"min"
+"mix"
+"modf"
+"normalize"
+"num_workgroups"
+"override"
+"pack2x16float"
+"pack2x16snorm"
+"pack2x16unorm"
+"pack4x8snorm"
+"pack4x8unorm"
+"perspective"
+"position"
+"pow"
"private"
"ptr"
+"quantizeToF16"
+"r"
+"r32float"
+"r32sint"
+"r32uint"
+"radians"
+"read"
+"read_write"
+"reflect"
+"refract"
"return"
+"reverseBits"
+"rg32float"
+"rg32sint"
+"rg32uint"
+"rgba16float"
+"rgba16sint"
+"rgba16uint"
+"rgba32float"
+"rgba32sint"
+"rgba32uint"
+"rgba8sint"
+"rgba8snorm"
+"rgba8uint"
+"rgba8unorm"
+"round"
+"sample"
+"sample_index"
+"sample_mask"
"sampler"
"sampler_comparison"
-"set"
+"saturate"
+"select"
+"sign"
+"sin"
+"sinh"
+"@size"
+"smoothstep"
+"sqrt"
+"staticAssert"
+"step"
"storage"
-"stage"
-"stride"
+"storageBarrier"
"struct"
"switch"
-"texture_depth_2d"
-"texture_depth_2d_array"
-"texture_depth_cube"
-"texture_depth_cube_array"
-"texture_depth_multisampled_2d"
-"texture_multisampled_2d"
-"texture_storage_1d"
-"texture_storage_2d_array"
-"texture_storage_2d"
-"texture_storage_2d_array"
-"texture_storage_3d"
+"tan"
+"tanh"
"texture_1d"
"texture_2d"
"texture_2d_array"
"texture_3d"
"texture_cube"
"texture_cube_array"
+"texture_depth_2d"
+"texture_depth_2d_array"
+"texture_depth_cube"
+"texture_depth_cube_array"
+"texture_depth_multisampled_2d"
+"textureDimensions"
+"textureGather"
+"textureGatherCompare"
+"textureLoad"
+"texture_multisampled_2d"
+"textureNumLayers"
+"textureNumLevels"
+"textureNumSamples"
+"textureSample"
+"textureSampleBias"
+"textureSampleCompare"
+"textureSampleCompareLevel"
+"textureSampleGrad"
+"textureSampleLevel"
+"texture_storage_1d"
+"texture_storage_2d"
+"texture_storage_2d_array"
+"texture_storage_3d"
+"textureStore"
+"transpose"
"true"
+"trunc"
"type"
"u32"
"uniform"
+"unpack2x16float"
+"unpack2x16snorm"
+"unpack2x16unorm"
+"unpack4x8snorm"
+"unpack4x8unorm"
"var"
"vec2"
"vec3"
"vec4"
-"vertex"
-"void"
+"@vertex"
+"vertex_index"
+"w"
+"while"
"workgroup"
-"workgroup_size"
+"workgroupBarrier"
+"workgroup_id"
+"@workgroup_size"
+"write"
+"x"
+"y"
+"z"
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/wrap_unary_operator.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/wrap_unary_operator.cc
index 8b52732..d6612f5 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/wrap_unary_operator.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/wrap_unary_operator.cc
@@ -18,6 +18,8 @@
#include <vector>
#include "src/tint/program_builder.h"
+#include "src/tint/sem/abstract_float.h"
+#include "src/tint/sem/abstract_int.h"
namespace tint::fuzzers::ast_fuzzer {
@@ -98,7 +100,8 @@
return {ast::UnaryOp::kNot};
}
- if (expr_type->is_signed_scalar_or_vector()) {
+ if (expr_type->is_signed_scalar_or_vector() ||
+ expr_type->is_abstract_integer_scalar_or_vector()) {
return {ast::UnaryOp::kNegation, ast::UnaryOp::kComplement};
}
@@ -106,7 +109,7 @@
return {ast::UnaryOp::kComplement};
}
- if (expr_type->is_float_scalar_or_vector()) {
+ if (expr_type->is_float_scalar_or_vector() || expr_type->is_abstract_float_scalar_or_vector()) {
return {ast::UnaryOp::kNegation};
}
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/fuzzer.cc b/src/tint/fuzzers/tint_regex_fuzzer/fuzzer.cc
index ac34684..ea1aea1 100644
--- a/src/tint/fuzzers/tint_regex_fuzzer/fuzzer.cc
+++ b/src/tint/fuzzers/tint_regex_fuzzer/fuzzer.cc
@@ -37,6 +37,9 @@
kReplaceIdentifier,
kReplaceLiteral,
kInsertReturnStatement,
+ kReplaceOperator,
+ kInsertBreakOrContinue,
+ kReplaceFunctionCallWithBuiltin,
kNumMutationKinds
};
@@ -64,43 +67,60 @@
MutationKind mutation_kind = static_cast<MutationKind>(
generator.GetUInt32(static_cast<uint32_t>(MutationKind::kNumMutationKinds)));
+ WgslMutator mutator(generator);
switch (mutation_kind) {
case MutationKind::kSwapIntervals:
- if (!SwapRandomIntervals(delimiter, wgsl_code, generator)) {
+ if (!mutator.SwapRandomIntervals(delimiter, wgsl_code)) {
return 0;
}
break;
case MutationKind::kDeleteInterval:
- if (!DeleteRandomInterval(delimiter, wgsl_code, generator)) {
+ if (!mutator.DeleteRandomInterval(delimiter, wgsl_code)) {
return 0;
}
break;
case MutationKind::kDuplicateInterval:
- if (!DuplicateRandomInterval(delimiter, wgsl_code, generator)) {
+ if (!mutator.DuplicateRandomInterval(delimiter, wgsl_code)) {
return 0;
}
break;
case MutationKind::kReplaceIdentifier:
- if (!ReplaceRandomIdentifier(wgsl_code, generator)) {
+ if (!mutator.ReplaceRandomIdentifier(wgsl_code)) {
return 0;
}
break;
case MutationKind::kReplaceLiteral:
- if (!ReplaceRandomIntLiteral(wgsl_code, generator)) {
+ if (!mutator.ReplaceRandomIntLiteral(wgsl_code)) {
return 0;
}
break;
case MutationKind::kInsertReturnStatement:
- if (!InsertReturnStatement(wgsl_code, generator)) {
+ if (!mutator.InsertReturnStatement(wgsl_code)) {
return 0;
}
break;
+ case MutationKind::kReplaceOperator:
+ if (!mutator.ReplaceRandomOperator(wgsl_code)) {
+ return 0;
+ }
+ break;
+
+ case MutationKind::kInsertBreakOrContinue:
+ if (!mutator.InsertBreakOrContinue(wgsl_code)) {
+ return 0;
+ }
+ break;
+ case MutationKind::kReplaceFunctionCallWithBuiltin:
+ if (!mutator.ReplaceFunctionCallWithBuiltin(wgsl_code)) {
+ return 0;
+ }
+ break;
default:
assert(false && "Unreachable");
return 0;
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/regex_fuzzer_tests.cc b/src/tint/fuzzers/tint_regex_fuzzer/regex_fuzzer_tests.cc
index 891a111..9e0b1a0 100644
--- a/src/tint/fuzzers/tint_regex_fuzzer/regex_fuzzer_tests.cc
+++ b/src/tint/fuzzers/tint_regex_fuzzer/regex_fuzzer_tests.cc
@@ -12,6 +12,7 @@
// See the License for the specific language governing permissions and
// limitations under the License.
+#include <optional>
#include <string>
#include "gtest/gtest.h"
@@ -21,26 +22,47 @@
namespace tint::fuzzers::regex_fuzzer {
namespace {
+class WgslMutatorTest : public WgslMutator {
+ public:
+ explicit WgslMutatorTest(RandomGenerator& generator) : WgslMutator(generator) {}
+
+ using WgslMutator::DeleteInterval;
+ using WgslMutator::DuplicateInterval;
+ using WgslMutator::FindClosingBrace;
+ using WgslMutator::FindOperatorOccurrence;
+ using WgslMutator::GetFunctionBodyPositions;
+ using WgslMutator::GetFunctionCallIdentifiers;
+ using WgslMutator::GetIdentifiers;
+ using WgslMutator::GetIntLiterals;
+ using WgslMutator::GetLoopBodyPositions;
+ using WgslMutator::ReplaceRegion;
+ using WgslMutator::SwapIntervals;
+};
+
// Swaps two non-consecutive regions in the edge
TEST(SwapRegionsTest, SwapIntervalsEdgeNonConsecutive) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;";
std::string all_regions = R1 + R2 + R3;
// this call should swap R1 with R3.
- SwapIntervals(0, R1.length(), R1.length() + R2.length(), R3.length(), all_regions);
+ mutator.SwapIntervals(0, R1.length(), R1.length() + R2.length(), R3.length(), all_regions);
ASSERT_EQ(R3 + R2 + R1, all_regions);
}
// Swaps two non-consecutive regions not in the edge
TEST(SwapRegionsTest, SwapIntervalsNonConsecutiveNonEdge) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// this call should swap R2 with R4.
- SwapIntervals(R1.length(), R2.length(), R1.length() + R2.length() + R3.length(), R4.length(),
- all_regions);
+ mutator.SwapIntervals(R1.length(), R2.length(), R1.length() + R2.length() + R3.length(),
+ R4.length(), all_regions);
ASSERT_EQ(R1 + R4 + R3 + R2 + R5, all_regions);
}
@@ -48,12 +70,15 @@
// Swaps two consecutive regions not in the edge (sorrounded by other
// regions)
TEST(SwapRegionsTest, SwapIntervalsConsecutiveEdge) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4;
// this call should swap R2 with R3.
- SwapIntervals(R1.length(), R2.length(), R1.length() + R2.length(), R3.length(), all_regions);
+ mutator.SwapIntervals(R1.length(), R2.length(), R1.length() + R2.length(), R3.length(),
+ all_regions);
ASSERT_EQ(R1 + R3 + R2 + R4, all_regions);
}
@@ -61,113 +86,137 @@
// Swaps two consecutive regions not in the edge (not sorrounded by other
// regions)
TEST(SwapRegionsTest, SwapIntervalsConsecutiveNonEdge) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// this call should swap R4 with R5.
- SwapIntervals(R1.length() + R2.length() + R3.length(), R4.length(),
- R1.length() + R2.length() + R3.length() + R4.length(), R5.length(), all_regions);
+ mutator.SwapIntervals(R1.length() + R2.length() + R3.length(), R4.length(),
+ R1.length() + R2.length() + R3.length() + R4.length(), R5.length(),
+ all_regions);
ASSERT_EQ(R1 + R2 + R3 + R5 + R4, all_regions);
}
// Deletes the first region.
TEST(DeleteRegionTest, DeleteFirstRegion) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// This call should delete R1.
- DeleteInterval(0, R1.length(), all_regions);
+ mutator.DeleteInterval(0, R1.length(), all_regions);
ASSERT_EQ(";" + R2 + R3 + R4 + R5, all_regions);
}
// Deletes the last region.
TEST(DeleteRegionTest, DeleteLastRegion) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// This call should delete R5.
- DeleteInterval(R1.length() + R2.length() + R3.length() + R4.length(), R5.length(), all_regions);
+ mutator.DeleteInterval(R1.length() + R2.length() + R3.length() + R4.length(), R5.length(),
+ all_regions);
ASSERT_EQ(R1 + R2 + R3 + R4 + ";", all_regions);
}
// Deletes the middle region.
TEST(DeleteRegionTest, DeleteMiddleRegion) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// This call should delete R3.
- DeleteInterval(R1.length() + R2.length(), R3.length(), all_regions);
+ mutator.DeleteInterval(R1.length() + R2.length(), R3.length(), all_regions);
ASSERT_EQ(R1 + R2 + ";" + R4 + R5, all_regions);
}
TEST(InsertRegionTest, InsertRegionTest1) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// This call should insert R2 after R4.
- DuplicateInterval(R1.length(), R2.length(),
- R1.length() + R2.length() + R3.length() + R4.length() - 1, all_regions);
+ mutator.DuplicateInterval(R1.length(), R2.length(),
+ R1.length() + R2.length() + R3.length() + R4.length() - 1,
+ all_regions);
ASSERT_EQ(R1 + R2 + R3 + R4 + R2.substr(1, R2.size() - 1) + R5, all_regions);
}
TEST(InsertRegionTest, InsertRegionTest2) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// This call should insert R3 after R1.
- DuplicateInterval(R1.length() + R2.length(), R3.length(), R1.length() - 1, all_regions);
+ mutator.DuplicateInterval(R1.length() + R2.length(), R3.length(), R1.length() - 1, all_regions);
ASSERT_EQ(R1 + R3.substr(1, R3.length() - 1) + R2 + R3 + R4 + R5, all_regions);
}
TEST(InsertRegionTest, InsertRegionTest3) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = ";region1;", R2 = ";regionregion2;", R3 = ";regionregionregion3;",
R4 = ";regionregionregionregion4;", R5 = ";regionregionregionregionregion5;";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// This call should insert R2 after R5.
- DuplicateInterval(R1.length(), R2.length(), all_regions.length() - 1, all_regions);
+ mutator.DuplicateInterval(R1.length(), R2.length(), all_regions.length() - 1, all_regions);
ASSERT_EQ(R1 + R2 + R3 + R4 + R5 + R2.substr(1, R2.length() - 1), all_regions);
}
TEST(ReplaceIdentifierTest, ReplaceIdentifierTest1) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = "|region1|", R2 = "; region2;", R3 = "---------region3---------",
R4 = "++region4++", R5 = "***region5***";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// Replaces R3 with R1.
- ReplaceRegion(0, R1.length(), R1.length() + R2.length(), R3.length(), all_regions);
+ mutator.ReplaceRegion(0, R1.length(), R1.length() + R2.length(), R3.length(), all_regions);
ASSERT_EQ(R1 + R2 + R1 + R4 + R5, all_regions);
}
TEST(ReplaceIdentifierTest, ReplaceIdentifierTest2) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string R1 = "|region1|", R2 = "; region2;", R3 = "---------region3---------",
R4 = "++region4++", R5 = "***region5***";
std::string all_regions = R1 + R2 + R3 + R4 + R5;
// Replaces R5 with R3.
- ReplaceRegion(R1.length() + R2.length(), R3.length(),
- R1.length() + R2.length() + R3.length() + R4.length(), R5.length(), all_regions);
+ mutator.ReplaceRegion(R1.length() + R2.length(), R3.length(),
+ R1.length() + R2.length() + R3.length() + R4.length(), R5.length(),
+ all_regions);
ASSERT_EQ(R1 + R2 + R3 + R4 + R3, all_regions);
}
TEST(GetIdentifierTest, GetIdentifierTest1) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string wgsl_code =
R"(fn clamp_0acf8f() {
var res: vec2<f32> = clamp(vec2<f32>(), vec2<f32>(), vec2<f32>());
@@ -187,23 +236,19 @@
clamp_0acf8f();
})";
- std::vector<std::pair<size_t, size_t>> identifiers_pos = GetIdentifiers(wgsl_code);
-
+ std::vector<std::pair<size_t, size_t>> identifiers_pos = mutator.GetIdentifiers(wgsl_code);
std::vector<std::pair<size_t, size_t>> ground_truth = {
- std::make_pair(3, 12), std::make_pair(28, 3), std::make_pair(37, 4),
- std::make_pair(49, 5), std::make_pair(60, 3), std::make_pair(68, 4),
- std::make_pair(81, 4), std::make_pair(110, 6), std::make_pair(123, 2),
- std::make_pair(133, 4), std::make_pair(144, 7), std::make_pair(162, 4),
- std::make_pair(183, 12), std::make_pair(209, 6), std::make_pair(221, 3),
- std::make_pair(244, 8), std::make_pair(259, 2), std::make_pair(271, 4),
- std::make_pair(288, 12), std::make_pair(319, 7), std::make_pair(328, 14),
- std::make_pair(352, 2), std::make_pair(363, 4), std::make_pair(381, 3),
- std::make_pair(394, 3), std::make_pair(399, 3), std::make_pair(418, 12)};
-
+ {0, 2}, {3, 12}, {28, 3}, {32, 3}, {37, 4}, {42, 3}, {49, 5}, {55, 4},
+ {60, 3}, {68, 4}, {73, 3}, {81, 4}, {86, 3}, {110, 6}, {123, 2}, {126, 11},
+ {144, 7}, {152, 8}, {162, 4}, {167, 3}, {183, 12}, {209, 6}, {216, 4}, {221, 3},
+ {244, 8}, {259, 2}, {262, 13}, {288, 12}, {319, 7}, {328, 14}, {352, 2}, {355, 12},
+ {381, 3}, {385, 7}, {394, 3}, {399, 3}, {418, 12}};
ASSERT_EQ(ground_truth, identifiers_pos);
}
TEST(TestGetLiteralsValues, TestGetLiteralsValues1) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string wgsl_code =
R"(fn clamp_0acf8f() {
var res: vec2<f32> = clamp(vec2<f32>(), vec2<f32>(), vec2<f32>());
@@ -227,7 +272,7 @@
foo_1 = 5 + 7;
var foo_3 : i32 = -20;)";
- std::vector<std::pair<size_t, size_t>> literals_pos = GetIntLiterals(wgsl_code);
+ std::vector<std::pair<size_t, size_t>> literals_pos = mutator.GetIntLiterals(wgsl_code);
std::vector<std::string> ground_truth = {"3", "10", "5", "7", "-20"};
@@ -241,6 +286,8 @@
}
TEST(InsertReturnTest, FindClosingBrace) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string wgsl_code =
R"(fn clamp_0acf8f() {
if(false){
@@ -269,7 +316,7 @@
var foo_3 : i32 = -20;
)";
size_t opening_bracket_pos = 18;
- size_t closing_bracket_pos = FindClosingBrace(opening_bracket_pos, wgsl_code);
+ size_t closing_bracket_pos = mutator.FindClosingBrace(opening_bracket_pos, wgsl_code);
// The -1 is needed since the function body starts after the left bracket.
std::string function_body =
@@ -286,6 +333,8 @@
}
TEST(InsertReturnTest, FindClosingBraceFailing) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string wgsl_code =
R"(fn clamp_0acf8f() {
// This comment } causes the test to fail.
@@ -314,7 +363,7 @@
foo_1 = 5 + 7;
var foo_3 : i32 = -20;)";
size_t opening_bracket_pos = 18;
- size_t closing_bracket_pos = FindClosingBrace(opening_bracket_pos, wgsl_code);
+ size_t closing_bracket_pos = mutator.FindClosingBrace(opening_bracket_pos, wgsl_code);
// The -1 is needed since the function body starts after the left bracket.
std::string function_body =
@@ -329,67 +378,9 @@
ASSERT_NE(expected, function_body);
}
-TEST(TestInsertReturn, TestInsertReturn1) {
- std::string wgsl_code =
- R"(fn clamp_0acf8f() {
- var res: vec2<f32> = clamp(vec2<f32>(), vec2<f32>(), vec2<f32>());
- }
- @vertex
- fn vertex_main() -> @builtin(position) vec4<f32> {
- clamp_0acf8f();
- var foo_1: i32 = 3;
- return vec4<f32>();
- }
- @fragment
- fn fragment_main() {
- clamp_0acf8f();
- }
- @compute @workgroup_size(1)
- fn compute_main() {
- var<private> foo: f32 = 0.0;
- var foo_2: i32 = 10;
- clamp_0acf8f();
- }
- foo_1 = 5 + 7;
- var foo_3 : i32 = -20;)";
-
- std::vector<size_t> semicolon_pos;
- for (size_t pos = wgsl_code.find(";", 0); pos != std::string::npos;
- pos = wgsl_code.find(";", pos + 1)) {
- semicolon_pos.push_back(pos);
- }
-
- // should insert a return true statement after the first semicolon of the
- // first function the the WGSL-like string above.
- wgsl_code.insert(semicolon_pos[0] + 1, "return true;");
-
- std::string expected_wgsl_code =
- R"(fn clamp_0acf8f() {
- var res: vec2<f32> = clamp(vec2<f32>(), vec2<f32>(), vec2<f32>());return true;
- }
- @vertex
- fn vertex_main() -> @builtin(position) vec4<f32> {
- clamp_0acf8f();
- var foo_1: i32 = 3;
- return vec4<f32>();
- }
- @fragment
- fn fragment_main() {
- clamp_0acf8f();
- }
- @compute @workgroup_size(1)
- fn compute_main() {
- var<private> foo: f32 = 0.0;
- var foo_2: i32 = 10;
- clamp_0acf8f();
- }
- foo_1 = 5 + 7;
- var foo_3 : i32 = -20;)";
-
- ASSERT_EQ(expected_wgsl_code, wgsl_code);
-}
-
-TEST(TestInsertReturn, TestFunctionPositions) {
+TEST(TestInsertReturn, TestFunctionPositions1) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string wgsl_code =
R"(fn clamp_0acf8f() {
var res: vec2<f32> = clamp(vec2<f32>(), vec2<f32>(), vec2<f32>());
@@ -418,12 +409,38 @@
foo_1 = 5 + 7;
var foo_3 : i32 = -20;)";
- std::vector<size_t> function_positions = GetFunctionBodyPositions(wgsl_code);
- std::vector<size_t> expected_positions = {180, 586};
+ std::vector<std::pair<size_t, bool>> function_positions =
+ mutator.GetFunctionBodyPositions(wgsl_code);
+ std::vector<std::pair<size_t, bool>> expected_positions = {
+ {18, false}, {180, true}, {323, false}, {423, false}, {586, true}};
+ ASSERT_EQ(expected_positions, function_positions);
+}
+
+TEST(TestInsertReturn, TestFunctionPositions2) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
+ std::string wgsl_code =
+ R"(fn some_loop_body() {
+}
+
+fn f() {
+ var j : i32; i = (i + 1)) {
+ some_loop_body(); ((i < 5) && (j < 10));
+ for(var i : i32 = 0;
+ j = (i * 30);
+ }
+}
+)";
+
+ std::vector<std::pair<size_t, bool>> function_positions =
+ mutator.GetFunctionBodyPositions(wgsl_code);
+ std::vector<std::pair<size_t, bool>> expected_positions = {{20, false}, {32, false}};
ASSERT_EQ(expected_positions, function_positions);
}
TEST(TestInsertReturn, TestMissingSemicolon) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
std::string wgsl_code =
R"(fn clamp_0acf8f() {
var res: vec2<f32> = clamp(vec2<f32>(), vec2<f32>(), vec2<f32>())
@@ -452,8 +469,7 @@
foo_1 = 5 + 7;
var foo_3 : i32 = -20;)";
- RandomGenerator generator(0);
- InsertReturnStatement(wgsl_code, generator);
+ mutator.InsertReturnStatement(wgsl_code);
// No semicolons found in the function's body, so wgsl_code
// should remain unchanged.
@@ -487,5 +503,141 @@
ASSERT_EQ(expected_wgsl_code, wgsl_code);
}
+TEST(TestReplaceOperator, TestIdentifyOperators) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
+ std::string code =
+ R"(
+x += 2;
+y = a + b;
+z = -a;
+x *= b / c;
+t = t && t | t || t;
+b = b > c ^ c <= d;
+a >>= b;
+b <<= a;
+a = a << 2;
+b = b >> 3;
+c = a % 3;
+d %= e;
+)";
+ // These are the operator occurrences that will be observed by going through the file character
+ // by character. This includes, for example, identifying the ">" operator if search starts after
+ // the first character of ">>".
+ std::vector<std::pair<uint32_t, uint32_t>> operator_occurrences = {
+ {3, 2}, {4, 1}, {11, 1}, {15, 1}, {22, 1}, {24, 1}, {30, 2}, {31, 1},
+ {35, 1}, {42, 1}, {46, 2}, {47, 1}, {51, 1}, {55, 2}, {56, 1}, {63, 1},
+ {67, 1}, {71, 1}, {75, 2}, {76, 1}, {83, 3}, {84, 2}, {85, 1}, {92, 3},
+ {93, 2}, {94, 1}, {101, 1}, {105, 2}, {106, 1}, {113, 1}, {117, 2}, {118, 1},
+ {125, 1}, {129, 1}, {136, 2}, {137, 1}, {3, 2}};
+ uint32_t operator_occurrence_index = 0;
+ for (size_t i = 0; i < code.length(); i++) {
+ // Move on to the next operator occurrence if the current index into the code string exceeds
+ // the index associated with that operator occurrence. Exception: stay with the last
+ // operator occurrence if search has already passed the last operator in the file.
+ if (i < code.length() - 2 && i > operator_occurrences[operator_occurrence_index].first) {
+ operator_occurrence_index =
+ (operator_occurrence_index + 1) % operator_occurrences.size();
+ }
+ ASSERT_EQ(operator_occurrences[operator_occurrence_index],
+ mutator.FindOperatorOccurrence(code, static_cast<uint32_t>(i)).value());
+ }
+}
+
+TEST(TestInsertBreakOrContinue, TestLoopPositions1) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
+ std::string wgsl_code = " loop { } loop { } loop { }";
+ std::vector<size_t> loop_positions = mutator.GetLoopBodyPositions(wgsl_code);
+ std::vector<size_t> expected_positions = {6, 15, 24};
+ ASSERT_EQ(expected_positions, loop_positions);
+}
+
+TEST(TestInsertBreakOrContinue, TestLoopPositions2) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
+ std::string wgsl_code = R"( loop { } loop
+{ } loop { })";
+ std::vector<size_t> loop_positions = mutator.GetLoopBodyPositions(wgsl_code);
+ std::vector<size_t> expected_positions = {6, 15, 24};
+ ASSERT_EQ(expected_positions, loop_positions);
+}
+
+TEST(TestInsertBreakOrContinue, TestLoopPositions3) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
+ // This WGSL-like code is not valid, but it suffices to test regex-based matching (which is
+ // intended to work well on semi-valid code).
+ std::string wgsl_code =
+ R"(fn compute_main() {
+ loop {
+ var twice: i32 = 2 * i;
+ i++;
+ if i == 5 { break; }
+ loop
+ {
+ var twice: i32 = 2 * i;
+ i++;
+ while (i < 100) { i++; }
+ if i == 5 { break; }
+ }
+ }
+ for (a = 0; a < 100; a++) {
+ if (a > 50) {
+ break;
+ }
+ while (i < 100) { i++; }
+ }
+})";
+
+ std::vector<size_t> loop_positions = mutator.GetLoopBodyPositions(wgsl_code);
+ std::vector<size_t> expected_positions = {27, 108, 173, 249, 310};
+ ASSERT_EQ(expected_positions, loop_positions);
+}
+
+TEST(TestInsertBreakOrContinue, TestLoopPositions4) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
+ std::string wgsl_code =
+ R"(fn clamp_0acf8f() {
+ var res: vec2<f32> = clamp(vec2<f32>(), vec2<f32>(), vec2<f32>());
+ }
+ @vertex
+ fn vertex_main() -> @builtin(position) vec4<f32> {
+ clamp_0acf8f();"
+ return vec4<f32>();
+ }
+ @fragment
+ fn fragment_main() {
+ clamp_0acf8f();
+ }
+ @compute @workgroup_size(1)
+ fn compute_main() {"
+ var<private> foo: f32 = 0.0;
+ clamp_0acf8f ();
+ })";
+
+ std::vector<size_t> loop_positions = mutator.GetLoopBodyPositions(wgsl_code);
+ ASSERT_TRUE(loop_positions.empty());
+}
+
+TEST(TestReplaceFunctionCallWithBuiltin, FindFunctionCalls) {
+ RandomGenerator generator(0);
+ WgslMutatorTest mutator(generator);
+ std::string function_body = R"({
+ var<private> foo: f32 = 0.0;
+ var foo_2: i32 = 10;
+ clamp_0acf8f ();
+ _0acf8f();
+ f
+();
+ j = (i * 30);
+ })";
+ std::vector<std::pair<size_t, size_t>> call_identifiers =
+ mutator.GetFunctionCallIdentifiers(function_body);
+ std::vector<std::pair<size_t, size_t>> ground_truth{{82, 12}, {110, 7}, {131, 1}};
+ ASSERT_EQ(ground_truth, call_identifiers);
+}
+
} // namespace
} // namespace tint::fuzzers::regex_fuzzer
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
index 677560b..a965613 100644
--- a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
+++ b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
@@ -26,8 +26,10 @@
namespace tint::fuzzers::regex_fuzzer {
-std::vector<size_t> FindDelimiterIndices(const std::string& delimiter,
- const std::string& wgsl_code) {
+WgslMutator::WgslMutator(RandomGenerator& generator) : generator_(generator) {}
+
+std::vector<size_t> WgslMutator::FindDelimiterIndices(const std::string& delimiter,
+ const std::string& wgsl_code) {
std::vector<size_t> result;
for (size_t pos = wgsl_code.find(delimiter, 0); pos != std::string::npos;
pos = wgsl_code.find(delimiter, pos + 1)) {
@@ -37,7 +39,7 @@
return result;
}
-std::vector<std::pair<size_t, size_t>> GetIdentifiers(const std::string& wgsl_code) {
+std::vector<std::pair<size_t, size_t>> WgslMutator::GetIdentifiers(const std::string& wgsl_code) {
std::vector<std::pair<size_t, size_t>> result;
// This regular expression works by looking for a character that
@@ -45,23 +47,37 @@
// by a character which cannot be part of a WGSL identifer. The regex
// for the WGSL identifier is obtained from:
// https://www.w3.org/TR/WGSL/#identifiers.
- std::regex wgsl_identifier_regex("[^a-zA-Z]([a-zA-Z][0-9a-zA-Z_]*)[^0-9a-zA-Z_]");
+ std::regex identifier_regex("[_a-zA-Z][0-9a-zA-Z_]*");
- std::smatch match;
+ auto identifiers_begin =
+ std::sregex_iterator(wgsl_code.begin(), wgsl_code.end(), identifier_regex);
+ auto identifiers_end = std::sregex_iterator();
- std::string::const_iterator search_start(wgsl_code.cbegin());
- std::string prefix;
-
- while (regex_search(search_start, wgsl_code.cend(), match, wgsl_identifier_regex) == true) {
- prefix += match.prefix();
- result.push_back(std::make_pair(prefix.size() + 1, match.str(1).size()));
- prefix += match.str(0);
- search_start = match.suffix().first;
+ for (std::sregex_iterator i = identifiers_begin; i != identifiers_end; ++i) {
+ result.push_back(
+ {static_cast<size_t>(i->prefix().second - wgsl_code.cbegin()), i->str().size()});
}
return result;
}
-std::vector<std::pair<size_t, size_t>> GetIntLiterals(const std::string& s) {
+std::vector<std::pair<size_t, size_t>> WgslMutator::GetFunctionCallIdentifiers(
+ const std::string& wgsl_code) {
+ std::vector<std::pair<size_t, size_t>> result;
+
+ std::regex call_regex("([_a-zA-Z][0-9a-zA-Z_]*)[ \\n]*\\(");
+
+ auto identifiers_begin = std::sregex_iterator(wgsl_code.begin(), wgsl_code.end(), call_regex);
+ auto identifiers_end = std::sregex_iterator();
+
+ for (std::sregex_iterator i = identifiers_begin; i != identifiers_end; ++i) {
+ auto submatch = (*i)[1];
+ result.push_back(
+ {static_cast<size_t>(submatch.first - wgsl_code.cbegin()), submatch.str().size()});
+ }
+ return result;
+}
+
+std::vector<std::pair<size_t, size_t>> WgslMutator::GetIntLiterals(const std::string& s) {
std::vector<std::pair<size_t, size_t>> result;
// Looks for integer literals in decimal or hexadecimal form.
@@ -83,7 +99,7 @@
return result;
}
-size_t FindClosingBrace(size_t opening_bracket_pos, const std::string& wgsl_code) {
+size_t WgslMutator::FindClosingBrace(size_t opening_bracket_pos, const std::string& wgsl_code) {
size_t open_bracket_count = 1;
size_t pos = opening_bracket_pos + 1;
while (open_bracket_count >= 1 && pos < wgsl_code.size()) {
@@ -97,33 +113,52 @@
return (pos == wgsl_code.size() && open_bracket_count >= 1) ? 0 : pos - 1;
}
-std::vector<size_t> GetFunctionBodyPositions(const std::string& wgsl_code) {
+std::vector<std::pair<size_t, bool>> WgslMutator::GetFunctionBodyPositions(
+ const std::string& wgsl_code) {
// Finds all the functions with a non-void return value.
- std::regex function_regex("fn.*?->.*?\\{");
- std::smatch match;
- std::vector<size_t> result;
+ std::regex function_regex("fn[^a-zA-Z_0-9][^\\{]*\\{");
+ std::vector<std::pair<size_t, bool>> result;
- auto search_start(wgsl_code.cbegin());
- std::string prefix = "";
+ auto functions_begin = std::sregex_iterator(wgsl_code.begin(), wgsl_code.end(), function_regex);
+ auto functions_end = std::sregex_iterator();
- while (std::regex_search(search_start, wgsl_code.cend(), match, function_regex)) {
- result.push_back(static_cast<size_t>(match.suffix().first - wgsl_code.cbegin() - 1L));
- search_start = match.suffix().first;
+ for (std::sregex_iterator i = functions_begin; i != functions_end; ++i) {
+ bool returns_value = i->str().find("->") != std::string::npos;
+ result.push_back(
+ {static_cast<size_t>(i->suffix().first - wgsl_code.cbegin() - 1), returns_value});
}
return result;
}
-bool InsertReturnStatement(std::string& wgsl_code, RandomGenerator& generator) {
- std::vector<size_t> function_body_positions = GetFunctionBodyPositions(wgsl_code);
+std::vector<size_t> WgslMutator::GetLoopBodyPositions(const std::string& wgsl_code) {
+ // Finds all loops.
+ std::regex loop_regex("[^a-zA-Z_0-9](for|while|loop)[^\\{]*\\{");
+ std::vector<size_t> result;
+
+ auto loops_begin = std::sregex_iterator(wgsl_code.begin(), wgsl_code.end(), loop_regex);
+ auto loops_end = std::sregex_iterator();
+
+ for (std::sregex_iterator i = loops_begin; i != loops_end; ++i) {
+ result.push_back(static_cast<size_t>(i->suffix().first - wgsl_code.cbegin() - 1));
+ }
+ return result;
+}
+
+bool WgslMutator::InsertReturnStatement(std::string& wgsl_code) {
+ std::vector<std::pair<size_t, bool>> function_body_positions =
+ GetFunctionBodyPositions(wgsl_code);
// No function was found in wgsl_code.
if (function_body_positions.empty()) {
return false;
}
- // Pick a random function's opening bracket, find the corresponding closing
- // bracket, and find a semi-colon within the function body.
- size_t left_bracket_pos = generator.GetRandomElement(function_body_positions);
+ // Pick a random function
+ auto function = generator_.GetRandomElement(function_body_positions);
+
+ // Find the corresponding closing bracket for the function, and find a semi-colon within the
+ // function body.
+ size_t left_bracket_pos = function.first;
size_t right_bracket_pos = FindClosingBrace(left_bracket_pos, wgsl_code);
@@ -141,27 +176,64 @@
return false;
}
- size_t semicolon_position = generator.GetRandomElement(semicolon_positions);
-
- // Get all identifiers and integer literals to use as potential return values.
- std::vector<std::pair<size_t, size_t>> identifiers = GetIdentifiers(wgsl_code);
- auto return_values = identifiers;
- std::vector<std::pair<size_t, size_t>> int_literals = GetIntLiterals(wgsl_code);
- return_values.insert(return_values.end(), int_literals.begin(), int_literals.end());
- std::pair<size_t, size_t> return_value = generator.GetRandomElement(return_values);
- std::string return_statement =
- "return " + wgsl_code.substr(return_value.first, return_value.second) + ";";
+ std::string return_statement = "return";
+ if (function.second) {
+ // The function returns a value. Get all identifiers and integer literals to use as
+ // potential return values.
+ std::vector<std::pair<size_t, size_t>> identifiers = GetIdentifiers(wgsl_code);
+ auto return_values = identifiers;
+ std::vector<std::pair<size_t, size_t>> int_literals = GetIntLiterals(wgsl_code);
+ return_values.insert(return_values.end(), int_literals.begin(), int_literals.end());
+ std::pair<size_t, size_t> return_value = generator_.GetRandomElement(return_values);
+ return_statement += " " + wgsl_code.substr(return_value.first, return_value.second);
+ }
+ return_statement += ";";
// Insert the return statement immediately after the semicolon.
- wgsl_code.insert(semicolon_position + 1, return_statement);
+ wgsl_code.insert(generator_.GetRandomElement(semicolon_positions) + 1, return_statement);
return true;
}
-void SwapIntervals(size_t idx1,
- size_t reg1_len,
- size_t idx2,
- size_t reg2_len,
- std::string& wgsl_code) {
+bool WgslMutator::InsertBreakOrContinue(std::string& wgsl_code) {
+ std::vector<size_t> loop_body_positions = GetLoopBodyPositions(wgsl_code);
+
+ // No loop was found in wgsl_code.
+ if (loop_body_positions.empty()) {
+ return false;
+ }
+
+ // Pick a random loop's opening bracket, find the corresponding closing
+ // bracket, and find a semi-colon within the loop body.
+ size_t left_bracket_pos = generator_.GetRandomElement(loop_body_positions);
+
+ size_t right_bracket_pos = FindClosingBrace(left_bracket_pos, wgsl_code);
+
+ if (right_bracket_pos == 0) {
+ return false;
+ }
+
+ std::vector<size_t> semicolon_positions;
+ for (size_t pos = wgsl_code.find(";", left_bracket_pos + 1); pos < right_bracket_pos;
+ pos = wgsl_code.find(";", pos + 1)) {
+ semicolon_positions.push_back(pos);
+ }
+
+ if (semicolon_positions.empty()) {
+ return false;
+ }
+
+ size_t semicolon_position = generator_.GetRandomElement(semicolon_positions);
+
+ // Insert a break or continue immediately after the semicolon.
+ wgsl_code.insert(semicolon_position + 1, generator_.GetBool() ? "break;" : "continue;");
+ return true;
+}
+
+void WgslMutator::SwapIntervals(size_t idx1,
+ size_t reg1_len,
+ size_t idx2,
+ size_t reg2_len,
+ std::string& wgsl_code) {
std::string region_1 = wgsl_code.substr(idx1 + 1, reg1_len - 1);
std::string region_2 = wgsl_code.substr(idx2 + 1, reg2_len - 1);
@@ -172,36 +244,37 @@
wgsl_code.replace(idx1 + 1, region_1.size(), region_2);
}
-void DeleteInterval(size_t idx1, size_t reg_len, std::string& wgsl_code) {
+void WgslMutator::DeleteInterval(size_t idx1, size_t reg_len, std::string& wgsl_code) {
wgsl_code.erase(idx1 + 1, reg_len - 1);
}
-void DuplicateInterval(size_t idx1, size_t reg1_len, size_t idx2, std::string& wgsl_code) {
+void WgslMutator::DuplicateInterval(size_t idx1,
+ size_t reg1_len,
+ size_t idx2,
+ std::string& wgsl_code) {
std::string region = wgsl_code.substr(idx1 + 1, reg1_len - 1);
wgsl_code.insert(idx2 + 1, region);
}
-void ReplaceRegion(size_t idx1,
- size_t id1_len,
- size_t idx2,
- size_t id2_len,
- std::string& wgsl_code) {
+void WgslMutator::ReplaceRegion(size_t idx1,
+ size_t id1_len,
+ size_t idx2,
+ size_t id2_len,
+ std::string& wgsl_code) {
std::string region_1 = wgsl_code.substr(idx1, id1_len);
std::string region_2 = wgsl_code.substr(idx2, id2_len);
wgsl_code.replace(idx2, region_2.size(), region_1);
}
-void ReplaceInterval(size_t start_index,
- size_t length,
- std::string replacement_text,
- std::string& wgsl_code) {
+void WgslMutator::ReplaceInterval(size_t start_index,
+ size_t length,
+ std::string replacement_text,
+ std::string& wgsl_code) {
std::string region_1 = wgsl_code.substr(start_index, length);
wgsl_code.replace(start_index, length, replacement_text);
}
-bool SwapRandomIntervals(const std::string& delimiter,
- std::string& wgsl_code,
- RandomGenerator& generator) {
+bool WgslMutator::SwapRandomIntervals(const std::string& delimiter, std::string& wgsl_code) {
std::vector<size_t> delimiter_positions = FindDelimiterIndices(delimiter, wgsl_code);
// Need to have at least 3 indices.
@@ -212,12 +285,12 @@
// Choose indices:
// interval_1_start < interval_1_end <= interval_2_start < interval_2_end
uint32_t interval_1_start =
- generator.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()) - 2u);
- uint32_t interval_1_end = generator.GetUInt32(
+ generator_.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()) - 2u);
+ uint32_t interval_1_end = generator_.GetUInt32(
interval_1_start + 1u, static_cast<uint32_t>(delimiter_positions.size()) - 1u);
- uint32_t interval_2_start =
- generator.GetUInt32(interval_1_end, static_cast<uint32_t>(delimiter_positions.size()) - 1u);
- uint32_t interval_2_end = generator.GetUInt32(
+ uint32_t interval_2_start = generator_.GetUInt32(
+ interval_1_end, static_cast<uint32_t>(delimiter_positions.size()) - 1u);
+ uint32_t interval_2_end = generator_.GetUInt32(
interval_2_start + 1u, static_cast<uint32_t>(delimiter_positions.size()));
SwapIntervals(delimiter_positions[interval_1_start],
@@ -229,9 +302,7 @@
return true;
}
-bool DeleteRandomInterval(const std::string& delimiter,
- std::string& wgsl_code,
- RandomGenerator& generator) {
+bool WgslMutator::DeleteRandomInterval(const std::string& delimiter, std::string& wgsl_code) {
std::vector<size_t> delimiter_positions = FindDelimiterIndices(delimiter, wgsl_code);
// Need to have at least 2 indices.
@@ -240,9 +311,9 @@
}
uint32_t interval_start =
- generator.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()) - 1u);
- uint32_t interval_end =
- generator.GetUInt32(interval_start + 1u, static_cast<uint32_t>(delimiter_positions.size()));
+ generator_.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()) - 1u);
+ uint32_t interval_end = generator_.GetUInt32(interval_start + 1u,
+ static_cast<uint32_t>(delimiter_positions.size()));
DeleteInterval(delimiter_positions[interval_start],
delimiter_positions[interval_end] - delimiter_positions[interval_start],
@@ -251,9 +322,7 @@
return true;
}
-bool DuplicateRandomInterval(const std::string& delimiter,
- std::string& wgsl_code,
- RandomGenerator& generator) {
+bool WgslMutator::DuplicateRandomInterval(const std::string& delimiter, std::string& wgsl_code) {
std::vector<size_t> delimiter_positions = FindDelimiterIndices(delimiter, wgsl_code);
// Need to have at least 2 indices
@@ -262,11 +331,11 @@
}
uint32_t interval_start =
- generator.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()) - 1u);
- uint32_t interval_end =
- generator.GetUInt32(interval_start + 1u, static_cast<uint32_t>(delimiter_positions.size()));
+ generator_.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()) - 1u);
+ uint32_t interval_end = generator_.GetUInt32(interval_start + 1u,
+ static_cast<uint32_t>(delimiter_positions.size()));
uint32_t duplication_point =
- generator.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()));
+ generator_.GetUInt32(static_cast<uint32_t>(delimiter_positions.size()));
DuplicateInterval(delimiter_positions[interval_start],
delimiter_positions[interval_end] - delimiter_positions[interval_start],
@@ -275,7 +344,7 @@
return true;
}
-bool ReplaceRandomIdentifier(std::string& wgsl_code, RandomGenerator& generator) {
+bool WgslMutator::ReplaceRandomIdentifier(std::string& wgsl_code) {
std::vector<std::pair<size_t, size_t>> identifiers = GetIdentifiers(wgsl_code);
// Need at least 2 identifiers
@@ -283,12 +352,12 @@
return false;
}
- uint32_t id1_index = generator.GetUInt32(static_cast<uint32_t>(identifiers.size()));
- uint32_t id2_index = generator.GetUInt32(static_cast<uint32_t>(identifiers.size()));
+ uint32_t id1_index = generator_.GetUInt32(static_cast<uint32_t>(identifiers.size()));
+ uint32_t id2_index = generator_.GetUInt32(static_cast<uint32_t>(identifiers.size()));
// The two identifiers must be different
while (id1_index == id2_index) {
- id2_index = generator.GetUInt32(static_cast<uint32_t>(identifiers.size()));
+ id2_index = generator_.GetUInt32(static_cast<uint32_t>(identifiers.size()));
}
ReplaceRegion(identifiers[id1_index].first, identifiers[id1_index].second,
@@ -297,7 +366,7 @@
return true;
}
-bool ReplaceRandomIntLiteral(std::string& wgsl_code, RandomGenerator& generator) {
+bool WgslMutator::ReplaceRandomIntLiteral(std::string& wgsl_code) {
std::vector<std::pair<size_t, size_t>> literals = GetIntLiterals(wgsl_code);
// Need at least one integer literal
@@ -305,13 +374,13 @@
return false;
}
- uint32_t literal_index = generator.GetUInt32(static_cast<uint32_t>(literals.size()));
+ uint32_t literal_index = generator_.GetUInt32(static_cast<uint32_t>(literals.size()));
// INT_MAX = 2147483647, INT_MIN = -2147483648
std::vector<std::string> boundary_values = {"2147483647", "-2147483648", "1",
"-1", "0", "4294967295"};
- uint32_t boundary_index = generator.GetUInt32(static_cast<uint32_t>(boundary_values.size()));
+ uint32_t boundary_index = generator_.GetUInt32(static_cast<uint32_t>(boundary_values.size()));
ReplaceInterval(literals[literal_index].first, literals[literal_index].second,
boundary_values[boundary_index], wgsl_code);
@@ -319,4 +388,268 @@
return true;
}
+std::string WgslMutator::ChooseRandomReplacementForOperator(const std::string& existing_operator) {
+ // Operators are partitioned into three classes: assignment, expression and increment. The regex
+ // mutator will swap operators in the same class. The hypothesis is that this should exercise a
+ // number of type-safe swaps (e.g. changing += to *=), as well as some badly-typed yet
+ // interesting swaps (e.g. changing + to ^ when the operators are matrices), while avoiding
+ // making totally nonsensical replacements (such as changing ++ too /).
+ std::vector<std::string> assignment_operators{
+ "=", "+=", "-=", "*=", "/=", "%=", "&=", "|=", "^=", "<<=", ">>="};
+ std::vector<std::string> expression_operators{"+", "-", "*", "/", "%", "&&", "||",
+ "&", "|", "^", "<<", ">>", "<", ">",
+ "<=", ">=", "!", "!=", "~"};
+ std::vector<std::string> increment_operators{"++", "--"};
+ for (auto operators : {assignment_operators, expression_operators, increment_operators}) {
+ auto it = std::find(operators.begin(), operators.end(), existing_operator);
+ if (it != operators.end()) {
+ // The operator falls into this category, so select another operator from the category.
+ operators.erase(it);
+ return generator_.GetRandomElement(operators);
+ }
+ }
+ assert(false && "Unknown operator");
+ return "";
+}
+
+bool WgslMutator::ReplaceRandomOperator(std::string& wgsl_code) {
+ // Choose an index into the code at random.
+ const uint32_t start_index = generator_.GetUInt32(static_cast<uint32_t>(wgsl_code.size()));
+ // Find the first operator occurrence from the chosen point, wrapping back to the start of the
+ // file if needed.
+ auto maybe_operator_occurrence = FindOperatorOccurrence(wgsl_code, start_index);
+ if (!maybe_operator_occurrence.has_value()) {
+ // It is unlikely that there will be *no* operators in the file, but if this is the case
+ // then this mutation cannot be applied.
+ return false;
+ }
+ std::string existing_operator =
+ wgsl_code.substr(maybe_operator_occurrence->first, maybe_operator_occurrence->second);
+ // Replace the identified operator with a randomly-chosen alternative.
+ wgsl_code.replace(maybe_operator_occurrence->first, maybe_operator_occurrence->second,
+ ChooseRandomReplacementForOperator(existing_operator));
+ return true;
+}
+
+std::optional<std::pair<uint32_t, uint32_t>> WgslMutator::FindOperatorOccurrence(
+ const std::string& wgsl_code,
+ uint32_t start_index) {
+ // Loops through the characters of the code in a wrap-around fashion, looking for the first
+ // encountered token that is a WGSL operator.
+
+ for (size_t i = 0; i < wgsl_code.size(); i++) {
+ uint32_t current_index = static_cast<uint32_t>((start_index + i) % wgsl_code.size());
+
+ // To cater for multi-character operator tokens, get the three consecutive characters from
+ // the code string starting at the current index. Use null characters to account for the
+ // case where search has reached the end of the code string.
+ char first_character = wgsl_code[current_index];
+ char second_character =
+ current_index == wgsl_code.size() - 1 ? '\0' : wgsl_code[current_index + 1];
+ char third_character =
+ current_index >= wgsl_code.size() - 2 ? '\0' : wgsl_code[current_index + 2];
+
+ // This uses the extracted characters to match for the various WGSL operators.
+ switch (first_character) {
+ case '!':
+ case '^':
+ switch (second_character) {
+ case '=':
+ return {{current_index, 2}};
+ default:
+ return {{current_index, 1}};
+ }
+ case '|':
+ case '&':
+ case '+':
+ case '-':
+ if (second_character == first_character || second_character == '=') {
+ return {{current_index, 2}};
+ }
+ return {{current_index, 1}};
+ case '*':
+ case '/':
+ case '%':
+ switch (second_character) {
+ case '=':
+ return {{current_index, 2}};
+ default:
+ return {{current_index, 1}};
+ }
+ case '=':
+ if (second_character == '=') {
+ return {{current_index, 2}};
+ }
+ return {{current_index, 1}};
+ case '<':
+ case '>':
+ if (second_character == '=') {
+ return {{current_index, 2}};
+ }
+ if (second_character == first_character) {
+ if (third_character == '=') {
+ return {{current_index, 3}};
+ }
+ return {{current_index, 2}};
+ }
+ return {{current_index, 1}};
+ case '~':
+ return {{current_index, 1}};
+ default:
+ break;
+ }
+ }
+ // No operator was found, so empty is returned.
+ return {};
+}
+
+bool WgslMutator::ReplaceFunctionCallWithBuiltin(std::string& wgsl_code) {
+ std::vector<std::pair<size_t, bool>> function_body_positions =
+ GetFunctionBodyPositions(wgsl_code);
+
+ // No function was found in wgsl_code.
+ if (function_body_positions.empty()) {
+ return false;
+ }
+
+ // Pick a random function
+ auto function = generator_.GetRandomElement(function_body_positions);
+
+ // Find the corresponding closing bracket for the function, and find a semi-colon within the
+ // function body.
+ size_t left_bracket_pos = function.first;
+
+ size_t right_bracket_pos = FindClosingBrace(left_bracket_pos, wgsl_code);
+
+ if (right_bracket_pos == 0) {
+ return false;
+ }
+
+ std::string function_body(
+ wgsl_code.substr(left_bracket_pos, right_bracket_pos - left_bracket_pos));
+
+ std::vector<std::pair<size_t, size_t>> function_call_identifiers =
+ GetFunctionCallIdentifiers(function_body);
+ if (function_call_identifiers.empty()) {
+ return false;
+ }
+ auto function_call_identifier = generator_.GetRandomElement(function_call_identifiers);
+
+ std::vector<std::string> builtin_functions{"all",
+ "any",
+ "select",
+ "arrayLength",
+ "abs",
+ "acos",
+ "acosh",
+ "asin",
+ "asinh",
+ "atan",
+ "atanh",
+ "atan2",
+ "ceil",
+ "clamp",
+ "cos",
+ "cosh",
+ "cross",
+ "degrees",
+ "distance",
+ "exp",
+ "exp2",
+ "faceForward",
+ "floor",
+ "fma",
+ "fract",
+ "frexp",
+ "inverseSqrt",
+ "ldexp",
+ "length",
+ "log",
+ "log2",
+ "max",
+ "min",
+ "mix",
+ "modf",
+ "normalize",
+ "pow",
+ "quantizeToF16",
+ "radians",
+ "reflect",
+ "refract",
+ "round",
+ "saturate",
+ "sign",
+ "sin",
+ "sinh",
+ "smoothstep",
+ "sqrt",
+ "step",
+ "tan",
+ "tanh",
+ "trunc",
+ "abs",
+ "clamp",
+ "countLeadingZeros",
+ "countOneBits",
+ "countTrailingZeros",
+ "extractBits",
+ "firstLeadingBit",
+ "firstTrailingBit",
+ "insertBits",
+ "max",
+ "min",
+ "reverseBits",
+ "determinant",
+ "transpose",
+ "dot",
+ "dpdx",
+ "dpdxCoarse",
+ "dpdxFine",
+ "dpdy",
+ "dpdyCoarse",
+ "dpdyFine",
+ "fwidth",
+ "fwidthCoarse",
+ "fwidthFine",
+ "textureDimensions",
+ "textureGather",
+ "textureGatherCompare",
+ "textureLoad",
+ "textureNumLayers",
+ "textureNumLevels",
+ "textureNumSamples",
+ "textureSample",
+ "textureSampleBias",
+ "textureSampleCompare",
+ "textureSampleCompareLevel",
+ "textureSampleGrad",
+ "textureSampleLevel",
+ "textureStore",
+ "atomicLoad",
+ "atomicStore",
+ "atomicAdd",
+ "atomicSub",
+ "atomicMax",
+ "atomicMin",
+ "atomicAnd",
+ "atomicOr",
+ "atomicXor",
+ "pack4x8snorm",
+ "pack4x8unorm",
+ "pack2x16snorm",
+ "pack2x16unorm",
+ "pack2x16float",
+ "unpack4x8snorm",
+ "unpack4x8unorm",
+ "unpack2x16snorm",
+ "unpack2x16unorm",
+ "unpack2x16float",
+ "storageBarrier",
+ "workgroupBarrier"};
+ wgsl_code.replace(left_bracket_pos + function_call_identifier.first,
+ function_call_identifier.second,
+ generator_.GetRandomElement(builtin_functions));
+ return true;
+}
+
} // namespace tint::fuzzers::regex_fuzzer
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.h b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.h
index 23c45bb..fde4611 100644
--- a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.h
+++ b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.h
@@ -15,6 +15,7 @@
#ifndef SRC_TINT_FUZZERS_TINT_REGEX_FUZZER_WGSL_MUTATOR_H_
#define SRC_TINT_FUZZERS_TINT_REGEX_FUZZER_WGSL_MUTATOR_H_
+#include <optional>
#include <string>
#include <utility>
#include <vector>
@@ -23,152 +24,207 @@
namespace tint::fuzzers::regex_fuzzer {
-/// A function that given a delimiter, returns a vector that contains
-/// all the positions of the delimiter in the WGSL code.
-/// @param delimiter - the delimiter of the enclosed region.
-/// @param wgsl_code - the initial string (WGSL code) that will be mutated.
-/// @return a vector with the positions of the delimiter in the WGSL code.
-std::vector<size_t> FindDelimiterIndices(const std::string& delimiter,
- const std::string& wgsl_code);
+/// Class encapsulating code for regex-based mutation of WGSL shaders.
+class WgslMutator {
+ public:
+ /// Constructor
+ /// @param generator - pseudo-random generator to use in mutator
+ explicit WgslMutator(RandomGenerator& generator);
-/// A function that finds all the identifiers in a WGSL-like string.
-/// @param wgsl_code - the WGSL-like string where the identifiers will be found.
-/// @return a vector with the positions and the length of all the
-/// identifiers in wgsl_code.
-std::vector<std::pair<size_t, size_t>> GetIdentifiers(const std::string& wgsl_code);
+ /// A function that, given WGSL-like string and a delimiter,
+ /// generates another WGSL-like string by picking two random regions
+ /// enclosed by the delimiter and swapping them.
+ /// @param delimiter - the delimiter that will be used to find enclosed regions.
+ /// @param wgsl_code - the initial string (WGSL code) that will be mutated.
+ /// @return true if a swap happened or false otherwise.
+ bool SwapRandomIntervals(const std::string& delimiter, std::string& wgsl_code);
-/// A function that returns returns the starting position
-/// and the length of all the integer literals in a WGSL-like string.
-/// @param wgsl_code - the WGSL-like string where the int literals
-/// will be found.
-/// @return a vector with the starting positions and the length
-/// of all the integer literals.
-std::vector<std::pair<size_t, size_t>> GetIntLiterals(const std::string& wgsl_code);
+ /// A function that, given a WGSL-like string and a delimiter,
+ /// generates another WGSL-like string by deleting a random
+ /// region enclosed by the delimiter.
+ /// @param delimiter - the delimiter that will be used to find enclosed regions.
+ /// @param wgsl_code - the initial string (WGSL code) that will be mutated.
+ /// @return true if a deletion happened or false otherwise.
+ bool DeleteRandomInterval(const std::string& delimiter, std::string& wgsl_code);
-/// Finds a possible closing brace corresponding to the opening
-/// brace at position opening_bracket_pos.
-/// @param opening_bracket_pos - the position of the opening brace.
-/// @param wgsl_code - the WGSL-like string where the closing brace.
-/// @return the position of the closing bracket or 0 if there is no closing
-/// brace.
-size_t FindClosingBrace(size_t opening_bracket_pos, const std::string& wgsl_code);
+ /// A function that, given a WGSL-like string and a delimiter,
+ /// generates another WGSL-like string by duplicating a random
+ /// region enclosed by the delimiter.
+ /// @param delimiter - the delimiter that will be used to find enclosed regions.
+ /// @param wgsl_code - the initial string (WGSL code) that will be mutated.
+ /// @return true if a duplication happened or false otherwise.
+ bool DuplicateRandomInterval(const std::string& delimiter, std::string& wgsl_code);
-/// Returns the starting_position of the bodies of the functions
-/// that follow the regular expression: fn.*?->.*?\\{, which searches for the
-/// keyword fn followed by the function name, its return type and opening brace.
-/// @param wgsl_code - the WGSL-like string where the functions will be
-/// searched.
-/// @return a vector with the starting position of the function bodies in
-/// wgsl_code.
-std::vector<size_t> GetFunctionBodyPositions(const std::string& wgsl_code);
+ /// Replaces a randomly-chosen identifier in wgsl_code.
+ /// @param wgsl_code - WGSL-like string where the replacement will occur.
+ /// @return true if a replacement happened or false otherwise.
+ bool ReplaceRandomIdentifier(std::string& wgsl_code);
-/// Given 4 indices, idx1, idx2, idx3 and idx4 it swaps the regions
-/// in the interval (idx1, idx2] with the region in the interval (idx3, idx4]
-/// in wgsl_text.
-/// @param idx1 - starting index of the first region.
-/// @param reg1_len - length of the first region.
-/// @param idx2 - starting index of the second region.
-/// @param reg2_len - length of the second region.
-/// @param wgsl_code - the string where the swap will occur.
-void SwapIntervals(size_t idx1,
- size_t reg1_len,
- size_t idx2,
- size_t reg2_len,
- std::string& wgsl_code);
+ /// Replaces the value of a randomly-chosen integer with one of
+ /// the values in the set {INT_MAX, INT_MIN, 0, -1}.
+ /// @param wgsl_code - WGSL-like string where the replacement will occur.
+ /// @return true if a replacement happened or false otherwise.
+ bool ReplaceRandomIntLiteral(std::string& wgsl_code);
-/// Given index idx1 it delets the region of length interval_len
-/// starting at index idx1;
-/// @param idx1 - starting index of the first region.
-/// @param reg_len - terminating index of the second region.
-/// @param wgsl_code - the string where the swap will occur.
-void DeleteInterval(size_t idx1, size_t reg_len, std::string& wgsl_code);
+ /// Inserts a return statement in a randomly chosen function of a
+ /// WGSL-like string. The return value is a randomly-chosen identifier
+ /// or literal in the string.
+ /// @param wgsl_code - WGSL-like string that will be mutated.
+ /// @return true if the mutation was succesful or false otherwise.
+ bool InsertReturnStatement(std::string& wgsl_code);
-/// Given 2 indices, idx1, idx2, it inserts the region of length
-/// reg1_len starting at idx1 after idx2.
-/// @param idx1 - starting index of region.
-/// @param reg1_len - length of the region.
-/// @param idx2 - the position where the region will be inserted.
-/// @param wgsl_code - the string where the swap will occur.
-void DuplicateInterval(size_t idx1, size_t reg1_len, size_t idx2, std::string& wgsl_code);
+ /// Inserts a break or continue statement in a randomly chosen loop of a WGSL-like string.
+ /// @param wgsl_code - WGSL-like string that will be mutated.
+ /// @return true if the mutation was succesful or false otherwise.
+ bool InsertBreakOrContinue(std::string& wgsl_code);
-/// Replaces a region of a WGSL-like string of length id2_len starting
-/// at position idx2 with a region of length id1_len starting at
-/// position idx1.
-/// @param idx1 - starting position of the first region.
-/// @param id1_len - length of the first region.
-/// @param idx2 - starting position of the second region.
-/// @param id2_len - length of the second region.
-/// @param wgsl_code - the string where the replacement will occur.
-void ReplaceRegion(size_t idx1,
- size_t id1_len,
- size_t idx2,
- size_t id2_len,
- std::string& wgsl_code);
+ /// A function that, given WGSL-like string, generates a new WGSL-like string by replacing one
+ /// randomly-chosen operator in the original string with another operator.
+ /// @param wgsl_code - the initial WGSL-like string that will be mutated.
+ /// @return true if an operator replacement happened or false otherwise.
+ bool ReplaceRandomOperator(std::string& wgsl_code);
-/// Replaces an interval of length `length` starting at start_index
-/// with the `replacement_text`.
-/// @param start_index - starting position of the interval to be replaced.
-/// @param length - length of the interval to be replaced.
-/// @param replacement_text - the interval that will be used as a replacement.
-/// @param wgsl_code - the WGSL-like string where the replacement will occur.
-void ReplaceInterval(size_t start_index,
- size_t length,
- std::string replacement_text,
- std::string& wgsl_code);
+ /// Given a WGSL-like string, replaces a random identifier that appears to be a function call
+ /// with the name of a built-in function. This will often lead to an invalid module, as the
+ /// mutation does not aim to check whether the original and replacement function have the same
+ /// number or types of arguments.
+ /// @param wgsl_code - the initial WGSL-like string that will be mutated.
+ /// @return true if a function call replacement happened or false otherwise.
+ bool ReplaceFunctionCallWithBuiltin(std::string& wgsl_code);
-/// A function that, given WGSL-like string and a delimiter,
-/// generates another WGSL-like string by picking two random regions
-/// enclosed by the delimiter and swapping them.
-/// @param delimiter - the delimiter that will be used to find enclosed regions.
-/// @param wgsl_code - the initial string (WGSL code) that will be mutated.
-/// @param generator - the random number generator.
-/// @return true if a swap happened or false otherwise.
-bool SwapRandomIntervals(const std::string& delimiter,
- std::string& wgsl_code,
- RandomGenerator& generator);
+ protected:
+ /// Given index idx1 it delets the region of length interval_len
+ /// starting at index idx1;
+ /// @param idx1 - starting index of the first region.
+ /// @param reg_len - terminating index of the second region.
+ /// @param wgsl_code - the string where the swap will occur.
+ void DeleteInterval(size_t idx1, size_t reg_len, std::string& wgsl_code);
-/// A function that, given a WGSL-like string and a delimiter,
-/// generates another WGSL-like string by deleting a random
-/// region enclosed by the delimiter.
-/// @param delimiter - the delimiter that will be used to find enclosed regions.
-/// @param wgsl_code - the initial string (WGSL code) that will be mutated.
-/// @param generator - the random number generator.
-/// @return true if a deletion happened or false otherwise.
-bool DeleteRandomInterval(const std::string& delimiter,
- std::string& wgsl_code,
- RandomGenerator& generator);
+ /// Given 2 indices, idx1, idx2, it inserts the region of length
+ /// reg1_len starting at idx1 after idx2.
+ /// @param idx1 - starting index of region.
+ /// @param reg1_len - length of the region.
+ /// @param idx2 - the position where the region will be inserted.
+ /// @param wgsl_code - the string where the swap will occur.
+ void DuplicateInterval(size_t idx1, size_t reg1_len, size_t idx2, std::string& wgsl_code);
-/// A function that, given a WGSL-like string and a delimiter,
-/// generates another WGSL-like string by duplicating a random
-/// region enclosed by the delimiter.
-/// @param delimiter - the delimiter that will be used to find enclosed regions.
-/// @param wgsl_code - the initial string (WGSL code) that will be mutated.
-/// @param generator - the random number generator.
-/// @return true if a duplication happened or false otherwise.
-bool DuplicateRandomInterval(const std::string& delimiter,
- std::string& wgsl_code,
- RandomGenerator& generator);
+ /// Finds a possible closing brace corresponding to the opening
+ /// brace at position opening_bracket_pos.
+ /// @param opening_bracket_pos - the position of the opening brace.
+ /// @param wgsl_code - the WGSL-like string where the closing brace.
+ /// @return the position of the closing bracket or 0 if there is no closing
+ /// brace.
+ size_t FindClosingBrace(size_t opening_bracket_pos, const std::string& wgsl_code);
-/// Replaces a randomly-chosen identifier in wgsl_code.
-/// @param wgsl_code - WGSL-like string where the replacement will occur.
-/// @param generator - the random number generator.
-/// @return true if a replacement happened or false otherwise.
-bool ReplaceRandomIdentifier(std::string& wgsl_code, RandomGenerator& generator);
+ /// Returns the starting position of the bodies of the functions identified by an appropriate
+ /// function, together with a boolean indicating whether the function returns a value or not.
+ /// @param wgsl_code - the WGSL-like string where the functions will be
+ /// searched.
+ /// @return a vector of pairs, where each pair provides the starting position of the function
+ /// body, and the value true if and only if the function returns a value.
+ std::vector<std::pair<size_t, bool>> GetFunctionBodyPositions(const std::string& wgsl_code);
-/// Replaces the value of a randomly-chosen integer with one of
-/// the values in the set {INT_MAX, INT_MIN, 0, -1}.
-/// @param wgsl_code - WGSL-like string where the replacement will occur.
-/// @param generator - the random number generator.
-/// @return true if a replacement happened or false otherwise.
-bool ReplaceRandomIntLiteral(std::string& wgsl_code, RandomGenerator& generator);
+ /// Returns the starting position of the bodies of the loops identified by an appropriate
+ /// regular expressions.
+ /// @param wgsl_code - the WGSL-like string in which loops will be searched for.
+ /// @return a vector with the starting position of the loop bodies in wgsl_code.
+ std::vector<size_t> GetLoopBodyPositions(const std::string& wgsl_code);
-/// Inserts a return statement in a randomly chosen function of a
-/// WGSL-like string. The return value is a randomly-chosen identifier
-/// or literal in the string.
-/// @param wgsl_code - WGSL-like string that will be mutated.
-/// @param generator - the random number generator.
-/// @return true if the mutation was succesful or false otherwise.
-bool InsertReturnStatement(std::string& wgsl_code, RandomGenerator& generator);
+ /// A function that finds all the identifiers in a WGSL-like string.
+ /// @param wgsl_code - the WGSL-like string where the identifiers will be found.
+ /// @return a vector with the positions and the length of all the
+ /// identifiers in wgsl_code.
+ std::vector<std::pair<size_t, size_t>> GetIdentifiers(const std::string& wgsl_code);
+
+ /// A function that finds the identifiers in a WGSL-like string that appear to be used as
+ /// function names in function call expressions.
+ /// @param wgsl_code - the WGSL-like string where the identifiers will be found.
+ /// @return a vector with the positions and the length of all the
+ /// identifiers in wgsl_code.
+ std::vector<std::pair<size_t, size_t>> GetFunctionCallIdentifiers(const std::string& wgsl_code);
+
+ /// A function that returns returns the starting position
+ /// and the length of all the integer literals in a WGSL-like string.
+ /// @param wgsl_code - the WGSL-like string where the int literals
+ /// will be found.
+ /// @return a vector with the starting positions and the length
+ /// of all the integer literals.
+ std::vector<std::pair<size_t, size_t>> GetIntLiterals(const std::string& wgsl_code);
+
+ /// Replaces a region of a WGSL-like string of length id2_len starting
+ /// at position idx2 with a region of length id1_len starting at
+ /// position idx1.
+ /// @param idx1 - starting position of the first region.
+ /// @param id1_len - length of the first region.
+ /// @param idx2 - starting position of the second region.
+ /// @param id2_len - length of the second region.
+ /// @param wgsl_code - the string where the replacement will occur.
+ void ReplaceRegion(size_t idx1,
+ size_t id1_len,
+ size_t idx2,
+ size_t id2_len,
+ std::string& wgsl_code);
+
+ /// Given 4 indices, idx1, idx2, idx3 and idx4 it swaps the regions
+ /// in the interval (idx1, idx2] with the region in the interval (idx3, idx4]
+ /// in wgsl_text.
+ /// @param idx1 - starting index of the first region.
+ /// @param reg1_len - length of the first region.
+ /// @param idx2 - starting index of the second region.
+ /// @param reg2_len - length of the second region.
+ /// @param wgsl_code - the string where the swap will occur.
+ void SwapIntervals(size_t idx1,
+ size_t reg1_len,
+ size_t idx2,
+ size_t reg2_len,
+ std::string& wgsl_code);
+
+ /// Finds the next occurrence of an operator in a WGSL-like string from a given starting
+ /// position, wrapping around to the start of the string if no operator is found before reaching
+ /// the end, and returning empty of no operator is found at all. There is no guarantee that the
+ /// result will correspond to a WGSL operator token, e.g. the identified characters could be
+ /// part of a comment, or e.g. the file might contain >>=, in which case the operator
+ /// >= will be identified should it happen that the starting index corresponds to the second >
+ /// character of this operator. Given that the regex mutator does not aim to guarantee
+ /// well-formed programs, these issues are acceptable.
+ /// @param wgsl_code - the WGSL-like string in which operator occurrences will be found.
+ /// @param start_index - the index at which search should start
+ /// @return empty if no operator is found, otherwise a pair comprising the index at which the
+ /// operator starts and the character length of the operator.
+ std::optional<std::pair<uint32_t, uint32_t>> FindOperatorOccurrence(
+ const std::string& wgsl_code,
+ uint32_t start_index);
+
+ private:
+ /// A function that given a delimiter, returns a vector that contains
+ /// all the positions of the delimiter in the WGSL code.
+ /// @param delimiter - the delimiter of the enclosed region.
+ /// @param wgsl_code - the initial string (WGSL code) that will be mutated.
+ /// @return a vector with the positions of the delimiter in the WGSL code.
+ std::vector<size_t> FindDelimiterIndices(const std::string& delimiter,
+ const std::string& wgsl_code);
+
+ /// Replaces an interval of length `length` starting at start_index
+ /// with the `replacement_text`.
+ /// @param start_index - starting position of the interval to be replaced.
+ /// @param length - length of the interval to be replaced.
+ /// @param replacement_text - the interval that will be used as a replacement.
+ /// @param wgsl_code - the WGSL-like string where the replacement will occur.
+ void ReplaceInterval(size_t start_index,
+ size_t length,
+ std::string replacement_text,
+ std::string& wgsl_code);
+
+ /// Given a string representing a WGSL operator, randomly returns a different WGSL operator in
+ /// the same category as the original, where the three categories are assignment operators (such
+ /// as = and +=), expression operators (such as + and ^) and increment operators (++ and --).
+ /// @param existing_operator - the characters comprising some WGSL operator
+ /// @return another WGSL operator falling into the same category.
+ std::string ChooseRandomReplacementForOperator(const std::string& existing_operator);
+
+ RandomGenerator& generator_;
+};
+
} // namespace tint::fuzzers::regex_fuzzer
#endif // SRC_TINT_FUZZERS_TINT_REGEX_FUZZER_WGSL_MUTATOR_H_
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index ab77532..16f2ff4 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -134,6 +134,7 @@
match fi32f16: f32 | f16 | i32
match iu32: i32 | u32
match aiu32: ai | i32 | u32
+match afi32f16: ai | af | f32 | i32 | f16
match scalar: f32 | f16 | i32 | u32 | bool
match abstract_or_scalar: ai | af | f32 | f16 | i32 | u32 | bool
match af_f32: af | f32
@@ -822,8 +823,8 @@
@const op ~ <T: aiu32>(T) -> T
@const op ~ <T: aiu32, N: num> (vec<N, T>) -> vec<N, T>
-op - <T: fi32f16>(T) -> T
-op - <T: fi32f16, N: num> (vec<N, T>) -> vec<N, T>
+@const op - <T: afi32f16>(T) -> T
+@const op - <T: afi32f16, N: num> (vec<N, T>) -> vec<N, T>
////////////////////////////////////////////////////////////////////////////////
// Binary Operators //
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index 4a8c097..e2b16b5 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -18,6 +18,7 @@
#include <limits>
#include <optional>
#include <string>
+#include <type_traits>
#include <unordered_map>
#include <utility>
@@ -45,14 +46,38 @@
namespace {
-/// TypeDispatch is a helper for calling the function `f`, passing a single zero-value argument of
-/// the C++ type that corresponds to the sem::Type `type`. For example, calling `TypeDispatch()`
-/// with a type of `sem::I32*` will call the function f with a single argument of `i32(0)`.
+/// Helper that calls 'f' passing in `c`'s value
+template <typename F>
+auto aiu32Dispatch(const sem::Constant* c, F&& f) {
+ return Switch(
+ c->Type(), [&](const sem::AbstractInt*) { return f(c->As<AInt>()); },
+ [&](const sem::I32*) { return f(c->As<i32>()); },
+ [&](const sem::U32*) { return f(c->As<u32>()); });
+}
+
+/// Helper that calls 'f' passing in `c`'s value
+template <typename F>
+auto afi32f16Dispatch(const sem::Constant* c, F&& f) {
+ return Switch(
+ c->Type(), [&](const sem::AbstractInt*) { return f(c->As<AInt>()); },
+ [&](const sem::AbstractFloat*) { return f(c->As<AFloat>()); },
+ [&](const sem::F32*) { return f(c->As<f32>()); },
+ [&](const sem::I32*) { return f(c->As<i32>()); },
+ [&](const sem::F16*) {
+ // TODO(crbug.com/tint/1502): Support const eval for f16
+ return nullptr;
+ });
+}
+
+/// ZeroTypeDispatch is a helper for calling the function `f`, passing a single zero-value argument
+/// of the C++ type that corresponds to the sem::Type `type`. For example, calling
+/// `ZeroTypeDispatch()` with a type of `sem::I32*` will call the function f with a single argument
+/// of `i32(0)`.
/// @returns the value returned by calling `f`.
/// @note `type` must be a scalar or abstract numeric type. Other types will not call `f`, and will
/// return the zero-initialized value of the return type for `f`.
template <typename F>
-auto TypeDispatch(const sem::Type* type, F&& f) {
+auto ZeroTypeDispatch(const sem::Type* type, F&& f) {
return Switch(
type, //
[&](const sem::AbstractInt*) { return f(AInt(0)); }, //
@@ -64,20 +89,6 @@
[&](const sem::Bool*) { return f(static_cast<bool>(0)); });
}
-/// IntegerDispatch is a helper for calling the function `f`, passing the integer value of the
-/// constant c.
-/// @returns the value returned by calling `f`.
-/// @note `c` must be of an integer type. Other types will not call `f`, and will return the
-/// zero-initialized value of the return type for `f`
-template <typename F>
-auto IntegerDispatch(const sem::Constant* c, F&& f) {
- return Switch(
- c->Type(), //
- [&](const sem::AbstractInt*) { return f(c->As<AInt>()); }, //
- [&](const sem::I32*) { return f(c->As<i32>()); }, //
- [&](const sem::U32*) { return f(c->As<u32>()); });
-}
-
/// @returns `value` if `T` is not a Number, otherwise ValueOf returns the inner value of the
/// Number.
template <typename T>
@@ -142,7 +153,7 @@
return this;
}
bool failed = false;
- auto* res = TypeDispatch(target_ty, [&](auto zero_to) -> const Constant* {
+ auto* res = ZeroTypeDispatch(target_ty, [&](auto zero_to) -> const Constant* {
// `T` is the source type, `value` is the source value.
// `TO` is the target type.
using TO = std::decay_t<decltype(zero_to)>;
@@ -333,7 +344,7 @@
return CreateComposite(builder, s, std::move(zeros));
},
[&](Default) -> const Constant* {
- return TypeDispatch(type, [&](auto zero) -> const Constant* {
+ return ZeroTypeDispatch(type, [&](auto zero) -> const Constant* {
return CreateElement(builder, type, zero);
});
});
@@ -643,12 +654,34 @@
sem::Expression const* const* args,
size_t) {
return TransformElements(builder, args[0]->ConstantValue(), [&](const sem::Constant* c) {
- return IntegerDispatch(c, [&](auto i) { //
+ return aiu32Dispatch(c, [&](auto i) { //
return CreateElement(builder, c->Type(), decltype(i)(~i.value));
});
});
}
+const sem::Constant* ConstEval::OpMinus(const sem::Type*,
+ sem::Expression const* const* args,
+ size_t) {
+ return TransformElements(builder, args[0]->ConstantValue(), [&](const sem::Constant* c) {
+ return afi32f16Dispatch(c, [&](auto i) { //
+ // For signed integrals, avoid C++ UB by not negating the smallest negative number. In
+ // WGSL, this operation is well defined to return the same value, see:
+ // https://gpuweb.github.io/gpuweb/wgsl/#arithmetic-expr.
+ using T = UnwrapNumber<decltype(i)>;
+ if constexpr (std::is_integral_v<T>) {
+ auto v = i.value;
+ if (v != std::numeric_limits<T>::min()) {
+ v = -v;
+ }
+ return CreateElement(builder, c->Type(), decltype(i)(v));
+ } else {
+ return CreateElement(builder, c->Type(), decltype(i)(-i.value));
+ }
+ });
+ });
+}
+
utils::Result<const sem::Constant*> ConstEval::Convert(const sem::Type* target_ty,
const sem::Constant* value,
const Source& source) {
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index b1e7352..6a85c85 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -187,6 +187,15 @@
sem::Expression const* const* args,
size_t num_args);
+ /// Minus operator '-'
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param num_args the number of input arguments (must be 1)
+ /// @return the result value, or null if the value cannot be calculated
+ const sem::Constant* OpMinus(const sem::Type* ty,
+ sem::Expression const* const* args,
+ size_t num_args);
+
private:
/// Adds the given error message to the diagnostics
void AddError(const std::string& msg, const Source& source) const;
diff --git a/src/tint/resolver/const_eval_test.cc b/src/tint/resolver/const_eval_test.cc
index dcc1604..7991f6f 100644
--- a/src/tint/resolver/const_eval_test.cc
+++ b/src/tint/resolver/const_eval_test.cc
@@ -13,6 +13,7 @@
// limitations under the License.
#include <cmath>
+#include <type_traits>
#include "gtest/gtest.h"
#include "src/tint/resolver/resolver_test_helper.h"
@@ -2930,13 +2931,36 @@
namespace unary_op {
template <typename T>
+auto Highest() {
+ return T(T::kHighest);
+}
+
+template <typename T>
+auto Lowest() {
+ return T(T::kLowest);
+}
+
+template <typename T>
+constexpr auto Negate(const Number<T>& v) {
+ // For signed integrals, avoid C++ UB by not negating the smallest negative number. In
+ // WGSL, this operation is well defined to return the same value, see:
+ // https://gpuweb.github.io/gpuweb/wgsl/#arithmetic-expr.
+ if constexpr (std::is_integral_v<T> && std::is_signed_v<T>) {
+ if (v == std::numeric_limits<T>::min()) {
+ return v;
+ }
+ }
+ return -v;
+}
+
+template <typename T>
struct Values {
T input;
T expect;
};
struct Case {
- std::variant<Values<AInt>, Values<u32>, Values<i32>> values;
+ std::variant<Values<AInt>, Values<AFloat>, Values<u32>, Values<i32>, Values<f32>> values;
};
static std::ostream& operator<<(std::ostream& o, const Case& c) {
@@ -2952,6 +2976,8 @@
using ResolverConstEvalUnaryOpTest = ResolverTestWithParam<std::tuple<ast::UnaryOp, Case>>;
TEST_P(ResolverConstEvalUnaryOpTest, Test) {
+ Enable(ast::Extension::kF16);
+
auto op = std::get<0>(GetParam());
auto c = std::get<1>(GetParam());
std::visit(
@@ -3000,6 +3026,60 @@
C(2_i, -3_i),
C(-3_i, 2_i),
})));
+
+INSTANTIATE_TEST_SUITE_P(Negation,
+ ResolverConstEvalUnaryOpTest,
+ testing::Combine(testing::Values(ast::UnaryOp::kNegation),
+ testing::ValuesIn({
+ // AInt
+ C(0_a, -0_a),
+ C(-0_a, 0_a),
+ C(1_a, -1_a),
+ C(-1_a, 1_a),
+ C(Highest<AInt>(), -Highest<AInt>()),
+ C(-Highest<AInt>(), Highest<AInt>()),
+ C(Lowest<AInt>(), Negate(Lowest<AInt>())),
+ C(Negate(Lowest<AInt>()), Lowest<AInt>()),
+ // i32
+ C(0_i, -0_i),
+ C(-0_i, 0_i),
+ C(1_i, -1_i),
+ C(-1_i, 1_i),
+ C(Highest<i32>(), -Highest<i32>()),
+ C(-Highest<i32>(), Highest<i32>()),
+ C(Lowest<i32>(), Negate(Lowest<i32>())),
+ C(Negate(Lowest<i32>()), Lowest<i32>()),
+ // AFloat
+ C(0.0_a, -0.0_a),
+ C(-0.0_a, 0.0_a),
+ C(1.0_a, -1.0_a),
+ C(-1.0_a, 1.0_a),
+ C(Highest<AFloat>(), -Highest<AFloat>()),
+ C(-Highest<AFloat>(), Highest<AFloat>()),
+ C(Lowest<AFloat>(), Negate(Lowest<AFloat>())),
+ C(Negate(Lowest<AFloat>()), Lowest<AFloat>()),
+ // f32
+ C(0.0_f, -0.0_f),
+ C(-0.0_f, 0.0_f),
+ C(1.0_f, -1.0_f),
+ C(-1.0_f, 1.0_f),
+ C(Highest<f32>(), -Highest<f32>()),
+ C(-Highest<f32>(), Highest<f32>()),
+ C(Lowest<f32>(), Negate(Lowest<f32>())),
+ C(Negate(Lowest<f32>()), Lowest<f32>()),
+ })));
+
+// Make sure UBSan doesn't trip on C++'s undefined behaviour of negating the smallest negative
+// number.
+TEST_F(ResolverConstEvalTest, UnaryNegateLowestAbstract) {
+ // const break_me = -(-9223372036854775808);
+ auto* c = GlobalConst("break_me", nullptr, Negation(Negation(Expr(9223372036854775808_a))));
+ (void)c;
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+ auto* sem = Sem().Get(c);
+ EXPECT_EQ(sem->ConstantValue()->As<AInt>(), 9223372036854775808_a);
+}
+
} // namespace unary_op
} // namespace
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index e7b44a3..811aa34 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -1832,8 +1832,52 @@
return ss.str();
}
-/// TypeMatcher for 'match scalar'
+/// TypeMatcher for 'match afi32f16'
/// @see src/tint/intrinsics.def:137:7
+class Afi32F16 : public TypeMatcher {
+ public:
+ /// Checks whether the given type matches the matcher rules, and returns the
+ /// expected, canonicalized type on success.
+ /// Match may define and refine the template types and numbers in state.
+ /// @param state the MatchState
+ /// @param type the type to match
+ /// @returns the canonicalized type on match, otherwise nullptr
+ const sem::Type* Match(MatchState& state,
+ const sem::Type* type) const override;
+ /// @param state the MatchState
+ /// @return a string representation of the matcher.
+ std::string String(MatchState* state) const override;
+};
+
+const sem::Type* Afi32F16::Match(MatchState& state, const sem::Type* ty) const {
+ if (match_af(ty)) {
+ return build_af(state);
+ }
+ if (match_ai(ty)) {
+ return build_ai(state);
+ }
+ if (match_i32(ty)) {
+ return build_i32(state);
+ }
+ if (match_f32(ty)) {
+ return build_f32(state);
+ }
+ if (match_f16(ty)) {
+ return build_f16(state);
+ }
+ return nullptr;
+}
+
+std::string Afi32F16::String(MatchState*) const {
+ std::stringstream ss;
+ // Note: We pass nullptr to the TypeMatcher::String() functions, as 'matcher's do not support
+ // template arguments, nor can they match sub-types. As such, they have no use for the MatchState.
+ ss << Ai().String(nullptr) << ", " << Af().String(nullptr) << ", " << F32().String(nullptr) << ", " << I32().String(nullptr) << " or " << F16().String(nullptr);
+ return ss.str();
+}
+
+/// TypeMatcher for 'match scalar'
+/// @see src/tint/intrinsics.def:138:7
class Scalar : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -1877,7 +1921,7 @@
}
/// TypeMatcher for 'match abstract_or_scalar'
-/// @see src/tint/intrinsics.def:138:7
+/// @see src/tint/intrinsics.def:139:7
class AbstractOrScalar : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -1927,7 +1971,7 @@
}
/// TypeMatcher for 'match af_f32'
-/// @see src/tint/intrinsics.def:139:7
+/// @see src/tint/intrinsics.def:140:7
class AfF32 : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -1962,7 +2006,7 @@
}
/// TypeMatcher for 'match af_f32f16'
-/// @see src/tint/intrinsics.def:140:7
+/// @see src/tint/intrinsics.def:141:7
class AfF32F16 : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -2000,7 +2044,7 @@
}
/// TypeMatcher for 'match scalar_no_f32'
-/// @see src/tint/intrinsics.def:141:7
+/// @see src/tint/intrinsics.def:142:7
class ScalarNoF32 : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -2041,7 +2085,7 @@
}
/// TypeMatcher for 'match scalar_no_f16'
-/// @see src/tint/intrinsics.def:142:7
+/// @see src/tint/intrinsics.def:143:7
class ScalarNoF16 : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -2082,7 +2126,7 @@
}
/// TypeMatcher for 'match scalar_no_i32'
-/// @see src/tint/intrinsics.def:143:7
+/// @see src/tint/intrinsics.def:144:7
class ScalarNoI32 : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -2123,7 +2167,7 @@
}
/// TypeMatcher for 'match scalar_no_u32'
-/// @see src/tint/intrinsics.def:144:7
+/// @see src/tint/intrinsics.def:145:7
class ScalarNoU32 : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -2164,7 +2208,7 @@
}
/// TypeMatcher for 'match scalar_no_bool'
-/// @see src/tint/intrinsics.def:145:7
+/// @see src/tint/intrinsics.def:146:7
class ScalarNoBool : public TypeMatcher {
public:
/// Checks whether the given type matches the matcher rules, and returns the
@@ -2205,7 +2249,7 @@
}
/// EnumMatcher for 'match f32_texel_format'
-/// @see src/tint/intrinsics.def:156:7
+/// @see src/tint/intrinsics.def:157:7
class F32TexelFormat : public NumberMatcher {
public:
/// Checks whether the given number matches the enum matcher rules.
@@ -2238,7 +2282,7 @@
}
/// EnumMatcher for 'match i32_texel_format'
-/// @see src/tint/intrinsics.def:158:7
+/// @see src/tint/intrinsics.def:159:7
class I32TexelFormat : public NumberMatcher {
public:
/// Checks whether the given number matches the enum matcher rules.
@@ -2270,7 +2314,7 @@
}
/// EnumMatcher for 'match u32_texel_format'
-/// @see src/tint/intrinsics.def:160:7
+/// @see src/tint/intrinsics.def:161:7
class U32TexelFormat : public NumberMatcher {
public:
/// Checks whether the given number matches the enum matcher rules.
@@ -2302,7 +2346,7 @@
}
/// EnumMatcher for 'match write_only'
-/// @see src/tint/intrinsics.def:163:7
+/// @see src/tint/intrinsics.def:164:7
class WriteOnly : public NumberMatcher {
public:
/// Checks whether the given number matches the enum matcher rules.
@@ -2328,7 +2372,7 @@
}
/// EnumMatcher for 'match function_private_workgroup'
-/// @see src/tint/intrinsics.def:165:7
+/// @see src/tint/intrinsics.def:166:7
class FunctionPrivateWorkgroup : public NumberMatcher {
public:
/// Checks whether the given number matches the enum matcher rules.
@@ -2358,7 +2402,7 @@
}
/// EnumMatcher for 'match workgroup_or_storage'
-/// @see src/tint/intrinsics.def:166:7
+/// @see src/tint/intrinsics.def:167:7
class WorkgroupOrStorage : public NumberMatcher {
public:
/// Checks whether the given number matches the enum matcher rules.
@@ -2524,6 +2568,7 @@
Fi32F16 Fi32F16_;
Iu32 Iu32_;
Aiu32 Aiu32_;
+ Afi32F16 Afi32F16_;
Scalar Scalar_;
AbstractOrScalar AbstractOrScalar_;
AfF32 AfF32_;
@@ -2550,7 +2595,7 @@
~Matchers();
/// The template types, types, and type matchers
- TypeMatcher const* const type[66] = {
+ TypeMatcher const* const type[67] = {
/* [0] */ &template_type_0_,
/* [1] */ &template_type_1_,
/* [2] */ &Bool_,
@@ -2608,15 +2653,16 @@
/* [54] */ &Fi32F16_,
/* [55] */ &Iu32_,
/* [56] */ &Aiu32_,
- /* [57] */ &Scalar_,
- /* [58] */ &AbstractOrScalar_,
- /* [59] */ &AfF32_,
- /* [60] */ &AfF32F16_,
- /* [61] */ &ScalarNoF32_,
- /* [62] */ &ScalarNoF16_,
- /* [63] */ &ScalarNoI32_,
- /* [64] */ &ScalarNoU32_,
- /* [65] */ &ScalarNoBool_,
+ /* [57] */ &Afi32F16_,
+ /* [58] */ &Scalar_,
+ /* [59] */ &AbstractOrScalar_,
+ /* [60] */ &AfF32_,
+ /* [61] */ &AfF32F16_,
+ /* [62] */ &ScalarNoF32_,
+ /* [63] */ &ScalarNoF16_,
+ /* [64] */ &ScalarNoI32_,
+ /* [65] */ &ScalarNoU32_,
+ /* [66] */ &ScalarNoBool_,
};
/// The template numbers, and number matchers
@@ -7914,7 +7960,7 @@
{
/* [1] */
/* name */ "U",
- /* matcher index */ 65,
+ /* matcher index */ 66,
},
{
/* [2] */
@@ -7924,7 +7970,7 @@
{
/* [3] */
/* name */ "U",
- /* matcher index */ 61,
+ /* matcher index */ 62,
},
{
/* [4] */
@@ -7934,7 +7980,7 @@
{
/* [5] */
/* name */ "U",
- /* matcher index */ 62,
+ /* matcher index */ 63,
},
{
/* [6] */
@@ -7944,7 +7990,7 @@
{
/* [7] */
/* name */ "U",
- /* matcher index */ 63,
+ /* matcher index */ 64,
},
{
/* [8] */
@@ -7954,7 +8000,7 @@
{
/* [9] */
/* name */ "U",
- /* matcher index */ 64,
+ /* matcher index */ 65,
},
{
/* [10] */
@@ -7964,7 +8010,7 @@
{
/* [11] */
/* name */ "T",
- /* matcher index */ 60,
+ /* matcher index */ 61,
},
{
/* [12] */
@@ -7979,12 +8025,12 @@
{
/* [14] */
/* name */ "T",
- /* matcher index */ 58,
+ /* matcher index */ 59,
},
{
/* [15] */
/* name */ "T",
- /* matcher index */ 57,
+ /* matcher index */ 58,
},
{
/* [16] */
@@ -7994,27 +8040,27 @@
{
/* [17] */
/* name */ "T",
- /* matcher index */ 65,
+ /* matcher index */ 66,
},
{
/* [18] */
/* name */ "T",
- /* matcher index */ 62,
+ /* matcher index */ 63,
},
{
/* [19] */
/* name */ "T",
- /* matcher index */ 61,
+ /* matcher index */ 62,
},
{
/* [20] */
/* name */ "T",
- /* matcher index */ 64,
+ /* matcher index */ 65,
},
{
/* [21] */
/* name */ "T",
- /* matcher index */ 63,
+ /* matcher index */ 64,
},
{
/* [22] */
@@ -8024,7 +8070,7 @@
{
/* [23] */
/* name */ "T",
- /* matcher index */ 54,
+ /* matcher index */ 57,
},
{
/* [24] */
@@ -13161,7 +13207,7 @@
/* parameters */ &kParameters[862],
/* return matcher indices */ &kMatcherIndices[1],
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::OpMinus,
},
{
/* [423] */
@@ -13173,7 +13219,7 @@
/* parameters */ &kParameters[863],
/* return matcher indices */ &kMatcherIndices[39],
/* flags */ OverloadFlags(OverloadFlag::kIsOperator, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::OpMinus,
},
{
/* [424] */
@@ -14521,8 +14567,8 @@
},
{
/* [2] */
- /* op -<T : fi32f16>(T) -> T */
- /* op -<T : fi32f16, N : num>(vec<N, T>) -> vec<N, T> */
+ /* op -<T : afi32f16>(T) -> T */
+ /* op -<T : afi32f16, N : num>(vec<N, T>) -> vec<N, T> */
/* num overloads */ 2,
/* overloads */ &kOverloads[422],
},
diff --git a/src/tint/resolver/intrinsic_table_test.cc b/src/tint/resolver/intrinsic_table_test.cc
index 30d7909..affcc76 100644
--- a/src/tint/resolver/intrinsic_table_test.cc
+++ b/src/tint/resolver/intrinsic_table_test.cc
@@ -604,8 +604,8 @@
EXPECT_EQ(Diagnostics().str(), R"(12:34 error: no matching overload for operator - (bool)
2 candidate operators:
- operator - (T) -> T where: T is f32, f16 or i32
- operator - (vecN<T>) -> vecN<T> where: T is f32, f16 or i32
+ operator - (T) -> T where: T is abstract-int, abstract-float, f32, i32 or f16
+ operator - (vecN<T>) -> vecN<T> where: T is abstract-int, abstract-float, f32, i32 or f16
)");
}
diff --git a/src/tint/resolver/storage_class_validation_test.cc b/src/tint/resolver/storage_class_validation_test.cc
index cb3ab16..f3e6f32 100644
--- a/src/tint/resolver/storage_class_validation_test.cc
+++ b/src/tint/resolver/storage_class_validation_test.cc
@@ -134,7 +134,7 @@
ASSERT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "56:78 error: using f16 types in 'uniform' or 'storage' storage class is not "
+ "56:78 error: using f16 types in 'storage' storage class is not "
"implemented yet");
}
@@ -153,7 +153,7 @@
ASSERT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "56:78 error: using f16 types in 'uniform' or 'storage' storage class is not "
+ "56:78 error: using f16 types in 'storage' storage class is not "
"implemented yet");
}
@@ -169,7 +169,7 @@
ASSERT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "56:78 error: using f16 types in 'uniform' or 'storage' storage class is not "
+ "56:78 error: using f16 types in 'storage' storage class is not "
"implemented yet");
}
@@ -188,9 +188,8 @@
ASSERT_FALSE(r()->Resolve());
- EXPECT_THAT(r()->error(),
- HasSubstr("56:78 error: using f16 types in 'uniform' or 'storage' storage "
- "class is not implemented yet"));
+ EXPECT_THAT(r()->error(), HasSubstr("56:78 error: using f16 types in 'storage' storage "
+ "class is not implemented yet"));
}
TEST_F(ResolverStorageClassValidationTest, StorageBufferStructF16_TemporallyBan) {
@@ -207,9 +206,8 @@
ASSERT_FALSE(r()->Resolve());
- EXPECT_THAT(r()->error(),
- HasSubstr("12:34 error: using f16 types in 'uniform' or 'storage' storage "
- "class is not implemented yet"));
+ EXPECT_THAT(r()->error(), HasSubstr("12:34 error: using f16 types in 'storage' storage "
+ "class is not implemented yet"));
}
TEST_F(ResolverStorageClassValidationTest, StorageBufferNoErrorStructF16Aliases_TemporallyBan) {
@@ -229,9 +227,8 @@
ASSERT_FALSE(r()->Resolve());
- EXPECT_THAT(r()->error(),
- HasSubstr("12:34 error: using f16 types in 'uniform' or 'storage' storage "
- "class is not implemented yet"));
+ EXPECT_THAT(r()->error(), HasSubstr("12:34 error: using f16 types in 'storage' storage "
+ "class is not implemented yet"));
}
TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) {
@@ -421,7 +418,7 @@
ASSERT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "56:78 error: using f16 types in 'uniform' or 'storage' storage class is not "
+ "56:78 error: using f16 types in 'uniform' storage class is not "
"implemented yet");
}
@@ -440,7 +437,7 @@
ASSERT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "56:78 error: using f16 types in 'uniform' or 'storage' storage class is not "
+ "56:78 error: using f16 types in 'uniform' storage class is not "
"implemented yet");
}
@@ -455,9 +452,8 @@
ASSERT_FALSE(r()->Resolve());
- EXPECT_THAT(r()->error(),
- HasSubstr("56:78 error: using f16 types in 'uniform' or 'storage' storage "
- "class is not implemented yet"));
+ EXPECT_THAT(r()->error(), HasSubstr("56:78 error: using f16 types in 'uniform' storage "
+ "class is not implemented yet"));
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferArrayF16_TemporallyBan) {
@@ -477,9 +473,8 @@
ASSERT_FALSE(r()->Resolve());
- EXPECT_THAT(r()->error(),
- HasSubstr("56:78 error: using f16 types in 'uniform' or 'storage' storage "
- "class is not implemented yet"));
+ EXPECT_THAT(r()->error(), HasSubstr("56:78 error: using f16 types in 'uniform' storage "
+ "class is not implemented yet"));
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferStructF16_TemporallyBan) {
@@ -496,9 +491,8 @@
ASSERT_FALSE(r()->Resolve());
- EXPECT_THAT(r()->error(),
- HasSubstr("12:34 error: using f16 types in 'uniform' or 'storage' storage "
- "class is not implemented yet"));
+ EXPECT_THAT(r()->error(), HasSubstr("12:34 error: using f16 types in 'uniform' storage "
+ "class is not implemented yet"));
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferStructF16Aliases_TemporallyBan) {
@@ -517,9 +511,8 @@
ASSERT_FALSE(r()->Resolve());
- EXPECT_THAT(r()->error(),
- HasSubstr("12:34 error: using f16 types in 'uniform' or 'storage' storage "
- "class is not implemented yet"));
+ EXPECT_THAT(r()->error(), HasSubstr("12:34 error: using f16 types in 'uniform' storage "
+ "class is not implemented yet"));
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferPointer) {
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index 0b8fb8c..f962e3b 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -396,7 +396,8 @@
// TODO(tint:1473, tint:1502): Remove this error after f16 is supported in "uniform" and
// "storage" storage class.
if (Is<sem::F16>(sem::Type::DeepestElementOf(store_ty))) {
- AddError("using f16 types in 'uniform' or 'storage' storage class is not implemented yet",
+ AddError("using f16 types in '" + std::string(ast::ToString(sc)) +
+ "' storage class is not implemented yet",
source);
return false;
}
diff --git a/src/tint/sem/f16.h b/src/tint/sem/f16.h
index 72984c1..87543ed 100644
--- a/src/tint/sem/f16.h
+++ b/src/tint/sem/f16.h
@@ -22,7 +22,7 @@
namespace tint::sem {
/// A float 16 type
-class F16 : public Castable<F16, Type> {
+class F16 final : public Castable<F16, Type> {
public:
/// Constructor
F16();
diff --git a/src/tint/sem/struct.h b/src/tint/sem/struct.h
index e026b11..5ee93e5 100644
--- a/src/tint/sem/struct.h
+++ b/src/tint/sem/struct.h
@@ -170,7 +170,7 @@
};
/// StructMember holds the semantic information for structure members.
-class StructMember : public Castable<StructMember, Node> {
+class StructMember final : public Castable<StructMember, Node> {
public:
/// Constructor
/// @param declaration the AST declaration node
diff --git a/src/tint/sem/type.cc b/src/tint/sem/type.cc
index 9d4c469..5b2baa0 100644
--- a/src/tint/sem/type.cc
+++ b/src/tint/sem/type.cc
@@ -136,6 +136,30 @@
return is_unsigned_scalar_or_vector() || is_signed_scalar_or_vector();
}
+bool Type::is_abstract_scalar_vector() const {
+ return Is([](const Vector* v) { return v->type()->Is<sem::AbstractNumeric>(); });
+}
+
+bool Type::is_abstract_integer_vector() const {
+ return Is([](const Vector* v) { return v->type()->Is<sem::AbstractInt>(); });
+}
+
+bool Type::is_abstract_float_vector() const {
+ return Is([](const Vector* v) { return v->type()->Is<sem::AbstractFloat>(); });
+}
+
+bool Type::is_abstract_scalar_or_vector() const {
+ return Is<sem::AbstractNumeric>() || is_abstract_scalar_vector();
+}
+
+bool Type::is_abstract_integer_scalar_or_vector() const {
+ return Is<sem::AbstractInt>() || is_abstract_integer_vector();
+}
+
+bool Type::is_abstract_float_scalar_or_vector() const {
+ return Is<sem::AbstractFloat>() || is_abstract_float_vector();
+}
+
bool Type::is_bool_vector() const {
return Is([](const Vector* v) { return v->type()->Is<Bool>(); });
}
diff --git a/src/tint/sem/type.h b/src/tint/sem/type.h
index 25f3a43..3866e8d 100644
--- a/src/tint/sem/type.h
+++ b/src/tint/sem/type.h
@@ -103,6 +103,18 @@
bool is_signed_scalar_or_vector() const;
/// @returns true if this type is an integer scalar or vector
bool is_integer_scalar_or_vector() const;
+ /// @returns true if this type is an abstract scalar vector
+ bool is_abstract_scalar_vector() const;
+ /// @returns true if this type is an abstract integer vector
+ bool is_abstract_integer_vector() const;
+ /// @returns true if this type is an abstract float vector
+ bool is_abstract_float_vector() const;
+ /// @returns true if this type is an abstract scalar or vector
+ bool is_abstract_scalar_or_vector() const;
+ /// @returns true if this type is an abstract integer scalar or vector
+ bool is_abstract_integer_scalar_or_vector() const;
+ /// @returns true if this type is an abstract float scalar or vector
+ bool is_abstract_float_scalar_or_vector() const;
/// @returns true if this type is a boolean vector
bool is_bool_vector() const;
/// @returns true if this type is boolean scalar or vector
diff --git a/src/tint/transform/decompose_memory_access.cc b/src/tint/transform/decompose_memory_access.cc
index c3d5c8e..a714581 100644
--- a/src/tint/transform/decompose_memory_access.cc
+++ b/src/tint/transform/decompose_memory_access.cc
@@ -73,7 +73,7 @@
/// OffsetLiteral is an implementation of Offset that constructs a u32 literal
/// value.
-struct OffsetLiteral : Castable<OffsetLiteral, Offset> {
+struct OffsetLiteral final : Castable<OffsetLiteral, Offset> {
uint32_t const literal = 0;
explicit OffsetLiteral(uint32_t lit) : literal(lit) {}
diff --git a/src/tint/transform/expand_compound_assignment.h b/src/tint/transform/expand_compound_assignment.h
index d38d297..1081df7 100644
--- a/src/tint/transform/expand_compound_assignment.h
+++ b/src/tint/transform/expand_compound_assignment.h
@@ -38,7 +38,7 @@
///
/// This transform also handles increment and decrement statements in the same
/// manner, by replacing `i++` with `i = i + 1`.
-class ExpandCompoundAssignment : public Castable<ExpandCompoundAssignment, Transform> {
+class ExpandCompoundAssignment final : public Castable<ExpandCompoundAssignment, Transform> {
public:
/// Constructor
ExpandCompoundAssignment();
diff --git a/src/tint/transform/localize_struct_array_assignment.h b/src/tint/transform/localize_struct_array_assignment.h
index 129c849..130f8cc 100644
--- a/src/tint/transform/localize_struct_array_assignment.h
+++ b/src/tint/transform/localize_struct_array_assignment.h
@@ -27,7 +27,8 @@
///
/// @note Depends on the following transforms to have been run first:
/// * SimplifyPointers
-class LocalizeStructArrayAssignment : public Castable<LocalizeStructArrayAssignment, Transform> {
+class LocalizeStructArrayAssignment final
+ : public Castable<LocalizeStructArrayAssignment, Transform> {
public:
/// Constructor
LocalizeStructArrayAssignment();
diff --git a/src/tint/transform/loop_to_for_loop.h b/src/tint/transform/loop_to_for_loop.h
index 0623d79..0e948c8 100644
--- a/src/tint/transform/loop_to_for_loop.h
+++ b/src/tint/transform/loop_to_for_loop.h
@@ -21,7 +21,7 @@
/// LoopToForLoop is a Transform that attempts to convert WGSL `loop {}`
/// statements into a for-loop statement.
-class LoopToForLoop : public Castable<LoopToForLoop, Transform> {
+class LoopToForLoop final : public Castable<LoopToForLoop, Transform> {
public:
/// Constructor
LoopToForLoop();
diff --git a/src/tint/transform/manager.h b/src/tint/transform/manager.h
index 9f5c6bc..04bf9fe 100644
--- a/src/tint/transform/manager.h
+++ b/src/tint/transform/manager.h
@@ -27,7 +27,7 @@
/// The inner transforms will execute in the appended order.
/// If any inner transform fails the manager will return immediately and
/// the error can be retrieved with the Output's diagnostics.
-class Manager : public Castable<Manager, Transform> {
+class Manager final : public Castable<Manager, Transform> {
public:
/// Constructor
Manager();
diff --git a/src/tint/transform/module_scope_var_to_entry_point_param.h b/src/tint/transform/module_scope_var_to_entry_point_param.h
index e3a50f4..40e6b7d 100644
--- a/src/tint/transform/module_scope_var_to_entry_point_param.h
+++ b/src/tint/transform/module_scope_var_to_entry_point_param.h
@@ -61,7 +61,7 @@
/// foo(&p, sptr);
/// }
/// ```
-class ModuleScopeVarToEntryPointParam
+class ModuleScopeVarToEntryPointParam final
: public Castable<ModuleScopeVarToEntryPointParam, Transform> {
public:
/// Constructor
diff --git a/src/tint/transform/multiplanar_external_texture.h b/src/tint/transform/multiplanar_external_texture.h
index fcb5156..afd15a1 100644
--- a/src/tint/transform/multiplanar_external_texture.h
+++ b/src/tint/transform/multiplanar_external_texture.h
@@ -50,7 +50,7 @@
/// decoding, gamut conversion, and gamma encoding steps. Specifically
// for BT.709 to SRGB conversion, it takes the fast path only doing the yuv->rgb
// step and skipping all other steps.
-class MultiplanarExternalTexture : public Castable<MultiplanarExternalTexture, Transform> {
+class MultiplanarExternalTexture final : public Castable<MultiplanarExternalTexture, Transform> {
public:
/// BindingsMap is a map where the key is the binding location of a
/// texture_external and the value is a struct containing the desired
@@ -60,7 +60,7 @@
/// NewBindingPoints is consumed by the MultiplanarExternalTexture transform.
/// Data holds information about location of each texture_external binding and
/// which binding slots it should expand into.
- struct NewBindingPoints : public Castable<Data, transform::Data> {
+ struct NewBindingPoints final : public Castable<Data, transform::Data> {
/// Constructor
/// @param bm a map to the new binding slots to use.
explicit NewBindingPoints(BindingsMap bm);
diff --git a/src/tint/transform/num_workgroups_from_uniform.h b/src/tint/transform/num_workgroups_from_uniform.h
index 0111ccc..292c823 100644
--- a/src/tint/transform/num_workgroups_from_uniform.h
+++ b/src/tint/transform/num_workgroups_from_uniform.h
@@ -44,7 +44,7 @@
///
/// @note Depends on the following transforms to have been run first:
/// * CanonicalizeEntryPointIO
-class NumWorkgroupsFromUniform : public Castable<NumWorkgroupsFromUniform, Transform> {
+class NumWorkgroupsFromUniform final : public Castable<NumWorkgroupsFromUniform, Transform> {
public:
/// Constructor
NumWorkgroupsFromUniform();
@@ -52,7 +52,7 @@
~NumWorkgroupsFromUniform() override;
/// Configuration options for the NumWorkgroupsFromUniform transform.
- struct Config : public Castable<Data, transform::Data> {
+ struct Config final : public Castable<Data, transform::Data> {
/// Constructor
/// @param ubo_bp the binding point to use for the generated uniform buffer. If ubo_bp
/// contains no value, a free binding point will be used to ensure the generated program is
diff --git a/src/tint/transform/promote_initializers_to_let.h b/src/tint/transform/promote_initializers_to_let.h
index 41f99d7..226c7d8 100644
--- a/src/tint/transform/promote_initializers_to_let.h
+++ b/src/tint/transform/promote_initializers_to_let.h
@@ -25,7 +25,7 @@
/// array or structure. For example, the following is not immediately expressable for HLSL:
/// `array<i32, 2>(1, 2)[0]`
/// @see crbug.com/tint/406
-class PromoteInitializersToLet : public Castable<PromoteInitializersToLet, Transform> {
+class PromoteInitializersToLet final : public Castable<PromoteInitializersToLet, Transform> {
public:
/// Constructor
PromoteInitializersToLet();
diff --git a/src/tint/transform/promote_side_effects_to_decl.h b/src/tint/transform/promote_side_effects_to_decl.h
index 1e629b3..d5d1126 100644
--- a/src/tint/transform/promote_side_effects_to_decl.h
+++ b/src/tint/transform/promote_side_effects_to_decl.h
@@ -23,7 +23,7 @@
/// declarations before the statement of usage with the goal of ensuring
/// left-to-right order of evaluation, while respecting short-circuit
/// evaluation.
-class PromoteSideEffectsToDecl : public Castable<PromoteSideEffectsToDecl, Transform> {
+class PromoteSideEffectsToDecl final : public Castable<PromoteSideEffectsToDecl, Transform> {
public:
/// Constructor
PromoteSideEffectsToDecl();
diff --git a/src/tint/transform/remove_continue_in_switch.h b/src/tint/transform/remove_continue_in_switch.h
index e706225..9e5a4d5 100644
--- a/src/tint/transform/remove_continue_in_switch.h
+++ b/src/tint/transform/remove_continue_in_switch.h
@@ -23,7 +23,7 @@
/// bool variable, and checking if the variable is set after the switch to
/// continue. It is necessary to work around FXC "error X3708: continue cannot
/// be used in a switch". See crbug.com/tint/1080.
-class RemoveContinueInSwitch : public Castable<RemoveContinueInSwitch, Transform> {
+class RemoveContinueInSwitch final : public Castable<RemoveContinueInSwitch, Transform> {
public:
/// Constructor
RemoveContinueInSwitch();
diff --git a/src/tint/transform/remove_phonies.h b/src/tint/transform/remove_phonies.h
index 20128a0..d04023b 100644
--- a/src/tint/transform/remove_phonies.h
+++ b/src/tint/transform/remove_phonies.h
@@ -25,7 +25,7 @@
/// RemovePhonies is a Transform that removes all phony-assignment statements,
/// while preserving function call expressions in the RHS of the assignment that
/// may have side-effects.
-class RemovePhonies : public Castable<RemovePhonies, Transform> {
+class RemovePhonies final : public Castable<RemovePhonies, Transform> {
public:
/// Constructor
RemovePhonies();
diff --git a/src/tint/transform/remove_unreachable_statements.h b/src/tint/transform/remove_unreachable_statements.h
index c75da3d..7f8b947 100644
--- a/src/tint/transform/remove_unreachable_statements.h
+++ b/src/tint/transform/remove_unreachable_statements.h
@@ -24,7 +24,7 @@
/// RemoveUnreachableStatements is a Transform that removes all statements
/// marked as unreachable.
-class RemoveUnreachableStatements : public Castable<RemoveUnreachableStatements, Transform> {
+class RemoveUnreachableStatements final : public Castable<RemoveUnreachableStatements, Transform> {
public:
/// Constructor
RemoveUnreachableStatements();
diff --git a/src/tint/transform/renamer.h b/src/tint/transform/renamer.h
index 354acda..000aee9 100644
--- a/src/tint/transform/renamer.h
+++ b/src/tint/transform/renamer.h
@@ -23,11 +23,11 @@
namespace tint::transform {
/// Renamer is a Transform that renames all the symbols in a program.
-class Renamer : public Castable<Renamer, Transform> {
+class Renamer final : public Castable<Renamer, Transform> {
public:
/// Data is outputted by the Renamer transform.
/// Data holds information about shader usage and constant buffer offsets.
- struct Data : public Castable<Data, transform::Data> {
+ struct Data final : public Castable<Data, transform::Data> {
/// Remappings is a map of old symbol name to new symbol name
using Remappings = std::unordered_map<std::string, std::string>;
@@ -59,7 +59,7 @@
/// Optional configuration options for the transform.
/// If omitted, then the renamer will use Target::kAll.
- struct Config : public Castable<Config, transform::Data> {
+ struct Config final : public Castable<Config, transform::Data> {
/// Constructor
/// @param tgt the targets to rename
/// @param keep_unicode if false, symbols with non-ascii code-points are
diff --git a/src/tint/transform/robustness.h b/src/tint/transform/robustness.h
index 138b48c..549b666 100644
--- a/src/tint/transform/robustness.h
+++ b/src/tint/transform/robustness.h
@@ -31,7 +31,7 @@
/// 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).
-class Robustness : public Castable<Robustness, Transform> {
+class Robustness final : public Castable<Robustness, Transform> {
public:
/// Storage class to be skipped in the transform
enum class StorageClass {
@@ -40,7 +40,7 @@
};
/// Configuration options for the transform
- struct Config : public Castable<Config, Data> {
+ struct Config final : public Castable<Config, Data> {
/// Constructor
Config();
diff --git a/src/tint/transform/simplify_pointers.h b/src/tint/transform/simplify_pointers.h
index 267b7b2..787c7d8 100644
--- a/src/tint/transform/simplify_pointers.h
+++ b/src/tint/transform/simplify_pointers.h
@@ -31,7 +31,7 @@
///
/// @note Depends on the following transforms to have been run first:
/// * Unshadow
-class SimplifyPointers : public Castable<SimplifyPointers, Transform> {
+class SimplifyPointers final : public Castable<SimplifyPointers, Transform> {
public:
/// Constructor
SimplifyPointers();
diff --git a/src/tint/transform/single_entry_point.h b/src/tint/transform/single_entry_point.h
index 0a922a7..59aa021 100644
--- a/src/tint/transform/single_entry_point.h
+++ b/src/tint/transform/single_entry_point.h
@@ -25,10 +25,10 @@
///
/// All module-scope variables, types, and functions that are not used by the
/// target entry point will also be removed.
-class SingleEntryPoint : public Castable<SingleEntryPoint, Transform> {
+class SingleEntryPoint final : public Castable<SingleEntryPoint, Transform> {
public:
/// Configuration options for the transform
- struct Config : public Castable<Config, Data> {
+ struct Config final : public Castable<Config, Data> {
/// Constructor
/// @param entry_point the name of the entry point to keep
explicit Config(std::string entry_point = "");
diff --git a/src/tint/transform/unshadow.h b/src/tint/transform/unshadow.h
index ce5e975..5ffe839 100644
--- a/src/tint/transform/unshadow.h
+++ b/src/tint/transform/unshadow.h
@@ -21,7 +21,7 @@
/// Unshadow is a Transform that renames any variables that shadow another
/// variable.
-class Unshadow : public Castable<Unshadow, Transform> {
+class Unshadow final : public Castable<Unshadow, Transform> {
public:
/// Constructor
Unshadow();
diff --git a/src/tint/transform/unwind_discard_functions.h b/src/tint/transform/unwind_discard_functions.h
index 3b1d838..105a9d8 100644
--- a/src/tint/transform/unwind_discard_functions.h
+++ b/src/tint/transform/unwind_discard_functions.h
@@ -36,7 +36,7 @@
///
/// @note Depends on the following transforms to have been run first:
/// * PromoteSideEffectsToDecl
-class UnwindDiscardFunctions : public Castable<UnwindDiscardFunctions, Transform> {
+class UnwindDiscardFunctions final : public Castable<UnwindDiscardFunctions, Transform> {
public:
/// Constructor
UnwindDiscardFunctions();
diff --git a/src/tint/transform/vectorize_scalar_matrix_constructors.h b/src/tint/transform/vectorize_scalar_matrix_constructors.h
index 83c4ce1..31c57f0 100644
--- a/src/tint/transform/vectorize_scalar_matrix_constructors.h
+++ b/src/tint/transform/vectorize_scalar_matrix_constructors.h
@@ -20,7 +20,7 @@
namespace tint::transform {
/// A transform that converts scalar matrix constructors to the vector form.
-class VectorizeScalarMatrixConstructors
+class VectorizeScalarMatrixConstructors final
: public Castable<VectorizeScalarMatrixConstructors, Transform> {
public:
/// Constructor
diff --git a/src/tint/transform/vertex_pulling.h b/src/tint/transform/vertex_pulling.h
index 7875600..92eb627 100644
--- a/src/tint/transform/vertex_pulling.h
+++ b/src/tint/transform/vertex_pulling.h
@@ -128,10 +128,10 @@
/// code, but these are types that the data may arrive as. We need to convert
/// these smaller types into the base types such as `f32` and `u32` for the
/// shader to use.
-class VertexPulling : public Castable<VertexPulling, Transform> {
+class VertexPulling final : public Castable<VertexPulling, Transform> {
public:
/// Configuration options for the transform
- struct Config : public Castable<Config, Data> {
+ struct Config final : public Castable<Config, Data> {
/// Constructor
Config();
diff --git a/src/tint/transform/zero_init_workgroup_memory.h b/src/tint/transform/zero_init_workgroup_memory.h
index c757725..07feaa8 100644
--- a/src/tint/transform/zero_init_workgroup_memory.h
+++ b/src/tint/transform/zero_init_workgroup_memory.h
@@ -22,7 +22,7 @@
/// ZeroInitWorkgroupMemory is a transform that injects code at the top of entry
/// points to zero-initialize workgroup memory used by that entry point (and all
/// transitive functions called by that entry point)
-class ZeroInitWorkgroupMemory : public Castable<ZeroInitWorkgroupMemory, Transform> {
+class ZeroInitWorkgroupMemory final : public Castable<ZeroInitWorkgroupMemory, Transform> {
public:
/// Constructor
ZeroInitWorkgroupMemory();
diff --git a/src/tint/writer/msl/generator_impl_unary_op_test.cc b/src/tint/writer/msl/generator_impl_unary_op_test.cc
index 843e2a4..21d6357 100644
--- a/src/tint/writer/msl/generator_impl_unary_op_test.cc
+++ b/src/tint/writer/msl/generator_impl_unary_op_test.cc
@@ -81,16 +81,15 @@
EXPECT_EQ(out.str(), "tint_unary_minus(expr)");
}
-TEST_F(MslUnaryOpTest, NegationOfIntMin) {
- auto* op = create<ast::UnaryOpExpression>(ast::UnaryOp::kNegation,
- Expr(i32(std::numeric_limits<int32_t>::min())));
+TEST_F(MslUnaryOpTest, IntMin) {
+ auto* op = Expr(i32(std::numeric_limits<int32_t>::min()));
WrapInFunction(op);
GeneratorImpl& gen = Build();
std::stringstream out;
ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
- EXPECT_EQ(out.str(), "tint_unary_minus((-2147483647 - 1))");
+ EXPECT_EQ(out.str(), "(-2147483647 - 1)");
}
} // namespace
