GLSL: implement push constants.
Bug: tint:2138
Change-Id: I238089989ab5e039f3da7331610dd8e9aede3183
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/169220
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: Stephen White <senorblanco@google.com>
Commit-Queue: Stephen White <senorblanco@chromium.org>
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index 2b0ec37..9a7c6ea 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -152,6 +152,8 @@
data.Add<ast::transform::SingleEntryPoint::Config>(entry_point);
}
+ data.Add<ast::transform::AddBlockAttribute::Config>(true);
+
manager.Add<ast::transform::PreservePadding>(); // Must come before DirectVariableAccess
manager.Add<ast::transform::Unshadow>(); // Must come before DirectVariableAccess
@@ -1897,9 +1899,7 @@
EmitIOVariable(sem);
return;
case core::AddressSpace::kPushConstant:
- diagnostics_.add_error(
- diag::System::Writer,
- "unhandled address space " + tint::ToString(sem->AddressSpace()));
+ EmitPushConstant(sem);
return;
default: {
TINT_ICE() << "unhandled address space " << sem->AddressSpace();
@@ -2092,6 +2092,17 @@
out << ";";
}
+void ASTPrinter::EmitPushConstant(const sem::GlobalVariable* var) {
+ auto* decl = var->Declaration();
+
+ auto out = Line();
+
+ auto name = decl->name->symbol.Name();
+ auto* type = var->Type()->UnwrapRef();
+ EmitTypeAndName(out, type, var->AddressSpace(), var->Access(), name);
+ out << ";";
+}
+
void ASTPrinter::EmitInterpolationQualifiers(StringStream& out,
VectorRef<const ast::Attribute*> attributes) {
for (auto* attr : attributes) {
@@ -2642,6 +2653,7 @@
break;
}
case core::AddressSpace::kUniform:
+ case core::AddressSpace::kPushConstant:
case core::AddressSpace::kHandle: {
out << "uniform ";
break;
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.h b/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
index f5cc633..cfaf716 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
@@ -310,6 +310,10 @@
/// @param var the global variable
void EmitIOVariable(const sem::GlobalVariable* var);
+ /// Handles emitting a global variable with the push_constant address space
+ /// @param var the global variable
+ void EmitPushConstant(const sem::GlobalVariable* var);
+
/// Handles emitting interpolation qualifiers
/// @param out the output of the expression stream
/// @param attrs the attributes
diff --git a/src/tint/lang/wgsl/ast/transform/add_block_attribute.cc b/src/tint/lang/wgsl/ast/transform/add_block_attribute.cc
index 82e0716..607bf54 100644
--- a/src/tint/lang/wgsl/ast/transform/add_block_attribute.cc
+++ b/src/tint/lang/wgsl/ast/transform/add_block_attribute.cc
@@ -39,6 +39,7 @@
TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::AddBlockAttribute);
TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::AddBlockAttribute::BlockAttribute);
+TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::AddBlockAttribute::Config);
namespace tint::ast::transform {
@@ -47,12 +48,13 @@
AddBlockAttribute::~AddBlockAttribute() = default;
Transform::ApplyResult AddBlockAttribute::Apply(const Program& src,
- const DataMap&,
+ const DataMap& inputs,
DataMap&) const {
ProgramBuilder b;
program::CloneContext ctx{&b, &src, /* auto_clone_symbols */ true};
auto& sem = src.Sem();
+ auto* cfg = inputs.Get<Config>();
// A map from a type in the source program to a block-decorated wrapper that contains it in the
// destination program.
@@ -80,6 +82,11 @@
bool needs_wrapping = !str || // Type is not a structure
str->HasFixedFootprint(); // Struct has a fixed footprint
+ if (cfg && cfg->skip_push_constants &&
+ var->AddressSpace() == core::AddressSpace::kPushConstant) {
+ continue;
+ }
+
if (needs_wrapping) {
const char* kMemberName = "inner";
@@ -129,4 +136,8 @@
ctx.dst->AllocateNodeID());
}
+AddBlockAttribute::Config::Config(bool skip_push_consts) : skip_push_constants(skip_push_consts) {}
+
+AddBlockAttribute::Config::~Config() = default;
+
} // namespace tint::ast::transform
diff --git a/src/tint/lang/wgsl/ast/transform/add_block_attribute.h b/src/tint/lang/wgsl/ast/transform/add_block_attribute.h
index 544e1ee..61a884d 100644
--- a/src/tint/lang/wgsl/ast/transform/add_block_attribute.h
+++ b/src/tint/lang/wgsl/ast/transform/add_block_attribute.h
@@ -66,6 +66,19 @@
/// Destructor
~AddBlockAttribute() override;
+ /// Transform configuration options
+ struct Config final : public Castable<Config, ast::transform::Data> {
+ /// Constructor
+ /// @param skip_push_const whether to skip push constants
+ explicit Config(bool skip_push_const);
+
+ /// Destructor
+ ~Config() override;
+
+ /// Whether to skip push constants
+ bool skip_push_constants;
+ };
+
/// @copydoc Transform::Apply
ApplyResult Apply(const Program& program,
const DataMap& inputs,
diff --git a/test/tint/var/uses/push_constant.wgsl.expected.glsl b/test/tint/var/uses/push_constant.wgsl.expected.glsl
index 00e2f76..6ed822e 100644
--- a/test/tint/var/uses/push_constant.wgsl.expected.glsl
+++ b/test/tint/var/uses/push_constant.wgsl.expected.glsl
@@ -1,128 +1,55 @@
-SKIP: FAILED
+#version 310 es
-
-enable chromium_experimental_push_constant;
-
-var<push_constant> a : i32;
-
-var<push_constant> b : i32;
-
-var<push_constant> c : i32;
-
-fn uses_a() {
- let foo = a;
+uniform int a;
+void uses_a() {
+ int foo = a;
}
-fn uses_uses_a() {
+void main1() {
uses_a();
}
-fn uses_b() {
- let foo = b;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ main1();
+ return;
+}
+#version 310 es
+
+uniform int a;
+void uses_a() {
+ int foo = a;
}
-@compute @workgroup_size(1)
-fn main1() {
+void uses_uses_a() {
uses_a();
}
-@compute @workgroup_size(1)
-fn main2() {
+void main2() {
uses_uses_a();
}
-@compute @workgroup_size(1)
-fn main3() {
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ main2();
+ return;
+}
+#version 310 es
+
+uniform int b;
+void uses_b() {
+ int foo = b;
+}
+
+void main3() {
uses_b();
}
-@compute @workgroup_size(1)
-fn main4() {
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ main3();
+ return;
}
-
-Failed to generate: error: unhandled address space push_constant
-
-enable chromium_experimental_push_constant;
-
-var<push_constant> a : i32;
-
-var<push_constant> b : i32;
-
-var<push_constant> c : i32;
-
-fn uses_a() {
- let foo = a;
-}
-
-fn uses_uses_a() {
- uses_a();
-}
-
-fn uses_b() {
- let foo = b;
-}
-
-@compute @workgroup_size(1)
-fn main1() {
- uses_a();
-}
-
-@compute @workgroup_size(1)
-fn main2() {
- uses_uses_a();
-}
-
-@compute @workgroup_size(1)
-fn main3() {
- uses_b();
-}
-
-@compute @workgroup_size(1)
-fn main4() {
-}
-
-Failed to generate: error: unhandled address space push_constant
-
-enable chromium_experimental_push_constant;
-
-var<push_constant> a : i32;
-
-var<push_constant> b : i32;
-
-var<push_constant> c : i32;
-
-fn uses_a() {
- let foo = a;
-}
-
-fn uses_uses_a() {
- uses_a();
-}
-
-fn uses_b() {
- let foo = b;
-}
-
-@compute @workgroup_size(1)
-fn main1() {
- uses_a();
-}
-
-@compute @workgroup_size(1)
-fn main2() {
- uses_uses_a();
-}
-
-@compute @workgroup_size(1)
-fn main3() {
- uses_b();
-}
-
-@compute @workgroup_size(1)
-fn main4() {
-}
-
-Failed to generate: error: unhandled address space push_constant
#version 310 es
void main4() {
