diff --git a/naga/src/back/msl/mod.rs b/naga/src/back/msl/mod.rs
index 376662a7d0..01ac1ac419 100644
--- a/naga/src/back/msl/mod.rs
+++ b/naga/src/back/msl/mod.rs
@@ -29,6 +29,20 @@ holding the result.
 [msl]: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
 [all-atom]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
 
+## Pointer-typed bounds-checked expressions and OOB locals
+
+MSL (unlike HLSL and GLSL) has native support for pointer-typed function
+arguments. When the [`BoundsCheckPolicy`] is `ReadZeroSkipWrite` and an
+out-of-bounds index expression is used for such an argument, our strategy is to
+pass a pointer to a dummy variable. These dummy variables are called "OOB
+locals". We emit at most one OOB local per function for each type, since all
+expressions producing a result of that type can share the same OOB local. (Note
+that the OOB local mechanism is not actually implementing "skip write", nor even
+"read zero" in some cases of read-after-write, but doing so would require
+additional effort and the difference is unlikely to matter.)
+
+[`BoundsCheckPolicy`]: crate::proc::BoundsCheckPolicy
+
 */
 
 use alloc::{
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 2935a6ad3c..2903ad80e9 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -18,7 +18,11 @@ use crate::{
     arena::{Handle, HandleSet},
     back::{self, Baked},
     common,
-    proc::{self, index, NameKey, TypeResolution},
+    proc::{
+        self,
+        index::{self, BoundsCheck},
+        NameKey, TypeResolution,
+    },
     valid, FastHashMap, FastHashSet,
 };
 
@@ -599,11 +603,34 @@ impl crate::Type {
     }
 }
 
+#[derive(Clone, Copy)]
 enum FunctionOrigin {
     Handle(Handle<crate::Function>),
     EntryPoint(proc::EntryPointIndex),
 }
 
+trait NameKeyExt {
+    fn local(origin: FunctionOrigin, local_handle: Handle<crate::LocalVariable>) -> NameKey {
+        match origin {
+            FunctionOrigin::Handle(handle) => NameKey::FunctionLocal(handle, local_handle),
+            FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointLocal(idx, local_handle),
+        }
+    }
+
+    /// Return the name key for a local variable used by ReadZeroSkipWrite bounds-check
+    /// policy when it needs to produce a pointer-typed result for an OOB access. These
+    /// are unique per accessed type, so the second argument is a type handle. See docs
+    /// for [`crate::back::msl`].
+    fn oob_local_for_type(origin: FunctionOrigin, ty: Handle<crate::Type>) -> NameKey {
+        match origin {
+            FunctionOrigin::Handle(handle) => NameKey::FunctionOobLocal(handle, ty),
+            FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointOobLocal(idx, ty),
+        }
+    }
+}
+
+impl NameKeyExt for NameKey {}
+
 /// A level of detail argument.
 ///
 /// When [`BoundsCheckPolicy::Restrict`] applies to an [`ImageLoad`] access, we
@@ -681,6 +708,7 @@ impl<'a> ExpressionContext<'a> {
             .choose_policy(pointer, &self.module.types, self.info)
     }
 
+    /// See docs for [`proc::index::access_needs_check`].
     fn access_needs_check(
         &self,
         base: Handle<crate::Expression>,
@@ -695,6 +723,19 @@ impl<'a> ExpressionContext<'a> {
         )
     }
 
+    /// See docs for [`proc::index::bounds_check_iter`].
+    fn bounds_check_iter(
+        &self,
+        chain: Handle<crate::Expression>,
+    ) -> impl Iterator<Item = BoundsCheck> + '_ {
+        index::bounds_check_iter(chain, self.module, self.function, self.info)
+    }
+
+    /// See docs for [`proc::index::oob_local_types`].
+    fn oob_local_types(&self) -> FastHashSet<Handle<crate::Type>> {
+        index::oob_local_types(self.module, self.function, self.info, self.policies)
+    }
+
     fn get_packed_vec_kind(&self, expr_handle: Handle<crate::Expression>) -> Option<crate::Scalar> {
         match self.function.expressions[expr_handle] {
             crate::Expression::AccessIndex { base, index } => {
@@ -902,6 +943,59 @@ impl<W: Write> Writer<W> {
         Ok(())
     }
 
+    /// Writes the local variables of the given function, as well as any extra
+    /// out-of-bounds locals that are needed.
+    ///
+    /// The names of the OOB locals are also added to `self.names` at the same
+    /// time.
+    fn put_locals(&mut self, context: &ExpressionContext) -> BackendResult {
+        let oob_local_types = context.oob_local_types();
+        for &ty in oob_local_types.iter() {
+            let name_key = NameKey::oob_local_for_type(context.origin, ty);
+            self.names.insert(name_key, self.namer.call("oob"));
+        }
+
+        for (name_key, ty, init) in context
+            .function
+            .local_variables
+            .iter()
+            .map(|(local_handle, local)| {
+                let name_key = NameKey::local(context.origin, local_handle);
+                (name_key, local.ty, local.init)
+            })
+            .chain(oob_local_types.iter().map(|&ty| {
+                let name_key = NameKey::oob_local_for_type(context.origin, ty);
+                (name_key, ty, None)
+            }))
+        {
+            let ty_name = TypeContext {
+                handle: ty,
+                gctx: context.module.to_ctx(),
+                names: &self.names,
+                access: crate::StorageAccess::empty(),
+                first_time: false,
+            };
+            write!(
+                self.out,
+                "{}{} {}",
+                back::INDENT,
+                ty_name,
+                self.names[&name_key]
+            )?;
+            match init {
+                Some(value) => {
+                    write!(self.out, " = ")?;
+                    self.put_expression(value, context, true)?;
+                }
+                None => {
+                    write!(self.out, " = {{}}")?;
+                }
+            };
+            writeln!(self.out, ";")?;
+        }
+        Ok(())
+    }
+
     fn put_level_of_detail(
         &mut self,
         level: LevelOfDetail,
@@ -1660,7 +1754,6 @@ impl<W: Write> Writer<W> {
         }
 
         let expression = &context.function.expressions[expr_handle];
-        log::trace!("expression {:?} = {:?}", expr_handle, expression);
         match *expression {
             crate::Expression::Literal(_)
             | crate::Expression::Constant(_)
@@ -1696,7 +1789,42 @@ impl<W: Write> Writer<W> {
                 {
                     write!(self.out, " ? ")?;
                     self.put_access_chain(expr_handle, policy, context)?;
-                    write!(self.out, " : DefaultConstructible()")?;
+                    write!(self.out, " : ")?;
+
+                    if context.resolve_type(base).pointer_space().is_some() {
+                        // We can't just use `DefaultConstructible` if this is a pointer.
+                        // Instead, we create a dummy local variable to serve as pointer
+                        // target if the access is out of bounds.
+                        let result_ty = context.info[expr_handle]
+                            .ty
+                            .inner_with(&context.module.types)
+                            .pointer_base_type();
+                        let result_ty_handle = match result_ty {
+                            Some(TypeResolution::Handle(handle)) => handle,
+                            Some(TypeResolution::Value(_)) => {
+                                // As long as the result of a pointer access expression is
+                                // passed to a function or stored in a let binding, the
+                                // type will be in the arena. If additional uses of
+                                // pointers become valid, this assumption might no longer
+                                // hold. Note that the LHS of a load or store doesn't
+                                // take this path -- there is dedicated code in `put_load`
+                                // and `put_store`.
+                                unreachable!(
+                                    "Expected type {result_ty:?} of access through pointer type {base:?} to be in the arena",
+                                );
+                            }
+                            None => {
+                                unreachable!(
+                                    "Expected access through pointer type {base:?} to return a pointer, but got {result_ty:?}",
+                                )
+                            }
+                        };
+                        let name_key =
+                            NameKey::oob_local_for_type(context.origin, result_ty_handle);
+                        self.out.write_str(&self.names[&name_key])?;
+                    } else {
+                        write!(self.out, "DefaultConstructible()")?;
+                    }
 
                     if !is_scoped {
                         write!(self.out, ")")?;
@@ -1736,14 +1864,7 @@ impl<W: Write> Writer<W> {
                 write!(self.out, "{name}")?;
             }
             crate::Expression::LocalVariable(handle) => {
-                let name_key = match context.origin {
-                    FunctionOrigin::Handle(fun_handle) => {
-                        NameKey::FunctionLocal(fun_handle, handle)
-                    }
-                    FunctionOrigin::EntryPoint(ep_index) => {
-                        NameKey::EntryPointLocal(ep_index, handle)
-                    }
-                };
+                let name_key = NameKey::local(context.origin, handle);
                 let name = &self.names[&name_key];
                 write!(self.out, "{name}")?;
             }
@@ -2647,68 +2768,44 @@ impl<W: Write> Writer<W> {
     #[allow(unused_variables)]
     fn put_bounds_checks(
         &mut self,
-        mut chain: Handle<crate::Expression>,
+        chain: Handle<crate::Expression>,
         context: &ExpressionContext,
         level: back::Level,
         prefix: &'static str,
     ) -> Result<bool, Error> {
         let mut check_written = false;
 
-        // Iterate over the access chain, handling each expression.
-        loop {
-            // Produce a `GuardedIndex`, so we can shared code between the
-            // `Access` and `AccessIndex` cases.
-            let (base, guarded_index) = match context.function.expressions[chain] {
-                crate::Expression::Access { base, index } => {
-                    (base, Some(index::GuardedIndex::Expression(index)))
-                }
-                crate::Expression::AccessIndex { base, index } => {
-                    // Don't try to check indices into structs. Validation already took
-                    // care of them, and index::needs_guard doesn't handle that case.
-                    let mut base_inner = context.resolve_type(base);
-                    if let crate::TypeInner::Pointer { base, .. } = *base_inner {
-                        base_inner = &context.module.types[base].inner;
-                    }
-                    match *base_inner {
-                        crate::TypeInner::Struct { .. } => (base, None),
-                        _ => (base, Some(index::GuardedIndex::Known(index))),
-                    }
-                }
-                _ => break,
-            };
+        // Iterate over the access chain, handling each required bounds check.
+        for item in context.bounds_check_iter(chain) {
+            let BoundsCheck {
+                base,
+                index,
+                length,
+            } = item;
 
-            if let Some(index) = guarded_index {
-                if let Some(length) = context.access_needs_check(base, index) {
-                    if check_written {
-                        write!(self.out, " && ")?;
-                    } else {
-                        write!(self.out, "{level}{prefix}")?;
-                        check_written = true;
-                    }
+            if check_written {
+                write!(self.out, " && ")?;
+            } else {
+                write!(self.out, "{level}{prefix}")?;
+                check_written = true;
+            }
 
-                    // Check that the index falls within bounds. Do this with a single
-                    // comparison, by casting the index to `uint` first, so that negative
-                    // indices become large positive values.
-                    write!(self.out, "uint(")?;
-                    self.put_index(index, context, true)?;
-                    self.out.write_str(") < ")?;
-                    match length {
-                        index::IndexableLength::Known(value) => write!(self.out, "{value}")?,
-                        index::IndexableLength::Dynamic => {
-                            let global =
-                                context.function.originating_global(base).ok_or_else(|| {
-                                    Error::GenericValidation(
-                                        "Could not find originating global".into(),
-                                    )
-                                })?;
-                            write!(self.out, "1 + ")?;
-                            self.put_dynamic_array_max_index(global, context)?
-                        }
-                    }
+            // Check that the index falls within bounds. Do this with a single
+            // comparison, by casting the index to `uint` first, so that negative
+            // indices become large positive values.
+            write!(self.out, "uint(")?;
+            self.put_index(index, context, true)?;
+            self.out.write_str(") < ")?;
+            match length {
+                index::IndexableLength::Known(value) => write!(self.out, "{value}")?,
+                index::IndexableLength::Dynamic => {
+                    let global = context.function.originating_global(base).ok_or_else(|| {
+                        Error::GenericValidation("Could not find originating global".into())
+                    })?;
+                    write!(self.out, "1 + ")?;
+                    self.put_dynamic_array_max_index(global, context)?
                 }
             }
-
-            chain = base
         }
 
         Ok(check_written)
@@ -5694,28 +5791,7 @@ template <typename A>
                 result_struct: None,
             };
 
-            for (local_handle, local) in fun.local_variables.iter() {
-                let ty_name = TypeContext {
-                    handle: local.ty,
-                    gctx: module.to_ctx(),
-                    names: &self.names,
-                    access: crate::StorageAccess::empty(),
-                    first_time: false,
-                };
-                let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)];
-                write!(self.out, "{}{} {}", back::INDENT, ty_name, local_name)?;
-                match local.init {
-                    Some(value) => {
-                        write!(self.out, " = ")?;
-                        self.put_expression(value, &context.expression, true)?;
-                    }
-                    None => {
-                        write!(self.out, " = {{}}")?;
-                    }
-                };
-                writeln!(self.out, ";")?;
-            }
-
+            self.put_locals(&context.expression)?;
             self.update_expressions_to_bake(fun, fun_info, &context.expression);
             self.put_block(back::Level(1), &fun.body, &context)?;
             writeln!(self.out, "}}")?;
@@ -6627,28 +6703,7 @@ template <typename A>
 
             // Finally, declare all the local variables that we need
             //TODO: we can postpone this till the relevant expressions are emitted
-            for (local_handle, local) in fun.local_variables.iter() {
-                let name = &self.names[&NameKey::EntryPointLocal(ep_index as _, local_handle)];
-                let ty_name = TypeContext {
-                    handle: local.ty,
-                    gctx: module.to_ctx(),
-                    names: &self.names,
-                    access: crate::StorageAccess::empty(),
-                    first_time: false,
-                };
-                write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?;
-                match local.init {
-                    Some(value) => {
-                        write!(self.out, " = ")?;
-                        self.put_expression(value, &context.expression, true)?;
-                    }
-                    None => {
-                        write!(self.out, " = {{}}")?;
-                    }
-                };
-                writeln!(self.out, ";")?;
-            }
-
+            self.put_locals(&context.expression)?;
             self.update_expressions_to_bake(fun, fun_info, &context.expression);
             self.put_block(back::Level(1), &fun.body, &context)?;
             writeln!(self.out, "}}")?;
diff --git a/naga/src/proc/index.rs b/naga/src/proc/index.rs
index 9f1c0ddb79..cf6a127ace 100644
--- a/naga/src/proc/index.rs
+++ b/naga/src/proc/index.rs
@@ -2,8 +2,10 @@
 Definitions for index bounds checking.
 */
 
+use core::iter::{self, zip};
+
 use crate::arena::{Handle, HandleSet, UniqueArena};
-use crate::valid;
+use crate::{valid, FastHashSet};
 
 /// How should code generated by Naga do bounds checks?
 ///
@@ -340,6 +342,128 @@ pub fn access_needs_check(
     Some(length)
 }
 
+/// Items returned by the [`bounds_check_iter`] iterator.
+#[cfg_attr(not(feature = "msl-out"), allow(dead_code))]
+pub(crate) struct BoundsCheck {
+    /// The base of the [`Access`] or [`AccessIndex`] expression.
+    ///
+    /// [`Access`]: crate::Expression::Access
+    /// [`AccessIndex`]: crate::Expression::AccessIndex
+    pub base: Handle<crate::Expression>,
+
+    /// The index being accessed.
+    pub index: GuardedIndex,
+
+    /// The length of `base`.
+    pub length: IndexableLength,
+}
+
+/// Returns an iterator of accesses within the chain of `Access` and
+/// `AccessIndex` expressions starting from `chain` that may need to be
+/// bounds-checked at runtime.
+///
+/// Items are yielded as [`BoundsCheck`] instances.
+///
+/// Accesses through a struct are omitted, since you never need a bounds check
+/// for accessing a struct field.
+///
+/// If `chain` isn't an `Access` or `AccessIndex` expression at all, the
+/// iterator is empty.
+pub(crate) fn bounds_check_iter<'a>(
+    mut chain: Handle<crate::Expression>,
+    module: &'a crate::Module,
+    function: &'a crate::Function,
+    info: &'a valid::FunctionInfo,
+) -> impl Iterator<Item = BoundsCheck> + 'a {
+    iter::from_fn(move || {
+        let (next_expr, result) = match function.expressions[chain] {
+            crate::Expression::Access { base, index } => {
+                (base, Some((base, GuardedIndex::Expression(index))))
+            }
+            crate::Expression::AccessIndex { base, index } => {
+                // Don't try to check indices into structs. Validation already took
+                // care of them, and access_needs_check doesn't handle that case.
+                let mut base_inner = info[base].ty.inner_with(&module.types);
+                if let crate::TypeInner::Pointer { base, .. } = *base_inner {
+                    base_inner = &module.types[base].inner;
+                }
+                match *base_inner {
+                    crate::TypeInner::Struct { .. } => (base, None),
+                    _ => (base, Some((base, GuardedIndex::Known(index)))),
+                }
+            }
+            _ => return None,
+        };
+        chain = next_expr;
+        Some(result)
+    })
+    .flatten()
+    .filter_map(|(base, index)| {
+        access_needs_check(base, index, module, &function.expressions, info).map(|length| {
+            BoundsCheck {
+                base,
+                index,
+                length,
+            }
+        })
+    })
+}
+
+/// Returns all the types which we need out-of-bounds locals for; that is,
+/// all of the types which the code might attempt to get an out-of-bounds
+/// pointer to, in which case we yield a pointer to the out-of-bounds local
+/// of the correct type.
+pub fn oob_local_types(
+    module: &crate::Module,
+    function: &crate::Function,
+    info: &valid::FunctionInfo,
+    policies: BoundsCheckPolicies,
+) -> FastHashSet<Handle<crate::Type>> {
+    let mut result = FastHashSet::default();
+
+    if policies.index != BoundsCheckPolicy::ReadZeroSkipWrite {
+        return result;
+    }
+
+    for statement in &function.body {
+        // The only situation in which we end up actually needing to create an
+        // out-of-bounds pointer is when passing one to a function.
+        //
+        // This is because pointers are never baked; they're just inlined everywhere
+        // they're used. That means that loads can just return 0, and stores can just do
+        // nothing; functions are the only case where you actually *have* to produce a
+        // pointer.
+        if let crate::Statement::Call {
+            function: callee,
+            ref arguments,
+            ..
+        } = *statement
+        {
+            // Now go through the arguments of the function looking for pointers which need bounds checks.
+            for (arg_info, &arg) in zip(&module.functions[callee].arguments, arguments) {
+                match module.types[arg_info.ty].inner {
+                    crate::TypeInner::ValuePointer { .. } => {
+                        // `ValuePointer`s should only ever be used when resolving the types of
+                        // expressions, since the arena can no longer be modified at that point; things
+                        // in the arena should always use proper `Pointer`s.
+                        unreachable!("`ValuePointer` found in arena")
+                    }
+                    crate::TypeInner::Pointer { base, .. } => {
+                        if bounds_check_iter(arg, module, function, info)
+                            .next()
+                            .is_some()
+                        {
+                            result.insert(base);
+                        }
+                    }
+                    _ => continue,
+                };
+            }
+        }
+    }
+    result
+}
+
 impl GuardedIndex {
     /// Make a `GuardedIndex::Known` from a `GuardedIndex::Expression` if possible.
     ///
diff --git a/naga/src/proc/namer.rs b/naga/src/proc/namer.rs
index 6b831cce79..05678242f8 100644
--- a/naga/src/proc/namer.rs
+++ b/naga/src/proc/namer.rs
@@ -21,9 +21,19 @@ pub enum NameKey {
     Function(Handle<crate::Function>),
     FunctionArgument(Handle<crate::Function>, u32),
     FunctionLocal(Handle<crate::Function>, Handle<crate::LocalVariable>),
+
+    /// A local variable used by ReadZeroSkipWrite bounds-check policy
+    /// when it needs to produce a pointer-typed result for an OOB access.
+    /// These are unique per accessed type, so the second element is a
+    /// type handle. See docs for [`crate::back::msl`].
+    FunctionOobLocal(Handle<crate::Function>, Handle<crate::Type>),
+
     EntryPoint(EntryPointIndex),
     EntryPointLocal(EntryPointIndex, Handle<crate::LocalVariable>),
     EntryPointArgument(EntryPointIndex, u32),
+
+    /// Entry point version of `FunctionOobLocal`.
+    EntryPointOobLocal(EntryPointIndex, Handle<crate::Type>),
 }
 
 /// This processor assigns names to all the things in a module
diff --git a/naga/tests/in/wgsl/pointer-function-arg-restrict.toml b/naga/tests/in/wgsl/pointer-function-arg-restrict.toml
new file mode 100644
index 0000000000..bc4112cddf
--- /dev/null
+++ b/naga/tests/in/wgsl/pointer-function-arg-restrict.toml
@@ -0,0 +1,4 @@
+targets = "METAL"
+
+[bounds_check_policies]
+index = "Restrict"
diff --git a/naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl b/naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl
new file mode 100644
index 0000000000..08ac388354
--- /dev/null
+++ b/naga/tests/in/wgsl/pointer-function-arg-restrict.wgsl
@@ -0,0 +1,61 @@
+fn takes_ptr(p: ptr<function, i32>) {}
+fn takes_array_ptr(p: ptr<function, array<i32, 4>>) {}
+fn takes_vec_ptr(p: ptr<function, vec2<i32>>) {}
+fn takes_mat_ptr(p: ptr<function, mat2x2<f32>>) {}
+
+fn local_var(i: u32) {
+    var arr = array(1, 2, 3, 4);
+    takes_ptr(&arr[i]);
+    takes_array_ptr(&arr);
+
+}
+
+fn mat_vec_ptrs(
+    pv: ptr<function, array<vec2<i32>, 4>>,
+    pm: ptr<function, array<mat2x2<f32>, 4>>,
+    i: u32,
+) {
+    takes_vec_ptr(&pv[i]);
+    takes_mat_ptr(&pm[i]);
+}
+
+fn argument(v: ptr<function, array<i32, 4>>, i: u32) {
+    takes_ptr(&v[i]);
+}
+
+fn argument_nested_x2(v: ptr<function, array<array<i32, 4>, 4>>, i: u32, j: u32) {
+    takes_ptr(&v[i][j]);
+
+    // Mixing compile and runtime bounds checks
+    takes_ptr(&v[i][0]);
+    takes_ptr(&v[0][j]);
+
+    takes_array_ptr(&v[i]);
+}
+
+fn argument_nested_x3(v: ptr<function, array<array<array<i32, 4>, 4>, 4>>, i: u32, j: u32) {
+    takes_ptr(&v[i][0][j]);
+    takes_ptr(&v[i][j][0]);
+    takes_ptr(&v[0][i][j]);
+}
+
+fn index_from_self(v: ptr<function, array<i32, 4>>, i: u32) {
+    takes_ptr(&v[v[i]]);
+}
+
+fn local_var_from_arg(a: array<i32, 4>, i: u32) {
+    var b = a;
+    takes_ptr(&b[i]);
+}
+
+fn let_binding(a: ptr<function, array<i32, 4>>, i: u32) {
+    let p0 = &a[i];
+    takes_ptr(p0);
+
+    let p1 = &a[0];
+    takes_ptr(p1);
+}
+
+// Runtime-sized arrays can only appear in storage buffers, while (in the base
+// language) pointers can only appear in function or private space, so there
+// is no interaction to test.
diff --git a/naga/tests/in/wgsl/pointer-function-arg-rzsw.toml b/naga/tests/in/wgsl/pointer-function-arg-rzsw.toml
new file mode 100644
index 0000000000..44773b7d04
--- /dev/null
+++ b/naga/tests/in/wgsl/pointer-function-arg-rzsw.toml
@@ -0,0 +1,4 @@
+targets = "METAL"
+
+[bounds_check_policies]
+index = "ReadZeroSkipWrite"
diff --git a/naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl b/naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl
new file mode 100644
index 0000000000..08ac388354
--- /dev/null
+++ b/naga/tests/in/wgsl/pointer-function-arg-rzsw.wgsl
@@ -0,0 +1,61 @@
+fn takes_ptr(p: ptr<function, i32>) {}
+fn takes_array_ptr(p: ptr<function, array<i32, 4>>) {}
+fn takes_vec_ptr(p: ptr<function, vec2<i32>>) {}
+fn takes_mat_ptr(p: ptr<function, mat2x2<f32>>) {}
+
+fn local_var(i: u32) {
+    var arr = array(1, 2, 3, 4);
+    takes_ptr(&arr[i]);
+    takes_array_ptr(&arr);
+
+}
+
+fn mat_vec_ptrs(
+    pv: ptr<function, array<vec2<i32>, 4>>,
+    pm: ptr<function, array<mat2x2<f32>, 4>>,
+    i: u32,
+) {
+    takes_vec_ptr(&pv[i]);
+    takes_mat_ptr(&pm[i]);
+}
+
+fn argument(v: ptr<function, array<i32, 4>>, i: u32) {
+    takes_ptr(&v[i]);
+}
+
+fn argument_nested_x2(v: ptr<function, array<array<i32, 4>, 4>>, i: u32, j: u32) {
+    takes_ptr(&v[i][j]);
+
+    // Mixing compile and runtime bounds checks
+    takes_ptr(&v[i][0]);
+    takes_ptr(&v[0][j]);
+
+    takes_array_ptr(&v[i]);
+}
+
+fn argument_nested_x3(v: ptr<function, array<array<array<i32, 4>, 4>, 4>>, i: u32, j: u32) {
+    takes_ptr(&v[i][0][j]);
+    takes_ptr(&v[i][j][0]);
+    takes_ptr(&v[0][i][j]);
+}
+
+fn index_from_self(v: ptr<function, array<i32, 4>>, i: u32) {
+    takes_ptr(&v[v[i]]);
+}
+
+fn local_var_from_arg(a: array<i32, 4>, i: u32) {
+    var b = a;
+    takes_ptr(&b[i]);
+}
+
+fn let_binding(a: ptr<function, array<i32, 4>>, i: u32) {
+    let p0 = &a[i];
+    takes_ptr(p0);
+
+    let p1 = &a[0];
+    takes_ptr(p1);
+}
+
+// Runtime-sized arrays can only appear in storage buffers, while (in the base
+// language) pointers can only appear in function or private space, so there
+// is no interaction to test.
diff --git a/naga/tests/in/wgsl/pointer-function-arg.toml b/naga/tests/in/wgsl/pointer-function-arg.toml
new file mode 100644
index 0000000000..e74ee3b97b
--- /dev/null
+++ b/naga/tests/in/wgsl/pointer-function-arg.toml
@@ -0,0 +1 @@
+targets = "METAL | GLSL | HLSL | WGSL"
diff --git a/naga/tests/in/wgsl/pointer-function-arg.wgsl b/naga/tests/in/wgsl/pointer-function-arg.wgsl
new file mode 100644
index 0000000000..606835eef7
--- /dev/null
+++ b/naga/tests/in/wgsl/pointer-function-arg.wgsl
@@ -0,0 +1,64 @@
+@compute @workgroup_size(1)
+fn main() {}
+
+fn takes_ptr(p: ptr<function, i32>) {}
+fn takes_array_ptr(p: ptr<function, array<i32, 4>>) {}
+fn takes_vec_ptr(p: ptr<function, vec2<i32>>) {}
+fn takes_mat_ptr(p: ptr<function, mat2x2<f32>>) {}
+
+fn local_var(i: u32) {
+    var arr = array(1, 2, 3, 4);
+    takes_ptr(&arr[i]);
+    takes_array_ptr(&arr);
+
+}
+
+fn mat_vec_ptrs(
+    pv: ptr<function, array<vec2<i32>, 4>>,
+    pm: ptr<function, array<mat2x2<f32>, 4>>,
+    i: u32,
+) {
+    takes_vec_ptr(&pv[i]);
+    takes_mat_ptr(&pm[i]);
+}
+
+fn argument(v: ptr<function, array<i32, 4>>, i: u32) {
+    takes_ptr(&v[i]);
+}
+
+fn argument_nested_x2(v: ptr<function, array<array<i32, 4>, 4>>, i: u32, j: u32) {
+    takes_ptr(&v[i][j]);
+
+    // Mixing compile and runtime bounds checks
+    takes_ptr(&v[i][0]);
+    takes_ptr(&v[0][j]);
+
+    takes_array_ptr(&v[i]);
+}
+
+fn argument_nested_x3(v: ptr<function, array<array<array<i32, 4>, 4>, 4>>, i: u32, j: u32) {
+    takes_ptr(&v[i][0][j]);
+    takes_ptr(&v[i][j][0]);
+    takes_ptr(&v[0][i][j]);
+}
+
+fn index_from_self(v: ptr<function, array<i32, 4>>, i: u32) {
+    takes_ptr(&v[v[i]]);
+}
+
+fn local_var_from_arg(a: array<i32, 4>, i: u32) {
+    var b = a;
+    takes_ptr(&b[i]);
+}
+
+fn let_binding(a: ptr<function, array<i32, 4>>, i: u32) {
+    let p0 = &a[i];
+    takes_ptr(p0);
+
+    let p1 = &a[0];
+    takes_ptr(p1);
+}
+
+// Runtime-sized arrays can only appear in storage buffers, while (in the base
+// language) pointers can only appear in function or private space, so there
+// is no interaction to test.
diff --git a/naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl b/naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl
new file mode 100644
index 0000000000..a3ffe7952e
--- /dev/null
+++ b/naga/tests/out/glsl/wgsl-pointer-function-arg.main.Compute.glsl
@@ -0,0 +1,80 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+
+void takes_ptr(inout int p) {
+    return;
+}
+
+void takes_array_ptr(inout int p_1[4]) {
+    return;
+}
+
+void takes_vec_ptr(inout ivec2 p_2) {
+    return;
+}
+
+void takes_mat_ptr(inout mat2x2 p_3) {
+    return;
+}
+
+void local_var(uint i) {
+    int arr[4] = int[4](1, 2, 3, 4);
+    takes_ptr(arr[i]);
+    takes_array_ptr(arr);
+    return;
+}
+
+void mat_vec_ptrs(inout ivec2 pv[4], inout mat2x2 pm[4], uint i_1) {
+    takes_vec_ptr(pv[i_1]);
+    takes_mat_ptr(pm[i_1]);
+    return;
+}
+
+void argument(inout int v[4], uint i_2) {
+    takes_ptr(v[i_2]);
+    return;
+}
+
+void argument_nested_x2_(inout int v_1[4][4], uint i_3, uint j) {
+    takes_ptr(v_1[i_3][j]);
+    takes_ptr(v_1[i_3][0]);
+    takes_ptr(v_1[0][j]);
+    takes_array_ptr(v_1[i_3]);
+    return;
+}
+
+void argument_nested_x3_(inout int v_2[4][4][4], uint i_4, uint j_1) {
+    takes_ptr(v_2[i_4][0][j_1]);
+    takes_ptr(v_2[i_4][j_1][0]);
+    takes_ptr(v_2[0][i_4][j_1]);
+    return;
+}
+
+void index_from_self(inout int v_3[4], uint i_5) {
+    int _e3 = v_3[i_5];
+    takes_ptr(v_3[_e3]);
+    return;
+}
+
+void local_var_from_arg(int a[4], uint i_6) {
+    int b[4] = int[4](0, 0, 0, 0);
+    b = a;
+    takes_ptr(b[i_6]);
+    return;
+}
+
+void let_binding(inout int a_1[4], uint i_7) {
+    takes_ptr(a_1[i_7]);
+    takes_ptr(a_1[0]);
+    return;
+}
+
+void main() {
+    return;
+}
+
diff --git a/naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl b/naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl
new file mode 100644
index 0000000000..1f4e76eeae
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-pointer-function-arg.hlsl
@@ -0,0 +1,93 @@
+void takes_ptr(inout int p)
+{
+    return;
+}
+
+void takes_array_ptr(inout int p_1[4])
+{
+    return;
+}
+
+void takes_vec_ptr(inout int2 p_2)
+{
+    return;
+}
+
+void takes_mat_ptr(inout float2x2 p_3)
+{
+    return;
+}
+
+typedef int ret_Constructarray4_int_[4];
+ret_Constructarray4_int_ Constructarray4_int_(int arg0, int arg1, int arg2, int arg3) {
+    int ret[4] = { arg0, arg1, arg2, arg3 };
+    return ret;
+}
+
+void local_var(uint i)
+{
+    int arr[4] = Constructarray4_int_(int(1), int(2), int(3), int(4));
+
+    takes_ptr(arr[min(uint(i), 3u)]);
+    takes_array_ptr(arr);
+    return;
+}
+
+void mat_vec_ptrs(inout int2 pv[4], inout float2x2 pm[4], uint i_1)
+{
+    takes_vec_ptr(pv[min(uint(i_1), 3u)]);
+    takes_mat_ptr(pm[min(uint(i_1), 3u)]);
+    return;
+}
+
+void argument(inout int v[4], uint i_2)
+{
+    takes_ptr(v[min(uint(i_2), 3u)]);
+    return;
+}
+
+void argument_nested_x2_(inout int v_1[4][4], uint i_3, uint j)
+{
+    takes_ptr(v_1[min(uint(i_3), 3u)][min(uint(j), 3u)]);
+    takes_ptr(v_1[min(uint(i_3), 3u)][0]);
+    takes_ptr(v_1[0][min(uint(j), 3u)]);
+    takes_array_ptr(v_1[min(uint(i_3), 3u)]);
+    return;
+}
+
+void argument_nested_x3_(inout int v_2[4][4][4], uint i_4, uint j_1)
+{
+    takes_ptr(v_2[min(uint(i_4), 3u)][0][min(uint(j_1), 3u)]);
+    takes_ptr(v_2[min(uint(i_4), 3u)][min(uint(j_1), 3u)][0]);
+    takes_ptr(v_2[0][min(uint(i_4), 3u)][min(uint(j_1), 3u)]);
+    return;
+}
+
+void index_from_self(inout int v_3[4], uint i_5)
+{
+    int _e3 = v_3[min(uint(i_5), 3u)];
+    takes_ptr(v_3[min(uint(_e3), 3u)]);
+    return;
+}
+
+void local_var_from_arg(int a[4], uint i_6)
+{
+    int b[4] = (int[4])0;
+
+    b = a;
+    takes_ptr(b[min(uint(i_6), 3u)]);
+    return;
+}
+
+void let_binding(inout int a_1[4], uint i_7)
+{
+    takes_ptr(a_1[min(uint(i_7), 3u)]);
+    takes_ptr(a_1[0]);
+    return;
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    return;
+}
diff --git a/naga/tests/out/hlsl/wgsl-pointer-function-arg.ron b/naga/tests/out/hlsl/wgsl-pointer-function-arg.ron
new file mode 100644
index 0000000000..a07b03300b
--- /dev/null
+++ b/naga/tests/out/hlsl/wgsl-pointer-function-arg.ron
@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)
diff --git a/naga/tests/out/msl/wgsl-pointer-function-arg-restrict.msl b/naga/tests/out/msl/wgsl-pointer-function-arg-restrict.msl
new file mode 100644
index 0000000000..66990660e0
--- /dev/null
+++ b/naga/tests/out/msl/wgsl-pointer-function-arg-restrict.msl
@@ -0,0 +1,123 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+struct type_2 {
+    int inner[4];
+};
+struct type_9 {
+    metal::int2 inner[4];
+};
+struct type_11 {
+    metal::float2x2 inner[4];
+};
+struct type_13 {
+    type_2 inner[4];
+};
+struct type_15 {
+    type_13 inner[4];
+};
+
+void takes_ptr(
+    thread int& p
+) {
+    return;
+}
+
+void takes_array_ptr(
+    thread type_2& p_1
+) {
+    return;
+}
+
+void takes_vec_ptr(
+    thread metal::int2& p_2
+) {
+    return;
+}
+
+void takes_mat_ptr(
+    thread metal::float2x2& p_3
+) {
+    return;
+}
+
+void local_var(
+    uint i
+) {
+    type_2 arr = type_2 {1, 2, 3, 4};
+    takes_ptr(arr.inner[metal::min(unsigned(i), 3u)]);
+    takes_array_ptr(arr);
+    return;
+}
+
+void mat_vec_ptrs(
+    thread type_9& pv,
+    thread type_11& pm,
+    uint i_1
+) {
+    takes_vec_ptr(pv.inner[metal::min(unsigned(i_1), 3u)]);
+    takes_mat_ptr(pm.inner[metal::min(unsigned(i_1), 3u)]);
+    return;
+}
+
+void argument(
+    thread type_2& v,
+    uint i_2
+) {
+    takes_ptr(v.inner[metal::min(unsigned(i_2), 3u)]);
+    return;
+}
+
+void argument_nested_x2_(
+    thread type_13& v_1,
+    uint i_3,
+    uint j
+) {
+    takes_ptr(v_1.inner[metal::min(unsigned(i_3), 3u)].inner[metal::min(unsigned(j), 3u)]);
+    takes_ptr(v_1.inner[metal::min(unsigned(i_3), 3u)].inner[0]);
+    takes_ptr(v_1.inner[0].inner[metal::min(unsigned(j), 3u)]);
+    takes_array_ptr(v_1.inner[metal::min(unsigned(i_3), 3u)]);
+    return;
+}
+
+void argument_nested_x3_(
+    thread type_15& v_2,
+    uint i_4,
+    uint j_1
+) {
+    takes_ptr(v_2.inner[metal::min(unsigned(i_4), 3u)].inner[0].inner[metal::min(unsigned(j_1), 3u)]);
+    takes_ptr(v_2.inner[metal::min(unsigned(i_4), 3u)].inner[metal::min(unsigned(j_1), 3u)].inner[0]);
+    takes_ptr(v_2.inner[0].inner[metal::min(unsigned(i_4), 3u)].inner[metal::min(unsigned(j_1), 3u)]);
+    return;
+}
+
+void index_from_self(
+    thread type_2& v_3,
+    uint i_5
+) {
+    int _e3 = v_3.inner[metal::min(unsigned(i_5), 3u)];
+    takes_ptr(v_3.inner[metal::min(unsigned(_e3), 3u)]);
+    return;
+}
+
+void local_var_from_arg(
+    type_2 a,
+    uint i_6
+) {
+    type_2 b = {};
+    b = a;
+    takes_ptr(b.inner[metal::min(unsigned(i_6), 3u)]);
+    return;
+}
+
+void let_binding(
+    thread type_2& a_1,
+    uint i_7
+) {
+    takes_ptr(a_1.inner[metal::min(unsigned(i_7), 3u)]);
+    takes_ptr(a_1.inner[0]);
+    return;
+}
diff --git a/naga/tests/out/msl/wgsl-pointer-function-arg-rzsw.msl b/naga/tests/out/msl/wgsl-pointer-function-arg-rzsw.msl
new file mode 100644
index 0000000000..8b76ec2505
--- /dev/null
+++ b/naga/tests/out/msl/wgsl-pointer-function-arg-rzsw.msl
@@ -0,0 +1,139 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+struct DefaultConstructible {
+    template<typename T>
+    operator T() && {
+        return T {};
+    }
+};
+
+struct type_2 {
+    int inner[4];
+};
+struct type_9 {
+    metal::int2 inner[4];
+};
+struct type_11 {
+    metal::float2x2 inner[4];
+};
+struct type_13 {
+    type_2 inner[4];
+};
+struct type_15 {
+    type_13 inner[4];
+};
+
+void takes_ptr(
+    thread int& p
+) {
+    return;
+}
+
+void takes_array_ptr(
+    thread type_2& p_1
+) {
+    return;
+}
+
+void takes_vec_ptr(
+    thread metal::int2& p_2
+) {
+    return;
+}
+
+void takes_mat_ptr(
+    thread metal::float2x2& p_3
+) {
+    return;
+}
+
+void local_var(
+    uint i
+) {
+    type_2 arr = type_2 {1, 2, 3, 4};
+    int oob = {};
+    takes_ptr(uint(i) < 4 ? arr.inner[i] : oob);
+    takes_array_ptr(arr);
+    return;
+}
+
+void mat_vec_ptrs(
+    thread type_9& pv,
+    thread type_11& pm,
+    uint i_1
+) {
+    metal::int2 oob_1 = {};
+    metal::float2x2 oob_2 = {};
+    takes_vec_ptr(uint(i_1) < 4 ? pv.inner[i_1] : oob_1);
+    takes_mat_ptr(uint(i_1) < 4 ? pm.inner[i_1] : oob_2);
+    return;
+}
+
+void argument(
+    thread type_2& v,
+    uint i_2
+) {
+    int oob_3 = {};
+    takes_ptr(uint(i_2) < 4 ? v.inner[i_2] : oob_3);
+    return;
+}
+
+void argument_nested_x2_(
+    thread type_13& v_1,
+    uint i_3,
+    uint j
+) {
+    int oob_4 = {};
+    type_2 oob_5 = {};
+    takes_ptr(uint(j) < 4 && uint(i_3) < 4 ? v_1.inner[i_3].inner[j] : oob_4);
+    takes_ptr(uint(i_3) < 4 ? v_1.inner[i_3].inner[0] : oob_4);
+    takes_ptr(uint(j) < 4 ? v_1.inner[0].inner[j] : oob_4);
+    takes_array_ptr(uint(i_3) < 4 ? v_1.inner[i_3] : oob_5);
+    return;
+}
+
+void argument_nested_x3_(
+    thread type_15& v_2,
+    uint i_4,
+    uint j_1
+) {
+    int oob_6 = {};
+    takes_ptr(uint(j_1) < 4 && uint(i_4) < 4 ? v_2.inner[i_4].inner[0].inner[j_1] : oob_6);
+    takes_ptr(uint(j_1) < 4 && uint(i_4) < 4 ? v_2.inner[i_4].inner[j_1].inner[0] : oob_6);
+    takes_ptr(uint(j_1) < 4 && uint(i_4) < 4 ? v_2.inner[0].inner[i_4].inner[j_1] : oob_6);
+    return;
+}
+
+void index_from_self(
+    thread type_2& v_3,
+    uint i_5
+) {
+    int oob_7 = {};
+    int _e3 = uint(i_5) < 4 ? v_3.inner[i_5] : DefaultConstructible();
+    takes_ptr(uint(_e3) < 4 ? v_3.inner[_e3] : oob_7);
+    return;
+}
+
+void local_var_from_arg(
+    type_2 a,
+    uint i_6
+) {
+    type_2 b = {};
+    int oob_8 = {};
+    b = a;
+    takes_ptr(uint(i_6) < 4 ? b.inner[i_6] : oob_8);
+    return;
+}
+
+void let_binding(
+    thread type_2& a_1,
+    uint i_7
+) {
+    int oob_9 = {};
+    takes_ptr(uint(i_7) < 4 ? a_1.inner[i_7] : oob_9);
+    takes_ptr(a_1.inner[0]);
+    return;
+}
diff --git a/naga/tests/out/msl/wgsl-pointer-function-arg.msl b/naga/tests/out/msl/wgsl-pointer-function-arg.msl
new file mode 100644
index 0000000000..c0d1d264fa
--- /dev/null
+++ b/naga/tests/out/msl/wgsl-pointer-function-arg.msl
@@ -0,0 +1,128 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+struct type_2 {
+    int inner[4];
+};
+struct type_9 {
+    metal::int2 inner[4];
+};
+struct type_11 {
+    metal::float2x2 inner[4];
+};
+struct type_13 {
+    type_2 inner[4];
+};
+struct type_15 {
+    type_13 inner[4];
+};
+
+void takes_ptr(
+    thread int& p
+) {
+    return;
+}
+
+void takes_array_ptr(
+    thread type_2& p_1
+) {
+    return;
+}
+
+void takes_vec_ptr(
+    thread metal::int2& p_2
+) {
+    return;
+}
+
+void takes_mat_ptr(
+    thread metal::float2x2& p_3
+) {
+    return;
+}
+
+void local_var(
+    uint i
+) {
+    type_2 arr = type_2 {1, 2, 3, 4};
+    takes_ptr(arr.inner[i]);
+    takes_array_ptr(arr);
+    return;
+}
+
+void mat_vec_ptrs(
+    thread type_9& pv,
+    thread type_11& pm,
+    uint i_1
+) {
+    takes_vec_ptr(pv.inner[i_1]);
+    takes_mat_ptr(pm.inner[i_1]);
+    return;
+}
+
+void argument(
+    thread type_2& v,
+    uint i_2
+) {
+    takes_ptr(v.inner[i_2]);
+    return;
+}
+
+void argument_nested_x2_(
+    thread type_13& v_1,
+    uint i_3,
+    uint j
+) {
+    takes_ptr(v_1.inner[i_3].inner[j]);
+    takes_ptr(v_1.inner[i_3].inner[0]);
+    takes_ptr(v_1.inner[0].inner[j]);
+    takes_array_ptr(v_1.inner[i_3]);
+    return;
+}
+
+void argument_nested_x3_(
+    thread type_15& v_2,
+    uint i_4,
+    uint j_1
+) {
+    takes_ptr(v_2.inner[i_4].inner[0].inner[j_1]);
+    takes_ptr(v_2.inner[i_4].inner[j_1].inner[0]);
+    takes_ptr(v_2.inner[0].inner[i_4].inner[j_1]);
+    return;
+}
+
+void index_from_self(
+    thread type_2& v_3,
+    uint i_5
+) {
+    int _e3 = v_3.inner[i_5];
+    takes_ptr(v_3.inner[_e3]);
+    return;
+}
+
+void local_var_from_arg(
+    type_2 a,
+    uint i_6
+) {
+    type_2 b = {};
+    b = a;
+    takes_ptr(b.inner[i_6]);
+    return;
+}
+
+void let_binding(
+    thread type_2& a_1,
+    uint i_7
+) {
+    takes_ptr(a_1.inner[i_7]);
+    takes_ptr(a_1.inner[0]);
+    return;
+}
+
+kernel void main_(
+) {
+    return;
+}
diff --git a/naga/tests/out/wgsl/wgsl-pointer-function-arg.wgsl b/naga/tests/out/wgsl/wgsl-pointer-function-arg.wgsl
new file mode 100644
index 0000000000..20a2349fcb
--- /dev/null
+++ b/naga/tests/out/wgsl/wgsl-pointer-function-arg.wgsl
@@ -0,0 +1,76 @@
+fn takes_ptr(p: ptr<function, i32>) {
+    return;
+}
+
+fn takes_array_ptr(p_1: ptr<function, array<i32, 4>>) {
+    return;
+}
+
+fn takes_vec_ptr(p_2: ptr<function, vec2<i32>>) {
+    return;
+}
+
+fn takes_mat_ptr(p_3: ptr<function, mat2x2<f32>>) {
+    return;
+}
+
+fn local_var(i: u32) {
+    var arr: array<i32, 4> = array<i32, 4>(1i, 2i, 3i, 4i);
+
+    takes_ptr((&arr[i]));
+    takes_array_ptr((&arr));
+    return;
+}
+
+fn mat_vec_ptrs(pv: ptr<function, array<vec2<i32>, 4>>, pm: ptr<function, array<mat2x2<f32>, 4>>, i_1: u32) {
+    takes_vec_ptr((&(*pv)[i_1]));
+    takes_mat_ptr((&(*pm)[i_1]));
+    return;
+}
+
+fn argument(v: ptr<function, array<i32, 4>>, i_2: u32) {
+    takes_ptr((&(*v)[i_2]));
+    return;
+}
+
+fn argument_nested_x2_(v_1: ptr<function, array<array<i32, 4>, 4>>, i_3: u32, j: u32) {
+    takes_ptr((&(*v_1)[i_3][j]));
+    takes_ptr((&(*v_1)[i_3][0]));
+    takes_ptr((&(*v_1)[0][j]));
+    takes_array_ptr((&(*v_1)[i_3]));
+    return;
+}
+
+fn argument_nested_x3_(v_2: ptr<function, array<array<array<i32, 4>, 4>, 4>>, i_4: u32, j_1: u32) {
+    takes_ptr((&(*v_2)[i_4][0][j_1]));
+    takes_ptr((&(*v_2)[i_4][j_1][0]));
+    takes_ptr((&(*v_2)[0][i_4][j_1]));
+    return;
+}
+
+fn index_from_self(v_3: ptr<function, array<i32, 4>>, i_5: u32) {
+    let _e3 = (*v_3)[i_5];
+    takes_ptr((&(*v_3)[_e3]));
+    return;
+}
+
+fn local_var_from_arg(a: array<i32, 4>, i_6: u32) {
+    var b: array<i32, 4>;
+
+    b = a;
+    takes_ptr((&b[i_6]));
+    return;
+}
+
+fn let_binding(a_1: ptr<function, array<i32, 4>>, i_7: u32) {
+    let p0_ = (&(*a_1)[i_7]);
+    takes_ptr(p0_);
+    let p1_ = (&(*a_1)[0]);
+    takes_ptr(p1_);
+    return;
+}
+
+@compute @workgroup_size(1, 1, 1) 
+fn main() {
+    return;
+}