diff --git a/src/tint/cmd/tint/main.cc b/src/tint/cmd/tint/main.cc
index 4d18bb9..1601a76 100644
--- a/src/tint/cmd/tint/main.cc
+++ b/src/tint/cmd/tint/main.cc
@@ -1080,11 +1080,18 @@
         gen_options.texture_builtins_from_uniform = std::move(textureBuiltinsFromUniform);
 
         auto entry_point = inspector.GetEntryPoint(entry_point_name);
+        uint32_t offset = entry_point.push_constant_size;
 
         if (entry_point.instance_index_used) {
             // Place the first_instance push constant member after user-defined push constants (if
             // any).
-            gen_options.first_instance_offset = entry_point.push_constant_size;
+            gen_options.first_instance_offset = offset;
+            offset += 4;
+        }
+        if (entry_point.frag_depth_used) {
+            gen_options.min_depth_offset = offset + 0;
+            gen_options.max_depth_offset = offset + 4;
+            offset += 8;
         }
 
         auto result = tint::glsl::writer::Generate(prg, gen_options, entry_point_name);
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 326c1a2..764b779 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -61,6 +61,7 @@
 #include "src/tint/lang/wgsl/ast/transform/binding_remapper.h"
 #include "src/tint/lang/wgsl/ast/transform/builtin_polyfill.h"
 #include "src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io.h"
+#include "src/tint/lang/wgsl/ast/transform/clamp_frag_depth.h"
 #include "src/tint/lang/wgsl/ast/transform/demote_to_helper.h"
 #include "src/tint/lang/wgsl/ast/transform/direct_variable_access.h"
 #include "src/tint/lang/wgsl/ast/transform/disable_uniformity_analysis.h"
@@ -206,6 +207,10 @@
 
     manager.Add<ast::transform::OffsetFirstIndex>();
 
+    // ClampFragDepth must come before CanonicalizeEntryPointIO, or the assignments to FragDepth are
+    // lost
+    manager.Add<ast::transform::ClampFragDepth>();
+
     // CanonicalizeEntryPointIO must come after Robustness
     manager.Add<ast::transform::CanonicalizeEntryPointIO>();
 
@@ -249,6 +254,8 @@
 
     data.Add<ast::transform::OffsetFirstIndex::Config>(std::nullopt, options.first_instance_offset);
 
+    data.Add<ast::transform::ClampFragDepth::Config>(options.min_depth_offset,
+                                                     options.max_depth_offset);
     SanitizedResult result;
     ast::transform::DataMap outputs;
     result.program = manager.Run(in, data, outputs);
diff --git a/src/tint/lang/glsl/writer/common/options.h b/src/tint/lang/glsl/writer/common/options.h
index 7545366..62cd80c 100644
--- a/src/tint/lang/glsl/writer/common/options.h
+++ b/src/tint/lang/glsl/writer/common/options.h
@@ -82,6 +82,12 @@
     /// Offset of the firstInstance push constant.
     std::optional<int32_t> first_instance_offset;
 
+    /// Offset of the minDepth push constant.
+    std::optional<uint32_t> min_depth_offset;
+
+    /// Offset of the maxDepth push constant.
+    std::optional<uint32_t> max_depth_offset;
+
     /// Options used to map WGSL textureNumLevels/textureNumSamples builtins to internal uniform
     /// buffer values. If not specified, emits corresponding GLSL builtins
     /// textureQueryLevels/textureSamples directly.
@@ -98,6 +104,8 @@
                  binding_remapper_options,
                  external_texture_options,
                  first_instance_offset,
+                 min_depth_offset,
+                 max_depth_offset,
                  texture_builtins_from_uniform);
 };
 
diff --git a/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc b/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
index 4272172..c782db5 100644
--- a/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
@@ -66,6 +66,7 @@
 
     if (options.clamp_frag_depth) {
         manager.Add<ast::transform::ClampFragDepth>();
+        data.Add<ast::transform::ClampFragDepth::Config>(0, 4);
     }
 
     manager.Add<ast::transform::DisableUniformityAnalysis>();
diff --git a/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.cc b/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.cc
index 59e6dd6..2b5e6a2 100644
--- a/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.cc
+++ b/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.cc
@@ -35,6 +35,7 @@
 #include "src/tint/lang/wgsl/ast/function.h"
 #include "src/tint/lang/wgsl/ast/module.h"
 #include "src/tint/lang/wgsl/ast/struct.h"
+#include "src/tint/lang/wgsl/ast/transform/push_constant_helper.h"
 #include "src/tint/lang/wgsl/program/clone_context.h"
 #include "src/tint/lang/wgsl/program/program_builder.h"
 #include "src/tint/lang/wgsl/resolver/resolve.h"
@@ -45,6 +46,7 @@
 #include "src/tint/utils/macros/scoped_assignment.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::ClampFragDepth);
+TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::ClampFragDepth::Config);
 
 namespace tint::ast::transform {
 
@@ -63,49 +65,37 @@
 
     /// Runs the transform
     /// @returns the new program or SkipTransform if the transform is not required
-    Transform::ApplyResult Run() {
-        // Abort on any use of push constants in the module.
-        for (auto* global : src.AST().GlobalVariables()) {
-            if (auto* var = global->As<ast::Var>()) {
-                auto* v = src.Sem().Get(var);
-                if (TINT_UNLIKELY(v->AddressSpace() == core::AddressSpace::kPushConstant)) {
-                    TINT_ICE()
-                        << "ClampFragDepth doesn't know how to handle module that already use push "
-                           "constants";
-                    return resolver::Resolve(b);
-                }
-            }
-        }
-
-        if (!ShouldRun()) {
+    Transform::ApplyResult Run(const DataMap& inputs) {
+        const Config* cfg = inputs.Get<Config>();
+        if (!cfg || !cfg->min_depth_offset.has_value() || !cfg->max_depth_offset.has_value()) {
             return SkipTransform;
         }
 
+        PushConstantHelper push_constant_helper(ctx);
+
         // At least one entry-point needs clamping. Add the following to the module:
         //
         //   enable chromium_experimental_push_constant;
         //
-        //   struct FragDepthClampArgs {
-        //       min : f32,
-        //       max : f32,
+        //   struct PushConstants {
+        //       min_depth : f32,
+        //       max_depth : f32,
         //   }
-        //   var<push_constant> frag_depth_clamp_args : FragDepthClampArgs;
+        //   var<push_constant> push_constants : PushConstants;
         //
         //   fn clamp_frag_depth(v : f32) -> f32 {
-        //       return clamp(v, frag_depth_clamp_args.min, frag_depth_clamp_args.max);
+        //       return clamp(v, push_constants.min, push_constants.max_depth);
         //   }
-        b.Enable(wgsl::Extension::kChromiumExperimentalPushConstant);
 
-        b.Structure(b.Symbols().New("FragDepthClampArgs"),
-                    Vector{b.Member("min", b.ty.f32()), b.Member("max", b.ty.f32())});
+        push_constant_helper.InsertMember("min_depth", b.ty.f32(), *cfg->min_depth_offset);
+        push_constant_helper.InsertMember("max_depth", b.ty.f32(), *cfg->max_depth_offset);
 
-        auto args_sym = b.Symbols().New("frag_depth_clamp_args");
-        b.GlobalVar(args_sym, b.ty("FragDepthClampArgs"), core::AddressSpace::kPushConstant);
+        Symbol buffer_name = push_constant_helper.Run();
 
         auto base_fn_sym = b.Symbols().New("clamp_frag_depth");
         b.Func(base_fn_sym, Vector{b.Param("v", b.ty.f32())}, b.ty.f32(),
-               Vector{b.Return(b.Call("clamp", "v", b.MemberAccessor(args_sym, "min"),
-                                      b.MemberAccessor(args_sym, "max")))});
+               Vector{b.Return(b.Call("clamp", "v", b.MemberAccessor(buffer_name, "min_depth"),
+                                      b.MemberAccessor(buffer_name, "max_depth")))});
 
         // If true, the currently cloned function returns frag depth directly as a scalar
         bool returns_frag_depth_as_value = false;
@@ -185,17 +175,6 @@
     }
 
   private:
-    /// @returns true if the transform should run
-    bool ShouldRun() {
-        for (auto* fn : src.AST().Functions()) {
-            if (fn->PipelineStage() == ast::PipelineStage::kFragment &&
-                (ReturnsFragDepthAsValue(fn) || ReturnsFragDepthInStruct(fn))) {
-                return true;
-            }
-        }
-
-        return false;
-    }
     /// @param attrs the attributes to examine
     /// @returns true if @p attrs contains a `@builtin(frag_depth)` attribute
     bool ContainsFragDepth(VectorRef<const ast::Attribute*> attrs) {
@@ -237,9 +216,15 @@
 ClampFragDepth::~ClampFragDepth() = default;
 
 ast::transform::Transform::ApplyResult ClampFragDepth::Apply(const Program& src,
-                                                             const ast::transform::DataMap&,
+                                                             const ast::transform::DataMap& inputs,
                                                              ast::transform::DataMap&) const {
-    return State{src}.Run();
+    return State{src}.Run(inputs);
 }
 
+ClampFragDepth::Config::Config(std::optional<uint32_t> min_depth_off,
+                               std::optional<uint32_t> max_depth_off)
+    : min_depth_offset(min_depth_off), max_depth_offset(max_depth_off) {}
+
+ClampFragDepth::Config::~Config() = default;
+
 }  // namespace tint::ast::transform
diff --git a/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.h b/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.h
index 5bcd3ea..599b91c 100644
--- a/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.h
+++ b/src/tint/lang/wgsl/ast/transform/clamp_frag_depth.h
@@ -69,6 +69,23 @@
     /// Destructor
     ~ClampFragDepth() override;
 
+    /// Transform configuration options
+    struct Config final : public Castable<Config, ast::transform::Data> {
+        /// Constructor
+        /// @param min_depth_off Offset of the minDepth push constant
+        /// @param max_depth_off Offset of the maxDepth push constant
+        Config(std::optional<uint32_t> min_depth_off, std::optional<uint32_t> max_depth_off);
+
+        /// Destructor
+        ~Config() override;
+
+        /// Offset of the min_depth push constant
+        std::optional<uint32_t> min_depth_offset;
+
+        /// Offset of the min_depth push constant
+        std::optional<uint32_t> max_depth_offset;
+    };
+
     /// @copydoc ast::transform::Transform::Apply
     ApplyResult Apply(const Program& program,
                       const ast::transform::DataMap& inputs,
diff --git a/src/tint/lang/wgsl/ast/transform/clamp_frag_depth_test.cc b/src/tint/lang/wgsl/ast/transform/clamp_frag_depth_test.cc
index cc09c6f..97f1f31 100644
--- a/src/tint/lang/wgsl/ast/transform/clamp_frag_depth_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/clamp_frag_depth_test.cc
@@ -40,48 +40,102 @@
     EXPECT_FALSE(ShouldRun<ClampFragDepth>(src));
 }
 
-TEST_F(ClampFragDepthTest, ShouldRunNoFragmentShader) {
+TEST_F(ClampFragDepthTest, ShouldRunNoConfig) {
     auto* src = R"(
-        fn f() -> f32 {
+        @fragment fn main() -> @builtin(frag_depth) f32 {
             return 0.0;
         }
-
-        @compute @workgroup_size(1) fn cs() {
-        }
-
-        @vertex fn vs() -> @builtin(position) vec4<f32> {
-            return vec4<f32>();
-        }
     )";
 
     EXPECT_FALSE(ShouldRun<ClampFragDepth>(src));
 }
 
-TEST_F(ClampFragDepthTest, ShouldRunFragmentShaderNoReturnType) {
+TEST_F(ClampFragDepthTest, ShouldRunNoMin) {
     auto* src = R"(
-        @fragment fn main() {
-        }
-    )";
-
-    EXPECT_FALSE(ShouldRun<ClampFragDepth>(src));
-}
-
-TEST_F(ClampFragDepthTest, ShouldRunFragmentShaderNoFragDepth) {
-    auto* src = R"(
-        @fragment fn main() -> @location(0) f32 {
+        @fragment fn main() -> @builtin(frag_depth) f32 {
             return 0.0;
         }
+    )";
 
-        struct S {
-            @location(0) a : f32,
-            @builtin(sample_mask) b : u32,
-        }
-        @fragment fn main2() -> S {
-            return S();
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(std::nullopt, 4);
+
+    EXPECT_FALSE(ShouldRun<ClampFragDepth>(src, config));
+}
+
+TEST_F(ClampFragDepthTest, ShouldRunNoMinNoMax) {
+    auto* src = R"(
+        @fragment fn main() -> @builtin(frag_depth) f32 {
+            return 0.0;
         }
     )";
 
-    EXPECT_FALSE(ShouldRun<ClampFragDepth>(src));
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, std::nullopt);
+
+    EXPECT_FALSE(ShouldRun<ClampFragDepth>(src, config));
+}
+
+TEST_F(ClampFragDepthTest, ShouldRun) {
+    auto* src = R"(
+        @fragment fn main() -> @builtin(frag_depth) f32 {
+            return 0.0;
+        }
+    )";
+
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+
+    EXPECT_TRUE(ShouldRun<ClampFragDepth>(src, config));
+}
+
+TEST_F(ClampFragDepthTest, ExistingPushConstant) {
+    auto* src = R"(
+        enable chromium_experimental_push_constant;
+
+        struct PushConstants {
+          a : f32,
+        }
+
+        var<push_constant> push_constants : PushConstants;
+        @fragment fn main() -> @builtin(frag_depth) f32 {
+            return push_constants.a;
+        }
+
+    )";
+
+    auto* expect = R"(
+enable chromium_experimental_push_constant;
+
+struct PushConstants_1 {
+  a : f32,
+  /* @offset(4) */
+  min_depth : f32,
+  /* @offset(8) */
+  max_depth : f32,
+}
+
+fn clamp_frag_depth(v : f32) -> f32 {
+  return clamp(v, push_constants.min_depth, push_constants.max_depth);
+}
+
+struct PushConstants {
+  a : f32,
+}
+
+var<push_constant> push_constants : PushConstants_1;
+
+@fragment
+fn main() -> @builtin(frag_depth) f32 {
+  return clamp_frag_depth(push_constants.a);
+}
+)";
+
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(4, 8);
+
+    auto got = Run<ClampFragDepth>(src, config);
+    EXPECT_EQ(expect, str(got));
 }
 
 TEST_F(ClampFragDepthTest, ShouldRunFragDepthAsDirectReturn) {
@@ -91,7 +145,10 @@
         }
     )";
 
-    EXPECT_TRUE(ShouldRun<ClampFragDepth>(src));
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+
+    EXPECT_TRUE(ShouldRun<ClampFragDepth>(src, config));
 }
 
 TEST_F(ClampFragDepthTest, ShouldRunFragDepthInStruct) {
@@ -106,7 +163,10 @@
         }
     )";
 
-    EXPECT_TRUE(ShouldRun<ClampFragDepth>(src));
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+
+    EXPECT_TRUE(ShouldRun<ClampFragDepth>(src, config));
 }
 
 TEST_F(ClampFragDepthTest, SingleReturnOfFragDepth) {
@@ -119,15 +179,17 @@
     auto* expect = R"(
 enable chromium_experimental_push_constant;
 
-struct FragDepthClampArgs {
-  min : f32,
-  max : f32,
+struct PushConstants {
+  /* @offset(0) */
+  min_depth : f32,
+  /* @offset(4) */
+  max_depth : f32,
 }
 
-var<push_constant> frag_depth_clamp_args : FragDepthClampArgs;
+var<push_constant> push_constants : PushConstants;
 
 fn clamp_frag_depth(v : f32) -> f32 {
-  return clamp(v, frag_depth_clamp_args.min, frag_depth_clamp_args.max);
+  return clamp(v, push_constants.min_depth, push_constants.max_depth);
 }
 
 @fragment
@@ -136,7 +198,9 @@
 }
 )";
 
-    auto got = Run<ClampFragDepth>(src);
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+    auto got = Run<ClampFragDepth>(src, config);
     EXPECT_EQ(expect, str(got));
 }
 
@@ -153,15 +217,17 @@
     auto* expect = R"(
 enable chromium_experimental_push_constant;
 
-struct FragDepthClampArgs {
-  min : f32,
-  max : f32,
+struct PushConstants {
+  /* @offset(0) */
+  min_depth : f32,
+  /* @offset(4) */
+  max_depth : f32,
 }
 
-var<push_constant> frag_depth_clamp_args : FragDepthClampArgs;
+var<push_constant> push_constants : PushConstants;
 
 fn clamp_frag_depth(v : f32) -> f32 {
-  return clamp(v, frag_depth_clamp_args.min, frag_depth_clamp_args.max);
+  return clamp(v, push_constants.min_depth, push_constants.max_depth);
 }
 
 @fragment
@@ -173,7 +239,9 @@
 }
 )";
 
-    auto got = Run<ClampFragDepth>(src);
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+    auto got = Run<ClampFragDepth>(src, config);
     EXPECT_EQ(expect, str(got));
 }
 
@@ -190,15 +258,17 @@
     auto* expect = R"(
 enable chromium_experimental_push_constant;
 
-struct FragDepthClampArgs {
-  min : f32,
-  max : f32,
+struct PushConstants {
+  /* @offset(0) */
+  min_depth : f32,
+  /* @offset(4) */
+  max_depth : f32,
 }
 
-var<push_constant> frag_depth_clamp_args : FragDepthClampArgs;
+var<push_constant> push_constants : PushConstants;
 
 fn clamp_frag_depth(v : f32) -> f32 {
-  return clamp(v, frag_depth_clamp_args.min, frag_depth_clamp_args.max);
+  return clamp(v, push_constants.min_depth, push_constants.max_depth);
 }
 
 @fragment
@@ -212,7 +282,9 @@
 }
 )";
 
-    auto got = Run<ClampFragDepth>(src);
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+    auto got = Run<ClampFragDepth>(src, config);
     EXPECT_EQ(expect, str(got));
 }
 
@@ -230,15 +302,17 @@
     auto* expect = R"(
 enable chromium_experimental_push_constant;
 
-struct FragDepthClampArgs {
-  min : f32,
-  max : f32,
+struct PushConstants {
+  /* @offset(0) */
+  min_depth : f32,
+  /* @offset(4) */
+  max_depth : f32,
 }
 
-var<push_constant> frag_depth_clamp_args : FragDepthClampArgs;
+var<push_constant> push_constants : PushConstants;
 
 fn clamp_frag_depth(v : f32) -> f32 {
-  return clamp(v, frag_depth_clamp_args.min, frag_depth_clamp_args.max);
+  return clamp(v, push_constants.min_depth, push_constants.max_depth);
 }
 
 struct S {
@@ -256,7 +330,9 @@
 }
 )";
 
-    auto got = Run<ClampFragDepth>(src);
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+    auto got = Run<ClampFragDepth>(src, config);
     EXPECT_EQ(expect, str(got));
 }
 
@@ -285,15 +361,17 @@
     auto* expect = R"(
 enable chromium_experimental_push_constant;
 
-struct FragDepthClampArgs {
-  min : f32,
-  max : f32,
+struct PushConstants {
+  /* @offset(0) */
+  min_depth : f32,
+  /* @offset(4) */
+  max_depth : f32,
 }
 
-var<push_constant> frag_depth_clamp_args : FragDepthClampArgs;
+var<push_constant> push_constants : PushConstants;
 
 fn clamp_frag_depth(v : f32) -> f32 {
-  return clamp(v, frag_depth_clamp_args.min, frag_depth_clamp_args.max);
+  return clamp(v, push_constants.min_depth, push_constants.max_depth);
 }
 
 struct S {
@@ -330,7 +408,9 @@
 }
 )";
 
-    auto got = Run<ClampFragDepth>(src);
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+    auto got = Run<ClampFragDepth>(src, config);
     EXPECT_EQ(expect, str(got));
 }
 
@@ -352,15 +432,17 @@
     auto* expect = R"(
 enable chromium_experimental_push_constant;
 
-struct FragDepthClampArgs {
-  min : f32,
-  max : f32,
+struct PushConstants {
+  /* @offset(0) */
+  min_depth : f32,
+  /* @offset(4) */
+  max_depth : f32,
 }
 
-var<push_constant> frag_depth_clamp_args : FragDepthClampArgs;
+var<push_constant> push_constants : PushConstants;
 
 fn clamp_frag_depth(v : f32) -> f32 {
-  return clamp(v, frag_depth_clamp_args.min, frag_depth_clamp_args.max);
+  return clamp(v, push_constants.min_depth, push_constants.max_depth);
 }
 
 struct S {
@@ -386,7 +468,9 @@
 }
 )";
 
-    auto got = Run<ClampFragDepth>(src);
+    DataMap config;
+    config.Add<ClampFragDepth::Config>(0, 4);
+    auto got = Run<ClampFragDepth>(src, config);
     EXPECT_EQ(expect, str(got));
 }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index 05505d1..6d4621a 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -450,13 +450,24 @@
     }
 
     void EmitAssignment(const ast::AssignmentStatement* stmt) {
-        // If assigning to a phony, just generate the RHS and we're done. Note that, because
-        // this isn't used, a subsequent transform could remove it due to it being dead code.
-        // This could then change the interface for the program (i.e. a global var no longer
-        // used). If that happens we have to either fix this to store to a phony value, or make
-        // sure we pull the interface before doing the dead code elimination.
+        // If assigning to a phony, and the expression evaluation stage is runtime or override, just
+        // generate the RHS and we're done. Note that, because this isn't used, a subsequent
+        // transform could remove it due to it being dead code. This could then change the interface
+        // for the program (i.e. a global var no longer used). If that happens we have to either fix
+        // this to store to a phony value, or make sure we pull the interface before doing the dead
+        // code elimination.
         if (stmt->lhs->Is<ast::PhonyExpression>()) {
-            (void)EmitValueExpression(stmt->rhs);
+            const auto* sem = program_.Sem().GetVal(stmt->rhs);
+            switch (sem->Stage()) {
+                case core::EvaluationStage::kRuntime:
+                case core::EvaluationStage::kOverride:
+                    EmitValueExpression(stmt->rhs);
+                    break;
+                case core::EvaluationStage::kNotEvaluated:
+                case core::EvaluationStage::kConstant:
+                    // Don't emit.
+                    break;
+            }
             return;
         }
 
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
index b47ee3a..9867cb4 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
@@ -26,6 +26,7 @@
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 #include "gmock/gmock.h"
+#include "src/tint/lang/core/builtin_fn.h"
 #include "src/tint/lang/core/constant/scalar.h"
 #include "src/tint/lang/core/fluent_types.h"
 #include "src/tint/lang/core/ir/block.h"
@@ -1164,5 +1165,27 @@
 )");
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// Bugs
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IR_FromProgramTest, BugChromium324466107) {
+    Func("f", Empty, ty.void_(),
+         Vector{
+             // Abstract type on the RHS - cannot be emitted.
+             Assign(Phony(), Call(core::BuiltinFn::kFrexp, Call(ty.vec2<Infer>(), 2.0_a))),
+         });
+
+    auto m = Build();
+    ASSERT_EQ(m, Success);
+
+    EXPECT_EQ(Disassemble(m.Get()),
+              R"(%f = func():void -> %b1 {
+  %b1 = block {
+    ret
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::wgsl::reader
diff --git a/src/tint/lang/wgsl/sem/function_expression.cc b/src/tint/lang/wgsl/sem/function_expression.cc
index d20ab00..d033140 100644
--- a/src/tint/lang/wgsl/sem/function_expression.cc
+++ b/src/tint/lang/wgsl/sem/function_expression.cc
@@ -26,16 +26,21 @@
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 #include "src/tint/lang/wgsl/sem/function_expression.h"
+#include "src/tint/lang/wgsl/ast/identifier_expression.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::FunctionExpression);
 
 namespace tint::sem {
 
-FunctionExpression::FunctionExpression(const ast::Expression* declaration,
+FunctionExpression::FunctionExpression(const ast::IdentifierExpression* declaration,
                                        const Statement* statement,
                                        const sem::Function* function)
     : Base(declaration, statement), function_(function) {}
 
 FunctionExpression::~FunctionExpression() = default;
 
+const ast::IdentifierExpression* FunctionExpression::Declaration() const {
+    return static_cast<const ast::IdentifierExpression*>(Base::Declaration());
+}
+
 }  // namespace tint::sem
diff --git a/src/tint/lang/wgsl/sem/function_expression.h b/src/tint/lang/wgsl/sem/function_expression.h
index 5b0625f..8e2c4de 100644
--- a/src/tint/lang/wgsl/sem/function_expression.h
+++ b/src/tint/lang/wgsl/sem/function_expression.h
@@ -31,6 +31,9 @@
 #include "src/tint/lang/wgsl/sem/expression.h"
 
 // Forward declarations
+namespace tint::ast {
+class IdentifierExpression;
+}  // namespace tint::ast
 namespace tint::sem {
 class Function;
 }  // namespace tint::sem
@@ -45,13 +48,16 @@
     /// @param declaration the AST node
     /// @param statement the statement that owns this expression
     /// @param function the function that the expression resolved to
-    FunctionExpression(const ast::Expression* declaration,
+    FunctionExpression(const ast::IdentifierExpression* declaration,
                        const Statement* statement,
                        const sem::Function* function);
 
     /// Destructor
     ~FunctionExpression() override;
 
+    /// @returns the AST node
+    const ast::IdentifierExpression* Declaration() const;
+
     /// @return the function that the expression resolved to
     const sem::Function* Function() const { return function_; }
 
diff --git a/src/tint/lang/wgsl/sem/member_accessor_expression.cc b/src/tint/lang/wgsl/sem/member_accessor_expression.cc
index 68bcd7a..8a10ac1 100644
--- a/src/tint/lang/wgsl/sem/member_accessor_expression.cc
+++ b/src/tint/lang/wgsl/sem/member_accessor_expression.cc
@@ -48,6 +48,10 @@
 
 MemberAccessorExpression::~MemberAccessorExpression() = default;
 
+const ast::MemberAccessorExpression* MemberAccessorExpression::Declaration() const {
+    return static_cast<const ast::MemberAccessorExpression*>(declaration_);
+}
+
 StructMemberAccess::StructMemberAccess(const ast::MemberAccessorExpression* declaration,
                                        const core::type::Type* type,
                                        const Statement* statement,
diff --git a/src/tint/lang/wgsl/sem/member_accessor_expression.h b/src/tint/lang/wgsl/sem/member_accessor_expression.h
index 67dcf08..bb6f388 100644
--- a/src/tint/lang/wgsl/sem/member_accessor_expression.h
+++ b/src/tint/lang/wgsl/sem/member_accessor_expression.h
@@ -48,6 +48,9 @@
     /// Destructor
     ~MemberAccessorExpression() override;
 
+    /// @returns the AST node
+    const ast::MemberAccessorExpression* Declaration() const;
+
   protected:
     /// Constructor
     /// @param declaration the AST node
diff --git a/src/tint/lang/wgsl/sem/variable.cc b/src/tint/lang/wgsl/sem/variable.cc
index 9d5e1b2..e056c77 100644
--- a/src/tint/lang/wgsl/sem/variable.cc
+++ b/src/tint/lang/wgsl/sem/variable.cc
@@ -94,4 +94,8 @@
 
 VariableUser::~VariableUser() = default;
 
+const ast::IdentifierExpression* VariableUser::Declaration() const {
+    return static_cast<const ast::IdentifierExpression*>(Base::Declaration());
+}
+
 }  // namespace tint::sem
diff --git a/src/tint/lang/wgsl/sem/variable.h b/src/tint/lang/wgsl/sem/variable.h
index 9809fcb..15ca981 100644
--- a/src/tint/lang/wgsl/sem/variable.h
+++ b/src/tint/lang/wgsl/sem/variable.h
@@ -292,6 +292,9 @@
                  sem::Variable* variable);
     ~VariableUser() override;
 
+    /// @returns the AST node
+    const ast::IdentifierExpression* Declaration() const;
+
     /// @returns the variable that this expression refers to
     const sem::Variable* Variable() const { return variable_; }