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