diff --git a/README.md b/README.md index 044d4d99ab..13bd5afec2 100644 --- a/README.md +++ b/README.md @@ -24,7 +24,7 @@ Back-end | Status | Feature | Notes | SPIR-V | :white_check_mark: | spv-out | | WGSL | :ok: | wgsl-out | | Metal | :white_check_mark: | msl-out | | -HLSL | :construction: | hlsl-out | Shader Model 5.0+ (DirectX 11+) | +HLSL | :ok: | hlsl-out | Shader Model 5.0+ (DirectX 11+) | GLSL | :ok: | glsl-out | | AIR | | | | DXIL/DXIR | | | | diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 483d4e6892..68841b5475 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -10,6 +10,7 @@ use crate::{ use std::{fmt, mem}; const LOCATION_SEMANTIC: &str = "LOC"; +const STORE_TEMP_NAME: &str = "_value"; /// Shorthand result used internally by the backend pub(super) type BackendResult = Result<(), Error>; @@ -28,6 +29,20 @@ enum SubAccess { }, } +enum StoreValue { + Expression(Handle), + TempIndex { + depth: usize, + index: u32, + ty: proc::TypeResolution, + }, + TempAccess { + depth: usize, + base: Handle, + member_index: u32, + }, +} + /// Structure contains information required for generating /// wrapped structure of all entry points arguments struct EntryPointBinding { @@ -521,11 +536,7 @@ impl<'a, W: fmt::Write> Writer<'a, W> { let size = module.constants[const_handle].to_array_length().unwrap(); write!(self.out, "{}", size)?; } - crate::ArraySize::Dynamic => { - //TODO: https://github.com/gfx-rs/naga/issues/1127 - log::warn!("Dynamically sized arrays are not properly supported yet"); - write!(self.out, "1")?; - } + crate::ArraySize::Dynamic => unreachable!(), } write!(self.out, "]")?; @@ -965,7 +976,14 @@ impl<'a, W: fmt::Write> Writer<'a, W> { .pointer_class() == Some(crate::StorageClass::Storage) { - return Err(Error::Unimplemented("Storage stores".to_string())); + let var_handle = self.fill_access_chain(module, pointer, func_ctx)?; + self.write_storage_store( + module, + var_handle, + StoreValue::Expression(value), + func_ctx, + indent, + )?; } else if let Some((const_handle, base_ty)) = array_info { let size = module.constants[const_handle].to_array_length().unwrap(); writeln!(self.out, "{}{{", INDENT.repeat(indent))?; @@ -1813,7 +1831,7 @@ impl<'a, W: fmt::Write> Writer<'a, W> { Ok(()) } - /// Helper function to write down the load operation on a `ByteAddressBuffer`. + /// Helper function to write down the Load operation on a `ByteAddressBuffer`. fn write_storage_load( &mut self, module: &Module, @@ -1895,6 +1913,203 @@ impl<'a, W: fmt::Write> Writer<'a, W> { Ok(()) } + fn write_store_value( + &mut self, + module: &Module, + value: &StoreValue, + func_ctx: &back::FunctionCtx, + ) -> BackendResult { + match *value { + StoreValue::Expression(expr) => self.write_expr(module, expr, &func_ctx)?, + StoreValue::TempIndex { + depth, + index, + ty: _, + } => write!(self.out, "{}{}[{}]", STORE_TEMP_NAME, depth, index)?, + StoreValue::TempAccess { + depth, + base, + member_index, + } => { + let name = &self.names[&NameKey::StructMember(base, member_index)]; + write!(self.out, "{}{}.{}", STORE_TEMP_NAME, depth, name)? + } + } + Ok(()) + } + + /// Helper function to write down the Store operation on a `ByteAddressBuffer`. + fn write_storage_store( + &mut self, + module: &Module, + var_handle: Handle, + value: StoreValue, + func_ctx: &back::FunctionCtx<'_>, + indent: usize, + ) -> BackendResult { + let temp_resolution; + let ty_resolution = match value { + StoreValue::Expression(expr) => &func_ctx.info[expr].ty, + StoreValue::TempIndex { + depth: _, + index: _, + ref ty, + } => ty, + StoreValue::TempAccess { + depth: _, + base, + member_index, + } => { + let ty_handle = match module.types[base].inner { + TypeInner::Struct { ref members, .. } => members[member_index as usize].ty, + _ => unreachable!(), + }; + temp_resolution = proc::TypeResolution::Handle(ty_handle); + &temp_resolution + } + }; + match *ty_resolution.inner_with(&module.types) { + TypeInner::Scalar { .. } => { + // working around the borrow checker in `self.write_expr` + let chain = mem::take(&mut self.temp_access_chain); + let var_name = &self.names[&NameKey::GlobalVariable(var_handle)]; + write!( + self.out, + "{}{}.Store(", + back::INDENT.repeat(indent), + var_name + )?; + self.write_storage_address(module, &chain, func_ctx)?; + write!(self.out, ", asuint(")?; + self.write_store_value(module, &value, func_ctx)?; + writeln!(self.out, "));")?; + self.temp_access_chain = chain; + } + TypeInner::Vector { size, .. } => { + // working around the borrow checker in `self.write_expr` + let chain = mem::take(&mut self.temp_access_chain); + let var_name = &self.names[&NameKey::GlobalVariable(var_handle)]; + write!( + self.out, + "{}{}.Store{}(", + back::INDENT.repeat(indent), + var_name, + size as u8 + )?; + self.write_storage_address(module, &chain, func_ctx)?; + write!(self.out, ", asuint(")?; + self.write_store_value(module, &value, func_ctx)?; + writeln!(self.out, "));")?; + self.temp_access_chain = chain; + } + TypeInner::Matrix { + columns, + rows, + width, + } => { + // first, assign the value to a temporary + writeln!(self.out, "{}{{", back::INDENT.repeat(indent))?; + let depth = indent + 1; + write!( + self.out, + "{}{}{}x{} {}{} = ", + back::INDENT.repeat(indent + 1), + crate::ScalarKind::Float.to_hlsl_str(width)?, + columns as u8, + rows as u8, + STORE_TEMP_NAME, + depth, + )?; + self.write_store_value(module, &value, func_ctx)?; + writeln!(self.out, ";")?; + // then iterate the stores + let row_stride = width as u32 * rows as u32; + for i in 0..columns as u32 { + self.temp_access_chain + .push(SubAccess::Offset(i * row_stride)); + let ty_inner = TypeInner::Vector { + size: rows, + kind: crate::ScalarKind::Float, + width, + }; + let sv = StoreValue::TempIndex { + depth, + index: i, + ty: proc::TypeResolution::Value(ty_inner), + }; + self.write_storage_store(module, var_handle, sv, func_ctx, indent + 1)?; + self.temp_access_chain.pop(); + } + // done + writeln!(self.out, "{}}}", back::INDENT.repeat(indent))?; + } + TypeInner::Array { + base, + size: crate::ArraySize::Constant(const_handle), + .. + } => { + // first, assign the value to a temporary + writeln!(self.out, "{}{{", back::INDENT.repeat(indent))?; + write!(self.out, "{}", back::INDENT.repeat(indent + 1))?; + self.write_value_type(module, &module.types[base].inner)?; + let depth = indent + 1; + write!(self.out, " {}{}", STORE_TEMP_NAME, depth)?; + self.write_array_size(module, crate::ArraySize::Constant(const_handle))?; + write!(self.out, " = ")?; + self.write_store_value(module, &value, func_ctx)?; + writeln!(self.out, ";")?; + // then iterate the stores + let count = module.constants[const_handle].to_array_length().unwrap(); + let stride = module.types[base].inner.span(&module.constants); + for i in 0..count { + self.temp_access_chain.push(SubAccess::Offset(i * stride)); + let sv = StoreValue::TempIndex { + depth, + index: i, + ty: proc::TypeResolution::Handle(base), + }; + self.write_storage_store(module, var_handle, sv, func_ctx, indent + 1)?; + self.temp_access_chain.pop(); + } + // done + writeln!(self.out, "{}}}", back::INDENT.repeat(indent))?; + } + TypeInner::Struct { ref members, .. } => { + // first, assign the value to a temporary + writeln!(self.out, "{}{{", back::INDENT.repeat(indent))?; + let depth = indent + 1; + let struct_ty = ty_resolution.handle().unwrap(); + let struct_name = &self.names[&NameKey::Type(struct_ty)]; + write!( + self.out, + "{}{} {}{} = ", + back::INDENT.repeat(indent + 1), + struct_name, + STORE_TEMP_NAME, + depth + )?; + self.write_store_value(module, &value, func_ctx)?; + writeln!(self.out, ";")?; + // then iterate the stores + for (i, member) in members.iter().enumerate() { + self.temp_access_chain + .push(SubAccess::Offset(member.offset)); + let sv = StoreValue::TempAccess { + depth, + base: struct_ty, + member_index: i as u32, + }; + self.write_storage_store(module, var_handle, sv, func_ctx, indent + 1)?; + self.temp_access_chain.pop(); + } + // done + writeln!(self.out, "{}}}", back::INDENT.repeat(indent))?; + } + _ => unreachable!(), + } + Ok(()) + } + fn fill_access_chain( &mut self, module: &Module, diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index f97d504de5..e6f07a2633 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -17,14 +17,19 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { let baz: f32 = foo; foo = 1.0; + // test storage loads let matrix = bar.matrix; let arr = bar.arr; - let index = 3u; let b = bar.matrix[index].x; - let a = bar.data[arrayLength(&bar.data) - 2u]; + // test storage stores + bar.matrix[1].z = 1.0; + bar.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); + bar.arr = array, 2>(vec2(0u), vec2(1u)); + + // test array indexing var c = array(a, i32(b), 3, 4, 5); c[vi + 1u] = 42; let value = c[vi]; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 8d299c9a89..0946fb8bfc 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -24,6 +24,19 @@ float4 foo(VertexInput_foo vertexinput_foo) : SV_Position float4 _expr13 = asfloat(bar.Load4(12+0)); float b = _expr13.x; int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 80) / 4) - 2u)*4+8)); + bar.Store(8+4+0, asuint(1.0)); + { + float4x4 _value2 = float4x4(float4(0.0.xxxx), float4(1.0.xxxx), float4(2.0.xxxx), float4(3.0.xxxx)); + bar.Store4(0+0, asuint(_value2[0])); + bar.Store4(0+16, asuint(_value2[1])); + bar.Store4(0+32, asuint(_value2[2])); + bar.Store4(0+48, asuint(_value2[3])); + } + { + uint2 _value2[2] = { uint2(0u.xx), uint2(1u.xx) }; + bar.Store2(4+0, asuint(_value2[0])); + bar.Store2(4+8, asuint(_value2[1])); + } { int _result[5]={ a, int(b), 3, 4, 5 }; for(int _i=0; _i<5; ++_i) c[_i] = _result[_i]; diff --git a/tests/out/hlsl/boids.hlsl b/tests/out/hlsl/boids.hlsl new file mode 100644 index 0000000000..c3e742ca75 --- /dev/null +++ b/tests/out/hlsl/boids.hlsl @@ -0,0 +1,147 @@ +static const uint NUM_PARTICLES = 1500; + +struct Particle { + float2 pos; + float2 vel; +}; + +struct SimParams { + float deltaT; + float rule1Distance; + float rule2Distance; + float rule3Distance; + float rule1Scale; + float rule2Scale; + float rule3Scale; +}; + +cbuffer params : register(b0) { SimParams params; } +ByteAddressBuffer particlesSrc : register(t1); +RWByteAddressBuffer particlesDst : register(u2); + +struct ComputeInput_main { + uint3 global_invocation_id1 : SV_DispatchThreadID; +}; + +[numthreads(64, 1, 1)] +void main(ComputeInput_main computeinput_main) +{ + float2 vPos = (float2)0; + float2 vVel = (float2)0; + float2 cMass = (float2)0; + float2 cVel = (float2)0; + float2 colVel = (float2)0; + int cMassCount = 0; + int cVelCount = 0; + float2 pos = (float2)0; + float2 vel = (float2)0; + uint i = 0u; + + uint index = computeinput_main.global_invocation_id1.x; + if ((index >= NUM_PARTICLES)) { + return; + } + float2 _expr10 = asfloat(particlesSrc.Load2(0+index*4+0)); + vPos = _expr10; + float2 _expr15 = asfloat(particlesSrc.Load2(4+index*4+0)); + vVel = _expr15; + cMass = float2(0.0, 0.0); + cVel = float2(0.0, 0.0); + colVel = float2(0.0, 0.0); + while(true) { + uint _expr37 = i; + if ((_expr37 >= NUM_PARTICLES)) { + break; + } + uint _expr39 = i; + if ((_expr39 == index)) { + continue; + } + uint _expr42 = i; + float2 _expr45 = asfloat(particlesSrc.Load2(0+_expr42*4+0)); + pos = _expr45; + uint _expr47 = i; + float2 _expr50 = asfloat(particlesSrc.Load2(4+_expr47*4+0)); + vel = _expr50; + float2 _expr51 = pos; + float2 _expr52 = vPos; + float _expr55 = params.rule1Distance; + if ((distance(_expr51, _expr52) < _expr55)) { + float2 _expr57 = cMass; + float2 _expr58 = pos; + cMass = (_expr57 + _expr58); + int _expr60 = cMassCount; + cMassCount = (_expr60 + 1); + } + float2 _expr63 = pos; + float2 _expr64 = vPos; + float _expr67 = params.rule2Distance; + if ((distance(_expr63, _expr64) < _expr67)) { + float2 _expr69 = colVel; + float2 _expr70 = pos; + float2 _expr71 = vPos; + colVel = (_expr69 - (_expr70 - _expr71)); + } + float2 _expr74 = pos; + float2 _expr75 = vPos; + float _expr78 = params.rule3Distance; + if ((distance(_expr74, _expr75) < _expr78)) { + float2 _expr80 = cVel; + float2 _expr81 = vel; + cVel = (_expr80 + _expr81); + int _expr83 = cVelCount; + cVelCount = (_expr83 + 1); + } + uint _expr86 = i; + i = (_expr86 + 1u); + } + int _expr89 = cMassCount; + if ((_expr89 > 0)) { + float2 _expr92 = cMass; + int _expr93 = cMassCount; + float2 _expr97 = vPos; + cMass = ((_expr92 / float2(float(_expr93).xx)) - _expr97); + } + int _expr99 = cVelCount; + if ((_expr99 > 0)) { + float2 _expr102 = cVel; + int _expr103 = cVelCount; + cVel = (_expr102 / float2(float(_expr103).xx)); + } + float2 _expr107 = vVel; + float2 _expr108 = cMass; + float _expr110 = params.rule1Scale; + float2 _expr113 = colVel; + float _expr115 = params.rule2Scale; + float2 _expr118 = cVel; + float _expr120 = params.rule3Scale; + vVel = (((_expr107 + (_expr108 * _expr110)) + (_expr113 * _expr115)) + (_expr118 * _expr120)); + float2 _expr123 = vVel; + float2 _expr125 = vVel; + vVel = (normalize(_expr123) * clamp(length(_expr125), 0.0, 0.1)); + float2 _expr131 = vPos; + float2 _expr132 = vVel; + float _expr134 = params.deltaT; + vPos = (_expr131 + (_expr132 * _expr134)); + float2 _expr137 = vPos; + if ((_expr137.x < -1.0)) { + vPos.x = 1.0; + } + float2 _expr143 = vPos; + if ((_expr143.x > 1.0)) { + vPos.x = -1.0; + } + float2 _expr149 = vPos; + if ((_expr149.y < -1.0)) { + vPos.y = 1.0; + } + float2 _expr155 = vPos; + if ((_expr155.y > 1.0)) { + vPos.y = -1.0; + } + float2 _expr164 = vPos; + particlesDst.Store2(0+index*4+0, asuint(_expr164)); + float2 _expr168 = vVel; + particlesDst.Store2(4+index*4+0, asuint(_expr168)); + return; +} diff --git a/tests/out/hlsl/boids.hlsl.config b/tests/out/hlsl/boids.hlsl.config new file mode 100644 index 0000000000..c0d04b8f2a --- /dev/null +++ b/tests/out/hlsl/boids.hlsl.config @@ -0,0 +1,3 @@ +vertex=() +fragment=() +compute=(main:cs_5_0 ) diff --git a/tests/out/hlsl/collatz.hlsl b/tests/out/hlsl/collatz.hlsl new file mode 100644 index 0000000000..3f85ea87e2 --- /dev/null +++ b/tests/out/hlsl/collatz.hlsl @@ -0,0 +1,41 @@ + +RWByteAddressBuffer v_indices : register(u0); + +struct ComputeInput_main { + uint3 global_id1 : SV_DispatchThreadID; +}; + +uint collatz_iterations(uint n_base) +{ + uint n = (uint)0; + uint i = 0u; + + n = n_base; + while(true) { + uint _expr5 = n; + if ((_expr5 <= 1u)) { + break; + } + uint _expr8 = n; + if (((_expr8 % 2u) == 0u)) { + uint _expr13 = n; + n = (_expr13 / 2u); + } else { + uint _expr17 = n; + n = ((3u * _expr17) + 1u); + } + uint _expr21 = i; + i = (_expr21 + 1u); + } + uint _expr24 = i; + return _expr24; +} + +[numthreads(1, 1, 1)] +void main(ComputeInput_main computeinput_main) +{ + uint _expr8 = asuint(v_indices.Load(computeinput_main.global_id1.x*4+0)); + const uint _e9 = collatz_iterations(_expr8); + v_indices.Store(computeinput_main.global_id1.x*4+0, asuint(_e9)); + return; +} diff --git a/tests/out/hlsl/collatz.hlsl.config b/tests/out/hlsl/collatz.hlsl.config new file mode 100644 index 0000000000..c0d04b8f2a --- /dev/null +++ b/tests/out/hlsl/collatz.hlsl.config @@ -0,0 +1,3 @@ +vertex=() +fragment=() +compute=(main:cs_5_0 ) diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index b04acf40b6..16bcedf555 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -38,6 +38,9 @@ vertex fooOutput foo( metal::float4 _e13 = bar.matrix[3]; float b = _e13.x; int a = bar.data[(1 + (_buffer_sizes.size0 - 80 - 4) / 4) - 2u]; + bar.matrix[1].z = 1.0; + bar.matrix = metal::float4x4(metal::float4(0.0), metal::float4(1.0), metal::float4(2.0), metal::float4(3.0)); + for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type2 {metal::uint2(0u), metal::uint2(1u)}.inner[_i]; for(int _i=0; _i<5; ++_i) c.inner[_i] = type8 {a, static_cast(b), 3, 4, 5}.inner[_i]; c.inner[vi + 1u] = 42; int value = c.inner[vi]; diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index 32c4fc27d8..19727e05fb 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,35 +1,35 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 67 +; Bound: 82 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %35 "foo" %30 %33 +OpEntryPoint Vertex %39 "foo" %34 %37 OpSource GLSL 450 -OpName %21 "Bar" -OpMemberName %21 0 "matrix" -OpMemberName %21 1 "arr" -OpMemberName %21 2 "data" -OpName %23 "bar" -OpName %25 "foo" -OpName %27 "c" -OpName %30 "vi" -OpName %35 "foo" -OpDecorate %19 ArrayStride 8 -OpDecorate %20 ArrayStride 4 -OpDecorate %21 Block -OpMemberDecorate %21 0 Offset 0 -OpMemberDecorate %21 0 ColMajor -OpMemberDecorate %21 0 MatrixStride 16 -OpMemberDecorate %21 1 Offset 64 -OpMemberDecorate %21 2 Offset 80 -OpDecorate %22 ArrayStride 4 -OpDecorate %23 DescriptorSet 0 -OpDecorate %23 Binding 0 -OpDecorate %30 BuiltIn VertexIndex -OpDecorate %33 BuiltIn Position +OpName %25 "Bar" +OpMemberName %25 0 "matrix" +OpMemberName %25 1 "arr" +OpMemberName %25 2 "data" +OpName %27 "bar" +OpName %29 "foo" +OpName %31 "c" +OpName %34 "vi" +OpName %39 "foo" +OpDecorate %23 ArrayStride 8 +OpDecorate %24 ArrayStride 4 +OpDecorate %25 Block +OpMemberDecorate %25 0 Offset 0 +OpMemberDecorate %25 0 ColMajor +OpMemberDecorate %25 0 MatrixStride 16 +OpMemberDecorate %25 1 Offset 64 +OpMemberDecorate %25 2 Offset 80 +OpDecorate %26 ArrayStride 4 +OpDecorate %27 DescriptorSet 0 +OpDecorate %27 Binding 0 +OpDecorate %34 BuiltIn VertexIndex +OpDecorate %37 BuiltIn Position %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 @@ -39,66 +39,84 @@ OpDecorate %33 BuiltIn Position %9 = OpTypeInt 32 0 %8 = OpConstant %9 3 %10 = OpConstant %9 2 -%11 = OpConstant %4 5 -%12 = OpConstant %4 3 -%13 = OpConstant %4 4 -%14 = OpConstant %9 1 -%15 = OpConstant %4 42 -%17 = OpTypeVector %6 4 -%16 = OpTypeMatrix %17 4 -%18 = OpTypeVector %9 2 -%19 = OpTypeArray %18 %3 -%20 = OpTypeRuntimeArray %4 -%21 = OpTypeStruct %16 %19 %20 -%22 = OpTypeArray %4 %11 -%24 = OpTypePointer StorageBuffer %21 -%23 = OpVariable %24 StorageBuffer -%26 = OpTypePointer Function %6 -%28 = OpTypePointer Function %22 -%31 = OpTypePointer Input %9 -%30 = OpVariable %31 Input -%34 = OpTypePointer Output %17 -%33 = OpVariable %34 Output -%36 = OpTypeFunction %2 -%39 = OpTypePointer StorageBuffer %16 -%40 = OpConstant %9 0 -%43 = OpTypePointer StorageBuffer %19 -%46 = OpTypePointer StorageBuffer %17 -%50 = OpTypePointer StorageBuffer %20 -%53 = OpTypePointer StorageBuffer %4 -%59 = OpTypePointer Function %4 -%63 = OpTypeVector %4 4 -%35 = OpFunction %2 None %36 -%29 = OpLabel -%25 = OpVariable %26 Function %5 -%27 = OpVariable %28 Function -%32 = OpLoad %9 %30 -OpBranch %37 -%37 = OpLabel -%38 = OpLoad %6 %25 -OpStore %25 %7 -%41 = OpAccessChain %39 %23 %40 -%42 = OpLoad %16 %41 -%44 = OpAccessChain %43 %23 %14 -%45 = OpLoad %19 %44 -%47 = OpAccessChain %46 %23 %40 %8 -%48 = OpLoad %17 %47 -%49 = OpCompositeExtract %6 %48 0 -%51 = OpArrayLength %9 %23 2 -%52 = OpISub %9 %51 %10 -%54 = OpAccessChain %53 %23 %10 %52 -%55 = OpLoad %4 %54 -%56 = OpConvertFToS %4 %49 -%57 = OpCompositeConstruct %22 %55 %56 %12 %13 %11 -OpStore %27 %57 -%58 = OpIAdd %9 %32 %14 -%60 = OpAccessChain %59 %27 %58 -OpStore %60 %15 -%61 = OpAccessChain %59 %27 %32 -%62 = OpLoad %4 %61 -%64 = OpCompositeConstruct %63 %62 %62 %62 %62 -%65 = OpConvertSToF %17 %64 -%66 = OpMatrixTimesVector %17 %42 %65 -OpStore %33 %66 +%11 = OpConstant %4 1 +%12 = OpConstant %6 2.0 +%13 = OpConstant %6 3.0 +%14 = OpConstant %9 0 +%15 = OpConstant %9 1 +%16 = OpConstant %4 5 +%17 = OpConstant %4 3 +%18 = OpConstant %4 4 +%19 = OpConstant %4 42 +%21 = OpTypeVector %6 4 +%20 = OpTypeMatrix %21 4 +%22 = OpTypeVector %9 2 +%23 = OpTypeArray %22 %3 +%24 = OpTypeRuntimeArray %4 +%25 = OpTypeStruct %20 %23 %24 +%26 = OpTypeArray %4 %16 +%28 = OpTypePointer StorageBuffer %25 +%27 = OpVariable %28 StorageBuffer +%30 = OpTypePointer Function %6 +%32 = OpTypePointer Function %26 +%35 = OpTypePointer Input %9 +%34 = OpVariable %35 Input +%38 = OpTypePointer Output %21 +%37 = OpVariable %38 Output +%40 = OpTypeFunction %2 +%43 = OpTypePointer StorageBuffer %20 +%46 = OpTypePointer StorageBuffer %23 +%49 = OpTypePointer StorageBuffer %21 +%53 = OpTypePointer StorageBuffer %24 +%56 = OpTypePointer StorageBuffer %4 +%59 = OpTypePointer StorageBuffer %6 +%74 = OpTypePointer Function %4 +%78 = OpTypeVector %4 4 +%39 = OpFunction %2 None %40 +%33 = OpLabel +%29 = OpVariable %30 Function %5 +%31 = OpVariable %32 Function +%36 = OpLoad %9 %34 +OpBranch %41 +%41 = OpLabel +%42 = OpLoad %6 %29 +OpStore %29 %7 +%44 = OpAccessChain %43 %27 %14 +%45 = OpLoad %20 %44 +%47 = OpAccessChain %46 %27 %15 +%48 = OpLoad %23 %47 +%50 = OpAccessChain %49 %27 %14 %8 +%51 = OpLoad %21 %50 +%52 = OpCompositeExtract %6 %51 0 +%54 = OpArrayLength %9 %27 2 +%55 = OpISub %9 %54 %10 +%57 = OpAccessChain %56 %27 %10 %55 +%58 = OpLoad %4 %57 +%60 = OpAccessChain %59 %27 %14 %15 %10 +OpStore %60 %7 +%61 = OpCompositeConstruct %21 %5 %5 %5 %5 +%62 = OpCompositeConstruct %21 %7 %7 %7 %7 +%63 = OpCompositeConstruct %21 %12 %12 %12 %12 +%64 = OpCompositeConstruct %21 %13 %13 %13 %13 +%65 = OpCompositeConstruct %20 %61 %62 %63 %64 +%66 = OpAccessChain %43 %27 %14 +OpStore %66 %65 +%67 = OpCompositeConstruct %22 %14 %14 +%68 = OpCompositeConstruct %22 %15 %15 +%69 = OpCompositeConstruct %23 %67 %68 +%70 = OpAccessChain %46 %27 %15 +OpStore %70 %69 +%71 = OpConvertFToS %4 %52 +%72 = OpCompositeConstruct %26 %58 %71 %17 %18 %16 +OpStore %31 %72 +%73 = OpIAdd %9 %36 %15 +%75 = OpAccessChain %74 %31 %73 +OpStore %75 %19 +%76 = OpAccessChain %74 %31 %36 +%77 = OpLoad %4 %76 +%79 = OpCompositeConstruct %78 %77 %77 %77 %77 +%80 = OpConvertSToF %21 %79 +%81 = OpMatrixTimesVector %21 %45 %80 +OpStore %37 %81 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index cecc1e04d7..05c398dfee 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -20,6 +20,9 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { let _e13: vec4 = bar.matrix[3]; let b: f32 = _e13.x; let a: i32 = bar.data[(arrayLength(&bar.data) - 2u)]; + bar.matrix[1][2] = 1.0; + bar.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); + bar.arr = array,2>(vec2(0u), vec2(1u)); c = array(a, i32(b), 3, 4, 5); c[(vi + 1u)] = 42; let value: i32 = c[vi]; diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 03edc5f7ac..1e6b6b5345 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -396,7 +396,7 @@ fn convert_wgsl() { ), ( "boids", - Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ), ( "skybox", @@ -404,7 +404,12 @@ fn convert_wgsl() { ), ( "collatz", - Targets::SPIRV | Targets::METAL | Targets::IR | Targets::ANALYSIS | Targets::WGSL, + Targets::SPIRV + | Targets::METAL + | Targets::IR + | Targets::ANALYSIS + | Targets::HLSL + | Targets::WGSL, ), ( "shadow",