main: Replace --dawn-validation with --validate
Performs output validation with spirv-val for SPIR-V (as before), HLSL validation with DXC, and MSL validation with the Metal Shader Compiler. Disable HLSL tests that fail to validate Bug: tint:812 Change-Id: If78c351b4e23c7fb50d333eacf9ee7cc81d18564 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51280 Reviewed-by: Antonio Maiorano <amaiorano@google.com> Reviewed-by: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
This commit is contained in:
parent
a7a23eaa9e
commit
a70c14dbce
|
@ -19,6 +19,7 @@ executable("tint") {
|
|||
sources = [ "main.cc" ]
|
||||
deps = [
|
||||
"${tint_root_dir}/src:libtint",
|
||||
"${tint_root_dir}/src:tint_val",
|
||||
"${tint_spirv_tools_dir}/:spvtools",
|
||||
"${tint_spirv_tools_dir}/:spvtools_opt",
|
||||
"${tint_spirv_tools_dir}/:spvtools_val",
|
||||
|
|
|
@ -19,7 +19,7 @@ set(TINT_SRCS
|
|||
## Tint executable
|
||||
add_executable(tint ${TINT_SRCS})
|
||||
tint_default_compile_options(tint)
|
||||
target_link_libraries(tint libtint)
|
||||
target_link_libraries(tint libtint tint_val)
|
||||
|
||||
if(${TINT_BUILD_SPV_READER} OR ${TINT_BUILD_SPV_WRITER})
|
||||
target_link_libraries(tint SPIRV-Tools)
|
||||
|
|
177
samples/main.cc
177
samples/main.cc
|
@ -24,6 +24,8 @@
|
|||
#include "spirv-tools/libspirv.hpp"
|
||||
#endif // TINT_BUILD_SPV_READER
|
||||
|
||||
#include "src/utils/io/command.h"
|
||||
#include "src/val/val.h"
|
||||
#include "tint/tint.h"
|
||||
|
||||
namespace {
|
||||
|
@ -52,7 +54,7 @@ struct Options {
|
|||
|
||||
bool parse_only = false;
|
||||
bool dump_ast = false;
|
||||
bool dawn_validation = false;
|
||||
bool validate = false;
|
||||
bool demangle = false;
|
||||
bool dump_inspector_bindings = false;
|
||||
|
||||
|
@ -62,6 +64,9 @@ struct Options {
|
|||
std::string ep_name;
|
||||
|
||||
std::vector<std::string> transforms;
|
||||
|
||||
std::string dxc_path;
|
||||
std::string xcrun_path;
|
||||
};
|
||||
|
||||
const char kUsage[] = R"(Usage: tint [options] <input-file>
|
||||
|
@ -86,12 +91,16 @@ const char kUsage[] = R"(Usage: tint [options] <input-file>
|
|||
renamer
|
||||
--parse-only -- Stop after parsing the input
|
||||
--dump-ast -- Dump the generated AST to stdout
|
||||
--dawn-validation -- SPIRV outputs are validated with the same flags
|
||||
as Dawn does. Has no effect on non-SPIRV outputs.
|
||||
--demangle -- Preserve original source names. Demangle them.
|
||||
Affects AST dumping, and text-based output languages.
|
||||
--dump-inspector-bindings -- Dump reflection data about bindins to stdout.
|
||||
-h -- This help text)";
|
||||
-h -- This help text
|
||||
--validate -- Validates generated SPIR-V with spirv-val.
|
||||
Has no effect on non-SPIRV outputs.
|
||||
--dxc -- Path to DXC executable, used to validate HLSL output.
|
||||
When specified, automatically enables --validate
|
||||
--xcrun -- Path to xcrun executable, used to validate MSL output.
|
||||
When specified, automatically enables --validate)";
|
||||
|
||||
#ifdef _MSC_VER
|
||||
#pragma warning(disable : 4068; suppress : 4100)
|
||||
|
@ -389,12 +398,28 @@ bool ParseArgs(const std::vector<std::string>& args, Options* opts) {
|
|||
opts->parse_only = true;
|
||||
} else if (arg == "--dump-ast") {
|
||||
opts->dump_ast = true;
|
||||
} else if (arg == "--dawn-validation") {
|
||||
opts->dawn_validation = true;
|
||||
} else if (arg == "--demangle") {
|
||||
opts->demangle = true;
|
||||
} else if (arg == "--dump-inspector-bindings") {
|
||||
opts->dump_inspector_bindings = true;
|
||||
} else if (arg == "--validate") {
|
||||
opts->validate = true;
|
||||
} else if (arg == "--dxc") {
|
||||
++i;
|
||||
if (i >= args.size()) {
|
||||
std::cerr << "Missing value for " << arg << std::endl;
|
||||
return false;
|
||||
}
|
||||
opts->dxc_path = args[i];
|
||||
opts->validate = true;
|
||||
} else if (arg == "--xcrun") {
|
||||
++i;
|
||||
if (i >= args.size()) {
|
||||
std::cerr << "Missing value for " << arg << std::endl;
|
||||
return false;
|
||||
}
|
||||
opts->xcrun_path = args[i];
|
||||
opts->validate = true;
|
||||
} else if (!arg.empty()) {
|
||||
if (arg[0] == '-') {
|
||||
std::cerr << "Unrecognized option: " << arg << std::endl;
|
||||
|
@ -772,33 +797,47 @@ int main(int argc, const char** argv) {
|
|||
|
||||
std::unique_ptr<tint::writer::Writer> writer;
|
||||
|
||||
switch (options.format) {
|
||||
case Format::kSpirv:
|
||||
case Format::kSpvAsm:
|
||||
#if TINT_BUILD_SPV_WRITER
|
||||
if (options.format == Format::kSpirv || options.format == Format::kSpvAsm) {
|
||||
writer = std::make_unique<tint::writer::spirv::Generator>(program.get());
|
||||
}
|
||||
writer = std::make_unique<tint::writer::spirv::Generator>(program.get());
|
||||
break;
|
||||
#else
|
||||
std::cerr << "SPIR-V writer not enabled in tint build" << std::endl;
|
||||
return 1;
|
||||
#endif // TINT_BUILD_SPV_WRITER
|
||||
|
||||
case Format::kWgsl:
|
||||
#if TINT_BUILD_WGSL_WRITER
|
||||
if (options.format == Format::kWgsl) {
|
||||
writer = std::make_unique<tint::writer::wgsl::Generator>(program.get());
|
||||
}
|
||||
writer = std::make_unique<tint::writer::wgsl::Generator>(program.get());
|
||||
break;
|
||||
#else
|
||||
std::cerr << "WGSL writer not enabled in tint build" << std::endl;
|
||||
return 1;
|
||||
#endif // TINT_BUILD_WGSL_WRITER
|
||||
|
||||
case Format::kMsl:
|
||||
#if TINT_BUILD_MSL_WRITER
|
||||
if (options.format == Format::kMsl) {
|
||||
writer = std::make_unique<tint::writer::msl::Generator>(program.get());
|
||||
}
|
||||
writer = std::make_unique<tint::writer::msl::Generator>(program.get());
|
||||
break;
|
||||
#else
|
||||
std::cerr << "MSL writer not enabled in tint build" << std::endl;
|
||||
return 1;
|
||||
#endif // TINT_BUILD_MSL_WRITER
|
||||
|
||||
case Format::kHlsl:
|
||||
#if TINT_BUILD_HLSL_WRITER
|
||||
if (options.format == Format::kHlsl) {
|
||||
writer = std::make_unique<tint::writer::hlsl::Generator>(program.get());
|
||||
}
|
||||
writer = std::make_unique<tint::writer::hlsl::Generator>(program.get());
|
||||
break;
|
||||
#else
|
||||
std::cerr << "HLSL writer not enabled in tint build" << std::endl;
|
||||
return 1;
|
||||
#endif // TINT_BUILD_HLSL_WRITER
|
||||
|
||||
if (!writer) {
|
||||
std::cerr << "Unknown output format specified" << std::endl;
|
||||
return 1;
|
||||
default:
|
||||
std::cerr << "Unknown output format specified" << std::endl;
|
||||
return 1;
|
||||
}
|
||||
|
||||
if (!writer->Generate()) {
|
||||
|
@ -807,27 +846,82 @@ int main(int argc, const char** argv) {
|
|||
return 1;
|
||||
}
|
||||
|
||||
#if TINT_BUILD_SPV_WRITER
|
||||
bool dawn_validation_failed = false;
|
||||
std::ostringstream stream;
|
||||
bool validation_failed = false;
|
||||
std::ostringstream validation_msgs;
|
||||
|
||||
if (options.dawn_validation &&
|
||||
(options.format == Format::kSpvAsm || options.format == Format::kSpirv)) {
|
||||
// Use Vulkan 1.1, since this is what Tint, internally, uses.
|
||||
spvtools::SpirvTools tools(SPV_ENV_VULKAN_1_1);
|
||||
tools.SetMessageConsumer([&stream](spv_message_level_t, const char*,
|
||||
const spv_position_t& pos,
|
||||
const char* msg) {
|
||||
stream << (pos.line + 1) << ":" << (pos.column + 1) << ": " << msg
|
||||
<< std::endl;
|
||||
});
|
||||
auto* w = static_cast<tint::writer::spirv::Generator*>(writer.get());
|
||||
if (!tools.Validate(w->result().data(), w->result().size(),
|
||||
spvtools::ValidatorOptions())) {
|
||||
dawn_validation_failed = true;
|
||||
if (options.validate) {
|
||||
switch (options.format) {
|
||||
#if TINT_BUILD_SPV_WRITER
|
||||
case Format::kSpirv:
|
||||
case Format::kSpvAsm: {
|
||||
// Use Vulkan 1.1, since this is what Tint, internally, uses.
|
||||
spvtools::SpirvTools tools(SPV_ENV_VULKAN_1_1);
|
||||
tools.SetMessageConsumer(
|
||||
[&validation_msgs](spv_message_level_t, const char*,
|
||||
const spv_position_t& pos, const char* msg) {
|
||||
validation_msgs << (pos.line + 1) << ":" << (pos.column + 1)
|
||||
<< ": " << msg << std::endl;
|
||||
});
|
||||
auto* w = static_cast<tint::writer::spirv::Generator*>(writer.get());
|
||||
if (!tools.Validate(w->result().data(), w->result().size(),
|
||||
spvtools::ValidatorOptions())) {
|
||||
validation_failed = true;
|
||||
}
|
||||
break;
|
||||
}
|
||||
#endif
|
||||
#if TINT_BUILD_HLSL_WRITER
|
||||
case Format::kHlsl: {
|
||||
auto dxc = tint::utils::Command::LookPath(
|
||||
options.dxc_path.empty() ? "dxc" : options.dxc_path);
|
||||
if (dxc.Found()) {
|
||||
auto* w = static_cast<tint::writer::Text*>(writer.get());
|
||||
auto hlsl = w->result();
|
||||
auto res = tint::val::Hlsl(dxc.Path(), hlsl, program.get());
|
||||
if (res.failed) {
|
||||
validation_failed = true;
|
||||
validation_msgs << res.source << std::endl;
|
||||
validation_msgs << res.output;
|
||||
}
|
||||
} else {
|
||||
validation_failed = true;
|
||||
validation_msgs << "DXC executable not found. Cannot validate";
|
||||
}
|
||||
break;
|
||||
}
|
||||
#endif // TINT_BUILD_HLSL_WRITER
|
||||
#if TINT_BUILD_MSL_WRITER
|
||||
case Format::kMsl: {
|
||||
#ifdef _WIN32
|
||||
const char* default_xcrun_exe = "metal.exe";
|
||||
#else
|
||||
const char* default_xcrun_exe = "xcrun";
|
||||
#endif
|
||||
auto xcrun = tint::utils::Command::LookPath(options.xcrun_path.empty()
|
||||
? default_xcrun_exe
|
||||
: options.xcrun_path);
|
||||
if (xcrun.Found()) {
|
||||
auto* w = static_cast<tint::writer::Text*>(writer.get());
|
||||
auto msl = w->result();
|
||||
auto res = tint::val::Msl(xcrun.Path(), msl);
|
||||
if (res.failed) {
|
||||
validation_failed = true;
|
||||
validation_msgs << res.source << std::endl;
|
||||
validation_msgs << res.output;
|
||||
}
|
||||
} else {
|
||||
validation_failed = true;
|
||||
validation_msgs << "xcrun executable not found. Cannot validate";
|
||||
}
|
||||
break;
|
||||
}
|
||||
#endif // TINT_BUILD_MSL_WRITER
|
||||
default:
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
#if TINT_BUILD_SPV_WRITER
|
||||
if (options.format == Format::kSpvAsm) {
|
||||
auto* w = static_cast<tint::writer::spirv::Generator*>(writer.get());
|
||||
auto str = Disassemble(w->result());
|
||||
|
@ -841,12 +935,13 @@ int main(int argc, const char** argv) {
|
|||
return 1;
|
||||
}
|
||||
}
|
||||
if (dawn_validation_failed) {
|
||||
#endif // TINT_BUILD_SPV_WRITER
|
||||
|
||||
if (validation_failed) {
|
||||
std::cerr << std::endl << std::endl << "Validation Failure:" << std::endl;
|
||||
std::cerr << stream.str();
|
||||
std::cerr << validation_msgs.str();
|
||||
return 1;
|
||||
}
|
||||
#endif // TINT_BUILD_SPV_WRITER
|
||||
|
||||
if (options.format != Format::kSpvAsm && options.format != Format::kSpirv) {
|
||||
auto* w = static_cast<tint::writer::Text*>(writer.get());
|
||||
|
|
39
src/BUILD.gn
39
src/BUILD.gn
|
@ -208,6 +208,43 @@ tint_language_header("cldebuginfo100") {
|
|||
grammar_file = "${tint_spirv_headers_dir}/include/spirv/unified1/extinst.opencl.debuginfo.100.grammar.json"
|
||||
}
|
||||
|
||||
###############################################################################
|
||||
# Helper library for IO operations
|
||||
# Only to be used by tests and sample executable
|
||||
###############################################################################
|
||||
source_set("tint_utils_io") {
|
||||
sources = [
|
||||
"utils/io/command.h",
|
||||
"utils/io/tmpfile.h",
|
||||
]
|
||||
|
||||
if (is_linux || is_mac) {
|
||||
sources += [ "utils/io/command_posix.cc" ]
|
||||
sources += [ "utils/io/tmpfile_posix.cc" ]
|
||||
} else if (is_win) {
|
||||
sources += [ "utils/io/command_windows.cc" ]
|
||||
sources += [ "utils/io/tmpfile_windows.cc" ]
|
||||
} else {
|
||||
sources += [ "utils/io/command_other.cc" ]
|
||||
sources += [ "utils/io/tmpfile_other.cc" ]
|
||||
}
|
||||
|
||||
public_deps = [ ":libtint_core_all_src" ]
|
||||
}
|
||||
|
||||
###############################################################################
|
||||
# Helper library for validating generated shaders
|
||||
# As this depends on tint_utils_io, this is only to be used by tests and sample
|
||||
# executable
|
||||
###############################################################################
|
||||
source_set("tint_val") {
|
||||
sources = [
|
||||
"val/val.cc",
|
||||
"val/val.h",
|
||||
]
|
||||
public_deps = [ ":tint_utils_io" ]
|
||||
}
|
||||
|
||||
###############################################################################
|
||||
# Library - Tint core and optional modules of libtint
|
||||
###############################################################################
|
||||
|
@ -464,9 +501,9 @@ libtint_source_set("libtint_core_all_src") {
|
|||
"sem/i32_type.cc",
|
||||
"sem/i32_type.h",
|
||||
"sem/info.h",
|
||||
"sem/intrinsic.h",
|
||||
"sem/intrinsic_type.cc",
|
||||
"sem/intrinsic_type.h",
|
||||
"sem/intrinsic.h",
|
||||
"sem/matrix_type.cc",
|
||||
"sem/matrix_type.h",
|
||||
"sem/multisampled_texture_type.cc",
|
||||
|
|
|
@ -449,6 +449,23 @@ if(${TINT_BUILD_HLSL_WRITER})
|
|||
)
|
||||
endif()
|
||||
|
||||
## Tint IO utilities. Used by tint_val.
|
||||
add_library(tint_utils_io
|
||||
utils/io/command_${TINT_OS_CC_SUFFIX}.cc
|
||||
utils/io/command.h
|
||||
utils/io/tmpfile_${TINT_OS_CC_SUFFIX}.cc
|
||||
utils/io/tmpfile.h
|
||||
)
|
||||
tint_default_compile_options(tint_utils_io)
|
||||
|
||||
## Tint validation utilities. Used by tests and the tint executable.
|
||||
add_library(tint_val
|
||||
val/val.cc
|
||||
val/val.h
|
||||
)
|
||||
tint_default_compile_options(tint_val)
|
||||
target_link_libraries(tint_val tint_utils_io)
|
||||
|
||||
## Tint library
|
||||
add_library(libtint ${TINT_LIB_SRCS})
|
||||
tint_default_compile_options(libtint)
|
||||
|
@ -603,16 +620,12 @@ if(${TINT_BUILD_TESTS})
|
|||
sem/type_manager_test.cc
|
||||
sem/u32_type_test.cc
|
||||
sem/vector_type_test.cc
|
||||
utils/command_${TINT_OS_CC_SUFFIX}.cc
|
||||
utils/command_test.cc
|
||||
utils/command.h
|
||||
utils/get_or_create_test.cc
|
||||
utils/hash_test.cc
|
||||
utils/io/command_test.cc
|
||||
utils/io/tmpfile_test.cc
|
||||
utils/math_test.cc
|
||||
utils/scoped_assignment_test.cc
|
||||
utils/tmpfile_${TINT_OS_CC_SUFFIX}.cc
|
||||
utils/tmpfile_test.cc
|
||||
utils/tmpfile.h
|
||||
utils/unique_vector_test.cc
|
||||
writer/append_vector_test.cc
|
||||
writer/float_to_string_test.cc
|
||||
|
@ -913,7 +926,7 @@ if(${TINT_BUILD_TESTS})
|
|||
## Test executable
|
||||
target_include_directories(
|
||||
tint_unittests PRIVATE ${gmock_SOURCE_DIR}/include)
|
||||
target_link_libraries(tint_unittests libtint gmock)
|
||||
target_link_libraries(tint_unittests libtint gmock tint_val)
|
||||
tint_default_compile_options(tint_unittests)
|
||||
|
||||
if(${TINT_BUILD_SPV_READER} OR ${TINT_BUILD_SPV_WRITER})
|
||||
|
|
|
@ -14,7 +14,7 @@
|
|||
|
||||
#include "gmock/gmock.h"
|
||||
#include "src/reader/spirv/parser_impl_test_helper.h"
|
||||
#include "src/utils/command.h"
|
||||
#include "src/utils/io/command.h"
|
||||
#include "src/writer/hlsl/test_helper.h"
|
||||
#include "src/writer/msl/test_helper.h"
|
||||
|
||||
|
|
|
@ -12,8 +12,8 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#ifndef SRC_UTILS_COMMAND_H_
|
||||
#define SRC_UTILS_COMMAND_H_
|
||||
#ifndef SRC_UTILS_IO_COMMAND_H_
|
||||
#define SRC_UTILS_IO_COMMAND_H_
|
||||
|
||||
#include <string>
|
||||
#include <utility>
|
||||
|
@ -81,4 +81,4 @@ class Command {
|
|||
} // namespace utils
|
||||
} // namespace tint
|
||||
|
||||
#endif // SRC_UTILS_COMMAND_H_
|
||||
#endif // SRC_UTILS_IO_COMMAND_H_
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/command.h"
|
||||
#include "src/utils/io/command.h"
|
||||
|
||||
namespace tint {
|
||||
namespace utils {
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/command.h"
|
||||
#include "src/utils/io/command.h"
|
||||
|
||||
#include <sys/poll.h>
|
||||
#include <sys/stat.h>
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/command.h"
|
||||
#include "src/utils/io/command.h"
|
||||
|
||||
#include "gtest/gtest.h"
|
||||
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/command.h"
|
||||
#include "src/utils/io/command.h"
|
||||
|
||||
#define WIN32_LEAN_AND_MEAN 1
|
||||
#include <Windows.h>
|
|
@ -12,8 +12,8 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#ifndef SRC_UTILS_TMPFILE_H_
|
||||
#define SRC_UTILS_TMPFILE_H_
|
||||
#ifndef SRC_UTILS_IO_TMPFILE_H_
|
||||
#define SRC_UTILS_IO_TMPFILE_H_
|
||||
|
||||
#include <sstream>
|
||||
#include <string>
|
||||
|
@ -73,4 +73,4 @@ class TmpFile {
|
|||
} // namespace utils
|
||||
} // namespace tint
|
||||
|
||||
#endif // SRC_UTILS_TMPFILE_H_
|
||||
#endif // SRC_UTILS_IO_TMPFILE_H_
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/tmpfile.h"
|
||||
#include "src/utils/io/tmpfile.h"
|
||||
|
||||
namespace tint {
|
||||
namespace utils {
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/tmpfile.h"
|
||||
#include "src/utils/io/tmpfile.h"
|
||||
|
||||
#include <unistd.h>
|
||||
#include <limits>
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/tmpfile.h"
|
||||
#include "src/utils/io/tmpfile.h"
|
||||
|
||||
#include <fstream>
|
||||
|
|
@ -12,7 +12,7 @@
|
|||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/utils/tmpfile.h"
|
||||
#include "src/utils/io/tmpfile.h"
|
||||
|
||||
#include <stdio.h>
|
||||
#include <cstdio>
|
|
@ -0,0 +1,133 @@
|
|||
// Copyright 2021 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/val/val.h"
|
||||
|
||||
#include "src/ast/module.h"
|
||||
#include "src/program.h"
|
||||
#include "src/utils/io/command.h"
|
||||
#include "src/utils/io/tmpfile.h"
|
||||
|
||||
namespace tint {
|
||||
namespace val {
|
||||
|
||||
Result Hlsl(const std::string& dxc_path,
|
||||
const std::string& source,
|
||||
Program* program) {
|
||||
Result result;
|
||||
|
||||
auto dxc = utils::Command(dxc_path);
|
||||
if (!dxc.Found()) {
|
||||
result.output = "DXC not found at '" + std::string(dxc_path) + "'";
|
||||
result.failed = true;
|
||||
return result;
|
||||
}
|
||||
|
||||
result.source = source;
|
||||
|
||||
utils::TmpFile file;
|
||||
file << source;
|
||||
|
||||
bool found_an_entrypoint = false;
|
||||
for (auto* func : program->AST().Functions()) {
|
||||
if (func->IsEntryPoint()) {
|
||||
found_an_entrypoint = true;
|
||||
|
||||
const char* profile = "";
|
||||
|
||||
switch (func->pipeline_stage()) {
|
||||
case ast::PipelineStage::kNone:
|
||||
result.output = "Invalid PipelineStage";
|
||||
result.failed = true;
|
||||
return result;
|
||||
case ast::PipelineStage::kVertex:
|
||||
profile = "-T vs_6_0";
|
||||
break;
|
||||
case ast::PipelineStage::kFragment:
|
||||
profile = "-T ps_6_0";
|
||||
break;
|
||||
case ast::PipelineStage::kCompute:
|
||||
profile = "-T cs_6_0";
|
||||
break;
|
||||
}
|
||||
|
||||
auto name = program->Symbols().NameFor(func->symbol());
|
||||
auto res = dxc(profile, "-E " + name, file.Path());
|
||||
if (!res.out.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += res.out;
|
||||
}
|
||||
if (!res.err.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += res.err;
|
||||
}
|
||||
result.failed = (res.error_code != 0);
|
||||
}
|
||||
}
|
||||
|
||||
if (!found_an_entrypoint) {
|
||||
result.output = "No entrypoint found";
|
||||
result.failed = true;
|
||||
return result;
|
||||
}
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
Result Msl(const std::string& xcrun_path, const std::string& source) {
|
||||
Result result;
|
||||
|
||||
auto xcrun = utils::Command(xcrun_path);
|
||||
if (!xcrun.Found()) {
|
||||
result.output = "xcrun not found at '" + std::string(xcrun_path) + "'";
|
||||
result.failed = true;
|
||||
return result;
|
||||
}
|
||||
|
||||
result.source = source;
|
||||
|
||||
utils::TmpFile file(".metal");
|
||||
file << result.source;
|
||||
|
||||
#ifdef _WIN32
|
||||
// On Windows, we should actually be running metal.exe from the Metal
|
||||
// Developer Tools for Windows
|
||||
auto res = xcrun("-x", "metal", "-c", "-o", "NUL", file.Path());
|
||||
#else
|
||||
auto res =
|
||||
xcrun("-sdk", "macosx", "metal", "-o", "/dev/null", "-c", file.Path());
|
||||
#endif
|
||||
if (!res.out.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += res.out;
|
||||
}
|
||||
if (!res.err.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += res.err;
|
||||
}
|
||||
result.failed = (res.error_code != 0);
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
} // namespace val
|
||||
} // namespace tint
|
|
@ -0,0 +1,58 @@
|
|||
// Copyright 2021 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_VAL_VAL_H_
|
||||
#define SRC_VAL_VAL_H_
|
||||
|
||||
#include <string>
|
||||
|
||||
// Forward declarations
|
||||
namespace tint {
|
||||
class Program;
|
||||
} // namespace tint
|
||||
|
||||
namespace tint {
|
||||
namespace val {
|
||||
|
||||
/// The return structure of Validate()
|
||||
struct Result {
|
||||
/// True if validation passed
|
||||
bool failed = false;
|
||||
/// Output of DXC.
|
||||
std::string output;
|
||||
/// The generated source that was compiled
|
||||
std::string source;
|
||||
};
|
||||
|
||||
/// Hlsl attempts to compile the shader with DXC, verifying that the shader
|
||||
/// compiles successfully.
|
||||
/// @param dxc_path path to DXC
|
||||
/// @param source the generated HLSL source
|
||||
/// @param program the HLSL program
|
||||
/// @return the result of the compile
|
||||
Result Hlsl(const std::string& dxc_path,
|
||||
const std::string& source,
|
||||
Program* program);
|
||||
|
||||
/// Msl attempts to compile the shader with the Metal Shader Compiler,
|
||||
/// verifying that the shader compiles successfully.
|
||||
/// @param xcrun_path path to xcrun
|
||||
/// @param source the generated MSL source
|
||||
/// @return the result of the compile
|
||||
Result Msl(const std::string& xcrun_path, const std::string& source);
|
||||
|
||||
} // namespace val
|
||||
} // namespace tint
|
||||
|
||||
#endif // SRC_VAL_VAL_H_
|
|
@ -14,9 +14,6 @@
|
|||
|
||||
#include "src/writer/hlsl/test_helper.h"
|
||||
|
||||
#include "src/utils/command.h"
|
||||
#include "src/utils/tmpfile.h"
|
||||
|
||||
namespace tint {
|
||||
namespace writer {
|
||||
namespace hlsl {
|
||||
|
@ -31,81 +28,16 @@ void EnableHLSLValidation(const char* dxc) {
|
|||
dxc_path = dxc;
|
||||
}
|
||||
|
||||
CompileResult Compile(Program* program, GeneratorImpl* generator) {
|
||||
CompileResult result;
|
||||
|
||||
val::Result Validate(Program* program, GeneratorImpl* generator) {
|
||||
if (!dxc_path) {
|
||||
result.status = CompileResult::Status::kVerificationNotEnabled;
|
||||
return result;
|
||||
}
|
||||
|
||||
auto dxc = utils::Command(dxc_path);
|
||||
if (!dxc.Found()) {
|
||||
result.output = "DXC not found at '" + std::string(dxc_path) + "'";
|
||||
result.status = CompileResult::Status::kFailed;
|
||||
return result;
|
||||
return val::Result{};
|
||||
}
|
||||
|
||||
std::ostringstream hlsl;
|
||||
if (!generator->Generate(hlsl)) {
|
||||
result.output = generator->error();
|
||||
result.status = CompileResult::Status::kFailed;
|
||||
return result;
|
||||
return {true, generator->error(), ""};
|
||||
}
|
||||
result.hlsl = hlsl.str();
|
||||
|
||||
utils::TmpFile file;
|
||||
file << result.hlsl;
|
||||
|
||||
bool found_an_entrypoint = false;
|
||||
for (auto* func : program->AST().Functions()) {
|
||||
if (func->IsEntryPoint()) {
|
||||
found_an_entrypoint = true;
|
||||
|
||||
const char* profile = "";
|
||||
|
||||
switch (func->pipeline_stage()) {
|
||||
case ast::PipelineStage::kNone:
|
||||
result.output = "Invalid PipelineStage";
|
||||
result.status = CompileResult::Status::kFailed;
|
||||
return result;
|
||||
case ast::PipelineStage::kVertex:
|
||||
profile = "-T vs_6_0";
|
||||
break;
|
||||
case ast::PipelineStage::kFragment:
|
||||
profile = "-T ps_6_0";
|
||||
break;
|
||||
case ast::PipelineStage::kCompute:
|
||||
profile = "-T cs_6_0";
|
||||
break;
|
||||
}
|
||||
|
||||
auto name = program->Symbols().NameFor(func->symbol());
|
||||
auto res = dxc(profile, "-E " + name, file.Path());
|
||||
if (!res.out.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += res.out;
|
||||
}
|
||||
if (!res.err.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += res.err;
|
||||
}
|
||||
result.status = (res.error_code == 0) ? CompileResult::Status::kSuccess
|
||||
: CompileResult::Status::kFailed;
|
||||
}
|
||||
}
|
||||
|
||||
if (!found_an_entrypoint) {
|
||||
result.output = "No entrypoint found";
|
||||
result.status = CompileResult::Status::kFailed;
|
||||
return result;
|
||||
}
|
||||
|
||||
return result;
|
||||
return val::Hlsl(dxc_path, hlsl.str(), program);
|
||||
}
|
||||
|
||||
} // namespace hlsl
|
||||
|
|
|
@ -23,6 +23,7 @@
|
|||
#include "src/transform/hlsl.h"
|
||||
#include "src/transform/manager.h"
|
||||
#include "src/transform/renamer.h"
|
||||
#include "src/val/val.h"
|
||||
#include "src/writer/hlsl/generator_impl.h"
|
||||
|
||||
namespace tint {
|
||||
|
@ -34,23 +35,11 @@ namespace hlsl {
|
|||
/// @param dxc_path the path to the DXC executable
|
||||
void EnableHLSLValidation(const char* dxc_path);
|
||||
|
||||
/// The return structure of Compile()
|
||||
struct CompileResult {
|
||||
/// Status is an enumerator of status codes from Compile()
|
||||
enum class Status { kSuccess, kFailed, kVerificationNotEnabled };
|
||||
/// The resulting status of the compile
|
||||
Status status;
|
||||
/// Output of DXC.
|
||||
std::string output;
|
||||
/// The HLSL source that was compiled
|
||||
std::string hlsl;
|
||||
};
|
||||
|
||||
/// Compile attempts to compile the shader with DXC if found on PATH.
|
||||
/// Validate attempts to compile the shader with DXC if found on PATH.
|
||||
/// @param program the HLSL program
|
||||
/// @param generator the HLSL generator
|
||||
/// @return the result of the compile
|
||||
CompileResult Compile(Program* program, GeneratorImpl* generator);
|
||||
val::Result Validate(Program* program, GeneratorImpl* generator);
|
||||
|
||||
/// Helper class for testing
|
||||
template <typename BODY>
|
||||
|
@ -122,10 +111,10 @@ class TestHelperBase : public BODY, public ProgramBuilder {
|
|||
/// If DXC finds problems the test will fail.
|
||||
/// If DXC is not on `PATH` then Validate() does nothing.
|
||||
void Validate() const {
|
||||
auto res = Compile(program.get(), gen_.get());
|
||||
if (res.status == CompileResult::Status::kFailed) {
|
||||
auto res = hlsl::Validate(program.get(), gen_.get());
|
||||
if (res.failed) {
|
||||
FAIL() << "HLSL Validation failed.\n\n"
|
||||
<< res.hlsl << "\n\n"
|
||||
<< res.source << "\n\n"
|
||||
<< res.output;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -14,8 +14,8 @@
|
|||
|
||||
#include "src/writer/msl/test_helper.h"
|
||||
|
||||
#include "src/utils/command.h"
|
||||
#include "src/utils/tmpfile.h"
|
||||
#include "src/utils/io/command.h"
|
||||
#include "src/utils/io/tmpfile.h"
|
||||
|
||||
namespace tint {
|
||||
namespace writer {
|
||||
|
@ -31,50 +31,16 @@ void EnableMSLValidation(const char* xcrun) {
|
|||
xcrun_path = xcrun;
|
||||
}
|
||||
|
||||
CompileResult Compile(Program* program) {
|
||||
CompileResult result;
|
||||
|
||||
val::Result Validate(Program* program) {
|
||||
if (!xcrun_path) {
|
||||
result.status = CompileResult::Status::kVerificationNotEnabled;
|
||||
return result;
|
||||
}
|
||||
|
||||
auto xcrun = utils::Command(xcrun_path);
|
||||
if (!xcrun.Found()) {
|
||||
result.output = "xcrun not found at '" + std::string(xcrun_path) + "'";
|
||||
result.status = CompileResult::Status::kFailed;
|
||||
return result;
|
||||
return val::Result{};
|
||||
}
|
||||
|
||||
auto gen = std::make_unique<GeneratorImpl>(program);
|
||||
if (!gen->Generate()) {
|
||||
result.output = gen->error();
|
||||
result.status = CompileResult::Status::kFailed;
|
||||
return result;
|
||||
return {true, gen->error(), ""};
|
||||
}
|
||||
result.msl = gen->result();
|
||||
|
||||
utils::TmpFile file(".metal");
|
||||
file << result.msl;
|
||||
|
||||
auto xcrun_res =
|
||||
xcrun("-sdk", "macosx", "metal", "-o", "/dev/null", "-c", file.Path());
|
||||
if (!xcrun_res.out.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += xcrun_res.out;
|
||||
}
|
||||
if (!xcrun_res.err.empty()) {
|
||||
if (!result.output.empty()) {
|
||||
result.output += "\n";
|
||||
}
|
||||
result.output += xcrun_res.err;
|
||||
}
|
||||
result.status = (xcrun_res.error_code == 0) ? CompileResult::Status::kSuccess
|
||||
: CompileResult::Status::kFailed;
|
||||
|
||||
return result;
|
||||
return val::Msl(xcrun_path, gen->result());
|
||||
}
|
||||
|
||||
} // namespace msl
|
||||
|
|
|
@ -22,6 +22,7 @@
|
|||
#include "gtest/gtest.h"
|
||||
#include "src/program_builder.h"
|
||||
#include "src/transform/msl.h"
|
||||
#include "src/val/val.h"
|
||||
#include "src/writer/msl/generator_impl.h"
|
||||
|
||||
namespace tint {
|
||||
|
@ -33,22 +34,10 @@ namespace msl {
|
|||
/// @param xcrun_path the path to the `xcrun` executable
|
||||
void EnableMSLValidation(const char* xcrun_path);
|
||||
|
||||
/// The return structure of Compile()
|
||||
struct CompileResult {
|
||||
/// Status is an enumerator of status codes from Compile()
|
||||
enum class Status { kSuccess, kFailed, kVerificationNotEnabled };
|
||||
/// The resulting status of the compilation
|
||||
Status status;
|
||||
/// Output of the Metal compiler
|
||||
std::string output;
|
||||
/// The MSL source that was compiled
|
||||
std::string msl;
|
||||
};
|
||||
|
||||
/// Compile attempts to compile the shader with xcrun if found on PATH.
|
||||
/// Validate attempts to compile the shader with DXC if found on PATH.
|
||||
/// @param program the MSL program
|
||||
/// @return the result of the compile
|
||||
CompileResult Compile(Program* program);
|
||||
val::Result Validate(Program* program);
|
||||
|
||||
/// Helper class for testing
|
||||
template <typename BASE>
|
||||
|
@ -115,9 +104,11 @@ class TestHelperBase : public BASE, public ProgramBuilder {
|
|||
/// This function does nothing, if the Metal compiler path has not been
|
||||
/// configured by calling `EnableMSLValidation()`.
|
||||
void Validate() {
|
||||
auto res = Compile(program.get());
|
||||
if (res.status == CompileResult::Status::kFailed) {
|
||||
FAIL() << "MSL Validation failed.\n\n" << res.msl << "\n\n" << res.output;
|
||||
auto res = msl::Validate(program.get());
|
||||
if (res.failed) {
|
||||
FAIL() << "MSL Validation failed.\n\n"
|
||||
<< res.source << "\n\n"
|
||||
<< res.output;
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -90,11 +90,11 @@ source_set("tint_unittests_main") {
|
|||
sources = [ "../src/test_main.cc" ]
|
||||
configs += [ ":tint_unittests_config" ]
|
||||
deps += [
|
||||
":tint_test_helpers",
|
||||
":tint_unittests_hlsl_writer_src",
|
||||
":tint_unittests_msl_writer_src",
|
||||
":tint_unittests_spv_reader_src",
|
||||
"${tint_root_dir}/src:libtint",
|
||||
"${tint_root_dir}/src:tint_val",
|
||||
]
|
||||
}
|
||||
}
|
||||
|
@ -114,36 +114,6 @@ config("tint_unittests_config") {
|
|||
]
|
||||
}
|
||||
|
||||
source_set("tint_test_helpers") {
|
||||
testonly = true
|
||||
|
||||
sources = [
|
||||
"../src/ast/intrinsic_texture_helper_test.cc",
|
||||
"../src/ast/intrinsic_texture_helper_test.h",
|
||||
"../src/transform/test_helper.h",
|
||||
"../src/utils/command.h",
|
||||
"../src/utils/tmpfile.h",
|
||||
]
|
||||
|
||||
if (is_linux || is_mac) {
|
||||
sources += [ "../src/utils/command_posix.cc" ]
|
||||
sources += [ "../src/utils/tmpfile_posix.cc" ]
|
||||
} else if (is_win) {
|
||||
sources += [ "../src/utils/command_windows.cc" ]
|
||||
sources += [ "../src/utils/tmpfile_windows.cc" ]
|
||||
} else {
|
||||
sources += [ "../src/utils/command_other.cc" ]
|
||||
sources += [ "../src/utils/tmpfile_other.cc" ]
|
||||
}
|
||||
|
||||
configs += [ ":tint_unittests_config" ]
|
||||
|
||||
public_deps = [
|
||||
":gmock_and_gtest",
|
||||
"${tint_root_dir}/src:libtint",
|
||||
]
|
||||
}
|
||||
|
||||
template("tint_unittests_source_set") {
|
||||
source_set(target_name) {
|
||||
forward_variables_from(invoker, "*", [ "configs" ])
|
||||
|
@ -162,7 +132,11 @@ template("tint_unittests_source_set") {
|
|||
if (!defined(invoker.deps)) {
|
||||
deps = []
|
||||
}
|
||||
deps += [ ":tint_test_helpers" ]
|
||||
deps += [
|
||||
":gmock_and_gtest",
|
||||
"${tint_root_dir}/src:libtint",
|
||||
"${tint_root_dir}/src:tint_val",
|
||||
]
|
||||
}
|
||||
}
|
||||
|
||||
|
@ -203,6 +177,8 @@ tint_unittests_source_set("tint_unittests_core_src") {
|
|||
"../src/ast/identifier_expression_test.cc",
|
||||
"../src/ast/if_statement_test.cc",
|
||||
"../src/ast/int_literal_test.cc",
|
||||
"../src/ast/intrinsic_texture_helper_test.cc",
|
||||
"../src/ast/intrinsic_texture_helper_test.h",
|
||||
"../src/ast/location_decoration_test.cc",
|
||||
"../src/ast/loop_statement_test.cc",
|
||||
"../src/ast/matrix_test.cc",
|
||||
|
@ -307,15 +283,16 @@ tint_unittests_source_set("tint_unittests_core_src") {
|
|||
"../src/transform/renamer_test.cc",
|
||||
"../src/transform/simplify_test.cc",
|
||||
"../src/transform/single_entry_point_test.cc",
|
||||
"../src/transform/test_helper.h",
|
||||
"../src/transform/transform_test.cc",
|
||||
"../src/transform/var_for_dynamic_index_test.cc",
|
||||
"../src/transform/vertex_pulling_test.cc",
|
||||
"../src/utils/command_test.cc",
|
||||
"../src/utils/get_or_create_test.cc",
|
||||
"../src/utils/hash_test.cc",
|
||||
"../src/utils/io/command_test.cc",
|
||||
"../src/utils/io/tmpfile_test.cc",
|
||||
"../src/utils/math_test.cc",
|
||||
"../src/utils/scoped_assignment_test.cc",
|
||||
"../src/utils/tmpfile_test.cc",
|
||||
"../src/utils/unique_vector_test.cc",
|
||||
"../src/writer/append_vector_test.cc",
|
||||
"../src/writer/float_to_string_test.cc",
|
||||
|
@ -365,7 +342,10 @@ tint_unittests_source_set("tint_unittests_spv_reader_src") {
|
|||
"../src/reader/spirv/usage_test.cc",
|
||||
]
|
||||
|
||||
deps = [ "${tint_root_dir}/src:libtint_spv_reader_src" ]
|
||||
deps = [
|
||||
":tint_unittests_core_src",
|
||||
"${tint_root_dir}/src:libtint_spv_reader_src",
|
||||
]
|
||||
}
|
||||
|
||||
tint_unittests_source_set("tint_unittests_spv_writer_src") {
|
||||
|
@ -406,6 +386,7 @@ tint_unittests_source_set("tint_unittests_spv_writer_src") {
|
|||
]
|
||||
|
||||
deps = [
|
||||
":tint_unittests_core_src",
|
||||
"${tint_root_dir}/src:libtint_spv_writer_src",
|
||||
"${tint_spirv_tools_dir}/:spvtools",
|
||||
]
|
||||
|
@ -489,7 +470,10 @@ tint_unittests_source_set("tint_unittests_wgsl_reader_src") {
|
|||
"../src/reader/wgsl/token_test.cc",
|
||||
]
|
||||
|
||||
deps = [ "${tint_root_dir}/src:libtint_wgsl_reader_src" ]
|
||||
deps = [
|
||||
":tint_unittests_core_src",
|
||||
"${tint_root_dir}/src:libtint_wgsl_reader_src",
|
||||
]
|
||||
}
|
||||
|
||||
tint_unittests_source_set("tint_unittests_wgsl_writer_src") {
|
||||
|
@ -525,7 +509,10 @@ tint_unittests_source_set("tint_unittests_wgsl_writer_src") {
|
|||
"../src/writer/wgsl/test_helper.h",
|
||||
]
|
||||
|
||||
deps = [ "${tint_root_dir}/src:libtint_wgsl_writer_src" ]
|
||||
deps = [
|
||||
":tint_unittests_core_src",
|
||||
"${tint_root_dir}/src:libtint_wgsl_writer_src",
|
||||
]
|
||||
}
|
||||
|
||||
tint_unittests_source_set("tint_unittests_msl_writer_src") {
|
||||
|
@ -563,7 +550,10 @@ tint_unittests_source_set("tint_unittests_msl_writer_src") {
|
|||
"../src/writer/msl/test_helper.h",
|
||||
]
|
||||
|
||||
deps = [ "${tint_root_dir}/src:libtint_msl_writer_src" ]
|
||||
deps = [
|
||||
":tint_unittests_core_src",
|
||||
"${tint_root_dir}/src:libtint_msl_writer_src",
|
||||
]
|
||||
}
|
||||
|
||||
tint_unittests_source_set("tint_unittests_hlsl_writer_src") {
|
||||
|
@ -603,7 +593,10 @@ tint_unittests_source_set("tint_unittests_hlsl_writer_src") {
|
|||
"../src/writer/hlsl/test_helper.h",
|
||||
]
|
||||
|
||||
deps = [ "${tint_root_dir}/src:libtint_hlsl_writer_src" ]
|
||||
deps = [
|
||||
":tint_unittests_core_src",
|
||||
"${tint_root_dir}/src:libtint_hlsl_writer_src",
|
||||
]
|
||||
}
|
||||
|
||||
source_set("tint_unittests_src") {
|
||||
|
|
|
@ -6,7 +6,7 @@ struct Buf{
|
|||
data : Arr;
|
||||
};
|
||||
|
||||
[[group(0), binding (0)]] var<storage> b : [[access(read)]] Buf;
|
||||
[[group(0), binding (0)]] var<storage> b : [[access(read_write)]] Buf;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main() {
|
||||
|
|
|
@ -1,5 +1,5 @@
|
|||
|
||||
ByteAddressBuffer b : register(t0, space0);
|
||||
RWByteAddressBuffer b : register(u0, space0);
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main() {
|
||||
|
|
|
@ -8,7 +8,7 @@ struct Buf {
|
|||
/* 0x0004 */ uint data[50];
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const device Buf& b [[buffer(0)]]) {
|
||||
kernel void tint_symbol(device Buf& b [[buffer(0)]]) {
|
||||
uint i = 0u;
|
||||
{
|
||||
bool tint_msl_is_first_1 = true;
|
||||
|
|
|
@ -17,7 +17,6 @@
|
|||
OpMemberDecorate %Buf 0 Offset 0
|
||||
OpMemberDecorate %Buf 1 Offset 4
|
||||
OpDecorate %_arr_uint_uint_50 ArrayStride 4
|
||||
OpDecorate %b NonWritable
|
||||
OpDecorate %b DescriptorSet 0
|
||||
OpDecorate %b Binding 0
|
||||
%uint = OpTypeInt 32 0
|
||||
|
|
|
@ -6,7 +6,7 @@ struct Buf {
|
|||
data : Arr;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<storage> b : [[access(read)]] Buf;
|
||||
[[group(0), binding(0)]] var<storage> b : [[access(read_write)]] Buf;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main() {
|
||||
|
|
|
@ -20,7 +20,7 @@ SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
|
|||
|
||||
function usage() {
|
||||
echo "test-all.sh is a simple wrapper around <tint>/tools/test-runner that"
|
||||
echo "injects the <tint>/tools directory as the second command line argument"
|
||||
echo "injects the <tint>/test directory as the second command line argument"
|
||||
echo
|
||||
echo "Usage of <tint>/tools/test-runner:"
|
||||
"${SCRIPT_DIR}/../tools/test-runner" --help
|
||||
|
@ -41,4 +41,4 @@ if [ ! -x "$TINT" ]; then
|
|||
exit 1
|
||||
fi
|
||||
|
||||
"${SCRIPT_DIR}/../tools/test-runner" ${@:2} "${TINT}" "${SCRIPT_DIR}"
|
||||
"${SCRIPT_DIR}/../tools/test-runner" "${@:2}" "${TINT}" "${SCRIPT_DIR}"
|
||||
|
|
|
@ -68,11 +68,13 @@ optional flags:`)
|
|||
}
|
||||
|
||||
func run() error {
|
||||
var formatList, filter string
|
||||
var formatList, filter, dxcPath, xcrunPath string
|
||||
numCPU := runtime.NumCPU()
|
||||
generateExpected := false
|
||||
flag.StringVar(&formatList, "format", "all", "comma separated list of formats to emit. Possible values are: all, wgsl, spvasm, msl, hlsl")
|
||||
flag.StringVar(&filter, "filter", "**.wgsl, **.spvasm, **.spv", "comma separated list of glob patterns for test files")
|
||||
flag.StringVar(&dxcPath, "dxc", "", "path to DXC executable for validating HLSL output")
|
||||
flag.StringVar(&xcrunPath, "xcrun", "", "path to xcrun executable for validating MSL output")
|
||||
flag.BoolVar(&generateExpected, "generate-expected", false, "create or update all expected outputs")
|
||||
flag.IntVar(&numCPU, "j", numCPU, "maximum number of concurrent threads to run tests")
|
||||
flag.Usage = showUsage
|
||||
|
@ -147,6 +149,46 @@ func run() error {
|
|||
}
|
||||
}
|
||||
|
||||
default_msl_exe := "xcrun"
|
||||
if runtime.GOOS == "windows" {
|
||||
default_msl_exe = "metal.exe"
|
||||
}
|
||||
|
||||
// If explicit verification compilers have been specified, check they exist.
|
||||
// Otherwise, look on PATH for them, but don't error if they cannot be found.
|
||||
for _, tool := range []struct {
|
||||
name string
|
||||
lang string
|
||||
path *string
|
||||
}{
|
||||
{"dxc", "hlsl", &dxcPath},
|
||||
{default_msl_exe, "msl", &xcrunPath},
|
||||
} {
|
||||
if *tool.path == "" {
|
||||
p, err := exec.LookPath(tool.name)
|
||||
if err == nil && fileutils.IsExe(p) {
|
||||
*tool.path = p
|
||||
}
|
||||
} else if !fileutils.IsExe(*tool.path) {
|
||||
return fmt.Errorf("%v not found at '%v'", tool.name, *tool.path)
|
||||
}
|
||||
|
||||
color.Set(color.FgCyan)
|
||||
fmt.Printf("%-4s", tool.lang)
|
||||
color.Unset()
|
||||
fmt.Printf(" validation ")
|
||||
if *tool.path == "" {
|
||||
color.Set(color.FgRed)
|
||||
fmt.Printf("DISABLED")
|
||||
} else {
|
||||
color.Set(color.FgGreen)
|
||||
fmt.Printf("ENABLED")
|
||||
}
|
||||
color.Unset()
|
||||
fmt.Println()
|
||||
}
|
||||
fmt.Println()
|
||||
|
||||
results := make([]map[outputFormat]*status, len(files))
|
||||
jobs := make(chan job, 256)
|
||||
|
||||
|
@ -157,7 +199,7 @@ func run() error {
|
|||
go func() {
|
||||
defer wg.Done()
|
||||
for job := range jobs {
|
||||
job.run(exe, generateExpected)
|
||||
job.run(exe, dxcPath, xcrunPath, generateExpected)
|
||||
}
|
||||
}()
|
||||
}
|
||||
|
@ -313,7 +355,7 @@ type job struct {
|
|||
result *status
|
||||
}
|
||||
|
||||
func (j job) run(exe string, generateExpected bool) {
|
||||
func (j job) run(exe, dxcPath, xcrunPath string, generateExpected bool) {
|
||||
// Is there an expected output?
|
||||
expected := loadExpectedFile(j.file, j.format)
|
||||
if strings.HasPrefix(expected, "SKIP") { // Special SKIP token
|
||||
|
@ -321,9 +363,28 @@ func (j job) run(exe string, generateExpected bool) {
|
|||
return
|
||||
}
|
||||
|
||||
expected = strings.ReplaceAll(expected, "\r\n", "\n")
|
||||
|
||||
args := []string{j.file, "--format", string(j.format)}
|
||||
|
||||
// Can we validate?
|
||||
switch j.format {
|
||||
case spvasm:
|
||||
args = append(args, "--validate") // spirv-val is statically linked, always available
|
||||
case hlsl:
|
||||
if dxcPath != "" {
|
||||
args = append(args, "--dxc", dxcPath)
|
||||
}
|
||||
case msl:
|
||||
if xcrunPath != "" {
|
||||
args = append(args, "--xcrun", xcrunPath)
|
||||
}
|
||||
}
|
||||
|
||||
// Invoke the compiler...
|
||||
var err error
|
||||
if ok, out := invoke(exe, j.file, "--format", string(j.format), "--dawn-validation"); ok {
|
||||
if ok, out := invoke(exe, args...); ok {
|
||||
out = strings.ReplaceAll(out, "\r\n", "\n")
|
||||
if generateExpected {
|
||||
// If --generate-expected was passed, write out the output
|
||||
err = saveExpectedFile(j.file, j.format, out)
|
||||
|
|
|
@ -0,0 +1,21 @@
|
|||
// Copyright 2021 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.
|
||||
|
||||
// Package fileutils contains utility functions for files
|
||||
package fileutils
|
||||
|
||||
// IsExe returns true if the file at path is an executable
|
||||
func IsExe(path string) bool {
|
||||
return true
|
||||
}
|
Loading…
Reference in New Issue