Import Tint changes from Dawn
Changes:
- c96f8183280aa15e26e7fd9e197c818b7d219cf8 [tint][fuzz][ir] Add BindingRemapper fuzzer by Ben Clayton <bclayton@google.com>
- 3e18d6aead7a9bfce0a44e1cc224d8aad23a2563 [tint][fuzz][ir] Add BinaryPolyfill fuzzer by Ben Clayton <bclayton@google.com>
- b5045a6c9881bf4113313d534eaf8abdb235659b [tint][fuzz][ir] Add Bgra8UnormPolyfill fuzzer by Ben Clayton <bclayton@google.com>
- 8801bee379627f4a258d278a8a7a2f196d1cca9c [tint][fuzz][ir] Add AddEmptyEntryPoint fuzzer by Ben Clayton <bclayton@google.com>
- 6391acef4b6007a827546c5f226a384af4d008cb [tint][fuzz][ast] Add AddEmptyEntryPoint fuzzer by Ben Clayton <bclayton@google.com>
- c4d9ee724ea25288a5248ceb469514c0273eea7b [tint][fuzz][ast] Add AddBlockAttribute fuzzer by Ben Clayton <bclayton@google.com>
- ba06f4775aa0daeccd5af239f01fbdcd86a65e38 [tint][fuzz] Limit the first_index_offset value by Ben Clayton <bclayton@google.com>
- bc99356b1de0296043aebab79b9451d025c273c5 [Compat] Fix a bug when textureDimension with 1 arg is fi... by Shrek Shao <shrekshao@google.com>
- d011f42aeced8eb0ac7f3e999e8819baea66117d [tint][fuzz] Add fuzzer command line flags by Ben Clayton <bclayton@google.com>
- 988805ac8637dc9d93887fe521f5a5805600d02d [tint][spirv-ast] Fix emission of discard followed by sta... by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: c96f8183280aa15e26e7fd9e197c818b7d219cf8
Change-Id: Ic3ab680cb92fe4c0427ddc35186ad4307e810f4e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/185320
Commit-Queue: dan sinclair <dsinclair@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: dan sinclair <dsinclair@google.com>
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.cmake b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
index 853fd5d..d51bd4d 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.cmake
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.cmake
@@ -48,6 +48,7 @@
tint_cmd_fuzz_ir_fuzz
tint_lang_core
tint_lang_core_constant
+ tint_lang_core_ir_transform_fuzz
tint_lang_core_type
tint_lang_wgsl
tint_lang_wgsl_ast
@@ -142,6 +143,7 @@
tint_lang_wgsl_features
tint_lang_wgsl_program
tint_lang_wgsl_sem
+ tint_lang_wgsl_writer_ir_to_program
tint_utils_bytes
tint_utils_containers
tint_utils_diagnostic
@@ -168,4 +170,10 @@
)
endif(TINT_BUILD_WGSL_READER)
+if(TINT_BUILD_WGSL_WRITER)
+ tint_target_add_dependencies(tint_cmd_fuzz_wgsl_fuzz fuzz
+ tint_lang_wgsl_writer
+ )
+endif(TINT_BUILD_WGSL_WRITER)
+
endif(TINT_BUILD_WGSL_READER)
\ No newline at end of file
diff --git a/src/tint/cmd/fuzz/wgsl/BUILD.gn b/src/tint/cmd/fuzz/wgsl/BUILD.gn
index 97cec04..389a98e 100644
--- a/src/tint/cmd/fuzz/wgsl/BUILD.gn
+++ b/src/tint/cmd/fuzz/wgsl/BUILD.gn
@@ -56,6 +56,7 @@
"${tint_src_dir}/lang/wgsl/features",
"${tint_src_dir}/lang/wgsl/program",
"${tint_src_dir}/lang/wgsl/sem",
+ "${tint_src_dir}/lang/wgsl/writer/ir_to_program",
"${tint_src_dir}/utils/bytes",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
@@ -75,6 +76,10 @@
if (tint_build_wgsl_reader) {
deps += [ "${tint_src_dir}/lang/wgsl/reader" ]
}
+
+ if (tint_build_wgsl_writer) {
+ deps += [ "${tint_src_dir}/lang/wgsl/writer" ]
+ }
}
}
if (tint_build_wgsl_reader) {
@@ -85,6 +90,7 @@
"${tint_src_dir}/cmd/fuzz/ir:fuzz",
"${tint_src_dir}/lang/core",
"${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/ir/transform:fuzz",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl:fuzz",
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.cc b/src/tint/cmd/fuzz/wgsl/fuzz.cc
index d6b90ec..6e4d9b8 100644
--- a/src/tint/cmd/fuzz/wgsl/fuzz.cc
+++ b/src/tint/cmd/fuzz/wgsl/fuzz.cc
@@ -35,6 +35,11 @@
#include "src/tint/utils/macros/defer.h"
#include "src/tint/utils/macros/static_init.h"
+#if TINT_BUILD_WGSL_WRITER
+#include "src/tint/lang/wgsl/program/program.h"
+#include "src/tint/lang/wgsl/writer/writer.h"
+#endif
+
namespace tint::fuzz::wgsl {
namespace {
@@ -63,6 +68,17 @@
void Run(std::string_view wgsl, Slice<const std::byte> data, const Options& options) {
tint::SetInternalCompilerErrorReporter(&TintInternalCompilerErrorReporter);
+#if TINT_BUILD_WGSL_WRITER
+ // Register the Program printer. This is used for debugging purposes.
+ tint::Program::printer = [](const tint::Program& program) {
+ auto result = tint::wgsl::writer::Generate(program, {});
+ if (result != Success) {
+ return result.Failure().reason.Str();
+ }
+ return result->wgsl;
+ };
+#endif
+
// Ensure that fuzzers are sorted. Without this, the fuzzers may be registered in any order,
// leading to non-determinism, which we must avoid.
TINT_STATIC_INIT(Fuzzers().Sort([](auto& a, auto& b) { return a.name < b.name; }));
@@ -80,13 +96,20 @@
if (options.run_concurrently) {
size_t n = Fuzzers().Length();
tint::Vector<std::thread, 32> threads;
- threads.Resize(n);
+ threads.Reserve(n);
for (size_t i = 0; i < n; i++) {
- threads[i] = std::thread([i, &program, &data] {
+ if (!options.filter.empty() &&
+ Fuzzers()[i].name.find(options.filter) == std::string::npos) {
+ continue;
+ }
+ threads.Push(std::thread([i, &program, &data, &options] {
auto& fuzzer = Fuzzers()[i];
currently_running = fuzzer.name;
+ if (options.verbose) {
+ std::cout << " • [" << i << "] Running: " << currently_running << std::endl;
+ }
fuzzer.fn(program, data);
- });
+ }));
}
for (auto& thread : threads) {
thread.join();
@@ -94,7 +117,14 @@
} else {
TINT_DEFER(currently_running = "");
for (auto& fuzzer : Fuzzers()) {
+ if (!options.filter.empty() && fuzzer.name.find(options.filter) == std::string::npos) {
+ continue;
+ }
+
currently_running = fuzzer.name;
+ if (options.verbose) {
+ std::cout << " • Running: " << currently_running << std::endl;
+ }
fuzzer.fn(program, data);
}
}
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.h b/src/tint/cmd/fuzz/wgsl/fuzz.h
index d54b3ec..c01455b 100644
--- a/src/tint/cmd/fuzz/wgsl/fuzz.h
+++ b/src/tint/cmd/fuzz/wgsl/fuzz.h
@@ -78,8 +78,12 @@
/// Options for Run()
struct Options {
+ /// If not empty, only run the fuzzers with the given substring.
+ std::string filter;
/// If true, the fuzzers will be run concurrently on separate threads.
bool run_concurrently = false;
+ /// If true, print the fuzzer name to stdout before running.
+ bool verbose = false;
};
/// Runs all the registered WGSL fuzzers with the supplied WGSL
diff --git a/src/tint/cmd/fuzz/wgsl/main_fuzz.cc b/src/tint/cmd/fuzz/wgsl/main_fuzz.cc
index 9762952..3420d40 100644
--- a/src/tint/cmd/fuzz/wgsl/main_fuzz.cc
+++ b/src/tint/cmd/fuzz/wgsl/main_fuzz.cc
@@ -70,8 +70,12 @@
};
auto& opt_help = opts.Add<tint::cli::BoolOption>("help", "shows the usage");
+ auto& opt_filter = opts.Add<tint::cli::StringOption>(
+ "filter", "runs only the fuzzers with the given substring");
auto& opt_concurrent =
opts.Add<tint::cli::BoolOption>("concurrent", "runs the fuzzers concurrently");
+ auto& opt_verbose =
+ opts.Add<tint::cli::BoolOption>("verbose", "prints the name of each fuzzer before running");
tint::cli::ParseOptions parse_opts;
parse_opts.ignore_unknown = true;
@@ -86,6 +90,8 @@
return 0;
}
+ options.filter = opt_filter.value.value_or("");
options.run_concurrently = opt_concurrent.value.value_or(false);
+ options.verbose = opt_verbose.value.value_or(false);
return 0;
}
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake
index 01120f4..223dbee 100644
--- a/src/tint/lang/core/ir/transform/BUILD.cmake
+++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -179,3 +179,33 @@
tint_lang_wgsl_writer
)
endif(TINT_BUILD_WGSL_WRITER)
+
+################################################################################
+# Target: tint_lang_core_ir_transform_fuzz
+# Kind: fuzz
+################################################################################
+tint_add_target(tint_lang_core_ir_transform_fuzz fuzz
+ lang/core/ir/transform/add_empty_entry_point_fuzz.cc
+ lang/core/ir/transform/bgra8unorm_polyfill_fuzz.cc
+ lang/core/ir/transform/binary_polyfill_fuzz.cc
+ lang/core/ir/transform/binding_remapper_fuzz.cc
+)
+
+tint_target_add_dependencies(tint_lang_core_ir_transform_fuzz fuzz
+ tint_api_common
+ tint_cmd_fuzz_ir_fuzz
+ tint_lang_core_ir
+ tint_lang_core_ir_transform
+ tint_utils_bytes
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_reflection
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_text
+ tint_utils_traits
+)
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn
index 56910e6..bd81416 100644
--- a/src/tint/lang/core/ir/transform/BUILD.gn
+++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -175,3 +175,30 @@
}
}
}
+
+tint_fuzz_source_set("fuzz") {
+ sources = [
+ "add_empty_entry_point_fuzz.cc",
+ "bgra8unorm_polyfill_fuzz.cc",
+ "binary_polyfill_fuzz.cc",
+ "binding_remapper_fuzz.cc",
+ ]
+ deps = [
+ "${tint_src_dir}/api/common",
+ "${tint_src_dir}/cmd/fuzz/ir:fuzz",
+ "${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/ir/transform",
+ "${tint_src_dir}/utils/bytes",
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/diagnostic",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/reflection",
+ "${tint_src_dir}/utils/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+}
diff --git a/src/tint/lang/core/ir/transform/add_empty_entry_point_fuzz.cc b/src/tint/lang/core/ir/transform/add_empty_entry_point_fuzz.cc
new file mode 100644
index 0000000..70056f6
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/add_empty_entry_point_fuzz.cc
@@ -0,0 +1,50 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/core/ir/transform/add_empty_entry_point.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+void AddEmptyEntryPointFuzzer(Module& module) {
+ if (auto res = AddEmptyEntryPoint(module); res != Success) {
+ return;
+ }
+
+ Capabilities capabilities;
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of AddEmptyEntryPoint failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
+
+TINT_IR_MODULE_FUZZER(tint::core::ir::transform::AddEmptyEntryPointFuzzer);
diff --git a/src/tint/lang/core/ir/transform/bgra8unorm_polyfill_fuzz.cc b/src/tint/lang/core/ir/transform/bgra8unorm_polyfill_fuzz.cc
new file mode 100644
index 0000000..cbb95fe
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/bgra8unorm_polyfill_fuzz.cc
@@ -0,0 +1,50 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/core/ir/transform/bgra8unorm_polyfill.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+void Bgra8UnormPolyfillFuzzer(Module& module) {
+ if (auto res = Bgra8UnormPolyfill(module); res != Success) {
+ return;
+ }
+
+ Capabilities capabilities;
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of Bgra8UnormPolyfill failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
+
+TINT_IR_MODULE_FUZZER(tint::core::ir::transform::Bgra8UnormPolyfillFuzzer);
diff --git a/src/tint/lang/core/ir/transform/binary_polyfill.h b/src/tint/lang/core/ir/transform/binary_polyfill.h
index 82fe8d5..1ff8efd 100644
--- a/src/tint/lang/core/ir/transform/binary_polyfill.h
+++ b/src/tint/lang/core/ir/transform/binary_polyfill.h
@@ -30,6 +30,7 @@
#include <string>
+#include "src/tint/utils/reflection/reflection.h"
#include "src/tint/utils/result/result.h"
// Forward declarations.
@@ -45,6 +46,9 @@
bool bitshift_modulo = false;
/// Should integer divide and modulo be polyfilled to avoid DBZ and integer overflow?
bool int_div_mod = false;
+
+ /// Reflection for this class
+ TINT_REFLECT(BinaryPolyfillConfig, bitshift_modulo, int_div_mod);
};
/// BinaryPolyfill is a transform that modifies binary instructions to prepare them for raising to
diff --git a/src/tint/lang/core/ir/transform/binary_polyfill_fuzz.cc b/src/tint/lang/core/ir/transform/binary_polyfill_fuzz.cc
new file mode 100644
index 0000000..13875ea
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/binary_polyfill_fuzz.cc
@@ -0,0 +1,50 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/core/ir/transform/binary_polyfill.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+void BinaryPolyfillFuzzer(Module& module, BinaryPolyfillConfig config) {
+ if (auto res = BinaryPolyfill(module, config); res != Success) {
+ return;
+ }
+
+ Capabilities capabilities;
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of BinaryPolyfill failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
+
+TINT_IR_MODULE_FUZZER(tint::core::ir::transform::BinaryPolyfillFuzzer);
diff --git a/src/tint/lang/core/ir/transform/binding_remapper_fuzz.cc b/src/tint/lang/core/ir/transform/binding_remapper_fuzz.cc
new file mode 100644
index 0000000..68ba651
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/binding_remapper_fuzz.cc
@@ -0,0 +1,51 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/core/ir/transform/binding_remapper.h"
+
+#include "src/tint/cmd/fuzz/ir/fuzz.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+void BindingRemapperFuzzer(Module& module,
+ const std::unordered_map<BindingPoint, BindingPoint>& binding_points) {
+ if (auto res = BindingRemapper(module, binding_points); res != Success) {
+ return;
+ }
+
+ Capabilities capabilities;
+ if (auto res = Validate(module, capabilities); res != Success) {
+ TINT_ICE() << "result of BindingRemapper failed IR validation\n" << res.Failure();
+ }
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
+
+TINT_IR_MODULE_FUZZER(tint::core::ir::transform::BindingRemapperFuzzer);
diff --git a/src/tint/lang/glsl/writer/writer_ast_fuzz.cc b/src/tint/lang/glsl/writer/writer_ast_fuzz.cc
index 096c259..577f4f2 100644
--- a/src/tint/lang/glsl/writer/writer_ast_fuzz.cc
+++ b/src/tint/lang/glsl/writer/writer_ast_fuzz.cc
@@ -40,6 +40,10 @@
return;
}
+ if (options.first_instance_offset > 0x10000) {
+ return; // Excessive values can cause OOM.
+ }
+
auto inspector = tint::inspector::Inspector(program);
auto entrypoints = inspector.GetEntryPoints();
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index 9673668..4720efc 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -588,10 +588,8 @@
current_label_id_ = current_function_.label_id();
TINT_DEFER(current_function_ = Function());
- for (auto* stmt : func_ast->body->statements) {
- if (!GenerateStatement(stmt)) {
- return false;
- }
+ if (!GenerateBlockStatementWithoutScoping(func_ast->body)) {
+ return false;
}
if (InsideBasicBlock()) {
@@ -2198,6 +2196,9 @@
bool Builder::GenerateBlockStatementWithoutScoping(const ast::BlockStatement* stmt) {
for (auto* block_stmt : stmt->statements) {
+ if (!InsideBasicBlock()) {
+ break;
+ }
if (!GenerateStatement(block_stmt)) {
return false;
}
diff --git a/src/tint/lang/wgsl/ast/transform/BUILD.cmake b/src/tint/lang/wgsl/ast/transform/BUILD.cmake
index 4f2817f..05dae10 100644
--- a/src/tint/lang/wgsl/ast/transform/BUILD.cmake
+++ b/src/tint/lang/wgsl/ast/transform/BUILD.cmake
@@ -248,6 +248,8 @@
# Condition: TINT_BUILD_WGSL_READER
################################################################################
tint_add_target(tint_lang_wgsl_ast_transform_fuzz fuzz
+ lang/wgsl/ast/transform/add_block_attribute_fuzz.cc
+ lang/wgsl/ast/transform/add_empty_entry_point_fuzz.cc
lang/wgsl/ast/transform/zero_init_workgroup_memory_fuzz.cc
)
diff --git a/src/tint/lang/wgsl/ast/transform/BUILD.gn b/src/tint/lang/wgsl/ast/transform/BUILD.gn
index f0626b0..7b2d224 100644
--- a/src/tint/lang/wgsl/ast/transform/BUILD.gn
+++ b/src/tint/lang/wgsl/ast/transform/BUILD.gn
@@ -237,7 +237,11 @@
}
if (tint_build_wgsl_reader) {
tint_fuzz_source_set("fuzz") {
- sources = [ "zero_init_workgroup_memory_fuzz.cc" ]
+ sources = [
+ "add_block_attribute_fuzz.cc",
+ "add_empty_entry_point_fuzz.cc",
+ "zero_init_workgroup_memory_fuzz.cc",
+ ]
deps = [
"${tint_src_dir}/lang/core",
"${tint_src_dir}/lang/core/constant",
diff --git a/src/tint/lang/wgsl/ast/transform/add_block_attribute.h b/src/tint/lang/wgsl/ast/transform/add_block_attribute.h
index 544e1ee..5e9892b 100644
--- a/src/tint/lang/wgsl/ast/transform/add_block_attribute.h
+++ b/src/tint/lang/wgsl/ast/transform/add_block_attribute.h
@@ -40,7 +40,7 @@
class AddBlockAttribute final : public Castable<AddBlockAttribute, Transform> {
public:
/// BlockAttribute is an InternalAttribute that is used to decorate a
- // structure that is used as a buffer in SPIR-V or GLSL.
+ /// structure that is used as a buffer in SPIR-V or GLSL.
class BlockAttribute final : public Castable<BlockAttribute, InternalAttribute> {
public:
/// Constructor
diff --git a/src/tint/lang/wgsl/ast/transform/add_block_attribute_fuzz.cc b/src/tint/lang/wgsl/ast/transform/add_block_attribute_fuzz.cc
new file mode 100644
index 0000000..7898243
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/transform/add_block_attribute_fuzz.cc
@@ -0,0 +1,46 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/cmd/fuzz/wgsl/fuzz.h"
+#include "src/tint/lang/wgsl/ast/transform/add_block_attribute.h"
+
+namespace tint::ast::transform {
+namespace {
+
+void AddBlockAttributeFuzzer(const Program& program) {
+ DataMap outputs;
+ if (auto result = AddBlockAttribute{}.Apply(program, DataMap{}, outputs)) {
+ if (!result->IsValid()) {
+ TINT_ICE() << "AddBlockAttribute returned invalid program:\n" << result->Diagnostics();
+ }
+ }
+}
+
+} // namespace
+} // namespace tint::ast::transform
+
+TINT_WGSL_PROGRAM_FUZZER(tint::ast::transform::AddBlockAttributeFuzzer);
diff --git a/src/tint/lang/wgsl/ast/transform/add_empty_entry_point_fuzz.cc b/src/tint/lang/wgsl/ast/transform/add_empty_entry_point_fuzz.cc
new file mode 100644
index 0000000..cee5702
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/transform/add_empty_entry_point_fuzz.cc
@@ -0,0 +1,46 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/cmd/fuzz/wgsl/fuzz.h"
+#include "src/tint/lang/wgsl/ast/transform/add_empty_entry_point.h"
+
+namespace tint::ast::transform {
+namespace {
+
+void AddEmptyEntryPointFuzzer(const Program& program) {
+ DataMap outputs;
+ if (auto result = AddEmptyEntryPoint{}.Apply(program, DataMap{}, outputs)) {
+ if (!result->IsValid()) {
+ TINT_ICE() << "AddEmptyEntryPoint returned invalid program:\n" << result->Diagnostics();
+ }
+ }
+}
+
+} // namespace
+} // namespace tint::ast::transform
+
+TINT_WGSL_PROGRAM_FUZZER(tint::ast::transform::AddEmptyEntryPointFuzzer);
diff --git a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_fuzz.cc b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_fuzz.cc
index 7fe5685..2e3650d 100644
--- a/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_fuzz.cc
+++ b/src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory_fuzz.cc
@@ -32,7 +32,7 @@
namespace tint::ast::transform {
-void ZeroInitWorkgroupMemoryFuzzer(const tint::Program& program) {
+void ZeroInitWorkgroupMemoryFuzzer(const Program& program) {
if (program.AST().HasOverrides()) {
return;
}
diff --git a/src/tint/lang/wgsl/inspector/inspector.cc b/src/tint/lang/wgsl/inspector/inspector.cc
index f00ab13..1cdbcac 100644
--- a/src/tint/lang/wgsl/inspector/inspector.cc
+++ b/src/tint/lang/wgsl/inspector/inspector.cc
@@ -1029,24 +1029,14 @@
std::unordered_set<BindingPoint> seen = {};
- auto sample_type_for_call_and_type = [](wgsl::BuiltinFn builtin, const core::type::Type* ty,
- const Vector<const ast::Expression*, 8>& args) {
- if (builtin == wgsl::BuiltinFn::kTextureNumLevels) {
+ auto sample_type_for_call_and_type = [](wgsl::BuiltinFn builtin) {
+ if (builtin == wgsl::BuiltinFn::kTextureNumLevels ||
+ builtin == wgsl::BuiltinFn::kTextureDimensions ||
+ builtin == wgsl::BuiltinFn::kTextureLoad) {
return TextureQueryType::kTextureNumLevels;
}
- if (builtin == wgsl::BuiltinFn::kTextureDimensions && args.Length() > 1) {
- // When textureDimension takes level as the input,
- // it requires calls to textureNumLevels to clamp mip levels.
- return TextureQueryType::kTextureNumLevels;
- }
- if (builtin == wgsl::BuiltinFn::kTextureLoad) {
- if (!ty->UnwrapRef()
- ->IsAnyOf<core::type::MultisampledTexture,
- core::type::DepthMultisampledTexture>()) {
- return TextureQueryType::kTextureNumLevels;
- }
- }
+ TINT_ASSERT(builtin == wgsl::BuiltinFn::kTextureNumSamples);
return TextureQueryType::kTextureNumSamples;
};
@@ -1101,19 +1091,27 @@
call->Target(),
[&](const sem::BuiltinFn* builtin) {
if (builtin->Fn() != wgsl::BuiltinFn::kTextureNumLevels &&
- builtin->Fn() != wgsl::BuiltinFn::kTextureDimensions &&
builtin->Fn() != wgsl::BuiltinFn::kTextureNumSamples &&
- builtin->Fn() != wgsl::BuiltinFn::kTextureLoad) {
+ builtin->Fn() != wgsl::BuiltinFn::kTextureLoad &&
+ // When textureDimension takes level as the input,
+ // it requires calls to textureNumLevels to clamp mip levels.
+ !(builtin->Fn() == wgsl::BuiltinFn::kTextureDimensions &&
+ call->Declaration()->args.Length() > 1)) {
return;
}
auto* texture_expr = call->Declaration()->args[0];
auto* texture_sem = sem.GetVal(texture_expr)->RootIdentifier();
TINT_ASSERT(texture_sem);
+ if (builtin->Fn() == wgsl::BuiltinFn::kTextureLoad &&
+ texture_sem->Type()
+ ->UnwrapRef()
+ ->IsAnyOf<core::type::MultisampledTexture,
+ core::type::DepthMultisampledTexture>()) {
+ return;
+ }
- auto type = sample_type_for_call_and_type(builtin->Fn(), texture_sem->Type(),
- call->Declaration()->args);
-
+ auto type = sample_type_for_call_and_type(builtin->Fn());
tint::Switch(
texture_sem, //
[&](const sem::GlobalVariable* global) { save_if_needed(global, type); },
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc
index 8d1a4b3..7468ccc 100644
--- a/src/tint/lang/wgsl/inspector/inspector_test.cc
+++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -3931,22 +3931,20 @@
Inspector& inspector = Initialize(shader);
auto info = inspector.GetTextureQueries("main");
- ASSERT_EQ(1u, info.size());
-
- EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[0].type);
- EXPECT_EQ(2u, info[0].group);
- EXPECT_EQ(3u, info[0].binding);
+ ASSERT_EQ(0u, info.size());
}
TEST_F(InspectorTextureTest, TextureLoadMultipleInEP) {
std::string shader = R"(
@group(2) @binding(3) var tex1: texture_2d<f32>;
@group(1) @binding(4) var tex2: texture_multisampled_2d<f32>;
+@group(0) @binding(1) var tex3: texture_2d<f32>;
@compute @workgroup_size(1)
fn main() {
let num1 = textureLoad(tex1, vec2(0, 0), 0);
let num2 = textureLoad(tex2, vec2(0, 0), 0);
+ let num3 = textureLoad(tex3, vec2(0, 0), 0);
})";
Inspector& inspector = Initialize(shader);
@@ -3957,9 +3955,9 @@
EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
EXPECT_EQ(2u, info[0].group);
EXPECT_EQ(3u, info[0].binding);
- EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[1].type);
- EXPECT_EQ(1u, info[1].group);
- EXPECT_EQ(4u, info[1].binding);
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[1].type);
+ EXPECT_EQ(0u, info[1].group);
+ EXPECT_EQ(1u, info[1].binding);
}
TEST_F(InspectorTextureTest, TextureInSubfunction) {
