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>

src/tint/cmd/main.cc

@@ -85,6 +85,7 @@ struct Options {
     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 @@ ${transforms}
   --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 @@ Format infer_format(const std::string& filename) {
     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 @@ bool ParseArgs(const std::vector<std::string>& args, Options* opts) {
                 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 @@ bool ParseArgs(const std::vector<std::string>& args, Options* opts) {
             }
             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 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) {
         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 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) {
             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";

src/tint/val/hlsl.cc

@@ -30,7 +30,8 @@ namespace tint::val {
 
 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 @@ Result HlslUsingDXC(const std::string& dxc_path,
             "/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 @@ Result HlslUsingDXC(const std::string& dxc_path,
 }
 
 #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 @@ Result HlslUsingFXC(const std::string& source, const EntryPointList& entry_point
         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

src/tint/val/val.h

@@ -43,18 +43,23 @@ struct Result {
 /// @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,

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)

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 @@ static const bool o = WGSL_SPEC_CONSTANT_0;
 void main() {
   return;
 }
-/tmp/tint_N41FiO:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-

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)

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 @@ static const float o = WGSL_SPEC_CONSTANT_0;
 void main() {
   return;
 }
-/tmp/tint_gYmCvK:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-

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)

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 @@ static const int o = WGSL_SPEC_CONSTANT_0;
 void main() {
   return;
 }
-/tmp/tint_kYjxDo:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-

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)

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 @@ static const uint o = WGSL_SPEC_CONSTANT_0;
 void main() {
   return;
 }
-/tmp/tint_7FSY6m:2:2: error: spec constant required for constant id 0
-#error spec constant required for constant id 0
- ^
-
-

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)

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 @@ static const bool o = WGSL_SPEC_CONSTANT_1234;
 void main() {
   return;
 }
-/tmp/tint_4uqZfq:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-

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)

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 @@ static const float o = WGSL_SPEC_CONSTANT_1234;
 void main() {
   return;
 }
-/tmp/tint_UepNPT:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-

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)

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 @@ static const int o = WGSL_SPEC_CONSTANT_1234;
 void main() {
   return;
 }
-/tmp/tint_mwzbqN:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-

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)

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 @@ static const uint o = WGSL_SPEC_CONSTANT_1234;
 void main() {
   return;
 }
-/tmp/tint_U0EZFZ:2:2: error: spec constant required for constant id 1234
-#error spec constant required for constant id 1234
- ^
-
-