[spv-out] implement array value indexing

src/back/spv/instructions.rs

@@ -433,12 +433,12 @@ impl super::Instruction {
     }
 
     pub(super) fn store(
-        pointer_type_id: Word,
+        pointer_id: Word,
         object_id: Word,
         memory_access: Option<spirv::MemoryAccess>,
     ) -> Self {
         let mut instruction = Self::new(Op::Store);
-        instruction.add_operand(pointer_type_id);
+        instruction.add_operand(pointer_id);
         instruction.add_operand(object_id);
 
         if let Some(memory_access) = memory_access {

src/back/spv/writer.rs

@@ -70,6 +70,7 @@ struct Function {
     signature: Option<Instruction>,
     parameters: Vec<Instruction>,
     variables: crate::FastHashMap<Handle<crate::LocalVariable>, LocalVariable>,
+    internal_variables: Vec<LocalVariable>,
     blocks: Vec<Block>,
     entry_point_context: Option<EntryPointContext>,
 }
@@ -86,6 +87,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);
@@ -339,6 +343,20 @@ impl Writer {
         }
     }
 
+    fn get_expression_type_id(
+        &mut self,
+        arena: &Arena<crate::Type>,
+        tr: &TypeResolution,
+    ) -> Result<Word, Error> {
+        let lookup_ty = match *tr {
+            TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
+            TypeResolution::Value(ref inner) => {
+                LookupType::Local(self.physical_layout.make_local(inner).unwrap())
+            }
+        };
+        self.get_type_id(arena, lookup_ty)
+    }
+
     fn get_pointer_id(
         &mut self,
         arena: &Arena<crate::Type>,
@@ -649,11 +667,6 @@ impl Writer {
         };
         self.check(exec_model.required_capabilities())?;
 
-        if self.flags.contains(WriterFlags::DEBUG) {
-            self.debugs
-                .push(Instruction::name(function_id, &entry_point.name));
-        }
-
         Ok(Instruction::entry_point(
             exec_model,
             function_id,
@@ -1288,23 +1301,74 @@ impl Writer {
         })
     }
 
+    #[allow(clippy::too_many_arguments)]
+    fn promote_access_expression_to_variable(
+        &mut self,
+        ir_types: &Arena<crate::Type>,
+        result_type_id: Word,
+        container_id: Word,
+        container_resolution: &TypeResolution,
+        index_id: Word,
+        element_ty: Handle<crate::Type>,
+        block: &mut Block,
+    ) -> Result<(Word, LocalVariable), Error> {
+        let container_type_id = self.get_expression_type_id(ir_types, container_resolution)?;
+        let pointer_type_id = self.id_gen.next();
+        Instruction::type_pointer(
+            pointer_type_id,
+            spirv::StorageClass::Function,
+            container_type_id,
+        )
+        .to_words(&mut self.logical_layout.declarations);
+
+        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))
+    }
+
     /// Cache an expression for a value.
-    fn cache_expression_value<'a>(
+    fn cache_expression_value(
         &mut self,
-        ir_module: &'a crate::Module,
+        ir_module: &crate::Module,
         ir_function: &crate::Function,
         fun_info: &FunctionInfo,
         expr_handle: Handle<crate::Expression>,
         block: &mut Block,
         function: &mut Function,
     ) -> Result<(), Error> {
-        let result_lookup_ty = match fun_info[expr_handle].ty {
-            TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
-            TypeResolution::Value(ref inner) => {
-                LookupType::Local(self.physical_layout.make_local(inner).unwrap())
-            }
-        };
-        let result_type_id = self.get_type_id(&ir_module.types, result_lookup_ty)?;
+        let result_type_id =
+            self.get_expression_type_id(&ir_module.types, &fun_info[expr_handle].ty)?;
 
         let id = match ir_function.expressions[expr_handle] {
             crate::Expression::Access { base, index } => {
@@ -1318,10 +1382,10 @@ impl Writer {
                     0
                 } else {
                     let index_id = self.cached[index];
+                    let base_id = self.cached[base];
                     match *fun_info[base].ty.inner_with(&ir_module.types) {
                         crate::TypeInner::Vector { .. } => {
                             let id = self.id_gen.next();
-                            let base_id = self.cached[base];
                             block.body.push(Instruction::vector_extract_dynamic(
                                 result_type_id,
                                 id,
@@ -1330,7 +1394,21 @@ impl Writer {
                             ));
                             id
                         }
-                        //TODO: support `crate::TypeInner::Array { .. }` ?
+                        crate::TypeInner::Array {
+                            base: ty_element, ..
+                        } => {
+                            let (id, variable) = self.promote_access_expression_to_variable(
+                                &ir_module.types,
+                                result_type_id,
+                                base_id,
+                                &fun_info[base].ty,
+                                index_id,
+                                ty_element,
+                                block,
+                            )?;
+                            function.internal_variables.push(variable);
+                            id
+                        }
                         ref other => {
                             log::error!("Unable to access {:?}", other);
                             return Err(Error::FeatureNotImplemented("access for type"));

tests/in/access.param.ron

@@ -0,0 +1,7 @@
+(
+	spv_version: (1, 1),
+	spv_capabilities: [ Shader, Image1D, Sampled1D ],
+	spv_debug: true,
+	spv_adjust_coordinate_space: false,
+	msl_custom: false,
+)

tests/in/access.wgsl

@@ -0,0 +1,8 @@
+// This snapshot tests accessing various containers, dereferencing pointers.
+
+[[stage(vertex)]]
+fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4<f32> {
+	let array = array<i32, 5>(1, 2, 3, 4, 5);
+	let value = array[vi];
+	return vec4<f32>(vec4<i32>(value));
+}

tests/in/interpolate.param.ron

@@ -4,5 +4,5 @@
 	spv_debug: true,
 	spv_adjust_coordinate_space: true,
 	msl_custom: false,
-        glsl_desktop_version: Some(400)
+	glsl_desktop_version: Some(400)
 )

tests/in/operators.wgsl

@@ -1,6 +1,6 @@
 [[stage(vertex)]]
 fn splat() -> [[builtin(position)]] vec4<f32> {
-	let a = (1.0 + vec2<f32>(2.0) - 3.0) / 4.0;
-	let b = vec4<i32>(5) % 2;
-	return a.xyxy + vec4<f32>(b);
+    let a = (1.0 + vec2<f32>(2.0) - 3.0) / 4.0;
+    let b = vec4<i32>(5) % 2;
+    return a.xyxy + vec4<f32>(b);
 }

tests/in/shadow.param.ron

@@ -1,8 +1,8 @@
 (
-    spv_flow_dump_prefix: "",
-    spv_version: (1, 2),
-    spv_capabilities: [ Shader ],
-    spv_debug: true,
-    spv_adjust_coordinate_space: true,
-    msl_custom: false,
+	spv_flow_dump_prefix: "",
+	spv_version: (1, 2),
+	spv_capabilities: [ Shader ],
+	spv_debug: true,
+	spv_adjust_coordinate_space: true,
+	msl_custom: false,
 )

tests/out/access.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+typedef int type3[5];
+
+struct fooInput {
+};
+struct fooOutput {
+    metal::float4 member [[position]];
+};
+vertex fooOutput foo(
+  metal::uint vi [[vertex_id]]
+) {
+    return fooOutput { static_cast<float4>(type3 {1, 2, 3, 4, 5}[vi]) };
+}

tests/out/access.spvasm

@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 31
+OpCapability Image1D
+OpCapability Shader
+OpCapability Sampled1D
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Vertex %19 "foo" %14 %17
+OpSource GLSL 450
+OpName %14 "vi"
+OpName %19 "foo"
+OpDecorate %12 ArrayStride 4
+OpDecorate %14 BuiltIn VertexIndex
+OpDecorate %17 BuiltIn Position
+%2 = OpTypeVoid
+%4 = OpTypeInt 32 1
+%3 = OpConstant  %4  5
+%5 = OpConstant  %4  1
+%6 = OpConstant  %4  2
+%7 = OpConstant  %4  3
+%8 = OpConstant  %4  4
+%9 = OpTypeInt 32 0
+%11 = OpTypeFloat 32
+%10 = OpTypeVector %11 4
+%12 = OpTypeArray %4 %3
+%15 = OpTypePointer Input %9
+%14 = OpVariable  %15  Input
+%18 = OpTypePointer Output %10
+%17 = OpVariable  %18  Output
+%20 = OpTypeFunction %2
+%23 = OpTypePointer Function %12
+%26 = OpTypePointer Function %4
+%28 = OpTypeVector %4 4
+%19 = OpFunction  %2  None %20
+%13 = OpLabel
+%24 = OpVariable  %23  Function
+%16 = OpLoad  %9  %14
+OpBranch %21
+%21 = OpLabel
+%22 = OpCompositeConstruct  %12  %5 %6 %7 %8 %3
+OpStore %24 %22
+%25 = OpAccessChain  %26  %24 %16
+%27 = OpLoad  %4  %25
+%29 = OpCompositeConstruct  %28  %27 %27 %27 %27
+%30 = OpConvertSToF  %10  %29
+OpStore %17 %30
+OpReturn
+OpFunctionEnd

tests/out/boids.spvasm

@@ -38,7 +38,6 @@ OpName %36 "vel"
 OpName %37 "i"
 OpName %40 "global_invocation_id"
 OpName %43 "main"
-OpName %43 "main"
 OpMemberDecorate %16 0 Offset 0
 OpMemberDecorate %16 1 Offset 8
 OpDecorate %17 Block

tests/out/collatz.spvasm

@@ -17,7 +17,6 @@ OpName %15 "i"
 OpName %18 "collatz_iterations"
 OpName %45 "global_id"
 OpName %48 "main"
-OpName %48 "main"
 OpDecorate %8 ArrayStride 4
 OpDecorate %9 Block
 OpMemberDecorate %9 0 Offset 0

tests/out/interpolate.spvasm

@@ -25,15 +25,13 @@ OpName %33 "centroid"
 OpName %35 "sample"
 OpName %37 "perspective"
 OpName %38 "main"
-OpName %38 "main"
 OpName %76 "position"
 OpName %79 "flat"
 OpName %82 "linear"
 OpName %85 "centroid"
 OpName %88 "sample"
 OpName %91 "perspective"
 OpName %93 "main"
-OpName %93 "main"
 OpMemberDecorate %23 0 Offset 0
 OpMemberDecorate %23 1 Offset 16
 OpMemberDecorate %23 2 Offset 20

tests/out/quad.spvasm

@@ -21,10 +21,8 @@ OpName %22 "uv"
 OpName %24 "uv"
 OpName %26 "position"
 OpName %28 "main"
-OpName %28 "main"
 OpName %48 "uv"
 OpName %51 "main"
-OpName %51 "main"
 OpMemberDecorate %9 0 Offset 0
 OpMemberDecorate %9 1 Offset 16
 OpDecorate %12 DescriptorSet 0

tests/out/shadow.spvasm

@@ -29,7 +29,6 @@ OpName %71 "i"
 OpName %74 "raw_normal"
 OpName %77 "position"
 OpName %82 "fs_main"
-OpName %82 "fs_main"
 OpDecorate %14 Block
 OpMemberDecorate %14 0 Offset 0
 OpMemberDecorate %17 0 Offset 0

tests/out/texture-array.spvasm

@@ -16,7 +16,6 @@ OpName %14 "sampler"
 OpName %16 "pc"
 OpName %19 "tex_coord"
 OpName %24 "main"
-OpName %24 "main"
 OpDecorate %8 Block
 OpMemberDecorate %8 0 Offset 0
 OpDecorate %11 DescriptorSet 0

tests/snapshots.rs

@@ -229,6 +229,7 @@ fn convert_wgsl() {
             "interpolate",
             Targets::SPIRV | Targets::METAL | Targets::GLSL,
         ),
+        ("access", Targets::SPIRV | Targets::METAL),
     ];
 
     for &(name, targets) in inputs.iter() {