From 84b48cf30cf42e8a74903d3dd28f0a28380ae076 Mon Sep 17 00:00:00 2001 From: Zhaoming Jiang Date: Wed, 1 Jun 2022 10:23:51 +0000 Subject: [PATCH] Tint: num_workgroups use free binding group if not specified In this patch NumWorkgroupsFromUniform::Config changed to storage std::optional, and if it has no value, NumWorkgroupsFromUniform will choose a free binding group, i.e. binding 0 of the largest used group plus 1 is used if at least one resource is bound, otherwise group 0 binding 0 is used. Tint CLI is also changed to provide a --hlsl-root-constant-binding-point option allowing user to specify the binding point for num_workgroups uniform buffer. Bug: tint:1566 Change-Id: I3b8c22a4276bab722d901f5b07d23a268786c417 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/91980 Kokoro: Kokoro Reviewed-by: Ben Clayton Commit-Queue: Zhaoming Jiang --- src/dawn/native/d3d12/ShaderModuleD3D12.cpp | 4 +- src/tint/cmd/main.cc | 55 ++++ .../transform/num_workgroups_from_uniform.cc | 28 +- .../transform/num_workgroups_from_uniform.h | 15 +- .../num_workgroups_from_uniform_test.cc | 256 ++++++++++++++++++ src/tint/writer/hlsl/generator.h | 3 +- 6 files changed, 351 insertions(+), 10 deletions(-) diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp index b9d4ef4299..7a23be5e7f 100644 --- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp +++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp @@ -645,8 +645,8 @@ ResultOrError TranslateToHLSL(dawn::platform::Platform* platform, tint::writer::hlsl::Options options; options.disable_workgroup_init = request.disableWorkgroupInit; if (request.usesNumWorkgroups) { - options.root_constant_binding_point.group = request.numWorkgroupsRegisterSpace; - options.root_constant_binding_point.binding = request.numWorkgroupsShaderRegister; + options.root_constant_binding_point = tint::sem::BindingPoint{ + request.numWorkgroupsRegisterSpace, request.numWorkgroupsShaderRegister}; } // TODO(dawn:549): HLSL generation outputs the indices into the // array_length_from_uniform buffer that were actually used. When the blob cache can diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc index b9c6d042e5..c067c39b14 100644 --- a/src/tint/cmd/main.cc +++ b/src/tint/cmd/main.cc @@ -15,7 +15,9 @@ #include #include #include +#include #include +#include // NOLINT(build/include_order) #include #include #include @@ -86,6 +88,7 @@ struct Options { std::string dxc_path; std::string xcrun_path; std::vector overrides; + std::optional hlsl_root_constant_binding_point; }; const char kUsage[] = R"(Usage: tint [options] @@ -112,6 +115,11 @@ ${transforms} Affects AST dumping, and text-based output languages. --dump-inspector-bindings -- Dump reflection data about bindins to stdout. -h -- This help text + --hlsl-root-constant-binding-point , -- Binding point for root constant. + Specify the binding point for generated uniform buffer + used for num_workgroups in HLSL. If not specified, then + default to binding 0 of the largest used group plus 1, + or group 0 if no resource bound. --validate -- Validates the generated shader --fxc -- Ask to validate HLSL output using FXC instead of DXC. When specified, automatically enables --validate @@ -220,6 +228,27 @@ std::vector split_on_comma(std::string list) { return res; } +std::optional parse_unsigned_number(std::string number) { + for (char c : number) { + if (!std::isdigit(c)) { + // Found a non-digital char, return nullopt + return std::nullopt; + } + } + + errno = 0; + char* p_end; + uint64_t result; + // std::strtoull will not throw exception. + result = std::strtoull(number.c_str(), &p_end, 10); + if ((errno != 0) || (static_cast(p_end - number.c_str()) != number.length())) { + // Unexpected conversion result + return std::nullopt; + } + + return result; +} + std::string TextureDimensionToString(tint::inspector::ResourceBinding::TextureDimension dim) { switch (dim) { case tint::inspector::ResourceBinding::TextureDimension::kNone: @@ -402,6 +431,31 @@ bool ParseArgs(const std::vector& args, Options* opts) { return false; } opts->overrides = split_on_comma(args[i]); + } else if (arg == "--hlsl-root-constant-binding-point") { + ++i; + if (i >= args.size()) { + std::cerr << "Missing value for " << arg << std::endl; + return false; + } + auto binding_points = split_on_comma(args[i]); + if (binding_points.size() != 2) { + std::cerr << "Invalid binding point for " << arg << ": " << args[i] << std::endl; + return false; + } + auto group = parse_unsigned_number(binding_points[0]); + if ((!group.has_value()) || (group.value() > std::numeric_limits::max())) { + std::cerr << "Invalid group for " << arg << ": " << binding_points[0] << std::endl; + return false; + } + auto binding = parse_unsigned_number(binding_points[1]); + if ((!binding.has_value()) || + (binding.value() > std::numeric_limits::max())) { + std::cerr << "Invalid binding for " << arg << ": " << binding_points[1] + << std::endl; + return false; + } + opts->hlsl_root_constant_binding_point = tint::sem::BindingPoint{ + static_cast(group.value()), static_cast(binding.value())}; } else if (!arg.empty()) { if (arg[0] == '-') { std::cerr << "Unrecognized option: " << arg << std::endl; @@ -723,6 +777,7 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) { tint::writer::hlsl::Options gen_options; gen_options.disable_workgroup_init = options.disable_workgroup_init; gen_options.generate_external_texture_bindings = true; + gen_options.root_constant_binding_point = options.hlsl_root_constant_binding_point; auto result = tint::writer::hlsl::Generate(program, gen_options); if (!result.success) { PrintWGSL(std::cerr, *program); diff --git a/src/tint/transform/num_workgroups_from_uniform.cc b/src/tint/transform/num_workgroups_from_uniform.cc index 17814df9f6..0bb1518544 100644 --- a/src/tint/transform/num_workgroups_from_uniform.cc +++ b/src/tint/transform/num_workgroups_from_uniform.cc @@ -122,10 +122,31 @@ void NumWorkgroupsFromUniform::Run(CloneContext& ctx, const DataMap& inputs, Dat auto* num_workgroups_struct = ctx.dst->Structure( ctx.dst->Sym(), {ctx.dst->Member(kNumWorkgroupsMemberName, ctx.dst->ty.vec3(ctx.dst->ty.u32()))}); + + uint32_t group, binding; + if (cfg->ubo_binding.has_value()) { + // If cfg->ubo_binding holds a value, use the specified binding point. + group = cfg->ubo_binding->group; + binding = cfg->ubo_binding->binding; + } else { + // If cfg->ubo_binding holds no value, use the binding 0 of the largest used group + // plus 1, or group 0 if no resource bound. + group = 0; + + for (auto* var : ctx.src->AST().GlobalVariables()) { + if (auto binding_point = var->BindingPoint()) { + if (binding_point.group->value >= group) { + group = binding_point.group->value + 1; + } + } + } + + binding = 0; + } + num_workgroups_ubo = ctx.dst->Global( ctx.dst->Sym(), ctx.dst->ty.Of(num_workgroups_struct), ast::StorageClass::kUniform, - ast::AttributeList{ - ctx.dst->GroupAndBinding(cfg->ubo_binding.group, cfg->ubo_binding.binding)}); + ast::AttributeList{ctx.dst->GroupAndBinding(group, binding)}); } return num_workgroups_ubo; }; @@ -151,7 +172,8 @@ void NumWorkgroupsFromUniform::Run(CloneContext& ctx, const DataMap& inputs, Dat ctx.Clone(); } -NumWorkgroupsFromUniform::Config::Config(sem::BindingPoint ubo_bp) : ubo_binding(ubo_bp) {} +NumWorkgroupsFromUniform::Config::Config(std::optional ubo_bp) + : ubo_binding(ubo_bp) {} NumWorkgroupsFromUniform::Config::Config(const Config&) = default; NumWorkgroupsFromUniform::Config::~Config() = default; diff --git a/src/tint/transform/num_workgroups_from_uniform.h b/src/tint/transform/num_workgroups_from_uniform.h index 93c4f15b61..9f0b6c1a7e 100644 --- a/src/tint/transform/num_workgroups_from_uniform.h +++ b/src/tint/transform/num_workgroups_from_uniform.h @@ -15,6 +15,8 @@ #ifndef SRC_TINT_TRANSFORM_NUM_WORKGROUPS_FROM_UNIFORM_H_ #define SRC_TINT_TRANSFORM_NUM_WORKGROUPS_FROM_UNIFORM_H_ +#include // NOLINT(build/include_order) + #include "src/tint/sem/binding_point.h" #include "src/tint/transform/transform.h" @@ -52,8 +54,11 @@ class NumWorkgroupsFromUniform : public Castable { /// Constructor - /// @param ubo_bp the binding point to use for the generated uniform buffer. - explicit Config(sem::BindingPoint ubo_bp); + /// @param ubo_bp the binding point to use for the generated uniform buffer. If ubo_bp + /// contains no value, a free binding point will be used to ensure the generated program is + /// valid. Specifically, binding 0 of the largest used group plus 1 is used if at least one + /// resource is bound, otherwise group 0 binding 0 is used. + explicit Config(std::optional ubo_bp); /// Copy constructor Config(const Config&); @@ -61,8 +66,10 @@ class NumWorkgroupsFromUniform : public Castable ubo_binding; }; /// @param program the program to inspect diff --git a/src/tint/transform/num_workgroups_from_uniform_test.cc b/src/tint/transform/num_workgroups_from_uniform_test.cc index de6c6652f1..ffc0ca8812 100644 --- a/src/tint/transform/num_workgroups_from_uniform_test.cc +++ b/src/tint/transform/num_workgroups_from_uniform_test.cc @@ -434,5 +434,261 @@ fn main(tint_symbol : tint_symbol_1) { EXPECT_EQ(expect, str(got)); } +// Test that group 0 binding 0 is used if no bound resource in the program and binding point is not +// specified in NumWorkgroupsFromUniform::Config. +TEST_F(NumWorkgroupsFromUniformTest, UnspecifiedBindingPoint_NoResourceBound) { + auto* src = R"( +struct Builtins1 { + @builtin(num_workgroups) num_wgs : vec3, +}; + +struct Builtins2 { + @builtin(global_invocation_id) gid : vec3, + @builtin(num_workgroups) num_wgs : vec3, + @builtin(workgroup_id) wgid : vec3, +}; + +@stage(compute) @workgroup_size(1) +fn main1(in : Builtins1) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main2(in : Builtins2) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main3(@builtin(num_workgroups) num_wgs : vec3) { + let groups_x = num_wgs.x; + let groups_y = num_wgs.y; + let groups_z = num_wgs.z; +} +)"; + + auto* expect = R"( +struct tint_symbol_6 { + num_workgroups : vec3, +} + +@group(0) @binding(0) var tint_symbol_7 : tint_symbol_6; + +struct Builtins1 { + num_wgs : vec3, +} + +struct Builtins2 { + gid : vec3, + num_wgs : vec3, + wgid : vec3, +} + +fn main1_inner(in : Builtins1) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main1() { + main1_inner(Builtins1(tint_symbol_7.num_workgroups)); +} + +struct tint_symbol_3 { + @builtin(global_invocation_id) + gid : vec3, + @builtin(workgroup_id) + wgid : vec3, +} + +fn main2_inner(in : Builtins2) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main2(tint_symbol_2 : tint_symbol_3) { + main2_inner(Builtins2(tint_symbol_2.gid, tint_symbol_7.num_workgroups, tint_symbol_2.wgid)); +} + +fn main3_inner(num_wgs : vec3) { + let groups_x = num_wgs.x; + let groups_y = num_wgs.y; + let groups_z = num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main3() { + main3_inner(tint_symbol_7.num_workgroups); +} +)"; + + DataMap data; + data.Add(CanonicalizeEntryPointIO::ShaderStyle::kHlsl); + // Make binding point unspecified. + data.Add(std::nullopt); + auto got = Run(src, data); + EXPECT_EQ(expect, str(got)); +} + +// Test that binding 0 of the largest used group plus 1 is used if at least one resource is bound in +// the program and binding point is not specified in NumWorkgroupsFromUniform::Config. +TEST_F(NumWorkgroupsFromUniformTest, UnspecifiedBindingPoint_MultipleResourceBound) { + auto* src = R"( +struct Builtins1 { + @builtin(num_workgroups) num_wgs : vec3, +}; + +struct Builtins2 { + @builtin(global_invocation_id) gid : vec3, + @builtin(num_workgroups) num_wgs : vec3, + @builtin(workgroup_id) wgid : vec3, +}; + +struct S0 { + @size(4) + m0 : u32, + m1 : array, +}; + +struct S1 { + @size(4) + m0 : u32, + m1 : array, +}; + +@group(0) @binding(0) var g2 : texture_2d; +@group(1) @binding(0) var g3 : texture_depth_2d; +@group(1) @binding(1) var g4 : texture_storage_2d; +@group(3) @binding(0) var g5 : texture_depth_cube_array; +@group(4) @binding(0) var g6 : texture_external; + +@group(0) @binding(1) var g8 : S0; +@group(1) @binding(3) var g9 : S0; +@group(3) @binding(2) var g10 : S0; + +@stage(compute) @workgroup_size(1) +fn main1(in : Builtins1) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; + g8.m0 = 1u; +} + +@stage(compute) @workgroup_size(1) +fn main2(in : Builtins2) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main3(@builtin(num_workgroups) num_wgs : vec3) { + let groups_x = num_wgs.x; + let groups_y = num_wgs.y; + let groups_z = num_wgs.z; +} +)"; + + auto* expect = R"( +struct tint_symbol_6 { + num_workgroups : vec3, +} + +@group(5) @binding(0) var tint_symbol_7 : tint_symbol_6; + +struct Builtins1 { + num_wgs : vec3, +} + +struct Builtins2 { + gid : vec3, + num_wgs : vec3, + wgid : vec3, +} + +struct S0 { + @size(4) + m0 : u32, + m1 : array, +} + +struct S1 { + @size(4) + m0 : u32, + m1 : array, +} + +@group(0) @binding(0) var g2 : texture_2d; + +@group(1) @binding(0) var g3 : texture_depth_2d; + +@group(1) @binding(1) var g4 : texture_storage_2d; + +@group(3) @binding(0) var g5 : texture_depth_cube_array; + +@group(4) @binding(0) var g6 : texture_external; + +@group(0) @binding(1) var g8 : S0; + +@group(1) @binding(3) var g9 : S0; + +@group(3) @binding(2) var g10 : S0; + +fn main1_inner(in : Builtins1) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; + g8.m0 = 1u; +} + +@stage(compute) @workgroup_size(1) +fn main1() { + main1_inner(Builtins1(tint_symbol_7.num_workgroups)); +} + +struct tint_symbol_3 { + @builtin(global_invocation_id) + gid : vec3, + @builtin(workgroup_id) + wgid : vec3, +} + +fn main2_inner(in : Builtins2) { + let groups_x = in.num_wgs.x; + let groups_y = in.num_wgs.y; + let groups_z = in.num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main2(tint_symbol_2 : tint_symbol_3) { + main2_inner(Builtins2(tint_symbol_2.gid, tint_symbol_7.num_workgroups, tint_symbol_2.wgid)); +} + +fn main3_inner(num_wgs : vec3) { + let groups_x = num_wgs.x; + let groups_y = num_wgs.y; + let groups_z = num_wgs.z; +} + +@stage(compute) @workgroup_size(1) +fn main3() { + main3_inner(tint_symbol_7.num_workgroups); +} +)"; + + DataMap data; + data.Add(CanonicalizeEntryPointIO::ShaderStyle::kHlsl); + // Make binding point unspecified. + data.Add(std::nullopt); + auto got = Run(src, data); + EXPECT_EQ(expect, str(got)); +} + } // namespace } // namespace tint::transform diff --git a/src/tint/writer/hlsl/generator.h b/src/tint/writer/hlsl/generator.h index f14da6d113..a18687a709 100644 --- a/src/tint/writer/hlsl/generator.h +++ b/src/tint/writer/hlsl/generator.h @@ -16,6 +16,7 @@ #define SRC_TINT_WRITER_HLSL_GENERATOR_H_ #include +#include // NOLINT(build/include_order) #include #include #include @@ -46,7 +47,7 @@ struct Options { Options& operator=(const Options&); /// The binding point to use for information passed via root constants. - sem::BindingPoint root_constant_binding_point; + std::optional root_constant_binding_point; /// Set to `true` to disable workgroup memory zero initialization bool disable_workgroup_init = false; /// Set to 'true' to generates binding mappings for external textures