benchmarks: Add a basic set of benchmarks
Add google benchmark to the DEPs.
Implement a basic set of benchmarks for each of the writers and the WGSL parser.
Add build rules for CMake. GN build rules TODO.
Add a simple go tool (ported from Marl) to diff two benchmarks. Less
noisy than the one provided by google benchmark.
Bug: tint:1378
Change-Id: I73cf92c5d9fd2d3bfac8f264864fd774afbd5d01
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/76840
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 36a743f..6049b1b 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -81,6 +81,7 @@
option_if_not_defined(TINT_BUILD_SPIRV_TOOLS_FUZZER "Build SPIRV-Tools fuzzer" OFF)
option_if_not_defined(TINT_BUILD_AST_FUZZER "Build AST fuzzer" OFF)
option_if_not_defined(TINT_BUILD_REGEX_FUZZER "Build regex fuzzer" OFF)
+option_if_not_defined(TINT_BUILD_BENCHMARKS "Build benchmarks" OFF)
option_if_not_defined(TINT_BUILD_TESTS "Build tests" ${TINT_BUILD_TESTS_DEFAULT})
option_if_not_defined(TINT_BUILD_AS_OTHER_OS "Override OS detection to force building of *_other.cc files" OFF)
option_if_not_defined(TINT_BUILD_REMOTE_COMPILE "Build the remote-compile tool for validating shaders on a remote machine" OFF)
@@ -111,6 +112,7 @@
message(STATUS "Tint build SPIRV-Tools fuzzer: ${TINT_BUILD_SPIRV_TOOLS_FUZZER}")
message(STATUS "Tint build AST fuzzer: ${TINT_BUILD_AST_FUZZER}")
message(STATUS "Tint build regex fuzzer: ${TINT_BUILD_REGEX_FUZZER}")
+message(STATUS "Tint build benchmarks: ${TINT_BUILD_BENCHMARKS}")
message(STATUS "Tint build tests: ${TINT_BUILD_TESTS}")
message(STATUS "Tint build with ASAN: ${TINT_ENABLE_ASAN}")
message(STATUS "Tint build with MSAN: ${TINT_ENABLE_MSAN}")
@@ -254,7 +256,7 @@
endif(DOXYGEN_FOUND)
endif()
-function(tint_default_compile_options TARGET)
+function(tint_core_compile_options TARGET)
target_include_directories(${TARGET} PUBLIC "${TINT_ROOT_SOURCE_DIR}")
target_include_directories(${TARGET} PUBLIC "${TINT_ROOT_SOURCE_DIR}/include")
@@ -263,20 +265,48 @@
"${TINT_ROOT_SOURCE_DIR}/third_party/spirv-headers/include")
endif()
- target_compile_definitions(${TARGET} PUBLIC
- -DTINT_BUILD_SPV_READER=$<BOOL:${TINT_BUILD_SPV_READER}>)
- target_compile_definitions(${TARGET} PUBLIC
- -DTINT_BUILD_WGSL_READER=$<BOOL:${TINT_BUILD_WGSL_READER}>)
- target_compile_definitions(${TARGET} PUBLIC
- -DTINT_BUILD_GLSL_WRITER=$<BOOL:${TINT_BUILD_GLSL_WRITER}>)
- target_compile_definitions(${TARGET} PUBLIC
- -DTINT_BUILD_HLSL_WRITER=$<BOOL:${TINT_BUILD_HLSL_WRITER}>)
- target_compile_definitions(${TARGET} PUBLIC
- -DTINT_BUILD_MSL_WRITER=$<BOOL:${TINT_BUILD_MSL_WRITER}>)
- target_compile_definitions(${TARGET} PUBLIC
- -DTINT_BUILD_SPV_WRITER=$<BOOL:${TINT_BUILD_SPV_WRITER}>)
- target_compile_definitions(${TARGET} PUBLIC
- -DTINT_BUILD_WGSL_WRITER=$<BOOL:${TINT_BUILD_WGSL_WRITER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_SPV_READER=$<BOOL:${TINT_BUILD_SPV_READER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_READER=$<BOOL:${TINT_BUILD_WGSL_READER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_GLSL_WRITER=$<BOOL:${TINT_BUILD_GLSL_WRITER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_HLSL_WRITER=$<BOOL:${TINT_BUILD_HLSL_WRITER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_MSL_WRITER=$<BOOL:${TINT_BUILD_MSL_WRITER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_SPV_WRITER=$<BOOL:${TINT_BUILD_SPV_WRITER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_WRITER=$<BOOL:${TINT_BUILD_WGSL_WRITER}>)
+
+ if (COMPILER_IS_LIKE_GNU)
+ target_compile_options(${TARGET} PRIVATE
+ -std=c++17
+ -fno-exceptions
+ -fno-rtti
+ )
+
+ if (${TINT_ENABLE_MSAN})
+ target_compile_options(${TARGET} PRIVATE -fsanitize=memory)
+ target_link_options(${TARGET} PRIVATE -fsanitize=memory)
+ elseif (${TINT_ENABLE_ASAN})
+ target_compile_options(${TARGET} PRIVATE -fsanitize=address)
+ target_link_options(${TARGET} PRIVATE -fsanitize=address)
+ elseif (${TINT_ENABLE_UBSAN})
+ target_compile_options(${TARGET} PRIVATE -fsanitize=undefined)
+ target_link_options(${TARGET} PRIVATE -fsanitize=undefined)
+ endif()
+ endif(COMPILER_IS_LIKE_GNU)
+
+ if (TINT_EMIT_COVERAGE)
+ if(CMAKE_CXX_COMPILER_ID MATCHES "GNU")
+ target_compile_options(${TARGET} PRIVATE "--coverage")
+ target_link_options(${TARGET} PRIVATE "gcov")
+ elseif(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
+ target_compile_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
+ target_link_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
+ else()
+ message(FATAL_ERROR "Coverage generation not supported for the ${CMAKE_CXX_COMPILER_ID} toolchain")
+ endif()
+ endif(TINT_EMIT_COVERAGE)
+endfunction()
+
+function(tint_default_compile_options TARGET)
+ tint_core_compile_options(${TARGET})
set(COMMON_GNU_OPTIONS
-Wall
@@ -299,11 +329,8 @@
-Weverything
)
- if (${COMPILER_IS_LIKE_GNU})
+ if (COMPILER_IS_LIKE_GNU)
target_compile_options(${TARGET} PRIVATE
- -std=c++17
- -fno-exceptions
- -fno-rtti
-pedantic-errors
${COMMON_GNU_OPTIONS}
)
@@ -314,30 +341,7 @@
${COMMON_CLANG_OPTIONS}
)
endif()
-
- if (${TINT_ENABLE_MSAN})
- target_compile_options(${TARGET} PRIVATE -fsanitize=memory)
- target_link_options(${TARGET} PRIVATE -fsanitize=memory)
- elseif (${TINT_ENABLE_ASAN})
- target_compile_options(${TARGET} PRIVATE -fsanitize=address)
- target_link_options(${TARGET} PRIVATE -fsanitize=address)
- elseif (${TINT_ENABLE_UBSAN})
- target_compile_options(${TARGET} PRIVATE -fsanitize=undefined)
- target_link_options(${TARGET} PRIVATE -fsanitize=undefined)
- endif()
- endif()
-
- if (${TINT_EMIT_COVERAGE})
- if(CMAKE_CXX_COMPILER_ID MATCHES "GNU")
- target_compile_options(${TARGET} PRIVATE "--coverage")
- target_link_options(${TARGET} PRIVATE "gcov")
- elseif(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
- target_compile_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
- target_link_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
- else()
- message(FATAL_ERROR "Coverage generation not supported for the ${CMAKE_CXX_COMPILER_ID} toolchain")
- endif()
- endif()
+ endif(COMPILER_IS_LIKE_GNU)
if (MSVC)
# Specify /EHs for exception handling.
@@ -379,7 +383,6 @@
)
endif()
endif()
-
endfunction()
add_subdirectory(third_party)
diff --git a/DEPS b/DEPS
index 6bb6635..fbd2af5 100644
--- a/DEPS
+++ b/DEPS
@@ -10,6 +10,7 @@
'chromium_git': 'https://chromium.googlesource.com',
'github': '/external/github.com',
+ 'benchmark_revision': 'e991355c02b93fe17713efe04cbc2e278e00fdbd',
'build_revision': '555c8b467c21e2c4b22d00e87e3faa0431df9ac2',
'buildtools_revision': 'f78b4b9f33bd8ef9944d5ce643daff1c31880189',
'catapult_revision': 'fa35beefb3429605035f98211ddb8750dee6a13d',
@@ -101,6 +102,9 @@
'third_party/catapult': Var('chromium_git') + '/catapult.git@' +
Var('catapult_revision'),
+ 'third_party/benchmark': Var('chromium_git') + Var('github') +
+ '/google/benchmark.git@' + Var('benchmark_revision'),
+
'third_party/googletest': Var('chromium_git') + Var('github') +
'/google/googletest.git@' + Var('googletest_revision'),
diff --git a/fuzzers/tint_black_box_fuzz_target.cc b/fuzzers/tint_black_box_fuzz_target.cc
index cefaa76..f2f2f68 100644
--- a/fuzzers/tint_black_box_fuzz_target.cc
+++ b/fuzzers/tint_black_box_fuzz_target.cc
@@ -56,13 +56,7 @@
}
fseek(file, 0, SEEK_END);
- auto tell_file_size = ftell(file);
- if (tell_file_size <= 0) {
- std::cerr << "Input file of incorrect size: " << input_file << std::endl;
- fclose(file);
- return {};
- }
- const auto file_size = static_cast<size_t>(tell_file_size);
+ const auto file_size = static_cast<size_t>(ftell(file));
if (0 != (file_size % sizeof(T))) {
std::cerr << "File " << input_file
<< " does not contain an integral number of objects: "
diff --git a/kokoro/linux/docker.sh b/kokoro/linux/docker.sh
index cef68d1..9818060 100755
--- a/kokoro/linux/docker.sh
+++ b/kokoro/linux/docker.sh
@@ -84,7 +84,8 @@
COMMON_CMAKE_FLAGS=""
COMMON_CMAKE_FLAGS+=" -DCMAKE_BUILD_TYPE=${BUILD_TYPE}"
- COMMON_CMAKE_FLAGS+=" -DTINT_DOCS_WARN_AS_ERROR=ON"
+ COMMON_CMAKE_FLAGS+=" -DTINT_DOCS_WARN_AS_ERROR=1"
+ COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_BENCHMARKS=1"
if [ "$BUILD_TOOLCHAIN" == "clang" ]; then
using clang-10.0.0
@@ -155,9 +156,9 @@
status "Checking disabling all readers and writers also builds"
show_cmds
- cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=OFF -DTINT_BUILD_SPV_WRITER=OFF -DTINT_BUILD_WGSL_READER=OFF -DTINT_BUILD_WGSL_WRITER=OFF -DTINT_BUILD_MSL_WRITER=OFF -DTINT_BUILD_HLSL_WRITER=OFF
+ cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=OFF -DTINT_BUILD_SPV_WRITER=OFF -DTINT_BUILD_WGSL_READER=OFF -DTINT_BUILD_WGSL_WRITER=OFF -DTINT_BUILD_MSL_WRITER=OFF -DTINT_BUILD_HLSL_WRITER=OFF -DTINT_BUILD_BENCHMARKS=OFF
cmake --build . -- --jobs=$(nproc)
- cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=ON -DTINT_BUILD_SPV_WRITER=ON -DTINT_BUILD_WGSL_READER=ON -DTINT_BUILD_WGSL_WRITER=ON -DTINT_BUILD_MSL_WRITER=ON -DTINT_BUILD_HLSL_WRITER=ON
+ cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=ON -DTINT_BUILD_SPV_WRITER=ON -DTINT_BUILD_WGSL_READER=ON -DTINT_BUILD_WGSL_WRITER=ON -DTINT_BUILD_MSL_WRITER=ON -DTINT_BUILD_HLSL_WRITER=ON -DTINT_BUILD_BENCHMARKS=ON
hide_cmds
else
status "Unsupported build system: $BUILD_SYSTEM"
diff --git a/kokoro/windows/build.bat b/kokoro/windows/build.bat
index a7d28ce..e22d5d9 100644
--- a/kokoro/windows/build.bat
+++ b/kokoro/windows/build.bat
@@ -116,7 +116,7 @@
@echo on
mkdir %BUILD_DIR%
cd /d %BUILD_DIR%
-set COMMON_CMAKE_FLAGS=-DTINT_BUILD_DOCS=O -DCMAKE_BUILD_TYPE=%BUILD_TYPE%
+set COMMON_CMAKE_FLAGS=-DTINT_BUILD_DOCS=O -DTINT_BUILD_BENCHMARKS=1 -DCMAKE_BUILD_TYPE=%BUILD_TYPE%
@echo off
call :status "Building tint"
diff --git a/samples/main.cc b/samples/main.cc
index 4348636..87f0b54 100644
--- a/samples/main.cc
+++ b/samples/main.cc
@@ -437,13 +437,7 @@
}
fseek(file, 0, SEEK_END);
- uint64_t tell_file_size = static_cast<uint64_t>(ftell(file));
- if (tell_file_size <= 0) {
- std::cerr << "Input file of incorrect size: " << input_file << std::endl;
- fclose(file);
- return {};
- }
- const auto file_size = static_cast<size_t>(tell_file_size);
+ const auto file_size = static_cast<size_t>(ftell(file));
if (0 != (file_size % sizeof(T))) {
std::cerr << "File " << input_file
<< " does not contain an integral number of objects: "
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index be42ae5..d760dcb 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -598,7 +598,10 @@
endif()
endif()
-if(${TINT_BUILD_TESTS})
+################################################################################
+# Tests
+################################################################################
+if(TINT_BUILD_TESTS)
set(TINT_TEST_SRCS
ast/alias_test.cc
ast/array_test.cc
@@ -1114,6 +1117,7 @@
endif()
add_executable(tint_unittests ${TINT_TEST_SRCS})
+ set_target_properties(${target} PROPERTIES FOLDER "Tests")
if(NOT MSVC)
target_compile_options(tint_unittests PRIVATE
@@ -1133,4 +1137,41 @@
endif()
add_test(NAME tint_unittests COMMAND tint_unittests)
-endif()
+endif(TINT_BUILD_TESTS)
+
+################################################################################
+# Benchmarks
+################################################################################
+if(TINT_BUILD_BENCHMARKS)
+ if(NOT TINT_BUILD_WGSL_READER)
+ message(FATAL_ERROR "TINT_BUILD_BENCHMARKS requires TINT_BUILD_WGSL_READER")
+ endif()
+
+ set(TINT_BENCHMARK_SRC
+ "benchmark/benchmark.cc"
+ "reader/wgsl/parser_bench.cc"
+ )
+
+ if (${TINT_BUILD_GLSL_WRITER})
+ list(APPEND TINT_BENCHMARK_SRC writer/glsl/generator_bench.cc)
+ endif()
+ if (${TINT_BUILD_HLSL_WRITER})
+ list(APPEND TINT_BENCHMARK_SRC writer/hlsl/generator_bench.cc)
+ endif()
+ if (${TINT_BUILD_MSL_WRITER})
+ list(APPEND TINT_BENCHMARK_SRC writer/msl/generator_bench.cc)
+ endif()
+ if (${TINT_BUILD_SPV_WRITER})
+ list(APPEND TINT_BENCHMARK_SRC writer/spirv/generator_bench.cc)
+ endif()
+ if (${TINT_BUILD_WGSL_WRITER})
+ list(APPEND TINT_BENCHMARK_SRC writer/wgsl/generator_bench.cc)
+ endif()
+
+ add_executable(tint-benchmark ${TINT_BENCHMARK_SRC})
+ set_target_properties(${target} PROPERTIES FOLDER "Benchmarks")
+
+ tint_core_compile_options(tint-benchmark)
+
+ target_link_libraries(tint-benchmark PRIVATE benchmark::benchmark libtint)
+endif(TINT_BUILD_BENCHMARKS)
diff --git a/src/benchmark/benchmark.cc b/src/benchmark/benchmark.cc
new file mode 100644
index 0000000..813af84
--- /dev/null
+++ b/src/benchmark/benchmark.cc
@@ -0,0 +1,123 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/benchmark/benchmark.h"
+
+#include <filesystem>
+#include <sstream>
+#include <utility>
+#include <vector>
+
+namespace tint::benchmark {
+namespace {
+
+std::filesystem::path kInputFileDir;
+
+/// Copies the content from the file named `input_file` to `buffer`,
+/// assuming each element in the file is of type `T`. If any error occurs,
+/// writes error messages to the standard error stream and returns false.
+/// Assumes the size of a `T` object is divisible by its required alignment.
+/// @returns true if we successfully read the file.
+template <typename T>
+std::variant<std::vector<T>, Error> ReadFile(const std::string& input_file) {
+ FILE* file = nullptr;
+#if defined(_MSC_VER)
+ fopen_s(&file, input_file.c_str(), "rb");
+#else
+ file = fopen(input_file.c_str(), "rb");
+#endif
+ if (!file) {
+ return Error{"Failed to open " + input_file};
+ }
+
+ fseek(file, 0, SEEK_END);
+ const auto file_size = static_cast<size_t>(ftell(file));
+ if (0 != (file_size % sizeof(T))) {
+ std::stringstream err;
+ err << "File " << input_file
+ << " does not contain an integral number of objects: " << file_size
+ << " bytes in the file, require " << sizeof(T) << " bytes per object";
+ fclose(file);
+ return Error{err.str()};
+ }
+ fseek(file, 0, SEEK_SET);
+
+ std::vector<T> buffer;
+ buffer.resize(file_size / sizeof(T));
+
+ size_t bytes_read = fread(buffer.data(), 1, file_size, file);
+ fclose(file);
+ if (bytes_read != file_size) {
+ return Error{"Failed to read " + input_file};
+ }
+
+ return buffer;
+}
+
+bool FindBenchmarkInputDir() {
+ // Attempt to find the benchmark input files by searching up from the current
+ // working directory.
+ auto path = std::filesystem::current_path();
+ while (std::filesystem::is_directory(path)) {
+ auto test = path / "test" / "benchmark";
+ if (std::filesystem::is_directory(test)) {
+ kInputFileDir = test;
+ return true;
+ }
+ auto parent = path.parent_path();
+ if (path == parent) {
+ break;
+ }
+ path = parent;
+ }
+ return false;
+}
+
+} // namespace
+
+std::variant<tint::Source::File, Error> LoadInputFile(std::string name) {
+ auto path = (kInputFileDir / name).string();
+ auto data = ReadFile<uint8_t>(path);
+ if (auto* buf = std::get_if<std::vector<uint8_t>>(&data)) {
+ return tint::Source::File(path, std::string(buf->begin(), buf->end()));
+ }
+ return std::get<Error>(data);
+}
+
+std::variant<ProgramAndFile, Error> LoadProgram(std::string name) {
+ auto res = benchmark::LoadInputFile(name);
+ if (auto err = std::get_if<benchmark::Error>(&res)) {
+ return *err;
+ }
+ auto& file = std::get<Source::File>(res);
+ auto program = reader::wgsl::Parse(&file);
+ if (program.Diagnostics().contains_errors()) {
+ return Error{program.Diagnostics().str()};
+ }
+ return ProgramAndFile{std::move(program), std::move(file)};
+}
+
+} // namespace tint::benchmark
+
+int main(int argc, char** argv) {
+ ::benchmark::Initialize(&argc, argv);
+ if (::benchmark::ReportUnrecognizedArguments(argc, argv)) {
+ return 1;
+ }
+ if (!tint::benchmark::FindBenchmarkInputDir()) {
+ std::cerr << "failed to locate benchmark input files" << std::endl;
+ return 1;
+ }
+ ::benchmark::RunSpecifiedBenchmarks();
+}
diff --git a/src/benchmark/benchmark.h b/src/benchmark/benchmark.h
new file mode 100644
index 0000000..40de31d
--- /dev/null
+++ b/src/benchmark/benchmark.h
@@ -0,0 +1,69 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_BENCHMARK_BENCHMARK_H_
+#define SRC_BENCHMARK_BENCHMARK_H_
+
+#include <memory>
+#include <string>
+#include <variant> // NOLINT: Found C system header after C++ system header.
+
+#include "benchmark/benchmark.h"
+#include "src/utils/concat.h"
+#include "tint/tint.h"
+
+namespace tint::benchmark {
+
+/// Error indicates an operation did not complete successfully.
+struct Error {
+ /// The error message.
+ std::string msg;
+};
+
+/// ProgramAndFile holds a Program and a Source::File.
+struct ProgramAndFile {
+ /// The tint program parsed from file.
+ Program program;
+ /// The source file
+ Source::File file;
+};
+
+/// LoadInputFile attempts to load a benchmark input file with the given file
+/// name.
+/// @param name the file name
+/// @returns either the loaded Source::File or an Error
+std::variant<Source::File, Error> LoadInputFile(std::string name);
+
+/// LoadInputFile attempts to load a benchmark input program with the given file
+/// name.
+/// @param name the file name
+/// @returns either the loaded Program or an Error
+std::variant<ProgramAndFile, Error> LoadProgram(std::string name);
+
+/// Declares a benchmark with the given function and WGSL file name
+#define TINT_BENCHMARK_WGSL_PROGRAM(FUNC, WGSL_NAME) \
+ BENCHMARK_CAPTURE(FUNC, WGSL_NAME, WGSL_NAME);
+
+/// Declares a set of benchmarks for the given function using a list of WGSL
+/// files in `<tint>/test/benchmark`.
+#define TINT_BENCHMARK_WGSL_PROGRAMS(FUNC) \
+ TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "empty.wgsl"); \
+ TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "particles.wgsl"); \
+ TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "simple_fragment.wgsl"); \
+ TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "simple_vertex.wgsl"); \
+ TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "simple_compute.wgsl");
+
+} // namespace tint::benchmark
+
+#endif // SRC_BENCHMARK_BENCHMARK_H_
diff --git a/src/reader/wgsl/parser_bench.cc b/src/reader/wgsl/parser_bench.cc
new file mode 100644
index 0000000..95fa446
--- /dev/null
+++ b/src/reader/wgsl/parser_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::reader::wgsl {
+namespace {
+
+void ParseWGSL(::benchmark::State& state, std::string input_name) {
+ auto res = benchmark::LoadInputFile(input_name);
+ if (auto err = std::get_if<benchmark::Error>(&res)) {
+ state.SkipWithError(err->msg.c_str());
+ return;
+ }
+ auto& file = std::get<Source::File>(res);
+ for (auto _ : state) {
+ auto res = Parse(&file);
+ if (res.Diagnostics().contains_errors()) {
+ state.SkipWithError(res.Diagnostics().str().c_str());
+ }
+ }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(ParseWGSL);
+
+} // namespace
+} // namespace tint::reader::wgsl
diff --git a/src/writer/glsl/generator_bench.cc b/src/writer/glsl/generator_bench.cc
new file mode 100644
index 0000000..2ac109f
--- /dev/null
+++ b/src/writer/glsl/generator_bench.cc
@@ -0,0 +1,50 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/ast/module.h"
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::glsl {
+namespace {
+
+void GenerateGLSL(::benchmark::State& state, std::string input_name) {
+ auto res = benchmark::LoadProgram(input_name);
+ if (auto err = std::get_if<benchmark::Error>(&res)) {
+ state.SkipWithError(err->msg.c_str());
+ return;
+ }
+ auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+ std::vector<std::string> entry_points;
+ for (auto& fn : program.AST().Functions()) {
+ if (fn->IsEntryPoint()) {
+ entry_points.emplace_back(program.Symbols().NameFor(fn->symbol));
+ }
+ }
+
+ for (auto _ : state) {
+ for (auto& ep : entry_points) {
+ auto res = Generate(&program, {}, ep);
+ if (!res.error.empty()) {
+ state.SkipWithError(res.error.c_str());
+ }
+ }
+ }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateGLSL);
+
+} // namespace
+} // namespace tint::writer::glsl
diff --git a/src/writer/hlsl/generator_bench.cc b/src/writer/hlsl/generator_bench.cc
new file mode 100644
index 0000000..e9168af
--- /dev/null
+++ b/src/writer/hlsl/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::hlsl {
+namespace {
+
+void GenerateHLSL(::benchmark::State& state, std::string input_name) {
+ auto res = benchmark::LoadProgram(input_name);
+ if (auto err = std::get_if<benchmark::Error>(&res)) {
+ state.SkipWithError(err->msg.c_str());
+ return;
+ }
+ auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+ for (auto _ : state) {
+ auto res = Generate(&program, {});
+ if (!res.error.empty()) {
+ state.SkipWithError(res.error.c_str());
+ }
+ }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateHLSL);
+
+} // namespace
+} // namespace tint::writer::hlsl
diff --git a/src/writer/msl/generator_bench.cc b/src/writer/msl/generator_bench.cc
new file mode 100644
index 0000000..d2b53df
--- /dev/null
+++ b/src/writer/msl/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::msl {
+namespace {
+
+void GenerateMSL(::benchmark::State& state, std::string input_name) {
+ auto res = benchmark::LoadProgram(input_name);
+ if (auto err = std::get_if<benchmark::Error>(&res)) {
+ state.SkipWithError(err->msg.c_str());
+ return;
+ }
+ auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+ for (auto _ : state) {
+ auto res = Generate(&program, {});
+ if (!res.error.empty()) {
+ state.SkipWithError(res.error.c_str());
+ }
+ }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateMSL);
+
+} // namespace
+} // namespace tint::writer::msl
diff --git a/src/writer/spirv/generator_bench.cc b/src/writer/spirv/generator_bench.cc
new file mode 100644
index 0000000..fb937d1
--- /dev/null
+++ b/src/writer/spirv/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::spirv {
+namespace {
+
+void GenerateSPIRV(::benchmark::State& state, std::string input_name) {
+ auto res = benchmark::LoadProgram(input_name);
+ if (auto err = std::get_if<benchmark::Error>(&res)) {
+ state.SkipWithError(err->msg.c_str());
+ return;
+ }
+ auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+ for (auto _ : state) {
+ auto res = Generate(&program, {});
+ if (!res.error.empty()) {
+ state.SkipWithError(res.error.c_str());
+ }
+ }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateSPIRV);
+
+} // namespace
+} // namespace tint::writer::spirv
diff --git a/src/writer/wgsl/generator_bench.cc b/src/writer/wgsl/generator_bench.cc
new file mode 100644
index 0000000..524d9cf
--- /dev/null
+++ b/src/writer/wgsl/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::wgsl {
+namespace {
+
+void GenerateWGSL(::benchmark::State& state, std::string input_name) {
+ auto res = benchmark::LoadProgram(input_name);
+ if (auto err = std::get_if<benchmark::Error>(&res)) {
+ state.SkipWithError(err->msg.c_str());
+ return;
+ }
+ auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+ for (auto _ : state) {
+ auto res = Generate(&program, {});
+ if (!res.error.empty()) {
+ state.SkipWithError(res.error.c_str());
+ }
+ }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateWGSL);
+
+} // namespace
+} // namespace tint::writer::wgsl
diff --git a/test/benchmark/empty.wgsl b/test/benchmark/empty.wgsl
new file mode 100644
index 0000000..e69de29
--- /dev/null
+++ b/test/benchmark/empty.wgsl
diff --git a/test/benchmark/empty.wgsl.expected.hlsl b/test/benchmark/empty.wgsl.expected.hlsl
new file mode 100644
index 0000000..da1baa5
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.hlsl
@@ -0,0 +1,4 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+ return;
+}
diff --git a/test/benchmark/empty.wgsl.expected.msl b/test/benchmark/empty.wgsl.expected.msl
new file mode 100644
index 0000000..466ceaa
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.msl
@@ -0,0 +1,3 @@
+#include <metal_stdlib>
+
+using namespace metal;
diff --git a/test/benchmark/empty.wgsl.expected.spvasm b/test/benchmark/empty.wgsl.expected.spvasm
new file mode 100644
index 0000000..65bef94
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.spvasm
@@ -0,0 +1,16 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 5
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+ OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpName %unused_entry_point "unused_entry_point"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+%unused_entry_point = OpFunction %void None %1
+ %4 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/benchmark/empty.wgsl.expected.wgsl b/test/benchmark/empty.wgsl.expected.wgsl
new file mode 100644
index 0000000..e69de29
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.wgsl
diff --git a/test/benchmark/particles.wgsl b/test/benchmark/particles.wgsl
new file mode 100644
index 0000000..4734a37
--- /dev/null
+++ b/test/benchmark/particles.wgsl
@@ -0,0 +1,183 @@
+////////////////////////////////////////////////////////////////////////////////
+// Utilities
+////////////////////////////////////////////////////////////////////////////////
+var<private> rand_seed : vec2<f32>;
+
+fn rand() -> f32 {
+ rand_seed.x = fract(cos(dot(rand_seed, vec2<f32>(23.14077926, 232.61690225))) * 136.8168);
+ rand_seed.y = fract(cos(dot(rand_seed, vec2<f32>(54.47856553, 345.84153136))) * 534.7645);
+ return rand_seed.y;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Vertex shader
+////////////////////////////////////////////////////////////////////////////////
+struct RenderParams {
+ modelViewProjectionMatrix : mat4x4<f32>;
+ right : vec3<f32>;
+ up : vec3<f32>;
+};
+[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+
+struct VertexInput {
+ [[location(0)]] position : vec3<f32>;
+ [[location(1)]] color : vec4<f32>;
+ [[location(2)]] quad_pos : vec2<f32>; // -1..+1
+};
+
+struct VertexOutput {
+ [[builtin(position)]] position : vec4<f32>;
+ [[location(0)]] color : vec4<f32>;
+ [[location(1)]] quad_pos : vec2<f32>; // -1..+1
+};
+
+[[stage(vertex)]]
+fn vs_main(in : VertexInput) -> VertexOutput {
+ var quad_pos = mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos;
+ var position = in.position + quad_pos * 0.01;
+ var out : VertexOutput;
+ out.position = render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0);
+ out.color = in.color;
+ out.quad_pos = in.quad_pos;
+ return out;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Fragment shader
+////////////////////////////////////////////////////////////////////////////////
+[[stage(fragment)]]
+fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+ var color = in.color;
+ // Apply a circular particle alpha mask
+ color.a = color.a * max(1.0 - length(in.quad_pos), 0.0);
+ return color;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Simulation Compute shader
+////////////////////////////////////////////////////////////////////////////////
+struct SimulationParams {
+ deltaTime : f32;
+ seed : vec4<f32>;
+};
+
+struct Particle {
+ position : vec3<f32>;
+ lifetime : f32;
+ color : vec4<f32>;
+ velocity : vec3<f32>;
+};
+
+struct Particles {
+ particles : array<Particle>;
+};
+
+[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
+[[binding(1), group(0)]] var<storage, read_write> data : Particles;
+[[binding(2), group(0)]] var texture : texture_2d<f32>;
+
+[[stage(compute), workgroup_size(64)]]
+fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ rand_seed = (sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw;
+
+ let idx = GlobalInvocationID.x;
+ var particle = data.particles[idx];
+
+ // Apply gravity
+ particle.velocity.z = particle.velocity.z - sim_params.deltaTime * 0.5;
+
+ // Basic velocity integration
+ particle.position = particle.position + sim_params.deltaTime * particle.velocity;
+
+ // Age each particle. Fade out before vanishing.
+ particle.lifetime = particle.lifetime - sim_params.deltaTime;
+ particle.color.a = smoothStep(0.0, 0.5, particle.lifetime);
+
+ // If the lifetime has gone negative, then the particle is dead and should be
+ // respawned.
+ if (particle.lifetime < 0.0) {
+ // Use the probability map to find where the particle should be spawned.
+ // Starting with the 1x1 mip level.
+ var coord = vec2<i32>(0, 0);
+ for (var level = textureNumLevels(texture) - 1; level > 0; level = level - 1) {
+ // Load the probability value from the mip-level
+ // Generate a random number and using the probabilty values, pick the
+ // next texel in the next largest mip level:
+ //
+ // 0.0 probabilites.r probabilites.g probabilites.b 1.0
+ // | | | | |
+ // | TOP-LEFT | TOP-RIGHT | BOTTOM-LEFT | BOTTOM_RIGHT |
+ //
+ let probabilites = textureLoad(texture, coord, level);
+ let value = vec4<f32>(rand());
+ let mask = (value >= vec4<f32>(0.0, probabilites.xyz)) & (value < probabilites);
+ coord = coord * 2;
+ coord.x = coord.x + select(0, 1, any(mask.yw)); // x y
+ coord.y = coord.y + select(0, 1, any(mask.zw)); // z w
+ }
+ let uv = vec2<f32>(coord) / vec2<f32>(textureDimensions(texture));
+ particle.position = vec3<f32>((uv - 0.5) * 3.0 * vec2<f32>(1.0, -1.0), 0.0);
+ particle.color = textureLoad(texture, coord, 0);
+ particle.velocity.x = (rand() - 0.5) * 0.1;
+ particle.velocity.y = (rand() - 0.5) * 0.1;
+ particle.velocity.z = rand() * 0.3;
+ particle.lifetime = 0.5 + rand() * 2.0;
+ }
+
+ // Store the new particle value
+ data.particles[idx] = particle;
+}
+
+struct UBO {
+ width : u32;
+};
+
+struct Buffer {
+ weights : array<f32>;
+};
+
+[[binding(3), group(0)]] var<uniform> ubo : UBO;
+[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
+[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
+[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
+[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+
+////////////////////////////////////////////////////////////////////////////////
+// import_level
+//
+// Loads the alpha channel from a texel of the source image, and writes it to
+// the buf_out.weights.
+////////////////////////////////////////////////////////////////////////////////
+[[stage(compute), workgroup_size(64)]]
+fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+ _ = &buf_in;
+ let offset = coord.x + coord.y * ubo.width;
+ buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// export_level
+//
+// Loads 4 f32 weight values from buf_in.weights, and stores summed value into
+// buf_out.weights, along with the calculated 'probabilty' vec4 values into the
+// mip level of tex_out. See simulate() in particle.wgsl to understand the
+// probability logic.
+////////////////////////////////////////////////////////////////////////////////
+[[stage(compute), workgroup_size(64)]]
+fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+ if (all(coord.xy < vec2<u32>(textureDimensions(tex_out)))) {
+ let dst_offset = coord.x + coord.y * ubo.width;
+ let src_offset = coord.x*2u + coord.y*2u * ubo.width;
+
+ let a = buf_in.weights[src_offset + 0u];
+ let b = buf_in.weights[src_offset + 1u];
+ let c = buf_in.weights[src_offset + 0u + ubo.width];
+ let d = buf_in.weights[src_offset + 1u + ubo.width];
+ let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
+
+ buf_out.weights[dst_offset] = sum / 4.0;
+
+ let probabilities = vec4<f32>(a, a+b, a+b+c, sum) / max(sum, 0.0001);
+ textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
+ }
+}
diff --git a/test/benchmark/particles.wgsl.expected.hlsl b/test/benchmark/particles.wgsl.expected.hlsl
new file mode 100644
index 0000000..2a91f43
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.hlsl
@@ -0,0 +1,204 @@
+static float2 rand_seed = float2(0.0f, 0.0f);
+
+float rand() {
+ rand_seed.x = frac((cos(dot(rand_seed, float2(23.140779495f, 232.616897583f))) * 136.816802979f));
+ rand_seed.y = frac((cos(dot(rand_seed, float2(54.478565216f, 345.841522217f))) * 534.764526367f));
+ return rand_seed.y;
+}
+
+cbuffer cbuffer_render_params : register(b0, space0) {
+ uint4 render_params[6];
+};
+
+struct VertexInput {
+ float3 position;
+ float4 color;
+ float2 quad_pos;
+};
+struct VertexOutput {
+ float4 position;
+ float4 color;
+ float2 quad_pos;
+};
+struct tint_symbol_5 {
+ float3 position : TEXCOORD0;
+ float4 color : TEXCOORD1;
+ float2 quad_pos : TEXCOORD2;
+};
+struct tint_symbol_6 {
+ float4 color : TEXCOORD0;
+ float2 quad_pos : TEXCOORD1;
+ float4 position : SV_Position;
+};
+
+float4x4 tint_symbol_17(uint4 buffer[6], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 16u)) / 4;
+ const uint scalar_offset_2 = ((offset + 32u)) / 4;
+ const uint scalar_offset_3 = ((offset + 48u)) / 4;
+ return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
+}
+
+VertexOutput vs_main_inner(VertexInput tint_symbol) {
+ float3 quad_pos = mul(tint_symbol.quad_pos, float2x3(asfloat(render_params[4].xyz), asfloat(render_params[5].xyz)));
+ float3 position = (tint_symbol.position + (quad_pos * 0.01f));
+ VertexOutput tint_symbol_1 = (VertexOutput)0;
+ tint_symbol_1.position = mul(float4(position, 1.0f), tint_symbol_17(render_params, 0u));
+ tint_symbol_1.color = tint_symbol.color;
+ tint_symbol_1.quad_pos = tint_symbol.quad_pos;
+ return tint_symbol_1;
+}
+
+tint_symbol_6 vs_main(tint_symbol_5 tint_symbol_4) {
+ const VertexInput tint_symbol_32 = {tint_symbol_4.position, tint_symbol_4.color, tint_symbol_4.quad_pos};
+ const VertexOutput inner_result = vs_main_inner(tint_symbol_32);
+ tint_symbol_6 wrapper_result = (tint_symbol_6)0;
+ wrapper_result.position = inner_result.position;
+ wrapper_result.color = inner_result.color;
+ wrapper_result.quad_pos = inner_result.quad_pos;
+ return wrapper_result;
+}
+
+struct tint_symbol_8 {
+ float4 color : TEXCOORD0;
+ float2 quad_pos : TEXCOORD1;
+ float4 position : SV_Position;
+};
+struct tint_symbol_9 {
+ float4 value : SV_Target0;
+};
+
+float4 fs_main_inner(VertexOutput tint_symbol) {
+ float4 color = tint_symbol.color;
+ color.a = (color.a * max((1.0f - length(tint_symbol.quad_pos)), 0.0f));
+ return color;
+}
+
+tint_symbol_9 fs_main(tint_symbol_8 tint_symbol_7) {
+ const VertexOutput tint_symbol_33 = {tint_symbol_7.position, tint_symbol_7.color, tint_symbol_7.quad_pos};
+ const float4 inner_result_1 = fs_main_inner(tint_symbol_33);
+ tint_symbol_9 wrapper_result_1 = (tint_symbol_9)0;
+ wrapper_result_1.value = inner_result_1;
+ return wrapper_result_1;
+}
+
+struct Particle {
+ float3 position;
+ float lifetime;
+ float4 color;
+ float3 velocity;
+};
+
+cbuffer cbuffer_sim_params : register(b0, space0) {
+ uint4 sim_params[2];
+};
+RWByteAddressBuffer data : register(u1, space0);
+Texture2D<float4> tint_symbol_2 : register(t2, space0);
+
+struct tint_symbol_11 {
+ uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+Particle tint_symbol_20(RWByteAddressBuffer buffer, uint offset) {
+ const Particle tint_symbol_34 = {asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load((offset + 12u))), asfloat(buffer.Load4((offset + 16u))), asfloat(buffer.Load3((offset + 32u)))};
+ return tint_symbol_34;
+}
+
+void tint_symbol_25(RWByteAddressBuffer buffer, uint offset, Particle value) {
+ buffer.Store3((offset + 0u), asuint(value.position));
+ buffer.Store((offset + 12u), asuint(value.lifetime));
+ buffer.Store4((offset + 16u), asuint(value.color));
+ buffer.Store3((offset + 32u), asuint(value.velocity));
+}
+
+void simulate_inner(uint3 GlobalInvocationID) {
+ rand_seed = ((asfloat(sim_params[1]).xy + float2(GlobalInvocationID.xy)) * asfloat(sim_params[1]).zw);
+ const uint idx = GlobalInvocationID.x;
+ Particle particle = tint_symbol_20(data, (48u * idx));
+ particle.velocity.z = (particle.velocity.z - (asfloat(sim_params[0].x) * 0.5f));
+ particle.position = (particle.position + (asfloat(sim_params[0].x) * particle.velocity));
+ particle.lifetime = (particle.lifetime - asfloat(sim_params[0].x));
+ particle.color.a = smoothstep(0.0f, 0.5f, particle.lifetime);
+ if ((particle.lifetime < 0.0f)) {
+ int2 coord = int2(0, 0);
+ {
+ int3 tint_tmp;
+ tint_symbol_2.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+ int level = (tint_tmp.z - 1);
+ [loop] for(; (level > 0); level = (level - 1)) {
+ const float4 probabilites = tint_symbol_2.Load(int3(coord, level));
+ const float4 value = float4((rand()).xxxx);
+ const bool4 mask = ((value >= float4(0.0f, probabilites.xyz)) & (value < probabilites));
+ coord = (coord * 2);
+ coord.x = (coord.x + (any(mask.yw) ? 1 : 0));
+ coord.y = (coord.y + (any(mask.zw) ? 1 : 0));
+ }
+ }
+ int2 tint_tmp_1;
+ tint_symbol_2.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
+ const float2 uv = (float2(coord) / float2(tint_tmp_1));
+ particle.position = float3((((uv - 0.5f) * 3.0f) * float2(1.0f, -1.0f)), 0.0f);
+ particle.color = tint_symbol_2.Load(int3(coord, 0));
+ particle.velocity.x = ((rand() - 0.5f) * 0.100000001f);
+ particle.velocity.y = ((rand() - 0.5f) * 0.100000001f);
+ particle.velocity.z = (rand() * 0.300000012f);
+ particle.lifetime = (0.5f + (rand() * 2.0f));
+ }
+ tint_symbol_25(data, (48u * idx), particle);
+}
+
+[numthreads(64, 1, 1)]
+void simulate(tint_symbol_11 tint_symbol_10) {
+ simulate_inner(tint_symbol_10.GlobalInvocationID);
+ return;
+}
+
+cbuffer cbuffer_ubo : register(b3, space0) {
+ uint4 ubo[1];
+};
+ByteAddressBuffer buf_in : register(t4, space0);
+RWByteAddressBuffer buf_out : register(u5, space0);
+Texture2D<float4> tex_in : register(t6, space0);
+RWTexture2D<float4> tex_out : register(u7, space0);
+
+struct tint_symbol_13 {
+ uint3 coord : SV_DispatchThreadID;
+};
+
+void import_level_inner(uint3 coord) {
+ const uint offset = (coord.x + (coord.y * ubo[0].x));
+ buf_out.Store((4u * offset), asuint(tex_in.Load(int3(int2(coord.xy), 0)).w));
+}
+
+[numthreads(64, 1, 1)]
+void import_level(tint_symbol_13 tint_symbol_12) {
+ import_level_inner(tint_symbol_12.coord);
+ return;
+}
+
+struct tint_symbol_15 {
+ uint3 coord : SV_DispatchThreadID;
+};
+
+void export_level_inner(uint3 coord) {
+ int2 tint_tmp_2;
+ tex_out.GetDimensions(tint_tmp_2.x, tint_tmp_2.y);
+ if (all((coord.xy < uint2(tint_tmp_2)))) {
+ const uint dst_offset = (coord.x + (coord.y * ubo[0].x));
+ const uint src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo[0].x));
+ const float a_1 = asfloat(buf_in.Load((4u * (src_offset + 0u))));
+ const float b = asfloat(buf_in.Load((4u * (src_offset + 1u))));
+ const float c = asfloat(buf_in.Load((4u * ((src_offset + 0u) + ubo[0].x))));
+ const float d = asfloat(buf_in.Load((4u * ((src_offset + 1u) + ubo[0].x))));
+ const float sum = dot(float4(a_1, b, c, d), float4((1.0f).xxxx));
+ buf_out.Store((4u * dst_offset), asuint((sum / 4.0f)));
+ const float4 probabilities = (float4(a_1, (a_1 + b), ((a_1 + b) + c), sum) / max(sum, 0.0001f));
+ tex_out[int2(coord.xy)] = probabilities;
+ }
+}
+
+[numthreads(64, 1, 1)]
+void export_level(tint_symbol_15 tint_symbol_14) {
+ export_level_inner(tint_symbol_14.coord);
+ return;
+}
diff --git a/test/benchmark/particles.wgsl.expected.msl b/test/benchmark/particles.wgsl.expected.msl
new file mode 100644
index 0000000..5644e7c
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.msl
@@ -0,0 +1,175 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+ return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+ return vec<T, M>(lhs) * rhs;
+}
+
+struct RenderParams {
+ /* 0x0000 */ float4x4 modelViewProjectionMatrix;
+ /* 0x0040 */ packed_float3 right;
+ /* 0x004c */ int8_t tint_pad[4];
+ /* 0x0050 */ packed_float3 up;
+ /* 0x005c */ int8_t tint_pad_1[4];
+};
+struct VertexInput {
+ float3 position;
+ float4 color;
+ float2 quad_pos;
+};
+struct VertexOutput {
+ float4 position;
+ float4 color;
+ float2 quad_pos;
+};
+struct tint_symbol_2 {
+ float3 position [[attribute(0)]];
+ float4 color [[attribute(1)]];
+ float2 quad_pos [[attribute(2)]];
+};
+struct tint_symbol_3 {
+ float4 color [[user(locn0)]];
+ float2 quad_pos [[user(locn1)]];
+ float4 position [[position]];
+};
+struct tint_symbol_5 {
+ float4 color [[user(locn0)]];
+ float2 quad_pos [[user(locn1)]];
+};
+struct tint_symbol_6 {
+ float4 value [[color(0)]];
+};
+struct SimulationParams {
+ /* 0x0000 */ float deltaTime;
+ /* 0x0004 */ int8_t tint_pad_2[12];
+ /* 0x0010 */ float4 seed;
+};
+struct Particle {
+ /* 0x0000 */ packed_float3 position;
+ /* 0x000c */ float lifetime;
+ /* 0x0010 */ float4 color;
+ /* 0x0020 */ packed_float3 velocity;
+ /* 0x002c */ int8_t tint_pad_3[4];
+};
+struct Particles {
+ /* 0x0000 */ Particle particles[1];
+};
+struct UBO {
+ /* 0x0000 */ uint width;
+};
+struct Buffer {
+ /* 0x0000 */ float weights[1];
+};
+
+float rand(thread float2* const tint_symbol_9) {
+ (*(tint_symbol_9))[0] = fract((cos(dot(*(tint_symbol_9), float2(23.140779495f, 232.616897583f))) * 136.816802979f));
+ (*(tint_symbol_9))[1] = fract((cos(dot(*(tint_symbol_9), float2(54.478565216f, 345.841522217f))) * 534.764526367f));
+ return (*(tint_symbol_9))[1];
+}
+
+VertexOutput vs_main_inner(VertexInput in, const constant RenderParams* const tint_symbol_10) {
+ float3 quad_pos = (float2x3((*(tint_symbol_10)).right, (*(tint_symbol_10)).up) * in.quad_pos);
+ float3 position = (in.position + (quad_pos * 0.01f));
+ VertexOutput out = {};
+ out.position = ((*(tint_symbol_10)).modelViewProjectionMatrix * float4(position, 1.0f));
+ out.color = in.color;
+ out.quad_pos = in.quad_pos;
+ return out;
+}
+
+vertex tint_symbol_3 vs_main(const constant RenderParams* tint_symbol_11 [[buffer(0)]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+ VertexInput const tint_symbol_7 = {.position=tint_symbol_1.position, .color=tint_symbol_1.color, .quad_pos=tint_symbol_1.quad_pos};
+ VertexOutput const inner_result = vs_main_inner(tint_symbol_7, tint_symbol_11);
+ tint_symbol_3 wrapper_result = {};
+ wrapper_result.position = inner_result.position;
+ wrapper_result.color = inner_result.color;
+ wrapper_result.quad_pos = inner_result.quad_pos;
+ return wrapper_result;
+}
+
+float4 fs_main_inner(VertexOutput in) {
+ float4 color = in.color;
+ color[3] = (color[3] * fmax((1.0f - length(in.quad_pos)), 0.0f));
+ return color;
+}
+
+fragment tint_symbol_6 fs_main(float4 position [[position]], tint_symbol_5 tint_symbol_4 [[stage_in]]) {
+ VertexOutput const tint_symbol_8 = {.position=position, .color=tint_symbol_4.color, .quad_pos=tint_symbol_4.quad_pos};
+ float4 const inner_result_1 = fs_main_inner(tint_symbol_8);
+ tint_symbol_6 wrapper_result_1 = {};
+ wrapper_result_1.value = inner_result_1;
+ return wrapper_result_1;
+}
+
+void simulate_inner(uint3 GlobalInvocationID, thread float2* const tint_symbol_12, const constant SimulationParams* const tint_symbol_13, device Particles* const tint_symbol_14, texture2d<float, access::sample> tint_symbol_15) {
+ *(tint_symbol_12) = ((float4((*(tint_symbol_13)).seed).xy + float2(uint3(GlobalInvocationID).xy)) * float4((*(tint_symbol_13)).seed).zw);
+ uint const idx = GlobalInvocationID[0];
+ Particle particle = (*(tint_symbol_14)).particles[idx];
+ particle.velocity[2] = (particle.velocity[2] - ((*(tint_symbol_13)).deltaTime * 0.5f));
+ particle.position = (particle.position + ((*(tint_symbol_13)).deltaTime * particle.velocity));
+ particle.lifetime = (particle.lifetime - (*(tint_symbol_13)).deltaTime);
+ particle.color[3] = smoothstep(0.0f, 0.5f, particle.lifetime);
+ if ((particle.lifetime < 0.0f)) {
+ int2 coord = int2(0, 0);
+ for(int level = as_type<int>((as_type<uint>(int(tint_symbol_15.get_num_mip_levels())) - as_type<uint>(1))); (level > 0); level = as_type<int>((as_type<uint>(level) - as_type<uint>(1)))) {
+ float4 const probabilites = tint_symbol_15.read(uint2(coord), level);
+ float4 const value = float4(rand(tint_symbol_12));
+ bool4 const mask = ((value >= float4(0.0f, float4(probabilites).xyz)) & (value < probabilites));
+ coord = as_type<int2>((as_type<uint2>(coord) * as_type<uint>(2)));
+ coord[0] = as_type<int>((as_type<uint>(coord[0]) + as_type<uint>(select(0, 1, any(bool4(mask).yw)))));
+ coord[1] = as_type<int>((as_type<uint>(coord[1]) + as_type<uint>(select(0, 1, any(bool4(mask).zw)))));
+ }
+ float2 const uv = (float2(coord) / float2(int2(tint_symbol_15.get_width(), tint_symbol_15.get_height())));
+ particle.position = float3((((uv - 0.5f) * 3.0f) * float2(1.0f, -1.0f)), 0.0f);
+ particle.color = tint_symbol_15.read(uint2(coord), 0);
+ particle.velocity[0] = ((rand(tint_symbol_12) - 0.5f) * 0.100000001f);
+ particle.velocity[1] = ((rand(tint_symbol_12) - 0.5f) * 0.100000001f);
+ particle.velocity[2] = (rand(tint_symbol_12) * 0.300000012f);
+ particle.lifetime = (0.5f + (rand(tint_symbol_12) * 2.0f));
+ }
+ (*(tint_symbol_14)).particles[idx] = particle;
+}
+
+kernel void simulate(const constant SimulationParams* tint_symbol_17 [[buffer(0)]], device Particles* tint_symbol_18 [[buffer(1)]], texture2d<float, access::sample> tint_symbol_19 [[texture(0)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
+ thread float2 tint_symbol_16 = 0.0f;
+ simulate_inner(GlobalInvocationID, &(tint_symbol_16), tint_symbol_17, tint_symbol_18, tint_symbol_19);
+ return;
+}
+
+void import_level_inner(uint3 coord, const constant UBO* const tint_symbol_20, device Buffer* const tint_symbol_21, texture2d<float, access::sample> tint_symbol_22) {
+ uint const offset = (coord[0] + (coord[1] * (*(tint_symbol_20)).width));
+ (*(tint_symbol_21)).weights[offset] = tint_symbol_22.read(uint2(int2(uint3(coord).xy)), 0)[3];
+}
+
+kernel void import_level(const constant UBO* tint_symbol_23 [[buffer(2)]], device Buffer* tint_symbol_24 [[buffer(3)]], texture2d<float, access::sample> tint_symbol_25 [[texture(1)]], uint3 coord [[thread_position_in_grid]]) {
+ import_level_inner(coord, tint_symbol_23, tint_symbol_24, tint_symbol_25);
+ return;
+}
+
+void export_level_inner(uint3 coord, texture2d<float, access::write> tint_symbol_26, const constant UBO* const tint_symbol_27, const device Buffer* const tint_symbol_28, device Buffer* const tint_symbol_29) {
+ if (all((uint3(coord).xy < uint2(int2(tint_symbol_26.get_width(), tint_symbol_26.get_height()))))) {
+ uint const dst_offset = (coord[0] + (coord[1] * (*(tint_symbol_27)).width));
+ uint const src_offset = ((coord[0] * 2u) + ((coord[1] * 2u) * (*(tint_symbol_27)).width));
+ float const a_1 = (*(tint_symbol_28)).weights[(src_offset + 0u)];
+ float const b = (*(tint_symbol_28)).weights[(src_offset + 1u)];
+ float const c = (*(tint_symbol_28)).weights[((src_offset + 0u) + (*(tint_symbol_27)).width)];
+ float const d = (*(tint_symbol_28)).weights[((src_offset + 1u) + (*(tint_symbol_27)).width)];
+ float const sum = dot(float4(a_1, b, c, d), float4(1.0f));
+ (*(tint_symbol_29)).weights[dst_offset] = (sum / 4.0f);
+ float4 const probabilities = (float4(a_1, (a_1 + b), ((a_1 + b) + c), sum) / fmax(sum, 0.0001f));
+ tint_symbol_26.write(probabilities, uint2(int2(uint3(coord).xy)));
+ }
+}
+
+kernel void export_level(texture2d<float, access::write> tint_symbol_30 [[texture(2)]], const constant UBO* tint_symbol_31 [[buffer(2)]], const device Buffer* tint_symbol_32 [[buffer(4)]], device Buffer* tint_symbol_33 [[buffer(3)]], uint3 coord [[thread_position_in_grid]]) {
+ export_level_inner(coord, tint_symbol_30, tint_symbol_31, tint_symbol_32, tint_symbol_33);
+ return;
+}
+
diff --git a/test/benchmark/particles.wgsl.expected.spvasm b/test/benchmark/particles.wgsl.expected.spvasm
new file mode 100644
index 0000000..fe13031
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.spvasm
@@ -0,0 +1,646 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 423
+; Schema: 0
+ OpCapability Shader
+ OpCapability ImageQuery
+ %67 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vs_main "vs_main" %position_1 %color_1 %quad_pos_1 %position_2 %color_2 %quad_pos_2 %vertex_point_size
+ OpEntryPoint Fragment %fs_main "fs_main" %position_3 %color_3 %quad_pos_3 %value_1
+ OpEntryPoint GLCompute %simulate "simulate" %GlobalInvocationID_1
+ OpEntryPoint GLCompute %import_level "import_level" %coord_1
+ OpEntryPoint GLCompute %export_level "export_level" %coord_2
+ OpExecutionMode %fs_main OriginUpperLeft
+ OpExecutionMode %simulate LocalSize 64 1 1
+ OpExecutionMode %import_level LocalSize 64 1 1
+ OpExecutionMode %export_level LocalSize 64 1 1
+ OpName %position_1 "position_1"
+ OpName %color_1 "color_1"
+ OpName %quad_pos_1 "quad_pos_1"
+ OpName %position_2 "position_2"
+ OpName %color_2 "color_2"
+ OpName %quad_pos_2 "quad_pos_2"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %position_3 "position_3"
+ OpName %color_3 "color_3"
+ OpName %quad_pos_3 "quad_pos_3"
+ OpName %value_1 "value_1"
+ OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
+ OpName %coord_1 "coord_1"
+ OpName %coord_2 "coord_2"
+ OpName %rand_seed "rand_seed"
+ OpName %RenderParams "RenderParams"
+ OpMemberName %RenderParams 0 "modelViewProjectionMatrix"
+ OpMemberName %RenderParams 1 "right"
+ OpMemberName %RenderParams 2 "up"
+ OpName %render_params "render_params"
+ OpName %SimulationParams "SimulationParams"
+ OpMemberName %SimulationParams 0 "deltaTime"
+ OpMemberName %SimulationParams 1 "seed"
+ OpName %sim_params "sim_params"
+ OpName %Particles "Particles"
+ OpMemberName %Particles 0 "particles"
+ OpName %Particle "Particle"
+ OpMemberName %Particle 0 "position"
+ OpMemberName %Particle 1 "lifetime"
+ OpMemberName %Particle 2 "color"
+ OpMemberName %Particle 3 "velocity"
+ OpName %data "data"
+ OpName %texture "texture"
+ OpName %UBO "UBO"
+ OpMemberName %UBO 0 "width"
+ OpName %ubo "ubo"
+ OpName %Buffer "Buffer"
+ OpMemberName %Buffer 0 "weights"
+ OpName %buf_in "buf_in"
+ OpName %buf_out "buf_out"
+ OpName %tex_in "tex_in"
+ OpName %tex_out "tex_out"
+ OpName %rand "rand"
+ OpName %VertexOutput "VertexOutput"
+ OpMemberName %VertexOutput 0 "position"
+ OpMemberName %VertexOutput 1 "color"
+ OpMemberName %VertexOutput 2 "quad_pos"
+ OpName %VertexInput "VertexInput"
+ OpMemberName %VertexInput 0 "position"
+ OpMemberName %VertexInput 1 "color"
+ OpMemberName %VertexInput 2 "quad_pos"
+ OpName %vs_main_inner "vs_main_inner"
+ OpName %in "in"
+ OpName %quad_pos "quad_pos"
+ OpName %position "position"
+ OpName %out "out"
+ OpName %vs_main "vs_main"
+ OpName %fs_main_inner "fs_main_inner"
+ OpName %in_0 "in"
+ OpName %color "color"
+ OpName %fs_main "fs_main"
+ OpName %simulate_inner "simulate_inner"
+ OpName %GlobalInvocationID "GlobalInvocationID"
+ OpName %particle "particle"
+ OpName %coord "coord"
+ OpName %level "level"
+ OpName %simulate "simulate"
+ OpName %import_level_inner "import_level_inner"
+ OpName %coord_0 "coord"
+ OpName %import_level "import_level"
+ OpName %export_level_inner "export_level_inner"
+ OpName %coord_3 "coord"
+ OpName %export_level "export_level"
+ OpDecorate %position_1 Location 0
+ OpDecorate %color_1 Location 1
+ OpDecorate %quad_pos_1 Location 2
+ OpDecorate %position_2 BuiltIn Position
+ OpDecorate %color_2 Location 0
+ OpDecorate %quad_pos_2 Location 1
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %position_3 BuiltIn FragCoord
+ OpDecorate %color_3 Location 0
+ OpDecorate %quad_pos_3 Location 1
+ OpDecorate %value_1 Location 0
+ OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
+ OpDecorate %coord_1 BuiltIn GlobalInvocationId
+ OpDecorate %coord_2 BuiltIn GlobalInvocationId
+ OpDecorate %RenderParams Block
+ OpMemberDecorate %RenderParams 0 Offset 0
+ OpMemberDecorate %RenderParams 0 ColMajor
+ OpMemberDecorate %RenderParams 0 MatrixStride 16
+ OpMemberDecorate %RenderParams 1 Offset 64
+ OpMemberDecorate %RenderParams 2 Offset 80
+ OpDecorate %render_params NonWritable
+ OpDecorate %render_params Binding 0
+ OpDecorate %render_params DescriptorSet 0
+ OpDecorate %SimulationParams Block
+ OpMemberDecorate %SimulationParams 0 Offset 0
+ OpMemberDecorate %SimulationParams 1 Offset 16
+ OpDecorate %sim_params NonWritable
+ OpDecorate %sim_params Binding 0
+ OpDecorate %sim_params DescriptorSet 0
+ OpDecorate %Particles Block
+ OpMemberDecorate %Particles 0 Offset 0
+ OpMemberDecorate %Particle 0 Offset 0
+ OpMemberDecorate %Particle 1 Offset 12
+ OpMemberDecorate %Particle 2 Offset 16
+ OpMemberDecorate %Particle 3 Offset 32
+ OpDecorate %_runtimearr_Particle ArrayStride 48
+ OpDecorate %data Binding 1
+ OpDecorate %data DescriptorSet 0
+ OpDecorate %texture Binding 2
+ OpDecorate %texture DescriptorSet 0
+ OpDecorate %UBO Block
+ OpMemberDecorate %UBO 0 Offset 0
+ OpDecorate %ubo NonWritable
+ OpDecorate %ubo Binding 3
+ OpDecorate %ubo DescriptorSet 0
+ OpDecorate %Buffer Block
+ OpMemberDecorate %Buffer 0 Offset 0
+ OpDecorate %_runtimearr_float ArrayStride 4
+ OpDecorate %buf_in NonWritable
+ OpDecorate %buf_in Binding 4
+ OpDecorate %buf_in DescriptorSet 0
+ OpDecorate %buf_out Binding 5
+ OpDecorate %buf_out DescriptorSet 0
+ OpDecorate %tex_in Binding 6
+ OpDecorate %tex_in DescriptorSet 0
+ OpDecorate %tex_out NonReadable
+ OpDecorate %tex_out Binding 7
+ OpDecorate %tex_out DescriptorSet 0
+ OpMemberDecorate %VertexOutput 0 Offset 0
+ OpMemberDecorate %VertexOutput 1 Offset 16
+ OpMemberDecorate %VertexOutput 2 Offset 32
+ OpMemberDecorate %VertexInput 0 Offset 0
+ OpMemberDecorate %VertexInput 1 Offset 16
+ OpMemberDecorate %VertexInput 2 Offset 32
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%_ptr_Input_v3float = OpTypePointer Input %v3float
+ %position_1 = OpVariable %_ptr_Input_v3float Input
+ %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+ %color_1 = OpVariable %_ptr_Input_v4float Input
+ %v2float = OpTypeVector %float 2
+%_ptr_Input_v2float = OpTypePointer Input %v2float
+ %quad_pos_1 = OpVariable %_ptr_Input_v2float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %13 = OpConstantNull %v4float
+ %position_2 = OpVariable %_ptr_Output_v4float Output %13
+ %color_2 = OpVariable %_ptr_Output_v4float Output %13
+%_ptr_Output_v2float = OpTypePointer Output %v2float
+ %17 = OpConstantNull %v2float
+ %quad_pos_2 = OpVariable %_ptr_Output_v2float Output %17
+%_ptr_Output_float = OpTypePointer Output %float
+ %20 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %20
+ %position_3 = OpVariable %_ptr_Input_v4float Input
+ %color_3 = OpVariable %_ptr_Input_v4float Input
+ %quad_pos_3 = OpVariable %_ptr_Input_v2float Input
+ %value_1 = OpVariable %_ptr_Output_v4float Output %13
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
+ %coord_1 = OpVariable %_ptr_Input_v3uint Input
+ %coord_2 = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Private_v2float = OpTypePointer Private %v2float
+ %rand_seed = OpVariable %_ptr_Private_v2float Private %17
+%mat4v4float = OpTypeMatrix %v4float 4
+%RenderParams = OpTypeStruct %mat4v4float %v3float %v3float
+%_ptr_Uniform_RenderParams = OpTypePointer Uniform %RenderParams
+%render_params = OpVariable %_ptr_Uniform_RenderParams Uniform
+%SimulationParams = OpTypeStruct %float %v4float
+%_ptr_Uniform_SimulationParams = OpTypePointer Uniform %SimulationParams
+ %sim_params = OpVariable %_ptr_Uniform_SimulationParams Uniform
+ %Particle = OpTypeStruct %v3float %float %v4float %v3float
+%_runtimearr_Particle = OpTypeRuntimeArray %Particle
+ %Particles = OpTypeStruct %_runtimearr_Particle
+%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles
+ %data = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
+ %47 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_47 = OpTypePointer UniformConstant %47
+ %texture = OpVariable %_ptr_UniformConstant_47 UniformConstant
+ %UBO = OpTypeStruct %uint
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+ %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+%_runtimearr_float = OpTypeRuntimeArray %float
+ %Buffer = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_Buffer = OpTypePointer StorageBuffer %Buffer
+ %buf_in = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
+ %buf_out = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
+ %tex_in = OpVariable %_ptr_UniformConstant_47 UniformConstant
+ %59 = OpTypeImage %float 2D 0 0 0 2 Rgba8
+%_ptr_UniformConstant_59 = OpTypePointer UniformConstant %59
+ %tex_out = OpVariable %_ptr_UniformConstant_59 UniformConstant
+ %60 = OpTypeFunction %float
+ %uint_0 = OpConstant %uint 0
+%_ptr_Private_float = OpTypePointer Private %float
+%float_23_1407795 = OpConstant %float 23.1407795
+%float_232_616898 = OpConstant %float 232.616898
+ %73 = OpConstantComposite %v2float %float_23_1407795 %float_232_616898
+%float_136_816803 = OpConstant %float 136.816803
+ %uint_1 = OpConstant %uint 1
+%float_54_4785652 = OpConstant %float 54.4785652
+%float_345_841522 = OpConstant %float 345.841522
+ %84 = OpConstantComposite %v2float %float_54_4785652 %float_345_841522
+%float_534_764526 = OpConstant %float 534.764526
+%VertexOutput = OpTypeStruct %v4float %v4float %v2float
+%VertexInput = OpTypeStruct %v3float %v4float %v2float
+ %89 = OpTypeFunction %VertexOutput %VertexInput
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+ %uint_2 = OpConstant %uint 2
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %107 = OpConstantNull %v3float
+%float_0_00999999978 = OpConstant %float 0.00999999978
+%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
+ %116 = OpConstantNull %VertexOutput
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+ %float_1 = OpConstant %float 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+ %void = OpTypeVoid
+ %135 = OpTypeFunction %void
+ %147 = OpTypeFunction %v4float %VertexOutput
+ %uint_3 = OpConstant %uint 3
+%_ptr_Function_float = OpTypePointer Function %float
+ %float_0 = OpConstant %float 0
+ %172 = OpTypeFunction %void %v3uint
+%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
+ %v2uint = OpTypeVector %uint 2
+%_ptr_StorageBuffer_Particle = OpTypePointer StorageBuffer %Particle
+%_ptr_Function_Particle = OpTypePointer Function %Particle
+ %194 = OpConstantNull %Particle
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %float_0_5 = OpConstant %float 0.5
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %int_0 = OpConstant %int 0
+ %232 = OpConstantComposite %v2int %int_0 %int_0
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+ %235 = OpConstantNull %v2int
+ %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+ %242 = OpConstantNull %int
+ %v4bool = OpTypeVector %bool 4
+ %int_2 = OpConstant %int 2
+ %v2bool = OpTypeVector %bool 2
+ %float_3 = OpConstant %float 3
+ %float_n1 = OpConstant %float -1
+ %302 = OpConstantComposite %v2float %float_1 %float_n1
+%float_0_100000001 = OpConstant %float 0.100000001
+%float_0_300000012 = OpConstant %float 0.300000012
+ %float_2 = OpConstant %float 2
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %402 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %float_4 = OpConstant %float 4
+%float_9_99999975en05 = OpConstant %float 9.99999975e-05
+ %rand = OpFunction %float None %60
+ %62 = OpLabel
+ %65 = OpAccessChain %_ptr_Private_float %rand_seed %uint_0
+ %70 = OpLoad %v2float %rand_seed
+ %69 = OpDot %float %70 %73
+ %68 = OpExtInst %float %67 Cos %69
+ %75 = OpFMul %float %68 %float_136_816803
+ %66 = OpExtInst %float %67 Fract %75
+ OpStore %65 %66
+ %77 = OpAccessChain %_ptr_Private_float %rand_seed %uint_1
+ %81 = OpLoad %v2float %rand_seed
+ %80 = OpDot %float %81 %84
+ %79 = OpExtInst %float %67 Cos %80
+ %86 = OpFMul %float %79 %float_534_764526
+ %78 = OpExtInst %float %67 Fract %86
+ OpStore %77 %78
+ %87 = OpAccessChain %_ptr_Private_float %rand_seed %uint_1
+ %88 = OpLoad %float %87
+ OpReturnValue %88
+ OpFunctionEnd
+%vs_main_inner = OpFunction %VertexOutput None %89
+ %in = OpFunctionParameter %VertexInput
+ %94 = OpLabel
+ %quad_pos = OpVariable %_ptr_Function_v3float Function %107
+ %position = OpVariable %_ptr_Function_v3float Function %107
+ %out = OpVariable %_ptr_Function_VertexOutput Function %116
+ %97 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_1
+ %98 = OpLoad %v3float %97
+ %100 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_2
+ %101 = OpLoad %v3float %100
+ %102 = OpCompositeConstruct %mat2v3float %98 %101
+ %103 = OpCompositeExtract %v2float %in 2
+ %104 = OpMatrixTimesVector %v3float %102 %103
+ OpStore %quad_pos %104
+ %108 = OpCompositeExtract %v3float %in 0
+ %109 = OpLoad %v3float %quad_pos
+ %111 = OpVectorTimesScalar %v3float %109 %float_0_00999999978
+ %112 = OpFAdd %v3float %108 %111
+ OpStore %position %112
+ %118 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+ %120 = OpAccessChain %_ptr_Uniform_mat4v4float %render_params %uint_0
+ %121 = OpLoad %mat4v4float %120
+ %122 = OpLoad %v3float %position
+ %123 = OpCompositeExtract %float %122 0
+ %124 = OpCompositeExtract %float %122 1
+ %125 = OpCompositeExtract %float %122 2
+ %127 = OpCompositeConstruct %v4float %123 %124 %125 %float_1
+ %128 = OpMatrixTimesVector %v4float %121 %127
+ OpStore %118 %128
+ %129 = OpAccessChain %_ptr_Function_v4float %out %uint_1
+ %130 = OpCompositeExtract %v4float %in 1
+ OpStore %129 %130
+ %132 = OpAccessChain %_ptr_Function_v2float %out %uint_2
+ %133 = OpCompositeExtract %v2float %in 2
+ OpStore %132 %133
+ %134 = OpLoad %VertexOutput %out
+ OpReturnValue %134
+ OpFunctionEnd
+ %vs_main = OpFunction %void None %135
+ %138 = OpLabel
+ %140 = OpLoad %v3float %position_1
+ %141 = OpLoad %v4float %color_1
+ %142 = OpLoad %v2float %quad_pos_1
+ %143 = OpCompositeConstruct %VertexInput %140 %141 %142
+ %139 = OpFunctionCall %VertexOutput %vs_main_inner %143
+ %144 = OpCompositeExtract %v4float %139 0
+ OpStore %position_2 %144
+ %145 = OpCompositeExtract %v4float %139 1
+ OpStore %color_2 %145
+ %146 = OpCompositeExtract %v2float %139 2
+ OpStore %quad_pos_2 %146
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fs_main_inner = OpFunction %v4float None %147
+ %in_0 = OpFunctionParameter %VertexOutput
+ %150 = OpLabel
+ %color = OpVariable %_ptr_Function_v4float Function %13
+ %151 = OpCompositeExtract %v4float %in_0 1
+ OpStore %color %151
+ %155 = OpAccessChain %_ptr_Function_float %color %uint_3
+ %156 = OpAccessChain %_ptr_Function_float %color %uint_3
+ %157 = OpLoad %float %156
+ %160 = OpCompositeExtract %v2float %in_0 2
+ %159 = OpExtInst %float %67 Length %160
+ %161 = OpFSub %float %float_1 %159
+ %158 = OpExtInst %float %67 NMax %161 %float_0
+ %163 = OpFMul %float %157 %158
+ OpStore %155 %163
+ %164 = OpLoad %v4float %color
+ OpReturnValue %164
+ OpFunctionEnd
+ %fs_main = OpFunction %void None %135
+ %166 = OpLabel
+ %168 = OpLoad %v4float %position_3
+ %169 = OpLoad %v4float %color_3
+ %170 = OpLoad %v2float %quad_pos_3
+ %171 = OpCompositeConstruct %VertexOutput %168 %169 %170
+ %167 = OpFunctionCall %v4float %fs_main_inner %171
+ OpStore %value_1 %167
+ OpReturn
+ OpFunctionEnd
+%simulate_inner = OpFunction %void None %172
+%GlobalInvocationID = OpFunctionParameter %v3uint
+ %175 = OpLabel
+ %particle = OpVariable %_ptr_Function_Particle Function %194
+ %coord = OpVariable %_ptr_Function_v2int Function %235
+ %level = OpVariable %_ptr_Function_int Function %242
+ %270 = OpVariable %_ptr_Function_v2int Function %235
+ %297 = OpVariable %_ptr_Function_v2float Function %17
+ %177 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_1
+ %178 = OpLoad %v4float %177
+ %179 = OpVectorShuffle %v2float %178 %178 0 1
+ %182 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
+ %180 = OpConvertUToF %v2float %182
+ %183 = OpFAdd %v2float %179 %180
+ %184 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_1
+ %185 = OpLoad %v4float %184
+ %186 = OpVectorShuffle %v2float %185 %185 2 3
+ %187 = OpFMul %v2float %183 %186
+ OpStore %rand_seed %187
+ %188 = OpCompositeExtract %uint %GlobalInvocationID 0
+ %190 = OpAccessChain %_ptr_StorageBuffer_Particle %data %uint_0 %188
+ %191 = OpLoad %Particle %190
+ OpStore %particle %191
+ %195 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+ %196 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+ %197 = OpLoad %float %196
+ %199 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+ %200 = OpLoad %float %199
+ %202 = OpFMul %float %200 %float_0_5
+ %203 = OpFSub %float %197 %202
+ OpStore %195 %203
+ %204 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+ %205 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+ %206 = OpLoad %v3float %205
+ %207 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+ %208 = OpLoad %float %207
+ %209 = OpAccessChain %_ptr_Function_v3float %particle %uint_3
+ %210 = OpLoad %v3float %209
+ %211 = OpVectorTimesScalar %v3float %210 %208
+ %212 = OpFAdd %v3float %206 %211
+ OpStore %204 %212
+ %213 = OpAccessChain %_ptr_Function_float %particle %uint_1
+ %214 = OpAccessChain %_ptr_Function_float %particle %uint_1
+ %215 = OpLoad %float %214
+ %216 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+ %217 = OpLoad %float %216
+ %218 = OpFSub %float %215 %217
+ OpStore %213 %218
+ %219 = OpAccessChain %_ptr_Function_float %particle %uint_2 %uint_3
+ %221 = OpAccessChain %_ptr_Function_float %particle %uint_1
+ %222 = OpLoad %float %221
+ %220 = OpExtInst %float %67 SmoothStep %float_0 %float_0_5 %222
+ OpStore %219 %220
+ %223 = OpAccessChain %_ptr_Function_float %particle %uint_1
+ %224 = OpLoad %float %223
+ %225 = OpFOrdLessThan %bool %224 %float_0
+ OpSelectionMerge %227 None
+ OpBranchConditional %225 %228 %227
+ %228 = OpLabel
+ OpStore %coord %232
+ %237 = OpLoad %47 %texture
+ %236 = OpImageQueryLevels %int %237
+ %239 = OpISub %int %236 %int_1
+ OpStore %level %239
+ OpBranch %243
+ %243 = OpLabel
+ OpLoopMerge %244 %245 None
+ OpBranch %246
+ %246 = OpLabel
+ %248 = OpLoad %int %level
+ %249 = OpSGreaterThan %bool %248 %int_0
+ %247 = OpLogicalNot %bool %249
+ OpSelectionMerge %250 None
+ OpBranchConditional %247 %251 %250
+ %251 = OpLabel
+ OpBranch %244
+ %250 = OpLabel
+ %253 = OpLoad %47 %texture
+ %254 = OpLoad %v2int %coord
+ %255 = OpLoad %int %level
+ %252 = OpImageFetch %v4float %253 %254 Lod %255
+ %256 = OpFunctionCall %float %rand
+ %257 = OpCompositeConstruct %v4float %256 %256 %256 %256
+ %258 = OpVectorShuffle %v3float %252 %252 0 1 2
+ %259 = OpCompositeExtract %float %258 0
+ %260 = OpCompositeExtract %float %258 1
+ %261 = OpCompositeExtract %float %258 2
+ %262 = OpCompositeConstruct %v4float %float_0 %259 %260 %261
+ %263 = OpFOrdGreaterThanEqual %v4bool %257 %262
+ %265 = OpFOrdLessThan %v4bool %257 %252
+ %266 = OpLogicalAnd %v4bool %263 %265
+ %267 = OpLoad %v2int %coord
+ %271 = OpCompositeConstruct %v2int %int_2 %int_2
+ %269 = OpIMul %v2int %267 %271
+ OpStore %coord %269
+ %272 = OpAccessChain %_ptr_Function_int %coord %uint_0
+ %273 = OpAccessChain %_ptr_Function_int %coord %uint_0
+ %274 = OpLoad %int %273
+ %278 = OpVectorShuffle %v2bool %266 %266 1 3
+ %276 = OpAny %bool %278
+ %275 = OpSelect %int %276 %int_1 %int_0
+ %279 = OpIAdd %int %274 %275
+ OpStore %272 %279
+ %280 = OpAccessChain %_ptr_Function_int %coord %uint_1
+ %281 = OpAccessChain %_ptr_Function_int %coord %uint_1
+ %282 = OpLoad %int %281
+ %285 = OpVectorShuffle %v2bool %266 %266 2 3
+ %284 = OpAny %bool %285
+ %283 = OpSelect %int %284 %int_1 %int_0
+ %286 = OpIAdd %int %282 %283
+ OpStore %280 %286
+ OpBranch %245
+ %245 = OpLabel
+ %287 = OpLoad %int %level
+ %288 = OpISub %int %287 %int_1
+ OpStore %level %288
+ OpBranch %243
+ %244 = OpLabel
+ %290 = OpLoad %v2int %coord
+ %289 = OpConvertSToF %v2float %290
+ %293 = OpLoad %47 %texture
+ %292 = OpImageQuerySizeLod %v2int %293 %int_0
+ %291 = OpConvertSToF %v2float %292
+ %294 = OpFDiv %v2float %289 %291
+ %295 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+ %298 = OpCompositeConstruct %v2float %float_0_5 %float_0_5
+ %296 = OpFSub %v2float %294 %298
+ %300 = OpVectorTimesScalar %v2float %296 %float_3
+ %303 = OpFMul %v2float %300 %302
+ %304 = OpCompositeExtract %float %303 0
+ %305 = OpCompositeExtract %float %303 1
+ %306 = OpCompositeConstruct %v3float %304 %305 %float_0
+ OpStore %295 %306
+ %307 = OpAccessChain %_ptr_Function_v4float %particle %uint_2
+ %309 = OpLoad %47 %texture
+ %310 = OpLoad %v2int %coord
+ %308 = OpImageFetch %v4float %309 %310 Lod %int_0
+ OpStore %307 %308
+ %311 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_0
+ %312 = OpFunctionCall %float %rand
+ %313 = OpFSub %float %312 %float_0_5
+ %315 = OpFMul %float %313 %float_0_100000001
+ OpStore %311 %315
+ %316 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_1
+ %317 = OpFunctionCall %float %rand
+ %318 = OpFSub %float %317 %float_0_5
+ %319 = OpFMul %float %318 %float_0_100000001
+ OpStore %316 %319
+ %320 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+ %321 = OpFunctionCall %float %rand
+ %323 = OpFMul %float %321 %float_0_300000012
+ OpStore %320 %323
+ %324 = OpAccessChain %_ptr_Function_float %particle %uint_1
+ %325 = OpFunctionCall %float %rand
+ %327 = OpFMul %float %325 %float_2
+ %328 = OpFAdd %float %float_0_5 %327
+ OpStore %324 %328
+ OpBranch %227
+ %227 = OpLabel
+ %329 = OpAccessChain %_ptr_StorageBuffer_Particle %data %uint_0 %188
+ %330 = OpLoad %Particle %particle
+ OpStore %329 %330
+ OpReturn
+ OpFunctionEnd
+ %simulate = OpFunction %void None %135
+ %332 = OpLabel
+ %334 = OpLoad %v3uint %GlobalInvocationID_1
+ %333 = OpFunctionCall %void %simulate_inner %334
+ OpReturn
+ OpFunctionEnd
+%import_level_inner = OpFunction %void None %172
+ %coord_0 = OpFunctionParameter %v3uint
+ %337 = OpLabel
+ %339 = OpCompositeExtract %uint %coord_0 0
+ %340 = OpCompositeExtract %uint %coord_0 1
+ %342 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+ %343 = OpLoad %uint %342
+ %344 = OpIMul %uint %340 %343
+ %345 = OpIAdd %uint %339 %344
+ %347 = OpAccessChain %_ptr_StorageBuffer_float %buf_out %uint_0 %345
+ %349 = OpLoad %47 %tex_in
+ %351 = OpVectorShuffle %v2uint %coord_0 %coord_0 0 1
+ %350 = OpBitcast %v2int %351
+ %348 = OpImageFetch %v4float %349 %350 Lod %int_0
+ %352 = OpCompositeExtract %float %348 3
+ OpStore %347 %352
+ OpReturn
+ OpFunctionEnd
+%import_level = OpFunction %void None %135
+ %354 = OpLabel
+ %356 = OpLoad %v3uint %coord_1
+ %355 = OpFunctionCall %void %import_level_inner %356
+ OpReturn
+ OpFunctionEnd
+%export_level_inner = OpFunction %void None %172
+ %coord_3 = OpFunctionParameter %v3uint
+ %359 = OpLabel
+ %413 = OpVariable %_ptr_Function_v4float Function %13
+ %361 = OpVectorShuffle %v2uint %coord_3 %coord_3 0 1
+ %364 = OpLoad %59 %tex_out
+ %363 = OpImageQuerySize %v2int %364
+ %362 = OpBitcast %v2uint %363
+ %365 = OpULessThan %v2bool %361 %362
+ %360 = OpAll %bool %365
+ OpSelectionMerge %366 None
+ OpBranchConditional %360 %367 %366
+ %367 = OpLabel
+ %368 = OpCompositeExtract %uint %coord_3 0
+ %369 = OpCompositeExtract %uint %coord_3 1
+ %370 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+ %371 = OpLoad %uint %370
+ %372 = OpIMul %uint %369 %371
+ %373 = OpIAdd %uint %368 %372
+ %374 = OpCompositeExtract %uint %coord_3 0
+ %375 = OpIMul %uint %374 %uint_2
+ %376 = OpCompositeExtract %uint %coord_3 1
+ %377 = OpIMul %uint %376 %uint_2
+ %378 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+ %379 = OpLoad %uint %378
+ %380 = OpIMul %uint %377 %379
+ %381 = OpIAdd %uint %375 %380
+ %382 = OpIAdd %uint %381 %uint_0
+ %383 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %382
+ %384 = OpLoad %float %383
+ %385 = OpIAdd %uint %381 %uint_1
+ %386 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %385
+ %387 = OpLoad %float %386
+ %388 = OpIAdd %uint %381 %uint_0
+ %389 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+ %390 = OpLoad %uint %389
+ %391 = OpIAdd %uint %388 %390
+ %392 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %391
+ %393 = OpLoad %float %392
+ %394 = OpIAdd %uint %381 %uint_1
+ %395 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+ %396 = OpLoad %uint %395
+ %397 = OpIAdd %uint %394 %396
+ %398 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %397
+ %399 = OpLoad %float %398
+ %401 = OpCompositeConstruct %v4float %384 %387 %393 %399
+ %400 = OpDot %float %401 %402
+ %403 = OpAccessChain %_ptr_StorageBuffer_float %buf_out %uint_0 %373
+ %405 = OpFDiv %float %400 %float_4
+ OpStore %403 %405
+ %406 = OpFAdd %float %384 %387
+ %407 = OpFAdd %float %384 %387
+ %408 = OpFAdd %float %407 %393
+ %409 = OpCompositeConstruct %v4float %384 %406 %408 %400
+ %410 = OpExtInst %float %67 NMax %400 %float_9_99999975en05
+ %414 = OpCompositeConstruct %v4float %410 %410 %410 %410
+ %412 = OpFDiv %v4float %409 %414
+ %416 = OpLoad %59 %tex_out
+ %418 = OpVectorShuffle %v2uint %coord_3 %coord_3 0 1
+ %417 = OpBitcast %v2int %418
+ OpImageWrite %416 %417 %412
+ OpBranch %366
+ %366 = OpLabel
+ OpReturn
+ OpFunctionEnd
+%export_level = OpFunction %void None %135
+ %420 = OpLabel
+ %422 = OpLoad %v3uint %coord_2
+ %421 = OpFunctionCall %void %export_level_inner %422
+ OpReturn
+ OpFunctionEnd
diff --git a/test/benchmark/particles.wgsl.expected.wgsl b/test/benchmark/particles.wgsl.expected.wgsl
new file mode 100644
index 0000000..856f496
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.wgsl
@@ -0,0 +1,144 @@
+var<private> rand_seed : vec2<f32>;
+
+fn rand() -> f32 {
+ rand_seed.x = fract((cos(dot(rand_seed, vec2<f32>(23.140779495, 232.616897583))) * 136.816802979));
+ rand_seed.y = fract((cos(dot(rand_seed, vec2<f32>(54.478565216, 345.841522217))) * 534.764526367));
+ return rand_seed.y;
+}
+
+struct RenderParams {
+ modelViewProjectionMatrix : mat4x4<f32>;
+ right : vec3<f32>;
+ up : vec3<f32>;
+};
+
+[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+
+struct VertexInput {
+ [[location(0)]]
+ position : vec3<f32>;
+ [[location(1)]]
+ color : vec4<f32>;
+ [[location(2)]]
+ quad_pos : vec2<f32>;
+};
+
+struct VertexOutput {
+ [[builtin(position)]]
+ position : vec4<f32>;
+ [[location(0)]]
+ color : vec4<f32>;
+ [[location(1)]]
+ quad_pos : vec2<f32>;
+};
+
+[[stage(vertex)]]
+fn vs_main(in : VertexInput) -> VertexOutput {
+ var quad_pos = (mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos);
+ var position = (in.position + (quad_pos * 0.01));
+ var out : VertexOutput;
+ out.position = (render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0));
+ out.color = in.color;
+ out.quad_pos = in.quad_pos;
+ return out;
+}
+
+[[stage(fragment)]]
+fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+ var color = in.color;
+ color.a = (color.a * max((1.0 - length(in.quad_pos)), 0.0));
+ return color;
+}
+
+struct SimulationParams {
+ deltaTime : f32;
+ seed : vec4<f32>;
+};
+
+struct Particle {
+ position : vec3<f32>;
+ lifetime : f32;
+ color : vec4<f32>;
+ velocity : vec3<f32>;
+};
+
+struct Particles {
+ particles : array<Particle>;
+};
+
+[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
+
+[[binding(1), group(0)]] var<storage, read_write> data : Particles;
+
+[[binding(2), group(0)]] var texture : texture_2d<f32>;
+
+[[stage(compute), workgroup_size(64)]]
+fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ rand_seed = ((sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw);
+ let idx = GlobalInvocationID.x;
+ var particle = data.particles[idx];
+ particle.velocity.z = (particle.velocity.z - (sim_params.deltaTime * 0.5));
+ particle.position = (particle.position + (sim_params.deltaTime * particle.velocity));
+ particle.lifetime = (particle.lifetime - sim_params.deltaTime);
+ particle.color.a = smoothStep(0.0, 0.5, particle.lifetime);
+ if ((particle.lifetime < 0.0)) {
+ var coord = vec2<i32>(0, 0);
+ for(var level = (textureNumLevels(texture) - 1); (level > 0); level = (level - 1)) {
+ let probabilites = textureLoad(texture, coord, level);
+ let value = vec4<f32>(rand());
+ let mask = ((value >= vec4<f32>(0.0, probabilites.xyz)) & (value < probabilites));
+ coord = (coord * 2);
+ coord.x = (coord.x + select(0, 1, any(mask.yw)));
+ coord.y = (coord.y + select(0, 1, any(mask.zw)));
+ }
+ let uv = (vec2<f32>(coord) / vec2<f32>(textureDimensions(texture)));
+ particle.position = vec3<f32>((((uv - 0.5) * 3.0) * vec2<f32>(1.0, -1.0)), 0.0);
+ particle.color = textureLoad(texture, coord, 0);
+ particle.velocity.x = ((rand() - 0.5) * 0.100000001);
+ particle.velocity.y = ((rand() - 0.5) * 0.100000001);
+ particle.velocity.z = (rand() * 0.300000012);
+ particle.lifetime = (0.5 + (rand() * 2.0));
+ }
+ data.particles[idx] = particle;
+}
+
+struct UBO {
+ width : u32;
+};
+
+struct Buffer {
+ weights : array<f32>;
+};
+
+[[binding(3), group(0)]] var<uniform> ubo : UBO;
+
+[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
+
+[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
+
+[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
+
+[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+
+[[stage(compute), workgroup_size(64)]]
+fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+ _ = &(buf_in);
+ let offset = (coord.x + (coord.y * ubo.width));
+ buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
+}
+
+[[stage(compute), workgroup_size(64)]]
+fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+ if (all((coord.xy < vec2<u32>(textureDimensions(tex_out))))) {
+ let dst_offset = (coord.x + (coord.y * ubo.width));
+ let src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo.width));
+ let a = buf_in.weights[(src_offset + 0u)];
+ let b = buf_in.weights[(src_offset + 1u)];
+ let c = buf_in.weights[((src_offset + 0u) + ubo.width)];
+ let d = buf_in.weights[((src_offset + 1u) + ubo.width)];
+ let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
+ buf_out.weights[dst_offset] = (sum / 4.0);
+ let probabilities = (vec4<f32>(a, (a + b), ((a + b) + c), sum) / max(sum, 0.0001));
+ textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
+ }
+}
diff --git a/test/benchmark/simple_compute.wgsl b/test/benchmark/simple_compute.wgsl
new file mode 100644
index 0000000..f5b883a
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl
@@ -0,0 +1,10 @@
+struct SB {
+ data : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+
+[[stage(compute), workgroup_size(1, 2, 3)]]
+fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+ buffer.data[id.x] = buffer.data[id.x] + 1;
+}
diff --git a/test/benchmark/simple_compute.wgsl.expected.hlsl b/test/benchmark/simple_compute.wgsl.expected.hlsl
new file mode 100644
index 0000000..92d72b0
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.hlsl
@@ -0,0 +1,15 @@
+RWByteAddressBuffer buffer : register(u0, space0);
+
+struct tint_symbol_1 {
+ uint3 id : SV_DispatchThreadID;
+};
+
+void main_inner(uint3 id) {
+ buffer.Store((4u * id.x), asuint((asint(buffer.Load((4u * id.x))) + 1)));
+}
+
+[numthreads(1, 2, 3)]
+void main(tint_symbol_1 tint_symbol) {
+ main_inner(tint_symbol.id);
+ return;
+}
diff --git a/test/benchmark/simple_compute.wgsl.expected.msl b/test/benchmark/simple_compute.wgsl.expected.msl
new file mode 100644
index 0000000..a2119e6
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.msl
@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB {
+ /* 0x0000 */ int data[1];
+};
+
+void tint_symbol_1_inner(uint3 id, device SB* const tint_symbol_2) {
+ (*(tint_symbol_2)).data[id[0]] = as_type<int>((as_type<uint>((*(tint_symbol_2)).data[id[0]]) + as_type<uint>(1)));
+}
+
+kernel void tint_symbol_1(device SB* tint_symbol_3 [[buffer(0)]], uint3 id [[thread_position_in_grid]]) {
+ tint_symbol_1_inner(id, tint_symbol_3);
+ return;
+}
+
diff --git a/test/benchmark/simple_compute.wgsl.expected.spvasm b/test/benchmark/simple_compute.wgsl.expected.spvasm
new file mode 100644
index 0000000..35c2109
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.spvasm
@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main" %id_1
+ OpExecutionMode %main LocalSize 1 2 3
+ OpName %id_1 "id_1"
+ OpName %SB "SB"
+ OpMemberName %SB 0 "data"
+ OpName %buffer "buffer"
+ OpName %main_inner "main_inner"
+ OpName %id "id"
+ OpName %main "main"
+ OpDecorate %id_1 BuiltIn GlobalInvocationId
+ OpDecorate %SB Block
+ OpMemberDecorate %SB 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %buffer DescriptorSet 0
+ OpDecorate %buffer Binding 0
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+ %id_1 = OpVariable %_ptr_Input_v3uint Input
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %SB = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_SB = OpTypePointer StorageBuffer %SB
+ %buffer = OpVariable %_ptr_StorageBuffer_SB StorageBuffer
+ %void = OpTypeVoid
+ %10 = OpTypeFunction %void %v3uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %int_1 = OpConstant %int 1
+ %24 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %10
+ %id = OpFunctionParameter %v3uint
+ %14 = OpLabel
+ %16 = OpCompositeExtract %uint %id 0
+ %18 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0 %16
+ %19 = OpCompositeExtract %uint %id 0
+ %20 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0 %19
+ %21 = OpLoad %int %20
+ %23 = OpIAdd %int %21 %int_1
+ OpStore %18 %23
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %24
+ %26 = OpLabel
+ %28 = OpLoad %v3uint %id_1
+ %27 = OpFunctionCall %void %main_inner %28
+ OpReturn
+ OpFunctionEnd
diff --git a/test/benchmark/simple_compute.wgsl.expected.wgsl b/test/benchmark/simple_compute.wgsl.expected.wgsl
new file mode 100644
index 0000000..ac98410
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+struct SB {
+ data : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+
+[[stage(compute), workgroup_size(1, 2, 3)]]
+fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+ buffer.data[id.x] = (buffer.data[id.x] + 1);
+}
diff --git a/test/benchmark/simple_fragment.wgsl b/test/benchmark/simple_fragment.wgsl
new file mode 100644
index 0000000..bb5ed94
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl
@@ -0,0 +1,12 @@
+struct Input {
+ [[location(0)]] color: vec4<f32>;
+};
+
+struct Output {
+ [[location(0)]] color: vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main(in : Input) -> Output {
+ return Output(in.color);
+}
diff --git a/test/benchmark/simple_fragment.wgsl.expected.hlsl b/test/benchmark/simple_fragment.wgsl.expected.hlsl
new file mode 100644
index 0000000..4a462c8
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.hlsl
@@ -0,0 +1,25 @@
+struct Input {
+ float4 color;
+};
+struct Output {
+ float4 color;
+};
+struct tint_symbol_2 {
+ float4 color : TEXCOORD0;
+};
+struct tint_symbol_3 {
+ float4 color : SV_Target0;
+};
+
+Output main_inner(Input tint_symbol) {
+ const Output tint_symbol_4 = {tint_symbol.color};
+ return tint_symbol_4;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+ const Input tint_symbol_5 = {tint_symbol_1.color};
+ const Output inner_result = main_inner(tint_symbol_5);
+ tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+ wrapper_result.color = inner_result.color;
+ return wrapper_result;
+}
diff --git a/test/benchmark/simple_fragment.wgsl.expected.msl b/test/benchmark/simple_fragment.wgsl.expected.msl
new file mode 100644
index 0000000..0bc1116
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.msl
@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Input {
+ float4 color;
+};
+struct Output {
+ float4 color;
+};
+struct tint_symbol_2 {
+ float4 color [[user(locn0)]];
+};
+struct tint_symbol_3 {
+ float4 color [[color(0)]];
+};
+
+Output tint_symbol_inner(Input in) {
+ Output const tint_symbol_4 = {.color=in.color};
+ return tint_symbol_4;
+}
+
+fragment tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+ Input const tint_symbol_5 = {.color=tint_symbol_1.color};
+ Output const inner_result = tint_symbol_inner(tint_symbol_5);
+ tint_symbol_3 wrapper_result = {};
+ wrapper_result.color = inner_result.color;
+ return wrapper_result;
+}
+
diff --git a/test/benchmark/simple_fragment.wgsl.expected.spvasm b/test/benchmark/simple_fragment.wgsl.expected.spvasm
new file mode 100644
index 0000000..531ed06
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.spvasm
@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %color_1 %color_2
+ OpExecutionMode %main OriginUpperLeft
+ OpName %color_1 "color_1"
+ OpName %color_2 "color_2"
+ OpName %Output "Output"
+ OpMemberName %Output 0 "color"
+ OpName %Input "Input"
+ OpMemberName %Input 0 "color"
+ OpName %main_inner "main_inner"
+ OpName %in "in"
+ OpName %main "main"
+ OpDecorate %color_1 Location 0
+ OpDecorate %color_2 Location 0
+ OpMemberDecorate %Output 0 Offset 0
+ OpMemberDecorate %Input 0 Offset 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+ %color_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %7 = OpConstantNull %v4float
+ %color_2 = OpVariable %_ptr_Output_v4float Output %7
+ %Output = OpTypeStruct %v4float
+ %Input = OpTypeStruct %v4float
+ %8 = OpTypeFunction %Output %Input
+ %void = OpTypeVoid
+ %16 = OpTypeFunction %void
+ %main_inner = OpFunction %Output None %8
+ %in = OpFunctionParameter %Input
+ %13 = OpLabel
+ %14 = OpCompositeExtract %v4float %in 0
+ %15 = OpCompositeConstruct %Output %14
+ OpReturnValue %15
+ OpFunctionEnd
+ %main = OpFunction %void None %16
+ %19 = OpLabel
+ %21 = OpLoad %v4float %color_1
+ %22 = OpCompositeConstruct %Input %21
+ %20 = OpFunctionCall %Output %main_inner %22
+ %23 = OpCompositeExtract %v4float %20 0
+ OpStore %color_2 %23
+ OpReturn
+ OpFunctionEnd
diff --git a/test/benchmark/simple_fragment.wgsl.expected.wgsl b/test/benchmark/simple_fragment.wgsl.expected.wgsl
new file mode 100644
index 0000000..1f4225e
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct Input {
+ [[location(0)]]
+ color : vec4<f32>;
+};
+
+struct Output {
+ [[location(0)]]
+ color : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main(in : Input) -> Output {
+ return Output(in.color);
+}
diff --git a/test/benchmark/simple_vertex.wgsl b/test/benchmark/simple_vertex.wgsl
new file mode 100644
index 0000000..8f9a07b
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl
@@ -0,0 +1,12 @@
+struct Input {
+ [[location(0)]] position: vec4<f32>;
+};
+
+struct Output {
+ [[builtin(position)]] position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn main(in : Input) -> Output {
+ return Output(in.position);
+}
diff --git a/test/benchmark/simple_vertex.wgsl.expected.hlsl b/test/benchmark/simple_vertex.wgsl.expected.hlsl
new file mode 100644
index 0000000..03b0d54
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.hlsl
@@ -0,0 +1,25 @@
+struct Input {
+ float4 position;
+};
+struct Output {
+ float4 position;
+};
+struct tint_symbol_2 {
+ float4 position : TEXCOORD0;
+};
+struct tint_symbol_3 {
+ float4 position : SV_Position;
+};
+
+Output main_inner(Input tint_symbol) {
+ const Output tint_symbol_4 = {tint_symbol.position};
+ return tint_symbol_4;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+ const Input tint_symbol_5 = {tint_symbol_1.position};
+ const Output inner_result = main_inner(tint_symbol_5);
+ tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+ wrapper_result.position = inner_result.position;
+ return wrapper_result;
+}
diff --git a/test/benchmark/simple_vertex.wgsl.expected.msl b/test/benchmark/simple_vertex.wgsl.expected.msl
new file mode 100644
index 0000000..3c18b5f
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.msl
@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Input {
+ float4 position;
+};
+struct Output {
+ float4 position;
+};
+struct tint_symbol_2 {
+ float4 position [[attribute(0)]];
+};
+struct tint_symbol_3 {
+ float4 position [[position]];
+};
+
+Output tint_symbol_inner(Input in) {
+ Output const tint_symbol_4 = {.position=in.position};
+ return tint_symbol_4;
+}
+
+vertex tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+ Input const tint_symbol_5 = {.position=tint_symbol_1.position};
+ Output const inner_result = tint_symbol_inner(tint_symbol_5);
+ tint_symbol_3 wrapper_result = {};
+ wrapper_result.position = inner_result.position;
+ return wrapper_result;
+}
+
diff --git a/test/benchmark/simple_vertex.wgsl.expected.spvasm b/test/benchmark/simple_vertex.wgsl.expected.spvasm
new file mode 100644
index 0000000..6490ac0
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.spvasm
@@ -0,0 +1,56 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %main "main" %position_1 %position_2 %vertex_point_size
+ OpName %position_1 "position_1"
+ OpName %position_2 "position_2"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %Output "Output"
+ OpMemberName %Output 0 "position"
+ OpName %Input "Input"
+ OpMemberName %Input 0 "position"
+ OpName %main_inner "main_inner"
+ OpName %in "in"
+ OpName %main "main"
+ OpDecorate %position_1 Location 0
+ OpDecorate %position_2 BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpMemberDecorate %Output 0 Offset 0
+ OpMemberDecorate %Input 0 Offset 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+ %position_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %7 = OpConstantNull %v4float
+ %position_2 = OpVariable %_ptr_Output_v4float Output %7
+%_ptr_Output_float = OpTypePointer Output %float
+ %10 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %10
+ %Output = OpTypeStruct %v4float
+ %Input = OpTypeStruct %v4float
+ %11 = OpTypeFunction %Output %Input
+ %void = OpTypeVoid
+ %19 = OpTypeFunction %void
+ %float_1 = OpConstant %float 1
+ %main_inner = OpFunction %Output None %11
+ %in = OpFunctionParameter %Input
+ %16 = OpLabel
+ %17 = OpCompositeExtract %v4float %in 0
+ %18 = OpCompositeConstruct %Output %17
+ OpReturnValue %18
+ OpFunctionEnd
+ %main = OpFunction %void None %19
+ %22 = OpLabel
+ %24 = OpLoad %v4float %position_1
+ %25 = OpCompositeConstruct %Input %24
+ %23 = OpFunctionCall %Output %main_inner %25
+ %26 = OpCompositeExtract %v4float %23 0
+ OpStore %position_2 %26
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/benchmark/simple_vertex.wgsl.expected.wgsl b/test/benchmark/simple_vertex.wgsl.expected.wgsl
new file mode 100644
index 0000000..7bc18b1
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct Input {
+ [[location(0)]]
+ position : vec4<f32>;
+};
+
+struct Output {
+ [[builtin(position)]]
+ position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn main(in : Input) -> Output {
+ return Output(in.position);
+}
diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt
index f54ed0a..47d6664 100644
--- a/third_party/CMakeLists.txt
+++ b/third_party/CMakeLists.txt
@@ -12,6 +12,11 @@
# See the License for the specific language governing permissions and
# limitations under the License.
+if (${TINT_BUILD_BENCHMARKS})
+ set(BENCHMARK_ENABLE_TESTING FALSE CACHE BOOL FALSE FORCE)
+ add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/benchmark EXCLUDE_FROM_ALL)
+endif()
+
if (${TINT_BUILD_TESTS} AND NOT TARGET gmock)
set(gtest_force_shared_crt ON CACHE BOOL "Controls whether a shared run-time library should be used even when Google Test is built as static library" FORCE)
add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/googletest EXCLUDE_FROM_ALL)
diff --git a/tools/benchdiff b/tools/benchdiff
new file mode 100755
index 0000000..3cc175f
--- /dev/null
+++ b/tools/benchdiff
@@ -0,0 +1,33 @@
+#!/usr/bin/env bash
+# Copyright 2022 The Tint Authors
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+set -e # Fail on any error.
+
+if [ ! -x "$(which go)" ] ; then
+ echo "error: go needs to be on \$PATH to use $0"
+ exit 1
+fi
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+ROOT_DIR="$( cd "${SCRIPT_DIR}/.." >/dev/null 2>&1 && pwd )"
+BINARY="${SCRIPT_DIR}/bin/benchdiff"
+
+# Rebuild the binary.
+# Note, go caches build artifacts, so this is quick for repeat calls
+pushd "${SCRIPT_DIR}/src/cmd/benchdiff" > /dev/null
+ go build -o "${BINARY}" main.go
+popd > /dev/null
+
+"${BINARY}" "$@"
diff --git a/tools/src/bench/bench.go b/tools/src/bench/bench.go
new file mode 100644
index 0000000..8bb1d99
--- /dev/null
+++ b/tools/src/bench/bench.go
@@ -0,0 +1,150 @@
+// Copyright 2022 The Tint Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+// Package bench provides types and methods for parsing Google benchmark results.
+package bench
+
+import (
+ "encoding/json"
+ "errors"
+ "fmt"
+ "regexp"
+ "strconv"
+ "strings"
+ "time"
+)
+
+// Test holds the results of a single benchmark test.
+type Test struct {
+ Name string
+ NumTasks uint
+ NumThreads uint
+ Duration time.Duration
+ Iterations uint
+}
+
+var testVarRE = regexp.MustCompile(`([\w])+:([0-9]+)`)
+
+func (t *Test) parseName() {
+ for _, match := range testVarRE.FindAllStringSubmatch(t.Name, -1) {
+ if len(match) != 3 {
+ continue
+ }
+ n, err := strconv.Atoi(match[2])
+ if err != nil {
+ continue
+ }
+ switch match[1] {
+ case "threads":
+ t.NumThreads = uint(n)
+ case "tasks":
+ t.NumTasks = uint(n)
+ }
+ }
+}
+
+// Benchmark holds a set of benchmark test results.
+type Benchmark struct {
+ Tests []Test
+}
+
+// Parse parses the benchmark results from the string s.
+// Parse will handle the json and 'console' formats.
+func Parse(s string) (Benchmark, error) {
+ type Parser = func(s string) (Benchmark, error)
+ for _, parser := range []Parser{parseConsole, parseJSON} {
+ b, err := parser(s)
+ switch err {
+ case nil:
+ return b, nil
+ case errWrongFormat:
+ default:
+ return Benchmark{}, err
+ }
+ }
+
+ return Benchmark{}, errors.New("Unrecognised file format")
+}
+
+var errWrongFormat = errors.New("Wrong format")
+var consoleLineRE = regexp.MustCompile(`([\w/:]+)\s+([0-9]+(?:.[0-9]+)?) ns\s+[0-9]+(?:.[0-9]+) ns\s+([0-9]+)`)
+
+func parseConsole(s string) (Benchmark, error) {
+ blocks := strings.Split(s, "------------------------------------------------------------------------------------------")
+ if len(blocks) != 3 {
+ return Benchmark{}, errWrongFormat
+ }
+
+ lines := strings.Split(blocks[2], "\n")
+ b := Benchmark{
+ Tests: make([]Test, 0, len(lines)),
+ }
+ for _, line := range lines {
+ if len(line) == 0 {
+ continue
+ }
+ matches := consoleLineRE.FindStringSubmatch(line)
+ if len(matches) != 4 {
+ return Benchmark{}, fmt.Errorf("Unable to parse the line:\n" + line)
+ }
+ ns, err := strconv.ParseFloat(matches[2], 64)
+ if err != nil {
+ return Benchmark{}, fmt.Errorf("Unable to parse the duration: " + matches[2])
+ }
+ iterations, err := strconv.Atoi(matches[3])
+ if err != nil {
+ return Benchmark{}, fmt.Errorf("Unable to parse the number of iterations: " + matches[3])
+ }
+
+ t := Test{
+ Name: matches[1],
+ Duration: time.Nanosecond * time.Duration(ns),
+ Iterations: uint(iterations),
+ }
+ t.parseName()
+ b.Tests = append(b.Tests, t)
+ }
+ return b, nil
+}
+
+func parseJSON(s string) (Benchmark, error) {
+ type T struct {
+ Name string `json:"name"`
+ Iterations uint `json:"iterations"`
+ Time float64 `json:"real_time"`
+ }
+ type B struct {
+ Tests []T `json:"benchmarks"`
+ }
+ b := B{}
+ d := json.NewDecoder(strings.NewReader(s))
+ if err := d.Decode(&b); err != nil {
+ return Benchmark{}, err
+ }
+
+ out := Benchmark{
+ Tests: make([]Test, len(b.Tests)),
+ }
+ for i, test := range b.Tests {
+ t := Test{
+ Name: test.Name,
+ Duration: time.Nanosecond * time.Duration(int64(test.Time)),
+ Iterations: test.Iterations,
+ }
+ t.parseName()
+ out.Tests[i] = t
+ }
+
+ return out, nil
+}
diff --git a/tools/src/cmd/benchdiff/main.go b/tools/src/cmd/benchdiff/main.go
new file mode 100644
index 0000000..14ec4f3
--- /dev/null
+++ b/tools/src/cmd/benchdiff/main.go
@@ -0,0 +1,177 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+// benchdiff is a tool that compares two Google benchmark results and displays
+// sorted performance differences.
+package main
+
+import (
+ "errors"
+ "flag"
+ "fmt"
+ "io/ioutil"
+ "os"
+ "path/filepath"
+ "sort"
+ "strings"
+ "text/tabwriter"
+ "time"
+
+ "dawn.googlesource.com/tint/tools/src/bench"
+)
+
+var (
+ minDiff = flag.Duration("min-diff", time.Microsecond*10, "Filter away time diffs less than this duration")
+ minRelDiff = flag.Float64("min-rel-diff", 0.01, "Filter away absolute relative diffs between [1, 1+x]")
+)
+
+func main() {
+ flag.ErrHelp = errors.New("benchdiff is a tool to compare two benchmark results")
+ flag.Parse()
+ flag.Usage = func() {
+ fmt.Fprintln(os.Stderr, "benchdiff <benchmark-a> <benchmark-b>")
+ flag.PrintDefaults()
+ }
+
+ args := flag.Args()
+ if len(args) < 2 {
+ flag.Usage()
+ os.Exit(1)
+ }
+
+ pathA, pathB := args[0], args[1]
+
+ if err := run(pathA, pathB); err != nil {
+ fmt.Fprintln(os.Stderr, err)
+ os.Exit(-1)
+ }
+}
+
+func run(pathA, pathB string) error {
+ fileA, err := ioutil.ReadFile(pathA)
+ if err != nil {
+ return err
+ }
+ benchA, err := bench.Parse(string(fileA))
+ if err != nil {
+ return err
+ }
+
+ fileB, err := ioutil.ReadFile(pathB)
+ if err != nil {
+ return err
+ }
+ benchB, err := bench.Parse(string(fileB))
+ if err != nil {
+ return err
+ }
+
+ compare(benchA, benchB, fileName(pathA), fileName(pathB))
+
+ return nil
+}
+
+func fileName(path string) string {
+ _, name := filepath.Split(path)
+ return name
+}
+
+func compare(benchA, benchB bench.Benchmark, nameA, nameB string) {
+ type times struct {
+ a time.Duration
+ b time.Duration
+ }
+ byName := map[string]times{}
+ for _, test := range benchA.Tests {
+ byName[test.Name] = times{a: test.Duration}
+ }
+ for _, test := range benchB.Tests {
+ t := byName[test.Name]
+ t.b = test.Duration
+ byName[test.Name] = t
+ }
+
+ type delta struct {
+ name string
+ times times
+ relDiff float64
+ absRelDiff float64
+ }
+ deltas := []delta{}
+ for name, times := range byName {
+ if times.a == 0 || times.b == 0 {
+ continue // Assuming test was missing from a or b
+ }
+ diff := times.b - times.a
+ absDiff := diff
+ if absDiff < 0 {
+ absDiff = -absDiff
+ }
+ if absDiff < *minDiff {
+ continue
+ }
+
+ relDiff := float64(times.b) / float64(times.a)
+ absRelDiff := relDiff
+ if absRelDiff < 1 {
+ absRelDiff = 1.0 / absRelDiff
+ }
+ if absRelDiff < (1.0 + *minRelDiff) {
+ continue
+ }
+
+ d := delta{
+ name: name,
+ times: times,
+ relDiff: relDiff,
+ absRelDiff: absRelDiff,
+ }
+ deltas = append(deltas, d)
+ }
+
+ sort.Slice(deltas, func(i, j int) bool { return deltas[j].relDiff < deltas[i].relDiff })
+
+ fmt.Println("A:", nameA)
+ fmt.Println("B:", nameB)
+ fmt.Println()
+
+ buf := strings.Builder{}
+ {
+ w := tabwriter.NewWriter(&buf, 1, 1, 0, ' ', 0)
+ fmt.Fprintln(w, "Test name\t | Δ (A → B)\t | % (A → B)\t | % (B → A)\t | × (A → B)\t | × (B → A)\t | A \t | B")
+ fmt.Fprintln(w, "\t-+\t-+\t-+\t-+\t-+\t-+\t-+\t-")
+ for _, delta := range deltas {
+ a2b := delta.times.b - delta.times.a
+ fmt.Fprintf(w, "%v \t | %v \t | %+2.1f%% \t | %+2.1f%% \t | %+.4f \t | %+.4f \t | %v \t | %v \t|\n",
+ delta.name,
+ a2b, // Δ (A → B)
+ 100*float64(a2b)/float64(delta.times.a), // % (A → B)
+ 100*float64(-a2b)/float64(delta.times.b), // % (B → A)
+ float64(delta.times.b)/float64(delta.times.a), // × (A → B)
+ float64(delta.times.a)/float64(delta.times.b), // × (B → A)
+ delta.times.a, // A
+ delta.times.b, // B
+ )
+ }
+ w.Flush()
+ }
+
+ // Split the table by line so we can add in a header line
+ lines := strings.Split(buf.String(), "\n")
+ fmt.Println(lines[0])
+ fmt.Println(strings.ReplaceAll(lines[1], " ", "-"))
+ for _, l := range lines[2:] {
+ fmt.Println(l)
+ }
+}
