Repository: gfx-rs/wgpu
Pull request #6346 (2024-09-30 21:02:58)

[naga spv-out] Internal doc fixes for `back::spv::index`.
In `naga::back::spv::index`, clarify documentation for:
- `BoundsCheckResult`
- `write_index_comparison`
- `write_bounds_check`



Patch:
diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs
index f0c3bfa84..97db1cb54 100644
--- a/naga/src/back/spv/block.rs
+++ b/naga/src/back/spv/block.rs
@@ -343,6 +343,32 @@ impl<'w> BlockContext<'w> {
 
                         load_id
                     }
+                    crate::TypeInner::Array {
+                        base: ty_element, ..
+                    } => {
+                        let index_id = self.cached[index];
+                        let base_id = self.cached[base];
+                        let base_ty = match self.fun_info[base].ty {
+                            TypeResolution::Handle(handle) => handle,
+                            TypeResolution::Value(_) => {
+                                return Err(Error::Validation(
+                                    "Array types should always be in the arena",
+                                ))
+                            }
+                        };
+                        let (id, variable) = self.writer.promote_access_expression_to_variable(
+                            &self.ir_module.types,
+                            result_type_id,
+                            base_id,
+                            base_ty,
+                            index_id,
+                            ty_element,
+                            block,
+                        )?;
+                        self.function.internal_variables.push(variable);
+                        id
+                    }
+                    // wgpu#4337: Support `crate::TypeInner::Matrix`
                     ref other => {
                         log::error!(
                             "Unable to access base {:?} of type {:?}",
@@ -350,7 +376,7 @@ impl<'w> BlockContext<'w> {
                             other
                         );
                         return Err(Error::Validation(
-                            "only vectors may be dynamically indexed by value",
+                            "only vectors and arrays may be dynamically indexed by value",
                         ));
                     }
                 }
diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs
index 0295d895b..832f806fd 100644
--- a/naga/src/back/spv/index.rs
+++ b/naga/src/back/spv/index.rs
@@ -11,16 +11,31 @@ use crate::{arena::Handle, proc::BoundsCheckPolicy};
 
 /// The results of performing a bounds check.
 ///
-/// On success, `write_bounds_check` returns a value of this type.
+/// On success, [`write_bounds_check`](BlockContext::write_bounds_check)
+/// returns a value of this type. The caller can assume that the right
+/// policy has been applied, and simply do what the variant says.
 pub(super) enum BoundsCheckResult {
     /// The index is statically known and in bounds, with the given value.
     KnownInBounds(u32),
 
     /// The given instruction computes the index to be used.
+    ///
+    /// When [`BoundsCheckPolicy::Restrict`] is in force, this is a
+    /// clamped version of the index the user supplied.
+    ///
+    /// When [`BoundsCheckPolicy::Unchecked`] is in force, this is
+    /// simply the index the user supplied. This variant indicates
+    /// that we couldn't prove statically that the index was in
+    /// bounds; otherwise we would have returned [`KnownInBounds`].
+    ///
+    /// [`KnownInBounds`]: BoundsCheckResult::KnownInBounds
     Computed(Word),
 
     /// The given instruction computes a boolean condition which is true
     /// if the index is in bounds.
+    ///
+    /// This is returned when [`BoundsCheckPolicy::ReadZeroSkipWrite`]
+    /// is in force.
     Conditional(Word),
 }
 
@@ -357,6 +372,8 @@ impl<'w> BlockContext<'w> {
 
     /// Write an index bounds comparison to `block`, if needed.
     ///
+    /// This is used to implement [`BoundsCheckPolicy::ReadZeroSkipWrite`].
+    ///
     /// If we're able to determine statically that `index` is in bounds for
     /// `sequence`, return `KnownInBounds(value)`, where `value` is the actual
     /// value of the index. (In principle, one could know that the index is in
@@ -477,11 +494,23 @@ impl<'w> BlockContext<'w> {
 
     /// Emit code for bounds checks for an array, vector, or matrix access.
     ///
-    /// This implements either `index_bounds_check_policy` or
-    /// `buffer_bounds_check_policy`, depending on the address space of the
-    /// pointer being accessed.
+    /// This tries to handle all the critical steps for bounds checks:
+    ///
+    /// - First, select the appropriate bounds check policy for `base`,
+    ///   depending on its address space.
+    ///
+    /// - Next, analyze `index` to see if its value is known at
+    ///   compile time, in which case we can decide statically whether
+    ///   the index is in bounds.
+    ///
+    /// - If the index's value is not known at compile time, emit code to:
+    ///
+    ///     - restrict its value (for [`BoundsCheckPolicy::Restrict`]), or
+    ///
+    ///     - check whether it's in bounds (for
+    ///       [`BoundsCheckPolicy::ReadZeroSkipWrite`]).
     ///
-    /// Return a `BoundsCheckResult` indicating how the index should be
+    /// Return a [`BoundsCheckResult`] indicating how the index should be
     /// consumed. See that type's documentation for details.
     pub(super) fn write_bounds_check(
         &mut self,
diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs
index 32bd1fcec..77f915abb 100644
--- a/naga/src/back/spv/mod.rs
+++ b/naga/src/back/spv/mod.rs
@@ -144,6 +144,7 @@ struct Function {
     signature: Option<Instruction>,
     parameters: Vec<FunctionArgument>,
     variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
+    internal_variables: Vec<LocalVariable>,
     blocks: Vec<TerminatedBlock>,
     entry_point_context: Option<EntryPointContext>,
 }
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index 678dcb424..7f41caf29 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -32,6 +32,9 @@ impl Function {
                 for local_var in self.variables.values() {
                     local_var.instruction.to_words(sink);
                 }
+                for internal_var in self.internal_variables.iter() {
+                    internal_var.instruction.to_words(sink);
+                }
             }
             for instruction in block.body.iter() {
                 instruction.to_words(sink);
@@ -135,6 +138,56 @@ impl Writer {
         self.capabilities_used.insert(spirv::Capability::Shader);
     }
 
+    #[allow(clippy::too_many_arguments)]
+    pub(super) fn promote_access_expression_to_variable(
+        &mut self,
+        ir_types: &UniqueArena<crate::Type>,
+        result_type_id: Word,
+        container_id: Word,
+        container_ty: Handle<crate::Type>,
+        index_id: Word,
+        element_ty: Handle<crate::Type>,
+        block: &mut Block,
+    ) -> Result<(Word, LocalVariable), Error> {
+        let pointer_type_id =
+            self.get_pointer_id(ir_types, container_ty, spirv::StorageClass::Function)?;
+
+        let variable = {
+            let id = self.id_gen.next();
+            LocalVariable {
+                id,
+                instruction: Instruction::variable(
+                    pointer_type_id,
+                    id,
+                    spirv::StorageClass::Function,
+                    None,
+                ),
+            }
+        };
+        block
+            .body
+            .push(Instruction::store(variable.id, container_id, None));
+
+        let element_pointer_id = self.id_gen.next();
+        let element_pointer_type_id =
+            self.get_pointer_id(ir_types, element_ty, spirv::StorageClass::Function)?;
+        block.body.push(Instruction::access_chain(
+            element_pointer_type_id,
+            element_pointer_id,
+            variable.id,
+            &[index_id],
+        ));
+        let id = self.id_gen.next();
+        block.body.push(Instruction::load(
+            result_type_id,
+            id,
+            element_pointer_id,
+            None,
+        ));
+
+        Ok((id, variable))
+    }
+
     /// Indicate that the code requires any one of the listed capabilities.
     ///
     /// If nothing in `capabilities` appears in the available capabilities
diff --git a/naga/src/lib.rs b/naga/src/lib.rs
index 85fd7a450..038e215a6 100644
--- a/naga/src/lib.rs
+++ b/naga/src/lib.rs
@@ -1402,21 +1402,20 @@ pub enum Expression {
     /// ## Dynamic indexing restrictions
     ///
     /// To accommodate restrictions in some of the shader languages that Naga
-    /// targets, it is not permitted to subscript a matrix or array with a
-    /// dynamically computed index unless that matrix or array appears behind a
-    /// pointer. In other words, if the inner type of `base` is [`Array`] or
-    /// [`Matrix`], then `index` must be a constant. But if the type of `base`
-    /// is a [`Pointer`] to an array or matrix or a [`ValuePointer`] with a
-    /// `size`, then the index may be any expression of integer type.
+    /// targets, it is not permitted to subscript a matrix with a dynamically
+    /// computed index unless that matrix appears behind a pointer. In other
+    /// words, if the inner type of `base` is [`Matrix`], then `index` must be a
+    /// constant. But if the type of `base` is a [`Pointer`] to an matrix, then
+    /// the index may be any expression of integer type.
     ///
     /// You can use the [`Expression::is_dynamic_index`] method to determine
-    /// whether a given index expression requires matrix or array base operands
-    /// to be behind a pointer.
+    /// whether a given index expression requires matrix base operands to be
+    /// behind a pointer.
     ///
     /// (It would be simpler to always require the use of `AccessIndex` when
-    /// subscripting arrays and matrices that are not behind pointers, but to
-    /// accommodate existing front ends, Naga also permits `Access`, with a
-    /// restricted `index`.)
+    /// subscripting matrices that are not behind pointers, but to accommodate
+    /// existing front ends, Naga also permits `Access`, with a restricted
+    /// `index`.)
     ///
     /// [`Vector`]: TypeInner::Vector
     /// [`Matrix`]: TypeInner::Matrix
diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs
index a5b3ea4e3..abbe0c7e4 100644
--- a/naga/src/proc/mod.rs
+++ b/naga/src/proc/mod.rs
@@ -521,12 +521,12 @@ impl crate::Expression {
         }
     }
 
-    /// Return true if this expression is a dynamic array index, for [`Access`].
+    /// Return true if this expression is a dynamic array/vector/matrix index,
+    /// for [`Access`].
     ///
     /// This method returns true if this expression is a dynamically computed
-    /// index, and as such can only be used to index matrices and arrays when
-    /// they appear behind a pointer. See the documentation for [`Access`] for
-    /// details.
+    /// index, and as such can only be used to index matrices when they appear
+    /// behind a pointer. See the documentation for [`Access`] for details.
     ///
     /// Note, this does not check the _type_ of the given expression. It's up to
     /// the caller to establish that the `Access` expression is well-typed
diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs
index d8af0cd23..f29ff40f0 100644
--- a/naga/src/proc/typifier.rs
+++ b/naga/src/proc/typifier.rs
@@ -92,6 +92,13 @@ pub enum TypeResolution {
     /// available in the associated arena. However, the `TypeInner` itself may
     /// contain `Handle<Type>` values referring to types from the arena.
     ///
+    /// The inner type must only be one of the following variants:
+    /// - TypeInner::Pointer
+    /// - TypeInner::ValuePointer
+    /// - TypeInner::Matrix (generated by matrix multiplication)
+    /// - TypeInner::Vector
+    /// - TypeInner::Scalar
+    ///
     /// [`TypeInner`]: crate::TypeInner
     Value(crate::TypeInner),
 }
diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs
index 0b0d115c5..2b479d3a7 100644
--- a/naga/src/valid/expression.rs
+++ b/naga/src/valid/expression.rs
@@ -240,9 +240,10 @@ impl super::Validator {
                 let base_type = &resolver[base];
                 // See the documentation for `Expression::Access`.
                 let dynamic_indexing_restricted = match *base_type {
-                    Ti::Vector { .. } => false,
-                    Ti::Matrix { .. } | Ti::Array { .. } => true,
-                    Ti::Pointer { .. }
+                    Ti::Matrix { .. } => true,
+                    Ti::Vector { .. }
+                    | Ti::Array { .. }
+                    | Ti::Pointer { .. }
                     | Ti::ValuePointer { size: Some(_), .. }
                     | Ti::BindingArray { .. } => false,
                     ref other => {
diff --git a/wgpu-core/src/binding_model.rs b/wgpu-core/src/binding_model.rs
index 2357a8c77..6f43ac523 100644
--- a/wgpu-core/src/binding_model.rs
+++ b/wgpu-core/src/binding_model.rs
@@ -373,10 +373,6 @@ impl BindingTypeMaxCountValidator {
             limits.max_sampled_textures_per_shader_stage,
             BindingTypeMaxCountErrorKind::SampledTextures,
         )?;
-        self.storage_buffers.validate(
-            limits.max_storage_buffers_per_shader_stage,
-            BindingTypeMaxCountErrorKind::StorageBuffers,
-        )?;
         self.samplers.validate(
             limits.max_samplers_per_shader_stage,
             BindingTypeMaxCountErrorKind::Samplers,
diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs
index 5de03917d..133a5af35 100644
--- a/wgpu-core/src/command/compute.rs
+++ b/wgpu-core/src/command/compute.rs
@@ -136,6 +136,8 @@ pub enum ComputePassErrorInner {
     BindGroupIndexOutOfRange { index: u32, max: u32 },
     #[error(transparent)]
     DestroyedResource(#[from] DestroyedResourceError),
+    #[error("Indirect buffer offset {0:?} is not a multiple of 4")]
+    UnalignedIndirectBufferOffset(BufferAddress),
     #[error("Indirect buffer uses bytes {offset}..{end_offset} which overruns indirect buffer of size {buffer_size}")]
     IndirectBufferOverrun {
         offset: u64,
@@ -495,8 +497,6 @@ impl Global {
             state.raw_encoder.begin_compute_pass(&hal_desc);
         }
 
-        // TODO: We should be draining the commands here, avoiding extra copies in the process.
-        //       (A command encoder can't be executed twice!)
         for command in base.commands {
             match command {
                 ArcComputeCommand::SetBindGroup {
@@ -780,7 +780,7 @@ fn set_push_constant(
         .binder
         .pipeline_layout
         .as_ref()
-        //TODO: don't error here, lazily update the push constants
+        // TODO: don't error here, lazily update the push constants using `state.push_constants`
         .ok_or(ComputePassErrorInner::Dispatch(
             DispatchError::MissingPipeline,
         ))?;
@@ -847,6 +847,10 @@ fn dispatch_indirect(
         .merge_single(&buffer, hal::BufferUses::INDIRECT)?;
     buffer.check_usage(wgt::BufferUsages::INDIRECT)?;
 
+    if offset % 4 != 0 {
+        return Err(ComputePassErrorInner::UnalignedIndirectBufferOffset(offset));
+    }
+
     let end_offset = offset + size_of::<wgt::DispatchIndirectArgs>() as u64;
     if end_offset > buffer.size {
         return Err(ComputePassErrorInner::IndirectBufferOverrun {
diff --git a/wgpu-core/src/command/render.rs b/wgpu-core/src/command/render.rs
index b6680333c..d22eb5f0d 100644
--- a/wgpu-core/src/command/render.rs
+++ b/wgpu-core/src/command/render.rs
@@ -638,6 +638,8 @@ pub enum RenderPassErrorInner {
     MissingFeatures(#[from] MissingFeatures),
     #[error(transparent)]
     MissingDownlevelFlags(#[from] MissingDownlevelFlags),
+    #[error("Indirect buffer offset {0:?} is not a multiple of 4")]
+    UnalignedIndirectBufferOffset(BufferAddress),
     #[error("Indirect draw uses bytes {offset}..{end_offset} {} which overruns indirect buffer of size {buffer_size}",
         count.map_or_else(String::new, |v| format!("(using count {v})")))]
     IndirectBufferOverrun {
@@ -2450,6 +2452,10 @@ fn multi_draw_indirect(
 
     let actual_count = count.map_or(1, |c| c.get());
 
+    if offset % 4 != 0 {
+        return Err(RenderPassErrorInner::UnalignedIndirectBufferOffset(offset));
+    }
+
     let end_offset = offset + stride as u64 * actual_count as u64;
     if end_offset > indirect_buffer.size {
         return Err(RenderPassErrorInner::IndirectBufferOverrun {
@@ -2534,6 +2540,10 @@ fn multi_draw_indirect_count(
     count_buffer.check_usage(BufferUsages::INDIRECT)?;
     let count_raw = count_buffer.try_raw(state.snatch_guard)?;
 
+    if offset % 4 != 0 {
+        return Err(RenderPassErrorInner::UnalignedIndirectBufferOffset(offset));
+    }
+
     let end_offset = offset + stride * max_count as u64;
     if end_offset > indirect_buffer.size {
         return Err(RenderPassErrorInner::IndirectBufferOverrun {
diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs
index 03b183e08..5a726ef70 100644
--- a/wgpu-core/src/device/resource.rs
+++ b/wgpu-core/src/device/resource.rs
@@ -38,7 +38,6 @@ use arrayvec::ArrayVec;
 use once_cell::sync::OnceCell;
 
 use smallvec::SmallVec;
-use thiserror::Error;
 use wgt::{
     math::align_to, DeviceLostReason, TextureFormat, TextureSampleType, TextureViewDimension,
 };
@@ -187,14 +186,6 @@ impl Drop for Device {
     }
 }
 
-#[derive(Clone, Debug, Error)]
-pub enum CreateDeviceError {
-    #[error("Not enough memory left to create device")]
-    OutOfMemory,
-    #[error("Failed to create internal buffer for initializing textures")]
-    FailedToCreateZeroBuffer(#[from] DeviceError),
-}
-
 impl Device {
     pub(crate) fn raw(&self) -> &dyn hal::DynDevice {
         self.raw.as_ref()
@@ -372,7 +363,7 @@ impl Device {
                     let Some(view) = view.upgrade() else {
                         continue;
                     };
-                    let Some(raw_view) = view.raw.snatch(self.snatchable_lock.write()) else {
+                    let Some(raw_view) = view.raw.snatch(&mut self.snatchable_lock.write()) else {
                         continue;
                     };
 
@@ -386,7 +377,8 @@ impl Device {
                     let Some(bind_group) = bind_group.upgrade() else {
                         continue;
                     };
-                    let Some(raw_bind_group) = bind_group.raw.snatch(self.snatchable_lock.write())
+                    let Some(raw_bind_group) =
+                        bind_group.raw.snatch(&mut self.snatchable_lock.write())
                     else {
                         continue;
                     };
@@ -437,13 +429,11 @@ impl Device {
                     .last_successful_submission_index
                     .load(Ordering::Acquire);
 
-                if let wgt::Maintain::WaitForSubmissionIndex(submission_index) = maintain {
-                    if submission_index > last_successful_submission_index {
-                        return Err(WaitIdleError::WrongSubmissionIndex(
-                            submission_index,
-                            last_successful_submission_index,
-                        ));
-                    }
+                if submission_index > last_successful_submission_index {
+                    return Err(WaitIdleError::WrongSubmissionIndex(
+                        submission_index,
+                        last_successful_submission_index,
+                    ));
                 }
 
                 submission_index
diff --git a/wgpu-core/src/present.rs b/wgpu-core/src/present.rs
index c9d0124bf..b12e35328 100644
--- a/wgpu-core/src/present.rs
+++ b/wgpu-core/src/present.rs
@@ -288,8 +288,11 @@ impl Global {
                     .textures
                     .remove(texture.tracker_index());
                 let suf = surface.raw(device.backend()).unwrap();
-                let exclusive_snatch_guard = device.snatchable_lock.write();
-                match texture.inner.snatch(exclusive_snatch_guard).unwrap() {
+                match texture
+                    .inner
+                    .snatch(&mut device.snatchable_lock.write())
+                    .unwrap()
+                {
                     resource::TextureInner::Surface { raw, parent_id } => {
                         if surface_id != parent_id {
                             log::error!("Presented frame is from a different surface");
@@ -359,8 +362,11 @@ impl Global {
                     .textures
                     .remove(texture.tracker_index());
                 let suf = surface.raw(device.backend());
-                let exclusive_snatch_guard = device.snatchable_lock.write();
-                match texture.inner.snatch(exclusive_snatch_guard).unwrap() {
+                match texture
+                    .inner
+                    .snatch(&mut device.snatchable_lock.write())
+                    .unwrap()
+                {
                     resource::TextureInner::Surface { raw, parent_id } => {
                         if surface_id == parent_id {
                             unsafe { suf.unwrap().discard_texture(raw) };
diff --git a/wgpu-core/src/resource.rs b/wgpu-core/src/resource.rs
index 5df285da5..32fde3dd5 100644
--- a/wgpu-core/src/resource.rs
+++ b/wgpu-core/src/resource.rs
@@ -737,8 +737,7 @@ impl Buffer {
         let device = &self.device;
 
         let temp = {
-            let snatch_guard = device.snatchable_lock.write();
-            let raw = match self.raw.snatch(snatch_guard) {
+            let raw = match self.raw.snatch(&mut device.snatchable_lock.write()) {
                 Some(raw) => raw,
                 None => {
                     return Err(DestroyError::AlreadyDestroyed);
@@ -1185,8 +1184,7 @@ impl Texture {
         let device = &self.device;
 
         let temp = {
-            let snatch_guard = device.snatchable_lock.write();
-            let raw = match self.inner.snatch(snatch_guard) {
+            let raw = match self.inner.snatch(&mut device.snatchable_lock.write()) {
                 Some(TextureInner::Native { raw }) => raw,
                 Some(TextureInner::Surface { .. }) => {
                     return Ok(());
diff --git a/wgpu-core/src/snatch.rs b/wgpu-core/src/snatch.rs
index 9866b7772..a817e2068 100644
--- a/wgpu-core/src/snatch.rs
+++ b/wgpu-core/src/snatch.rs
@@ -38,7 +38,7 @@ impl<T> Snatchable<T> {
     }
 
     /// Take the value. Requires a the snatchable lock's write guard.
-    pub fn snatch(&self, _guard: ExclusiveSnatchGuard) -> Option<T> {
+    pub fn snatch(&self, _guard: &mut ExclusiveSnatchGuard) -> Option<T> {
         unsafe { (*self.value.get()).take() }
     }
 


Test patch:
diff --git a/naga/tests/in/access.wgsl b/naga/tests/in/access.wgsl
index 956a694aa..3336522fd 100644
--- a/naga/tests/in/access.wgsl
+++ b/naga/tests/in/access.wgsl
@@ -167,3 +167,14 @@ fn assign_through_ptr() {
 	var arr = array<vec4<f32>, 2>(vec4(6.0), vec4(7.0));
     assign_array_through_ptr_fn(&arr);
 }
+
+@vertex
+fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
+	let arr = array<i32, 5>(1, 2, 3, 4, 5);
+	let value = arr[vi];
+	return vec4<f32>(vec4<i32>(value));
+}
+
+fn array_by_value(a: array<i32, 5>, i: i32) -> i32 {
+    return a[i];
+}
diff --git a/naga/tests/out/analysis/access.info.ron b/naga/tests/out/analysis/access.info.ron
index 308bb1a8b..830831cb1 100644
--- a/naga/tests/out/analysis/access.info.ron
+++ b/naga/tests/out/analysis/access.info.ron
@@ -2735,6 +2735,54 @@
             sampling: [],
             dual_source_blending: false,
         ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            uniformity: (
+                non_uniform_result: Some(0),
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                (""),
+                (""),
+                (""),
+                (""),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(25),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(1),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+        ),
     ],
     entry_points: [
         (
@@ -3981,6 +4029,144 @@
             sampling: [],
             dual_source_blending: false,
         ),
+        (
+            flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"),
+            available_stages: ("VERTEX | FRAGMENT | COMPUTE"),
+            uniformity: (
+                non_uniform_result: Some(0),
+                requirements: (""),
+            ),
+            may_kill: false,
+            sampling_set: [],
+            global_uses: [
+                (""),
+                (""),
+                (""),
+                (""),
+                (""),
+            ],
+            expressions: [
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(0),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Scalar((
+                        kind: Sint,
+                        width: 4,
+                    ))),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: None,
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(25),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Handle(2),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Vector(
+                        size: Quad,
+                        scalar: (
+                            kind: Sint,
+                            width: 4,
+                        ),
+                    )),
+                ),
+                (
+                    uniformity: (
+                        non_uniform_result: Some(0),
+                        requirements: (""),
+                    ),
+                    ref_count: 1,
+                    assignable_global: None,
+                    ty: Value(Vector(
+                        size: Quad,
+                        scalar: (
+                            kind: Float,
+                            width: 4,
+                        ),
+                    )),
+                ),
+            ],
+            sampling: [],
+            dual_source_blending: false,
+        ),
     ],
     const_expression_types: [
         Value(Scalar((
diff --git a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
index 2e51bbde6..4a4791c99 100644
--- a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
+++ b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
@@ -20,8 +20,8 @@ struct MatCx2InArray {
     mat4x2 am[2];
 };
 
-float read_from_private(inout float foo_1) {
-    float _e1 = foo_1;
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -34,11 +34,15 @@ void assign_through_ptr_fn(inout uint p) {
     return;
 }
 
-void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
-    foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
     return;
 }
 
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
 void main() {
     uint val = 33u;
     vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));
diff --git a/naga/tests/out/glsl/access.foo.Vertex.glsl b/naga/tests/out/glsl/access.foo.Vertex.glsl
new file mode 100644
index 000000000..e1f313840
--- /dev/null
+++ b/naga/tests/out/glsl/access.foo.Vertex.glsl
@@ -0,0 +1,52 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+struct GlobalConst {
+    uint a;
+    uvec3 b;
+    int c;
+};
+struct AlignedWrapper {
+    int value;
+};
+struct Baz {
+    mat3x2 m;
+};
+struct MatCx2InArray {
+    mat4x2 am[2];
+};
+
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
+    return _e1;
+}
+
+float test_arr_as_arg(float a[5][10]) {
+    return a[4][9];
+}
+
+void assign_through_ptr_fn(inout uint p) {
+    p = 42u;
+    return;
+}
+
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
+    return;
+}
+
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
+void main() {
+    uint vi_1 = uint(gl_VertexID);
+    int arr_1[5] = int[5](1, 2, 3, 4, 5);
+    int value = arr_1[vi_1];
+    gl_Position = vec4(ivec4(value));
+    gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
+    return;
+}
+
diff --git a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
index aacdda013..eca6bc54c 100644
--- a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
+++ b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
@@ -30,8 +30,8 @@ layout(std430) buffer type_13_block_1Fragment { ivec2 _group_0_binding_2_fs; };
 
 layout(location = 0) out vec4 _fs2p_location0;
 
-float read_from_private(inout float foo_1) {
-    float _e1 = foo_1;
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -44,11 +44,15 @@ void assign_through_ptr_fn(inout uint p) {
     return;
 }
 
-void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
-    foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
     return;
 }
 
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
 void main() {
     _group_0_binding_0_fs._matrix[1][2] = 1.0;
     _group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0));
diff --git a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
index d4a9b9294..a926eadf7 100644
--- a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
+++ b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
@@ -103,8 +103,8 @@ void test_matrix_within_array_within_struct_accesses() {
     return;
 }
 
-float read_from_private(inout float foo_1) {
-    float _e1 = foo_1;
+float read_from_private(inout float foo_2) {
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -117,11 +117,15 @@ void assign_through_ptr_fn(inout uint p) {
     return;
 }
 
-void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
-    foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+void assign_array_through_ptr_fn(inout vec4 foo_3[2]) {
+    foo_3 = vec4[2](vec4(1.0), vec4(2.0));
     return;
 }
 
+int array_by_value(int a_1[5], int i) {
+    return a_1[i];
+}
+
 void main() {
     uint vi = uint(gl_VertexID);
     float foo = 0.0;
@@ -133,10 +137,10 @@ void main() {
     mat4x3 _matrix = _group_0_binding_0_vs._matrix;
     uvec2 arr_1[2] = _group_0_binding_0_vs.arr;
     float b = _group_0_binding_0_vs._matrix[3u][0];
-    int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
+    int a_2 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
     ivec2 c = _group_0_binding_2_vs;
     float _e33 = read_from_private(foo);
-    c2_ = int[5](a_1, int(b), 3, 4, 5);
+    c2_ = int[5](a_2, int(b), 3, 4, 5);
     c2_[(vi + 1u)] = 42;
     int value = c2_[vi];
     float _e47 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl
index 142083be6..543b84196 100644
--- a/naga/tests/out/hlsl/access.hlsl
+++ b/naga/tests/out/hlsl/access.hlsl
@@ -201,9 +201,9 @@ void test_matrix_within_array_within_struct_accesses()
     return;
 }
 
-float read_from_private(inout float foo_1)
+float read_from_private(inout float foo_2)
 {
-    float _e1 = foo_1;
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -224,12 +224,17 @@ ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) {
     return ret;
 }
 
-void assign_array_through_ptr_fn(inout float4 foo_2[2])
+void assign_array_through_ptr_fn(inout float4 foo_3[2])
 {
-    foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
+    foo_3 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
     return;
 }
 
+int array_by_value(int a_1[5], int i)
+{
+    return a_1[i];
+}
+
 typedef int ret_Constructarray5_int_[5];
 ret_Constructarray5_int_ Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4) {
     int ret[5] = { arg0, arg1, arg2, arg3, arg4 };
@@ -266,10 +271,10 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
     float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48)));
     uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8)));
     float b = asfloat(bar.Load(0+3u*16+0));
-    int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
+    int a_2 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
     int2 c = asint(qux.Load2(0));
     const float _e33 = read_from_private(foo);
-    c2_ = Constructarray5_int_(a_1, int(b), 3, 4, 5);
+    c2_ = Constructarray5_int_(a_2, int(b), 3, 4, 5);
     c2_[(vi + 1u)] = 42;
     int value = c2_[vi];
     const float _e47 = test_arr_as_arg(ZeroValuearray5_array10_float__());
@@ -310,3 +315,10 @@ void assign_through_ptr()
     assign_array_through_ptr_fn(arr);
     return;
 }
+
+float4 foo_1(uint vi_1 : SV_VertexID) : SV_Position
+{
+    int arr_2[5] = Constructarray5_int_(1, 2, 3, 4, 5);
+    int value_1 = arr_2[vi_1];
+    return float4((value_1).xxxx);
+}
diff --git a/naga/tests/out/hlsl/access.ron b/naga/tests/out/hlsl/access.ron
index 73c9e4444..8960a612e 100644
--- a/naga/tests/out/hlsl/access.ron
+++ b/naga/tests/out/hlsl/access.ron
@@ -4,6 +4,10 @@
             entry_point:"foo_vert",
             target_profile:"vs_5_1",
         ),
+        (
+            entry_point:"foo_1",
+            target_profile:"vs_5_1",
+        ),
     ],
     fragment:[
         (
diff --git a/naga/tests/out/ir/access.compact.ron b/naga/tests/out/ir/access.compact.ron
index 1b95742ff..2d066b8ff 100644
--- a/naga/tests/out/ir/access.compact.ron
+++ b/naga/tests/out/ir/access.compact.ron
@@ -1655,6 +1655,47 @@
                 ),
             ],
         ),
+        (
+            name: Some("array_by_value"),
+            arguments: [
+                (
+                    name: Some("a"),
+                    ty: 25,
+                    binding: None,
+                ),
+                (
+                    name: Some("i"),
+                    ty: 2,
+                    binding: None,
+                ),
+            ],
+            result: Some((
+                ty: 2,
+                binding: None,
+            )),
+            local_variables: [],
+            expressions: [
+                FunctionArgument(0),
+                FunctionArgument(1),
+                Access(
+                    base: 0,
+                    index: 1,
+                ),
+            ],
+            named_expressions: {
+                0: "a",
+                1: "i",
+            },
+            body: [
+                Emit((
+                    start: 2,
+                    end: 3,
+                )),
+                Return(
+                    value: Some(2),
+                ),
+            ],
+        ),
     ],
     entry_points: [
         (
@@ -2230,5 +2271,81 @@
                 ],
             ),
         ),
+        (
+            name: "foo",
+            stage: Vertex,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            function: (
+                name: Some("foo"),
+                arguments: [
+                    (
+                        name: Some("vi"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexIndex)),
+                    ),
+                ],
+                result: Some((
+                    ty: 24,
+                    binding: Some(BuiltIn(Position(
+                        invariant: false,
+                    ))),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    Literal(I32(3)),
+                    Literal(I32(4)),
+                    Literal(I32(5)),
+                    Compose(
+                        ty: 25,
+                        components: [
+                            1,
+                            2,
+                            3,
+                            4,
+                            5,
+                        ],
+                    ),
+                    Access(
+                        base: 6,
+                        index: 0,
+                    ),
+                    Splat(
+                        size: Quad,
+                        value: 7,
+                    ),
+                    As(
+                        expr: 8,
+                        kind: Float,
+                        convert: Some(4),
+                    ),
+                ],
+                named_expressions: {
+                    0: "vi",
+                    6: "arr",
+                    7: "value",
+                },
+                body: [
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Emit((
+                        start: 7,
+                        end: 8,
+                    )),
+                    Emit((
+                        start: 8,
+                        end: 10,
+                    )),
+                    Return(
+                        value: Some(9),
+                    ),
+                ],
+            ),
+        ),
     ],
 )
\ No newline at end of file
diff --git a/naga/tests/out/ir/access.ron b/naga/tests/out/ir/access.ron
index 1b95742ff..2d066b8ff 100644
--- a/naga/tests/out/ir/access.ron
+++ b/naga/tests/out/ir/access.ron
@@ -1655,6 +1655,47 @@
                 ),
             ],
         ),
+        (
+            name: Some("array_by_value"),
+            arguments: [
+                (
+                    name: Some("a"),
+                    ty: 25,
+                    binding: None,
+                ),
+                (
+                    name: Some("i"),
+                    ty: 2,
+                    binding: None,
+                ),
+            ],
+            result: Some((
+                ty: 2,
+                binding: None,
+            )),
+            local_variables: [],
+            expressions: [
+                FunctionArgument(0),
+                FunctionArgument(1),
+                Access(
+                    base: 0,
+                    index: 1,
+                ),
+            ],
+            named_expressions: {
+                0: "a",
+                1: "i",
+            },
+            body: [
+                Emit((
+                    start: 2,
+                    end: 3,
+                )),
+                Return(
+                    value: Some(2),
+                ),
+            ],
+        ),
     ],
     entry_points: [
         (
@@ -2230,5 +2271,81 @@
                 ],
             ),
         ),
+        (
+            name: "foo",
+            stage: Vertex,
+            early_depth_test: None,
+            workgroup_size: (0, 0, 0),
+            function: (
+                name: Some("foo"),
+                arguments: [
+                    (
+                        name: Some("vi"),
+                        ty: 0,
+                        binding: Some(BuiltIn(VertexIndex)),
+                    ),
+                ],
+                result: Some((
+                    ty: 24,
+                    binding: Some(BuiltIn(Position(
+                        invariant: false,
+                    ))),
+                )),
+                local_variables: [],
+                expressions: [
+                    FunctionArgument(0),
+                    Literal(I32(1)),
+                    Literal(I32(2)),
+                    Literal(I32(3)),
+                    Literal(I32(4)),
+                    Literal(I32(5)),
+                    Compose(
+                        ty: 25,
+                        components: [
+                            1,
+                            2,
+                            3,
+                            4,
+                            5,
+                        ],
+                    ),
+                    Access(
+                        base: 6,
+                        index: 0,
+                    ),
+                    Splat(
+                        size: Quad,
+                        value: 7,
+                    ),
+                    As(
+                        expr: 8,
+                        kind: Float,
+                        convert: Some(4),
+                    ),
+                ],
+                named_expressions: {
+                    0: "vi",
+                    6: "arr",
+                    7: "value",
+                },
+                body: [
+                    Emit((
+                        start: 6,
+                        end: 7,
+                    )),
+                    Emit((
+                        start: 7,
+                        end: 8,
+                    )),
+                    Emit((
+                        start: 8,
+                        end: 10,
+                    )),
+                    Return(
+                        value: Some(9),
+                    ),
+                ],
+            ),
+        ),
     ],
 )
\ No newline at end of file
diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl
index 65dba4910..924b604e4 100644
--- a/naga/tests/out/msl/access.msl
+++ b/naga/tests/out/msl/access.msl
@@ -133,9 +133,9 @@ void test_matrix_within_array_within_struct_accesses(
 }
 
 float read_from_private(
-    thread float& foo_1
+    thread float& foo_2
 ) {
-    float _e1 = foo_1;
+    float _e1 = foo_2;
     return _e1;
 }
 
@@ -153,12 +153,19 @@ void assign_through_ptr_fn(
 }
 
 void assign_array_through_ptr_fn(
-    thread type_22& foo_2
+    thread type_22& foo_3
 ) {
-    foo_2 = type_22 {metal::float4(1.0), metal::float4(2.0)};
+    foo_3 = type_22 {metal::float4(1.0), metal::float4(2.0)};
     return;
 }
 
+int array_by_value(
+    type_20 a_1,
+    int i
+) {
+    return a_1.inner[i];
+}
+
 struct foo_vertInput {
 };
 struct foo_vertOutput {
@@ -181,10 +188,10 @@ vertex foo_vertOutput foo_vert(
     metal::float4x3 _matrix = bar._matrix;
     type_10 arr_1 = bar.arr;
     float b = bar._matrix[3u].x;
-    int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
+    int a_2 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
     metal::int2 c = qux;
     float _e33 = read_from_private(foo);
-    c2_ = type_20 {a_1, static_cast<int>(b), 3, 4, 5};
+    c2_ = type_20 {a_2, static_cast<int>(b), 3, 4, 5};
     c2_.inner[vi + 1u] = 42;
     int value = c2_.inner[vi];
     float _e47 = test_arr_as_arg(type_18 {});
@@ -217,3 +224,17 @@ kernel void assign_through_ptr(
     assign_array_through_ptr_fn(arr);
     return;
 }
+
+
+struct foo_1Input {
+};
+struct foo_1Output {
+    metal::float4 member_3 [[position]];
+};
+vertex foo_1Output foo_1(
+  uint vi_1 [[vertex_id]]
+) {
+    type_20 arr_2 = type_20 {1, 2, 3, 4, 5};
+    int value_1 = arr_2.inner[vi_1];
+    return foo_1Output { static_cast<metal::float4>(metal::int4(value_1)) };
+}
diff --git a/naga/tests/out/spv/access.spvasm b/naga/tests/out/spv/access.spvasm
index ab0112870..6cb87a3b1 100644
--- a/naga/tests/out/spv/access.spvasm
+++ b/naga/tests/out/spv/access.spvasm
@@ -1,16 +1,17 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 301
+; Bound: 323
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Vertex %219 "foo_vert" %214 %217
-OpEntryPoint Fragment %273 "foo_frag" %272
-OpEntryPoint GLCompute %291 "assign_through_ptr"
-OpExecutionMode %273 OriginUpperLeft
-OpExecutionMode %291 LocalSize 1 1 1
+OpEntryPoint Vertex %229 "foo_vert" %224 %227
+OpEntryPoint Fragment %282 "foo_frag" %281
+OpEntryPoint GLCompute %300 "assign_through_ptr"
+OpEntryPoint Vertex %314 "foo" %311 %313
+OpExecutionMode %282 OriginUpperLeft
+OpExecutionMode %300 LocalSize 1 1 1
 OpMemberName %6 0 "a"
 OpMemberName %6 1 "b"
 OpMemberName %6 2 "c"
@@ -47,14 +48,19 @@ OpName %200 "p"
 OpName %201 "assign_through_ptr_fn"
 OpName %206 "foo"
 OpName %207 "assign_array_through_ptr_fn"
-OpName %214 "vi"
-OpName %219 "foo_vert"
-OpName %231 "foo"
-OpName %232 "c2"
-OpName %273 "foo_frag"
-OpName %291 "assign_through_ptr"
-OpName %296 "val"
-OpName %297 "arr"
+OpName %214 "a"
+OpName %215 "i"
+OpName %216 "array_by_value"
+OpName %224 "vi"
+OpName %229 "foo_vert"
+OpName %241 "foo"
+OpName %242 "c2"
+OpName %282 "foo_frag"
+OpName %300 "assign_through_ptr"
+OpName %305 "val"
+OpName %306 "arr"
+OpName %311 "vi"
+OpName %314 "foo"
 OpMemberDecorate %6 0 Offset 0
 OpMemberDecorate %6 1 Offset 16
 OpMemberDecorate %6 2 Offset 28
@@ -99,9 +105,11 @@ OpDecorate %50 DescriptorSet 0
 OpDecorate %50 Binding 3
 OpDecorate %51 Block
 OpMemberDecorate %51 0 Offset 0
-OpDecorate %214 BuiltIn VertexIndex
-OpDecorate %217 BuiltIn Position
-OpDecorate %272 Location 0
+OpDecorate %224 BuiltIn VertexIndex
+OpDecorate %227 BuiltIn Position
+OpDecorate %281 Location 0
+OpDecorate %311 BuiltIn VertexIndex
+OpDecorate %313 BuiltIn Position
 %2 = OpTypeVoid
 %3 = OpTypeInt 32 0
 %4 = OpTypeVector %3 3
@@ -209,44 +217,49 @@ OpDecorate %272 Location 0
 %209 = OpConstantComposite  %31  %59 %59 %59 %59
 %210 = OpConstantComposite  %31  %61 %61 %61 %61
 %211 = OpConstantComposite  %34  %209 %210
-%215 = OpTypePointer Input %3
-%214 = OpVariable  %215  Input
-%218 = OpTypePointer Output %31
-%217 = OpVariable  %218  Output
-%221 = OpTypePointer StorageBuffer %23
-%224 = OpConstant  %8  0.0
-%225 = OpConstant  %3  3
-%226 = OpConstant  %5  3
-%227 = OpConstant  %5  4
-%228 = OpConstant  %5  5
-%229 = OpConstant  %5  42
-%230 = OpConstantNull  %29
-%233 = OpTypePointer Function %32
-%234 = OpConstantNull  %32
-%239 = OpTypePointer StorageBuffer %9
-%242 = OpTypePointer StorageBuffer %18
-%243 = OpConstant  %3  4
-%246 = OpTypePointer StorageBuffer %10
-%247 = OpTypePointer StorageBuffer %8
-%250 = OpTypePointer StorageBuffer %19
-%253 = OpTypePointer StorageBuffer %7
-%254 = OpTypePointer StorageBuffer %5
-%266 = OpTypeVector %5 4
-%272 = OpVariable  %218  Output
-%275 = OpConstantComposite  %10  %224 %224 %224
-%276 = OpConstantComposite  %10  %59 %59 %59
-%277 = OpConstantComposite  %10  %61 %61 %61
-%278 = OpConstantComposite  %10  %63 %63 %63
-%279 = OpConstantComposite  %9  %275 %276 %277 %278
-%280 = OpConstantComposite  %17  %36 %36
-%281 = OpConstantComposite  %17  %99 %99
-%282 = OpConstantComposite  %18  %280 %281
-%283 = OpConstantNull  %23
-%284 = OpConstantComposite  %31  %224 %224 %224 %224
-%292 = OpConstant  %3  33
-%293 = OpConstantComposite  %31  %67 %67 %67 %67
-%294 = OpConstantComposite  %31  %137 %137 %137 %137
-%295 = OpConstantComposite  %34  %293 %294
+%217 = OpTypeFunction %5 %32 %5
+%219 = OpTypePointer Function %32
+%225 = OpTypePointer Input %3
+%224 = OpVariable  %225  Input
+%228 = OpTypePointer Output %31
+%227 = OpVariable  %228  Output
+%231 = OpTypePointer StorageBuffer %23
+%234 = OpConstant  %8  0.0
+%235 = OpConstant  %3  3
+%236 = OpConstant  %5  3
+%237 = OpConstant  %5  4
+%238 = OpConstant  %5  5
+%239 = OpConstant  %5  42
+%240 = OpConstantNull  %29
+%243 = OpConstantNull  %32
+%248 = OpTypePointer StorageBuffer %9
+%251 = OpTypePointer StorageBuffer %18
+%252 = OpConstant  %3  4
+%255 = OpTypePointer StorageBuffer %10
+%256 = OpTypePointer StorageBuffer %8
+%259 = OpTypePointer StorageBuffer %19
+%262 = OpTypePointer StorageBuffer %7
+%263 = OpTypePointer StorageBuffer %5
+%275 = OpTypeVector %5 4
+%281 = OpVariable  %228  Output
+%284 = OpConstantComposite  %10  %234 %234 %234
+%285 = OpConstantComposite  %10  %59 %59 %59
+%286 = OpConstantComposite  %10  %61 %61 %61
+%287 = OpConstantComposite  %10  %63 %63 %63
+%288 = OpConstantComposite  %9  %284 %285 %286 %287
+%289 = OpConstantComposite  %17  %36 %36
+%290 = OpConstantComposite  %17  %99 %99
+%291 = OpConstantComposite  %18  %289 %290
+%292 = OpConstantNull  %23
+%293 = OpConstantComposite  %31  %234 %234 %234 %234
+%301 = OpConstant  %3  33
+%302 = OpConstantComposite  %31  %67 %67 %67 %67
+%303 = OpConstantComposite  %31  %137 %137 %137 %137
+%304 = OpConstantComposite  %34  %302 %303
+%311 = OpVariable  %225  Input
+%313 = OpVariable  %228  Output
+%315 = OpConstant  %5  2
+%316 = OpConstantComposite  %32  %58 %315 %236 %237 %238
 %54 = OpFunction  %2  None %55
 %53 = OpLabel
 %82 = OpVariable  %83  Function %58
@@ -390,72 +403,98 @@ OpBranch %212
 OpStore %206 %211
 OpReturn
 OpFunctionEnd
-%219 = OpFunction  %2  None %55
+%216 = OpFunction  %5  None %217
+%214 = OpFunctionParameter  %32
+%215 = OpFunctionParameter  %5
 %213 = OpLabel
-%231 = OpVariable  %27  Function %224
-%232 = OpVariable  %233  Function %234
-%216 = OpLoad  %3  %214
-%220 = OpAccessChain  %56  %44 %36
-%222 = OpAccessChain  %221  %47 %36
-%223 = OpAccessChain  %131  %50 %36
-OpBranch %235
-%235 = OpLabel
-%236 = OpLoad  %8  %231
-OpStore %231 %59
-%237 = OpFunctionCall  %2  %54
-%238 = OpFunctionCall  %2  %130
-%240 = OpAccessChain  %239  %42 %36
-%241 = OpLoad  %9  %240
-%244 = OpAccessChain  %242  %42 %243
-%245 = OpLoad  %18  %244
-%248 = OpAccessChain  %247  %42 %36 %225 %36
-%249 = OpLoad  %8  %248
-%251 = OpArrayLength  %3  %42 5
-%252 = OpISub  %3  %251 %14
-%255 = OpAccessChain  %254  %42 %30 %252 %36
-%256 = OpLoad  %5  %255
-%257 = OpLoad  %23  %222
-%258 = OpFunctionCall  %8  %188 %231
-%259 = OpConvertFToS  %5  %249
-%260 = OpCompositeConstruct  %32  %256 %259 %226 %227 %228
-OpStore %232 %260
-%261 = OpIAdd  %3  %216 %99
-%262 = OpAccessChain  %83  %232 %261
-OpStore %262 %229
-%263 = OpAccessChain  %83  %232 %216
-%264 = OpLoad  %5  %263
-%265 = OpFunctionCall  %8  %194 %230
-%267 = OpCompositeConstruct  %266  %264 %264 %264 %264
-%268 = OpConvertSToF  %31  %267
-%269 = OpMatrixTimesVector  %10  %241 %268
-%270 = OpCompositeConstruct  %31  %269 %61
-OpStore %217 %270
+%220 = OpVariable  %219  Function
+OpBranch %218
+%218 = OpLabel
+OpStore %220 %214
+%221 = OpAccessChain  %83  %220 %215
+%222 = OpLoad  %5  %221
+OpReturnValue %222
+OpFunctionEnd
+%229 = OpFunction  %2  None %55
+%223 = OpLabel
+%241 = OpVariable  %27  Function %234
+%242 = OpVariable  %219  Function %243
+%226 = OpLoad  %3  %224
+%230 = OpAccessChain  %56  %44 %36
+%232 = OpAccessChain  %231  %47 %36
+%233 = OpAccessChain  %131  %50 %36
+OpBranch %244
+%244 = OpLabel
+%245 = OpLoad  %8  %241
+OpStore %241 %59
+%246 = OpFunctionCall  %2  %54
+%247 = OpFunctionCall  %2  %130
+%249 = OpAccessChain  %248  %42 %36
+%250 = OpLoad  %9  %249
+%253 = OpAccessChain  %251  %42 %252
+%254 = OpLoad  %18  %253
+%257 = OpAccessChain  %256  %42 %36 %235 %36
+%258 = OpLoad  %8  %257
+%260 = OpArrayLength  %3  %42 5
+%261 = OpISub  %3  %260 %14
+%264 = OpAccessChain  %263  %42 %30 %261 %36
+%265 = OpLoad  %5  %264
+%266 = OpLoad  %23  %232
+%267 = OpFunctionCall  %8  %188 %241
+%268 = OpConvertFToS  %5  %258
+%269 = OpCompositeConstruct  %32  %265 %268 %236 %237 %238
+OpStore %242 %269
+%270 = OpIAdd  %3  %226 %99
+%271 = OpAccessChain  %83  %242 %270
+OpStore %271 %239
+%272 = OpAccessChain  %83  %242 %226
+%273 = OpLoad  %5  %272
+%274 = OpFunctionCall  %8  %194 %240
+%276 = OpCompositeConstruct  %275  %273 %273 %273 %273
+%277 = OpConvertSToF  %31  %276
+%278 = OpMatrixTimesVector  %10  %250 %277
+%279 = OpCompositeConstruct  %31  %278 %61
+OpStore %227 %279
+OpReturn
+OpFunctionEnd
+%282 = OpFunction  %2  None %55
+%280 = OpLabel
+%283 = OpAccessChain  %231  %47 %36
+OpBranch %294
+%294 = OpLabel
+%295 = OpAccessChain  %256  %42 %36 %99 %14
+OpStore %295 %59
+%296 = OpAccessChain  %248  %42 %36
+OpStore %296 %288
+%297 = OpAccessChain  %251  %42 %252
+OpStore %297 %291
+%298 = OpAccessChain  %263  %42 %30 %99 %36
+OpStore %298 %58
+OpStore %283 %292
+OpStore %281 %293
 OpReturn
 OpFunctionEnd
-%273 = OpFunction  %2  None %55
-%271 = OpLabel
-%274 = OpAccessChain  %221  %47 %36
-OpBranch %285
-%285 = OpLabel
-%286 = OpAccessChain  %247  %42 %36 %99 %14
-OpStore %286 %59
-%287 = OpAccessChain  %239  %42 %36
-OpStore %287 %279
-%288 = OpAccessChain  %242  %42 %243
-OpStore %288 %282
-%289 = OpAccessChain  %254  %42 %30 %99 %36
-OpStore %289 %58
-OpStore %274 %283
-OpStore %272 %284
+%300 = OpFunction  %2  None %55
+%299 = OpLabel
+%305 = OpVariable  %33  Function %301
+%306 = OpVariable  %35  Function %304
+OpBranch %307
+%307 = OpLabel
+%308 = OpFunctionCall  %2  %201 %305
+%309 = OpFunctionCall  %2  %207 %306
 OpReturn
 OpFunctionEnd
-%291 = OpFunction  %2  None %55
-%290 = OpLabel
-%296 = OpVariable  %33  Function %292
-%297 = OpVariable  %35  Function %295
-OpBranch %298
-%298 = OpLabel
-%299 = OpFunctionCall  %2  %201 %296
-%300 = OpFunctionCall  %2  %207 %297
+%314 = OpFunction  %2  None %55
+%310 = OpLabel
+%318 = OpVariable  %219  Function
+%312 = OpLoad  %3  %311
+OpBranch %317
+%317 = OpLabel
+OpStore %318 %316
+%319 = OpAccessChain  %83  %318 %312
+%320 = OpLoad  %5  %319
+%321 = OpCompositeConstruct  %275  %320 %320 %320 %320
+%322 = OpConvertSToF  %31  %321
+OpStore %313 %322
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/access.wgsl b/naga/tests/out/wgsl/access.wgsl
index 1409e80b1..0c29e5d60 100644
--- a/naga/tests/out/wgsl/access.wgsl
+++ b/naga/tests/out/wgsl/access.wgsl
@@ -107,8 +107,8 @@ fn test_matrix_within_array_within_struct_accesses() {
     return;
 }
 
-fn read_from_private(foo_1: ptr<function, f32>) -> f32 {
-    let _e1 = (*foo_1);
+fn read_from_private(foo_2: ptr<function, f32>) -> f32 {
+    let _e1 = (*foo_2);
     return _e1;
 }
 
@@ -121,11 +121,15 @@ fn assign_through_ptr_fn(p: ptr<function, u32>) {
     return;
 }
 
-fn assign_array_through_ptr_fn(foo_2: ptr<function, array<vec4<f32>, 2>>) {
-    (*foo_2) = array<vec4<f32>, 2>(vec4(1f), vec4(2f));
+fn assign_array_through_ptr_fn(foo_3: ptr<function, array<vec4<f32>, 2>>) {
+    (*foo_3) = array<vec4<f32>, 2>(vec4(1f), vec4(2f));
     return;
 }
 
+fn array_by_value(a_1: array<i32, 5>, i: i32) -> i32 {
+    return a_1[i];
+}
+
 @vertex 
 fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
     var foo: f32 = 0f;
@@ -138,11 +142,11 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
     let _matrix = bar._matrix;
     let arr_1 = bar.arr;
     let b = bar._matrix[3u][0];
-    let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value;
+    let a_2 = bar.data[(arrayLength((&bar.data)) - 2u)].value;
     let c = qux;
     let data_pointer = (&bar.data[0].value);
     let _e33 = read_from_private((&foo));
-    c2_ = array<i32, 5>(a_1, i32(b), 3i, 4i, 5i);
+    c2_ = array<i32, 5>(a_2, i32(b), 3i, 4i, 5i);
     c2_[(vi + 1u)] = 42i;
     let value = c2_[vi];
     let _e47 = test_arr_as_arg(array<array<f32, 10>, 5>());
@@ -168,3 +172,10 @@ fn assign_through_ptr() {
     assign_array_through_ptr_fn((&arr));
     return;
 }
+
+@vertex 
+fn foo_1(@builtin(vertex_index) vi_1: u32) -> @builtin(position) vec4<f32> {
+    const arr_2 = array<i32, 5>(1i, 2i, 3i, 4i, 5i);
+    let value_1 = arr_2[vi_1];
+    return vec4<f32>(vec4(value_1));
+}
diff --git a/naga/tests/wgsl_errors.rs b/naga/tests/wgsl_errors.rs
index 2d91ba01c..e5fb77644 100644
--- a/naga/tests/wgsl_errors.rs
+++ b/naga/tests/wgsl_errors.rs
@@ -1359,11 +1359,6 @@ fn missing_bindings2() {
 #[test]
 fn invalid_access() {
     check_validation! {
-        "
-        fn array_by_value(a: array<i32, 5>, i: i32) -> i32 {
-            return a[i];
-        }
-        ",
         "
         fn matrix_by_value(m: mat4x4<f32>, i: i32) -> vec4<f32> {
             return m[i];