diff --git a/src/back/spv/instructions.rs b/src/back/spv/instructions.rs index 18de62be04..90f8b0128e 100644 --- a/src/back/spv/instructions.rs +++ b/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, ) -> 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 { diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index 2b467bc3d9..e35a91ca53 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -70,6 +70,7 @@ struct Function { signature: Option, parameters: Vec, variables: crate::FastHashMap, LocalVariable>, + internal_variables: Vec, blocks: Vec, entry_point_context: Option, } @@ -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, + tr: &TypeResolution, + ) -> Result { + 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, @@ -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 { }) } - /// Cache an expression for a value. - fn cache_expression_value<'a>( + #[allow(clippy::too_many_arguments)] + fn promote_access_expression_to_variable( &mut self, - ir_module: &'a crate::Module, + ir_types: &Arena, + result_type_id: Word, + container_id: Word, + container_resolution: &TypeResolution, + index_id: Word, + element_ty: Handle, + 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( + &mut self, + ir_module: &crate::Module, ir_function: &crate::Function, fun_info: &FunctionInfo, expr_handle: Handle, 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")); diff --git a/tests/in/access.param.ron b/tests/in/access.param.ron new file mode 100644 index 0000000000..52f49f360f --- /dev/null +++ b/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, +) diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl new file mode 100644 index 0000000000..c5f6c27e74 --- /dev/null +++ b/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 { + let array = array(1, 2, 3, 4, 5); + let value = array[vi]; + return vec4(vec4(value)); +} diff --git a/tests/in/interpolate.param.ron b/tests/in/interpolate.param.ron index dfd15dc16b..804665025c 100644 --- a/tests/in/interpolate.param.ron +++ b/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) ) diff --git a/tests/in/operators.wgsl b/tests/in/operators.wgsl index 958924c3e5..7808e6e7c5 100644 --- a/tests/in/operators.wgsl +++ b/tests/in/operators.wgsl @@ -1,6 +1,6 @@ [[stage(vertex)]] fn splat() -> [[builtin(position)]] vec4 { - let a = (1.0 + vec2(2.0) - 3.0) / 4.0; - let b = vec4(5) % 2; - return a.xyxy + vec4(b); + let a = (1.0 + vec2(2.0) - 3.0) / 4.0; + let b = vec4(5) % 2; + return a.xyxy + vec4(b); } diff --git a/tests/in/shadow.param.ron b/tests/in/shadow.param.ron index dc51d3ff45..dfa9978c3d 100644 --- a/tests/in/shadow.param.ron +++ 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, ) diff --git a/tests/out/access.msl b/tests/out/access.msl new file mode 100644 index 0000000000..8c9d1c2260 --- /dev/null +++ b/tests/out/access.msl @@ -0,0 +1,15 @@ +#include +#include + +typedef int type3[5]; + +struct fooInput { +}; +struct fooOutput { + metal::float4 member [[position]]; +}; +vertex fooOutput foo( + metal::uint vi [[vertex_id]] +) { + return fooOutput { static_cast(type3 {1, 2, 3, 4, 5}[vi]) }; +} diff --git a/tests/out/access.spvasm b/tests/out/access.spvasm new file mode 100644 index 0000000000..985b9a8408 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/tests/out/boids.spvasm b/tests/out/boids.spvasm index 3f28ef288e..c9a05f398a 100644 --- a/tests/out/boids.spvasm +++ b/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 diff --git a/tests/out/collatz.spvasm b/tests/out/collatz.spvasm index cc4698fa40..929a4343d3 100644 --- a/tests/out/collatz.spvasm +++ b/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 diff --git a/tests/out/interpolate.spvasm b/tests/out/interpolate.spvasm index 063b86f7fc..112bd1f0f5 100644 --- a/tests/out/interpolate.spvasm +++ b/tests/out/interpolate.spvasm @@ -25,7 +25,6 @@ OpName %33 "centroid" OpName %35 "sample" OpName %37 "perspective" OpName %38 "main" -OpName %38 "main" OpName %76 "position" OpName %79 "flat" OpName %82 "linear" @@ -33,7 +32,6 @@ 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 diff --git a/tests/out/quad.spvasm b/tests/out/quad.spvasm index 676aee1160..7f5af3c54e 100644 --- a/tests/out/quad.spvasm +++ b/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 diff --git a/tests/out/shadow.spvasm b/tests/out/shadow.spvasm index 8189e93eed..1016e20605 100644 --- a/tests/out/shadow.spvasm +++ b/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 diff --git a/tests/out/texture-array.spvasm b/tests/out/texture-array.spvasm index dfc2b1f4db..1f4ba03f46 100644 --- a/tests/out/texture-array.spvasm +++ b/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 diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 80b7ecf212..f90272620d 100644 --- a/tests/snapshots.rs +++ b/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() {