tint: add --overrides flag to specify pipeline overrides
And use this to define constants for override vars in unit tests that
would fail for HLSL.
Bug: tint:1519
Change-Id: I4fd15c517868694d2bcd81d563399f817ed74ae6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/88882
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc
index fdb625b..52dfd70 100644
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@@ -85,6 +85,7 @@
bool use_fxc = false;
std::string dxc_path;
std::string xcrun_path;
+ std::vector<std::string> overrides;
};
const char kUsage[] = R"(Usage: tint [options] <input-file>
@@ -117,7 +118,8 @@
--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)";
+ When specified, automatically enables --validate
+ --overrides -- Pipeline overrides as NAME=VALUE, comma-separated.)";
Format parse_format(const std::string& fmt) {
(void)fmt;
@@ -200,7 +202,7 @@
return Format::kNone;
}
-std::vector<std::string> split_transform_names(std::string list) {
+std::vector<std::string> split_on_comma(std::string list) {
std::vector<std::string> res;
std::stringstream str(list);
@@ -357,7 +359,7 @@
std::cerr << "Missing value for " << arg << std::endl;
return false;
}
- opts->transforms = split_transform_names(args[i]);
+ opts->transforms = split_on_comma(args[i]);
} else if (arg == "--parse-only") {
opts->parse_only = true;
} else if (arg == "--disable-workgroup-init") {
@@ -387,6 +389,13 @@
}
opts->xcrun_path = args[i];
opts->validate = true;
+ } else if (arg == "--overrides") {
+ ++i;
+ if (i >= args.size()) {
+ std::cerr << "Missing value for " << arg << std::endl;
+ return false;
+ }
+ opts->overrides = split_on_comma(args[i]);
} else if (!arg.empty()) {
if (arg[0] == '-') {
std::cerr << "Unrecognized option: " << arg << std::endl;
@@ -723,7 +732,7 @@
tint::val::Result res;
if (options.use_fxc) {
#ifdef _WIN32
- res = tint::val::HlslUsingFXC(result.hlsl, result.entry_points);
+ res = tint::val::HlslUsingFXC(result.hlsl, result.entry_points, options.overrides);
#else
res.failed = true;
res.output = "FXC can only be used on Windows. Sorry :X";
@@ -732,7 +741,8 @@
auto dxc =
tint::utils::Command::LookPath(options.dxc_path.empty() ? "dxc" : options.dxc_path);
if (dxc.Found()) {
- res = tint::val::HlslUsingDXC(dxc.Path(), result.hlsl, result.entry_points);
+ res = tint::val::HlslUsingDXC(dxc.Path(), result.hlsl, result.entry_points,
+ options.overrides);
} else {
res.failed = true;
res.output = "DXC executable not found. Cannot validate";
diff --git a/src/tint/val/hlsl.cc b/src/tint/val/hlsl.cc
index 164152d..49c3850 100644
--- a/src/tint/val/hlsl.cc
+++ b/src/tint/val/hlsl.cc
@@ -30,7 +30,8 @@
Result HlslUsingDXC(const std::string& dxc_path,
const std::string& source,
- const EntryPointList& entry_points) {
+ const EntryPointList& entry_points,
+ const std::vector<std::string>& overrides) {
Result result;
auto dxc = utils::Command(dxc_path);
@@ -69,7 +70,13 @@
"/Zpr " // D3DCOMPILE_PACK_MATRIX_ROW_MAJOR
"/Gis"; // D3DCOMPILE_IEEE_STRICTNESS
- auto res = dxc(profile, "-E " + ep.first, compileFlags, file.Path());
+ std::string defs;
+ defs.reserve(overrides.size() * 20);
+ for (auto& o : overrides) {
+ defs += "/D" + o + " ";
+ }
+
+ auto res = dxc(profile, "-E " + ep.first, compileFlags, file.Path(), defs);
if (!res.out.empty()) {
if (!result.output.empty()) {
result.output += "\n";
@@ -95,7 +102,9 @@
}
#ifdef _WIN32
-Result HlslUsingFXC(const std::string& source, const EntryPointList& entry_points) {
+Result HlslUsingFXC(const std::string& source,
+ const EntryPointList& entry_points,
+ const std::vector<std::string>& overrides) {
Result result;
// This library leaks if an error happens in this function, but it is ok
@@ -139,12 +148,26 @@
UINT compileFlags = D3DCOMPILE_OPTIMIZATION_LEVEL0 | D3DCOMPILE_PACK_MATRIX_ROW_MAJOR |
D3DCOMPILE_IEEE_STRICTNESS;
+ auto overrides_copy = overrides; // Copy so that we can replace '=' with '\0'
+ std::vector<D3D_SHADER_MACRO> macros;
+ macros.reserve(overrides_copy.size() * 2);
+ for (auto& o : overrides_copy) {
+ if (auto sep = o.find_first_of('='); sep != std::string::npos) {
+ // Replace '=' with '\0' so we can point directly into the allocated string buffer
+ o[sep] = '\0';
+ macros.push_back(D3D_SHADER_MACRO{&o[0], &o[sep + 1]});
+ } else {
+ macros.emplace_back(D3D_SHADER_MACRO{o.c_str(), NULL});
+ }
+ }
+ macros.emplace_back(D3D_SHADER_MACRO{NULL, NULL});
+
ComPtr<ID3DBlob> compiledShader;
ComPtr<ID3DBlob> errors;
HRESULT cr = d3dCompile(source.c_str(), // pSrcData
source.length(), // SrcDataSize
nullptr, // pSourceName
- nullptr, // pDefines
+ macros.data(), // pDefines
nullptr, // pInclude
ep.first.c_str(), // pEntrypoint
profile, // pTarget
diff --git a/src/tint/val/val.h b/src/tint/val/val.h
index 4b3fff9..c869efb 100644
--- a/src/tint/val/val.h
+++ b/src/tint/val/val.h
@@ -43,18 +43,23 @@
/// @param dxc_path path to DXC
/// @param source the generated HLSL source
/// @param entry_points the list of entry points to validate
+/// @param overrides optional list of pipeline overrides
/// @return the result of the compile
Result HlslUsingDXC(const std::string& dxc_path,
const std::string& source,
- const EntryPointList& entry_points);
+ const EntryPointList& entry_points,
+ const std::vector<std::string>& overrides);
#ifdef _WIN32
/// Hlsl attempts to compile the shader with FXC, verifying that the shader
/// compiles successfully.
/// @param source the generated HLSL source
/// @param entry_points the list of entry points to validate
+/// @param overrides optional list of pipeline overrides
/// @return the result of the compile
-Result HlslUsingFXC(const std::string& source, const EntryPointList& entry_points);
+Result HlslUsingFXC(const std::string& source,
+ const EntryPointList& entry_points,
+ const std::vector<std::string>& overrides);
#endif // _WIN32
/// Msl attempts to compile the shader with the Metal Shader Compiler,
diff --git a/test/tint/var/override/named/no_init/bool.wgsl b/test/tint/var/override/named/no_init/bool.wgsl
index 04ca499..9c595f2 100644
--- a/test/tint/var/override/named/no_init/bool.wgsl
+++ b/test/tint/var/override/named/no_init/bool.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_0=0
override o : bool;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl
index 977045c..3ea88ba 100644
--- a/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl
+++ b/test/tint/var/override/named/no_init/bool.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_0
#error spec constant required for constant id 0
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_N41FiO:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-
diff --git a/test/tint/var/override/named/no_init/f32.wgsl b/test/tint/var/override/named/no_init/f32.wgsl
index 50dee2f..75b0841 100644
--- a/test/tint/var/override/named/no_init/f32.wgsl
+++ b/test/tint/var/override/named/no_init/f32.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_0=0
override o : f32;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl
index 2ae175b..34bec48 100644
--- a/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl
+++ b/test/tint/var/override/named/no_init/f32.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_0
#error spec constant required for constant id 0
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_gYmCvK:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-
diff --git a/test/tint/var/override/named/no_init/i32.wgsl b/test/tint/var/override/named/no_init/i32.wgsl
index d91c2f2..847fb27 100644
--- a/test/tint/var/override/named/no_init/i32.wgsl
+++ b/test/tint/var/override/named/no_init/i32.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_0=0
override o : i32;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl
index 7685972..034990e 100644
--- a/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl
+++ b/test/tint/var/override/named/no_init/i32.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_0
#error spec constant required for constant id 0
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_kYjxDo:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-
diff --git a/test/tint/var/override/named/no_init/u32.wgsl b/test/tint/var/override/named/no_init/u32.wgsl
index 5c11bc9..a56cecb 100644
--- a/test/tint/var/override/named/no_init/u32.wgsl
+++ b/test/tint/var/override/named/no_init/u32.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_0=0
override o : u32;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl b/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl
index 4d5ed7a..fb48bba 100644
--- a/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl
+++ b/test/tint/var/override/named/no_init/u32.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_0
#error spec constant required for constant id 0
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_7FSY6m:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-
diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl b/test/tint/var/override/numbered/no_init/bool.wgsl
index 98ad2ec..34818a1 100644
--- a/test/tint/var/override/numbered/no_init/bool.wgsl
+++ b/test/tint/var/override/numbered/no_init/bool.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_1234=0
@id(1234) override o : bool;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl
index eca5977..fd25726 100644
--- a/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl
+++ b/test/tint/var/override/numbered/no_init/bool.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_1234
#error spec constant required for constant id 1234
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_4uqZfq:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-
diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl b/test/tint/var/override/numbered/no_init/f32.wgsl
index d15b4ef..dde41d9 100644
--- a/test/tint/var/override/numbered/no_init/f32.wgsl
+++ b/test/tint/var/override/numbered/no_init/f32.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_1234=0
@id(1234) override o : f32;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl
index f5c5a75..2ae7042 100644
--- a/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl
+++ b/test/tint/var/override/numbered/no_init/f32.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_1234
#error spec constant required for constant id 1234
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_UepNPT:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-
diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl b/test/tint/var/override/numbered/no_init/i32.wgsl
index dca18e2..afc07ef 100644
--- a/test/tint/var/override/numbered/no_init/i32.wgsl
+++ b/test/tint/var/override/numbered/no_init/i32.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_1234=0
@id(1234) override o : i32;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl
index 0b62c77..a02d224 100644
--- a/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl
+++ b/test/tint/var/override/numbered/no_init/i32.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_1234
#error spec constant required for constant id 1234
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_mwzbqN:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-
diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl b/test/tint/var/override/numbered/no_init/u32.wgsl
index 5cdb598..be20c98 100644
--- a/test/tint/var/override/numbered/no_init/u32.wgsl
+++ b/test/tint/var/override/numbered/no_init/u32.wgsl
@@ -1,3 +1,4 @@
+// flags: --overrides WGSL_SPEC_CONSTANT_1234=0
@id(1234) override o : u32;
@stage(compute) @workgroup_size(1)
diff --git a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl
index 443a2b1..4a3a2e1 100644
--- a/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl
+++ b/test/tint/var/override/numbered/no_init/u32.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#ifndef WGSL_SPEC_CONSTANT_1234
#error spec constant required for constant id 1234
#endif
@@ -9,8 +7,3 @@
void main() {
return;
}
-/tmp/tint_U0EZFZ:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-
