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:
Ben Clayton 2021-06-02 21:02:34 +00:00 committed by Tint LUCI CQ
parent a7a23eaa9e
commit a70c14dbce
32 changed files with 550 additions and 261 deletions

View File

@ -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",

View File

@ -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)

View File

@ -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,31 +797,45 @@ 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());
}
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());
}
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());
}
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());
}
break;
#else
std::cerr << "HLSL writer not enabled in tint build" << std::endl;
return 1;
#endif // TINT_BUILD_HLSL_WRITER
if (!writer) {
default:
std::cerr << "Unknown output format specified" << std::endl;
return 1;
}
@ -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)) {
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([&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;
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())) {
dawn_validation_failed = true;
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());

View File

@ -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",

View File

@ -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})

View File

@ -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"

View File

@ -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_

View File

@ -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 {

View File

@ -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>

View File

@ -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"

View File

@ -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>

View File

@ -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_

View File

@ -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 {

View File

@ -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>

View File

@ -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>

View File

@ -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>

133
src/val/val.cc Normal file
View File

@ -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

58
src/val/val.h Normal file
View File

@ -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_

View File

@ -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

View File

@ -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;
}
}

View File

@ -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

View File

@ -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;
}
}

View File

@ -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") {

View File

@ -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() {

View File

@ -1,5 +1,5 @@
ByteAddressBuffer b : register(t0, space0);
RWByteAddressBuffer b : register(u0, space0);
[numthreads(1, 1, 1)]
void main() {

View File

@ -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;

View File

@ -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

View File

@ -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() {

View File

@ -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}"

View File

@ -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)

View File

@ -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
}