diff --git a/CHANGELOG.md b/CHANGELOG.md index 93ab09e9ef..446cb4464e 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,5 +1,11 @@ # Change Log +## v0.9 (TBD) + - WGSL: + - attributes are declared with `@attrib` instead of `[[attrib]]` + - `stride` attribute is removed + - block comments are supported + ## v0.8 (2021-12-18) - development release for wgpu-0.12 - lots of fixes in all parts diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 60598fe904..64a512c939 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -17,7 +17,6 @@ enum Attribute { Interpolate(Option, Option), Location(u32), Stage(ShaderStage), - Stride(u32), WorkGroupSize([u32; 3]), } @@ -349,7 +348,6 @@ impl Writer { }; write!(self.out, "@stage({}) ", stage_str)?; } - Attribute::Stride(stride) => write!(self.out, "@stride({}) ", stride)?, Attribute::WorkGroupSize(size) => { write!( self.out, @@ -420,15 +418,6 @@ impl Writer { // Write struct member name and type let member_name = &self.names[&NameKey::StructMember(handle, index as u32)]; write!(self.out, "{}: ", member_name)?; - // Write stride attribute for array struct member - if let TypeInner::Array { - base: _, - size: _, - stride, - } = module.types[member.ty].inner - { - self.write_attributes(&[Attribute::Stride(stride)])?; - } self.write_type(module, member.ty)?; write!(self.out, ";")?; writeln!(self.out)?; @@ -523,7 +512,11 @@ impl Writer { TypeInner::Atomic { kind, .. } => { write!(self.out, "atomic<{}>", scalar_kind_str(kind))?; } - TypeInner::Array { base, size, .. } => { + TypeInner::Array { + base, + size, + stride: _, + } => { // More info https://gpuweb.github.io/gpuweb/wgsl/#array-types // array -- Constant array // array -- Dynamic array diff --git a/src/front/glsl/mod.rs b/src/front/glsl/mod.rs index 1d4c6c1a94..0712f4bd50 100644 --- a/src/front/glsl/mod.rs +++ b/src/front/glsl/mod.rs @@ -10,7 +10,7 @@ pub use ast::{Precision, Profile}; pub use error::{Error, ErrorKind, ExpectedToken}; pub use token::TokenValue; -use crate::{FastHashMap, FastHashSet, Handle, Module, ShaderStage, Span, Type}; +use crate::{proc::Layouter, FastHashMap, FastHashSet, Handle, Module, ShaderStage, Span, Type}; use ast::{EntryArg, FunctionDeclaration, GlobalLookup}; use parser::ParsingContext; @@ -166,6 +166,8 @@ pub struct Parser { entry_args: Vec, + layouter: Layouter, + errors: Vec, module: Module, @@ -179,6 +181,7 @@ impl Parser { self.lookup_type.clear(); self.global_variables.clear(); self.entry_args.clear(); + self.layouter.clear(); // This is necessary because if the last parsing errored out, the module // wouldn't have been swapped diff --git a/src/front/glsl/types.rs b/src/front/glsl/types.rs index aee13c77f2..becc53ad6c 100644 --- a/src/front/glsl/types.rs +++ b/src/front/glsl/types.rs @@ -259,10 +259,13 @@ impl Parser { mut meta: Span, array_specifier: Option<(ArraySize, Span)>, ) -> Handle { - array_specifier - .map(|(size, size_meta)| { + match array_specifier { + Some((size, size_meta)) => { meta.subsume(size_meta); - let stride = self.module.types[base].inner.size(&self.module.constants); + self.layouter + .update(&self.module.types, &self.module.constants) + .unwrap(); + let stride = self.layouter[base].to_stride(); self.module.types.insert( Type { name: None, @@ -270,7 +273,8 @@ impl Parser { }, meta, ) - }) - .unwrap_or(base) + } + None => base, + } } } diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index f7d65afc22..756ce38497 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -158,7 +158,6 @@ pub enum Error<'a> { UnknownType(Span), UnknownStorageFormat(Span), UnknownConservativeDepth(Span), - ZeroStride(Span), ZeroSizeOrAlign(Span), InconsistentBinding(Span), UnknownLocalFunction(Span), @@ -232,7 +231,7 @@ impl<'a> Error<'a> { ExpectedToken::Constant => "constant".to_string(), ExpectedToken::PrimaryExpression => "expression".to_string(), ExpectedToken::FieldName => "field name or a closing curly bracket to signify the end of the struct".to_string(), - ExpectedToken::TypeAttribute => "type attribute ('stride')".to_string(), + ExpectedToken::TypeAttribute => "type attribute".to_string(), ExpectedToken::Statement => "statement".to_string(), ExpectedToken::SwitchItem => "switch item ('case' or 'default') or a closing curly bracket to signify the end of the switch statement ('}')".to_string(), ExpectedToken::WorkgroupSizeSeparator => "workgroup size separator (',') or a closing parenthesis".to_string(), @@ -393,11 +392,6 @@ impl<'a> Error<'a> { labels: vec![(bad_span.clone(), "unknown type".into())], notes: vec![], }, - Error::ZeroStride(ref bad_span) => ParseError { - message: "array stride must not be zero".to_string(), - labels: vec![(bad_span.clone(), "array stride must not be zero".into())], - notes: vec![], - }, Error::ZeroSizeOrAlign(ref bad_span) => ParseError { message: "struct member size or alignment must not be 0".to_string(), labels: vec![(bad_span.clone(), "struct member size or alignment must not be 0".into())], @@ -1048,7 +1042,9 @@ impl Composition { #[derive(Default)] struct TypeAttributes { - stride: Option, + // Although WGSL nas no type attributes at the moment, it had them in the past +// (`[[stride]]`) and may as well acquire some again in the future. +// Therefore, we are leaving the plumbing in for now. } #[derive(Clone, Debug, PartialEq)] @@ -2888,7 +2884,7 @@ impl Parser { fn parse_type_decl_impl<'a>( &mut self, lexer: &mut Lexer<'a>, - attribute: TypeAttributes, + _attribute: TypeAttributes, word: &'a str, type_arena: &mut UniqueArena, const_arena: &mut Arena, @@ -3024,15 +3020,11 @@ impl Parser { crate::ArraySize::Dynamic }; lexer.expect_generic_paren('>')?; - let stride = match attribute.stride { - Some(stride) => stride.get(), - None => { - self.layouter.update(type_arena, const_arena).unwrap(); - let layout = &self.layouter[base]; - Layouter::round_up(layout.alignment, layout.size) - } - }; + let stride = { + self.layouter.update(type_arena, const_arena).unwrap(); + self.layouter[base].to_stride() + }; crate::TypeInner::Array { base, size, stride } } "sampler" => crate::TypeInner::Sampler { comparison: false }, @@ -3240,22 +3232,11 @@ impl Parser { const_arena: &mut Arena, ) -> Result<(Handle, crate::StorageAccess), Error<'a>> { self.push_scope(Scope::TypeDecl, lexer); - let mut attribute = TypeAttributes::default(); + let attribute = TypeAttributes::default(); - while lexer.skip(Token::Attribute) { - self.push_scope(Scope::Attribute, lexer); - match lexer.next() { - (Token::Word("stride"), _) => { - lexer.expect(Token::Paren('('))?; - let (stride, span) = - lexer.capture_span(|lexer| parse_non_negative_sint_literal(lexer, 4))?; - attribute.stride = - Some(NonZeroU32::new(stride).ok_or(Error::ZeroStride(span))?); - lexer.expect(Token::Paren(')'))?; - } - other => return Err(Error::Unexpected(other, ExpectedToken::TypeAttribute)), - } - self.pop_scope(lexer); + if lexer.skip(Token::Attribute) { + let other = lexer.next(); + return Err(Error::Unexpected(other, ExpectedToken::TypeAttribute)); } let storage_access = crate::StorageAccess::default(); diff --git a/src/front/wgsl/tests.rs b/src/front/wgsl/tests.rs index bb415fa7cd..4d5e047a6a 100644 --- a/src/front/wgsl/tests.rs +++ b/src/front/wgsl/tests.rs @@ -398,7 +398,7 @@ fn parse_array_length() { parse_str( " struct Foo { - data: @stride(4) array; + data: array; }; // this is used as both input and output for convenience @group(0) @binding(0) diff --git a/src/proc/layouter.rs b/src/proc/layouter.rs index 2ce334492f..7cbe15cedc 100644 --- a/src/proc/layouter.rs +++ b/src/proc/layouter.rs @@ -12,6 +12,13 @@ pub struct TypeLayout { pub alignment: Alignment, } +impl TypeLayout { + /// Produce the stride as if this type is a base of an array. + pub fn to_stride(&self) -> u32 { + Layouter::round_up(self.alignment, self.size) + } +} + /// Helper processor that derives the sizes of all types. /// It uses the default layout algorithm/table, described in /// diff --git a/src/valid/type.rs b/src/valid/type.rs index 7fc1a6466f..db68eac4c6 100644 --- a/src/valid/type.rs +++ b/src/valid/type.rs @@ -97,8 +97,8 @@ pub enum TypeError { UnsupportedSpecializedArrayLength(Handle), #[error("Array type {0:?} must have a length of one or more")] NonPositiveArrayLength(Handle), - #[error("Array stride {stride} is smaller than the base element size {base_size}")] - InsufficientArrayStride { stride: u32, base_size: u32 }, + #[error("Array stride {stride} does not match the expected {expected}")] + InvalidArrayStride { stride: u32, expected: u32 }, #[error("Field '{0}' can't be dynamically-sized, has type {1:?}")] InvalidDynamicArray(String, Handle), #[error("Structure member[{index}] at {offset} overlaps the previous member")] @@ -328,19 +328,20 @@ impl super::Validator { return Err(TypeError::InvalidArrayBaseType(base)); } - //Note: `unwrap()` is fine, since `Layouter` goes first and calls it - let base_size = types[base].inner.size(constants); - if stride < base_size { - return Err(TypeError::InsufficientArrayStride { stride, base_size }); + let base_layout = self.layouter[base]; + let expected_stride = base_layout.to_stride(); + if stride != expected_stride { + return Err(TypeError::InvalidArrayStride { + stride, + expected: expected_stride, + }); } - let general_alignment = self.layouter[base].alignment; + let general_alignment = base_layout.alignment.get(); let uniform_layout = match base_info.uniform_layout { Ok(base_alignment) => { // combine the alignment requirements - let align = ((base_alignment.unwrap().get() - 1) - | (general_alignment.get() - 1)) - + 1; + let align = base_alignment.unwrap().get().max(general_alignment); if stride % align != 0 { Err(( handle, @@ -357,9 +358,7 @@ impl super::Validator { }; let storage_layout = match base_info.storage_layout { Ok(base_alignment) => { - let align = ((base_alignment.unwrap().get() - 1) - | (general_alignment.get() - 1)) - + 1; + let align = base_alignment.unwrap().get().max(general_alignment); if stride % align != 0 { Err(( handle, diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index ce26cb8a4c..90c707359c 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -1,11 +1,15 @@ // This snapshot tests accessing various containers, dereferencing pointers. +struct AlignedWrapper { + @align(8) value: i32; +}; + struct Bar { matrix: mat4x4; matrix_array: array, 2>; atom: atomic; - arr: @stride(8) array, 2>; - data: @stride(8) array; + arr: array, 2>; + data: array; }; @group(0) @binding(0) @@ -27,17 +31,17 @@ fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { let arr = bar.arr; let index = 3u; let b = bar.matrix[index].x; - let a = bar.data[arrayLength(&bar.data) - 2u]; + let a = bar.data[arrayLength(&bar.data) - 2u].value; // test pointer types - let data_pointer: ptr = &bar.data[0]; + let data_pointer: ptr = &bar.data[0].value; let foo_value = read_from_private(&foo); // 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)); - bar.data[1] = 1; + bar.data[1].value = 1; // test array indexing var c = array(a, i32(b), 3, 4, 5); diff --git a/tests/in/boids.wgsl b/tests/in/boids.wgsl index 1cc54cfc48..6f44752c64 100644 --- a/tests/in/boids.wgsl +++ b/tests/in/boids.wgsl @@ -16,7 +16,7 @@ struct SimParams { }; struct Particles { - particles : @stride(16) array; + particles : array; }; @group(0) @binding(0) var params : SimParams; diff --git a/tests/in/collatz.wgsl b/tests/in/collatz.wgsl index f4341c24a6..cf972a8ffd 100644 --- a/tests/in/collatz.wgsl +++ b/tests/in/collatz.wgsl @@ -1,5 +1,5 @@ struct PrimeIndices { - data: @stride(4) array; + data: array; }; // this is used as both input and output for convenience @group(0) @binding(0) diff --git a/tests/in/shadow.wgsl b/tests/in/shadow.wgsl index 367762d517..bbd0a3f687 100644 --- a/tests/in/shadow.wgsl +++ b/tests/in/shadow.wgsl @@ -12,7 +12,7 @@ struct Light { }; struct Lights { - data: @stride(96) array; + data: array; }; @group(0) @binding(1) diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index cca755cd68..632fa26666 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -5,12 +5,15 @@ precision highp int; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +struct AlignedWrapper { + int value; +}; layout(std430) buffer Bar_block_0Compute { mat4x4 matrix; mat2x2 matrix_array[2]; int atom; uvec2 arr[2]; - int data[]; + AlignedWrapper data[]; } _group_0_binding_0_cs; diff --git a/tests/out/glsl/access.foo.Vertex.glsl b/tests/out/glsl/access.foo.Vertex.glsl index efac041484..3508015cf1 100644 --- a/tests/out/glsl/access.foo.Vertex.glsl +++ b/tests/out/glsl/access.foo.Vertex.glsl @@ -3,12 +3,15 @@ precision highp float; precision highp int; +struct AlignedWrapper { + int value; +}; layout(std430) buffer Bar_block_0Vertex { mat4x4 matrix; mat2x2 matrix_array[2]; int atom; uvec2 arr[2]; - int data[]; + AlignedWrapper data[]; } _group_0_binding_0_vs; @@ -26,12 +29,12 @@ void main() { mat4x4 matrix = _group_0_binding_0_vs.matrix; uvec2 arr[2] = _group_0_binding_0_vs.arr; float b = _group_0_binding_0_vs.matrix[3][0]; - int a = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)]; - float _e25 = read_from_private(foo_1); + int a = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; + float _e27 = read_from_private(foo_1); _group_0_binding_0_vs.matrix[1][2] = 1.0; _group_0_binding_0_vs.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); _group_0_binding_0_vs.arr = uvec2[2](uvec2(0u), uvec2(1u)); - _group_0_binding_0_vs.data[1] = 1; + _group_0_binding_0_vs.data[1].value = 1; c = int[5](a, int(b), 3, 4, 5); c[(vi + 1u)] = 42; int value = c[vi]; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index aa89e7f901..f79a6798bc 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -1,4 +1,8 @@ +struct AlignedWrapper { + int value; +}; + RWByteAddressBuffer bar : register(u0); float read_from_private(inout float foo_2) @@ -24,8 +28,8 @@ float4 foo(uint vi : SV_VertexID) : SV_Position float4x4 matrix_ = float4x4(asfloat(bar.Load4(0+0)), asfloat(bar.Load4(0+16)), asfloat(bar.Load4(0+32)), asfloat(bar.Load4(0+48))); uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))}; float b = asfloat(bar.Load(0+48+0)); - int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120)); - const float _e25 = read_from_private(foo_1); + int a = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120)); + const float _e27 = read_from_private(foo_1); bar.Store(8+16+0, asuint(1.0)); { float4x4 _value2 = float4x4(float4(0.0.xxxx), float4(1.0.xxxx), float4(2.0.xxxx), float4(3.0.xxxx)); @@ -39,7 +43,7 @@ float4 foo(uint vi : SV_VertexID) : SV_Position bar.Store2(104+0, asuint(_value2[0])); bar.Store2(104+8, asuint(_value2[1])); } - bar.Store(8+120, asuint(1)); + bar.Store(0+8+120, asuint(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/msl/access.msl b/tests/out/msl/access.msl index 53897e1adf..39033f3fdd 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -6,19 +6,22 @@ struct _mslBufferSizes { metal::uint size0; }; -struct type_2 { +struct AlignedWrapper { + int value; +}; +struct type_3 { metal::float2x2 inner[2]; }; -struct type_5 { +struct type_6 { metal::uint2 inner[2]; }; -typedef int type_7[1]; +typedef AlignedWrapper type_7[1]; struct Bar { metal::float4x4 matrix; - type_2 matrix_array; + type_3 matrix_array; metal::atomic_int atom; char _pad3[4]; - type_5 arr; + type_6 arr; type_7 data; }; struct type_13 { @@ -47,14 +50,14 @@ vertex fooOutput foo( float baz = foo_1; foo_1 = 1.0; metal::float4x4 matrix = bar.matrix; - type_5 arr = bar.arr; + type_6 arr = bar.arr; float b = bar.matrix[3].x; - int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 4) / 8) - 2u]; - float _e25 = read_from_private(foo_1); + int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value; + float _e27 = read_from_private(foo_1); 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] = type_5 {metal::uint2(0u), metal::uint2(1u)}.inner[_i]; - bar.data[1] = 1; + for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type_6 {metal::uint2(0u), metal::uint2(1u)}.inner[_i]; + bar.data[1].value = 1; for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {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 6a599ad51a..bdd6160406 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,48 +1,51 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 119 +; Bound: 121 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %50 "foo" %45 %48 -OpEntryPoint GLCompute %96 "atomics" -OpExecutionMode %96 LocalSize 1 1 1 +OpEntryPoint Vertex %51 "foo" %46 %49 +OpEntryPoint GLCompute %98 "atomics" +OpExecutionMode %98 LocalSize 1 1 1 OpSource GLSL 450 -OpMemberName %29 0 "matrix" -OpMemberName %29 1 "matrix_array" -OpMemberName %29 2 "atom" -OpMemberName %29 3 "arr" -OpMemberName %29 4 "data" -OpName %29 "Bar" -OpName %33 "bar" -OpName %36 "foo" -OpName %37 "read_from_private" -OpName %41 "foo" -OpName %42 "c" -OpName %45 "vi" -OpName %50 "foo" -OpName %94 "tmp" -OpName %96 "atomics" -OpDecorate %25 ArrayStride 16 -OpDecorate %27 ArrayStride 8 +OpMemberName %21 0 "value" +OpName %21 "AlignedWrapper" +OpMemberName %30 0 "matrix" +OpMemberName %30 1 "matrix_array" +OpMemberName %30 2 "atom" +OpMemberName %30 3 "arr" +OpMemberName %30 4 "data" +OpName %30 "Bar" +OpName %34 "bar" +OpName %37 "foo" +OpName %38 "read_from_private" +OpName %42 "foo" +OpName %43 "c" +OpName %46 "vi" +OpName %51 "foo" +OpName %96 "tmp" +OpName %98 "atomics" +OpMemberDecorate %21 0 Offset 0 +OpDecorate %26 ArrayStride 16 OpDecorate %28 ArrayStride 8 -OpMemberDecorate %29 0 Offset 0 -OpMemberDecorate %29 0 ColMajor -OpMemberDecorate %29 0 MatrixStride 16 -OpMemberDecorate %29 1 Offset 64 -OpMemberDecorate %29 1 ColMajor -OpMemberDecorate %29 1 MatrixStride 8 -OpMemberDecorate %29 2 Offset 96 -OpMemberDecorate %29 3 Offset 104 -OpMemberDecorate %29 4 Offset 120 -OpDecorate %32 ArrayStride 4 -OpDecorate %33 DescriptorSet 0 -OpDecorate %33 Binding 0 -OpDecorate %29 Block -OpDecorate %45 BuiltIn VertexIndex -OpDecorate %48 BuiltIn Position +OpDecorate %29 ArrayStride 8 +OpMemberDecorate %30 0 Offset 0 +OpMemberDecorate %30 0 ColMajor +OpMemberDecorate %30 0 MatrixStride 16 +OpMemberDecorate %30 1 Offset 64 +OpMemberDecorate %30 1 ColMajor +OpMemberDecorate %30 1 MatrixStride 8 +OpMemberDecorate %30 2 Offset 96 +OpMemberDecorate %30 3 Offset 104 +OpMemberDecorate %30 4 Offset 120 +OpDecorate %33 ArrayStride 4 +OpDecorate %34 DescriptorSet 0 +OpDecorate %34 Binding 0 +OpDecorate %30 Block +OpDecorate %46 BuiltIn VertexIndex +OpDecorate %49 BuiltIn Position %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 @@ -62,127 +65,129 @@ OpDecorate %48 BuiltIn Position %18 = OpConstant %4 3 %19 = OpConstant %4 4 %20 = OpConstant %4 42 -%22 = OpTypeVector %6 4 -%21 = OpTypeMatrix %22 4 -%24 = OpTypeVector %6 2 -%23 = OpTypeMatrix %24 2 -%25 = OpTypeArray %23 %3 -%26 = OpTypeVector %9 2 -%27 = OpTypeArray %26 %3 -%28 = OpTypeRuntimeArray %4 -%29 = OpTypeStruct %21 %25 %4 %27 %28 -%30 = OpTypePointer Function %6 -%31 = OpTypePointer StorageBuffer %4 -%32 = OpTypeArray %4 %17 -%34 = OpTypePointer StorageBuffer %29 -%33 = OpVariable %34 StorageBuffer -%38 = OpTypeFunction %6 %30 -%43 = OpTypePointer Function %32 -%46 = OpTypePointer Input %9 -%45 = OpVariable %46 Input -%49 = OpTypePointer Output %22 -%48 = OpVariable %49 Output -%51 = OpTypeFunction %2 -%54 = OpTypePointer StorageBuffer %21 -%57 = OpTypePointer StorageBuffer %27 -%60 = OpTypePointer StorageBuffer %22 -%61 = OpTypePointer StorageBuffer %6 -%64 = OpTypePointer StorageBuffer %28 -%67 = OpConstant %9 4 -%86 = OpTypePointer Function %4 -%90 = OpTypeVector %4 4 -%98 = OpTypePointer StorageBuffer %4 -%101 = OpConstant %9 64 -%37 = OpFunction %6 None %38 -%36 = OpFunctionParameter %30 -%35 = OpLabel -OpBranch %39 -%39 = OpLabel -%40 = OpLoad %6 %36 -OpReturnValue %40 +%21 = OpTypeStruct %4 +%23 = OpTypeVector %6 4 +%22 = OpTypeMatrix %23 4 +%25 = OpTypeVector %6 2 +%24 = OpTypeMatrix %25 2 +%26 = OpTypeArray %24 %3 +%27 = OpTypeVector %9 2 +%28 = OpTypeArray %27 %3 +%29 = OpTypeRuntimeArray %21 +%30 = OpTypeStruct %22 %26 %4 %28 %29 +%31 = OpTypePointer Function %6 +%32 = OpTypePointer StorageBuffer %4 +%33 = OpTypeArray %4 %17 +%35 = OpTypePointer StorageBuffer %30 +%34 = OpVariable %35 StorageBuffer +%39 = OpTypeFunction %6 %31 +%44 = OpTypePointer Function %33 +%47 = OpTypePointer Input %9 +%46 = OpVariable %47 Input +%50 = OpTypePointer Output %23 +%49 = OpVariable %50 Output +%52 = OpTypeFunction %2 +%55 = OpTypePointer StorageBuffer %22 +%58 = OpTypePointer StorageBuffer %28 +%61 = OpTypePointer StorageBuffer %23 +%62 = OpTypePointer StorageBuffer %6 +%65 = OpTypePointer StorageBuffer %29 +%68 = OpTypePointer StorageBuffer %21 +%69 = OpConstant %9 4 +%88 = OpTypePointer Function %4 +%92 = OpTypeVector %4 4 +%100 = OpTypePointer StorageBuffer %4 +%103 = OpConstant %9 64 +%38 = OpFunction %6 None %39 +%37 = OpFunctionParameter %31 +%36 = OpLabel +OpBranch %40 +%40 = OpLabel +%41 = OpLoad %6 %37 +OpReturnValue %41 OpFunctionEnd -%50 = OpFunction %2 None %51 -%44 = OpLabel -%41 = OpVariable %30 Function %5 -%42 = OpVariable %43 Function -%47 = OpLoad %9 %45 -OpBranch %52 -%52 = OpLabel -%53 = OpLoad %6 %41 -OpStore %41 %7 -%55 = OpAccessChain %54 %33 %15 -%56 = OpLoad %21 %55 -%58 = OpAccessChain %57 %33 %8 -%59 = OpLoad %27 %58 -%62 = OpAccessChain %61 %33 %15 %8 %15 -%63 = OpLoad %6 %62 -%65 = OpArrayLength %9 %33 4 -%66 = OpISub %9 %65 %10 -%68 = OpAccessChain %31 %33 %67 %66 -%69 = OpLoad %4 %68 -%70 = OpFunctionCall %6 %37 %41 -%71 = OpAccessChain %61 %33 %15 %16 %10 -OpStore %71 %7 -%72 = OpCompositeConstruct %22 %5 %5 %5 %5 -%73 = OpCompositeConstruct %22 %7 %7 %7 %7 -%74 = OpCompositeConstruct %22 %13 %13 %13 %13 -%75 = OpCompositeConstruct %22 %14 %14 %14 %14 -%76 = OpCompositeConstruct %21 %72 %73 %74 %75 -%77 = OpAccessChain %54 %33 %15 -OpStore %77 %76 -%78 = OpCompositeConstruct %26 %15 %15 -%79 = OpCompositeConstruct %26 %16 %16 -%80 = OpCompositeConstruct %27 %78 %79 -%81 = OpAccessChain %57 %33 %8 -OpStore %81 %80 -%82 = OpAccessChain %31 %33 %67 %16 -OpStore %82 %12 -%83 = OpConvertFToS %4 %63 -%84 = OpCompositeConstruct %32 %69 %83 %18 %19 %17 -OpStore %42 %84 -%85 = OpIAdd %9 %47 %16 -%87 = OpAccessChain %86 %42 %85 -OpStore %87 %20 -%88 = OpAccessChain %86 %42 %47 -%89 = OpLoad %4 %88 -%91 = OpCompositeConstruct %90 %89 %89 %89 %89 -%92 = OpConvertSToF %22 %91 -%93 = OpMatrixTimesVector %22 %56 %92 -OpStore %48 %93 +%51 = OpFunction %2 None %52 +%45 = OpLabel +%42 = OpVariable %31 Function %5 +%43 = OpVariable %44 Function +%48 = OpLoad %9 %46 +OpBranch %53 +%53 = OpLabel +%54 = OpLoad %6 %42 +OpStore %42 %7 +%56 = OpAccessChain %55 %34 %15 +%57 = OpLoad %22 %56 +%59 = OpAccessChain %58 %34 %8 +%60 = OpLoad %28 %59 +%63 = OpAccessChain %62 %34 %15 %8 %15 +%64 = OpLoad %6 %63 +%66 = OpArrayLength %9 %34 4 +%67 = OpISub %9 %66 %10 +%70 = OpAccessChain %32 %34 %69 %67 %15 +%71 = OpLoad %4 %70 +%72 = OpFunctionCall %6 %38 %42 +%73 = OpAccessChain %62 %34 %15 %16 %10 +OpStore %73 %7 +%74 = OpCompositeConstruct %23 %5 %5 %5 %5 +%75 = OpCompositeConstruct %23 %7 %7 %7 %7 +%76 = OpCompositeConstruct %23 %13 %13 %13 %13 +%77 = OpCompositeConstruct %23 %14 %14 %14 %14 +%78 = OpCompositeConstruct %22 %74 %75 %76 %77 +%79 = OpAccessChain %55 %34 %15 +OpStore %79 %78 +%80 = OpCompositeConstruct %27 %15 %15 +%81 = OpCompositeConstruct %27 %16 %16 +%82 = OpCompositeConstruct %28 %80 %81 +%83 = OpAccessChain %58 %34 %8 +OpStore %83 %82 +%84 = OpAccessChain %32 %34 %69 %16 %15 +OpStore %84 %12 +%85 = OpConvertFToS %4 %64 +%86 = OpCompositeConstruct %33 %71 %85 %18 %19 %17 +OpStore %43 %86 +%87 = OpIAdd %9 %48 %16 +%89 = OpAccessChain %88 %43 %87 +OpStore %89 %20 +%90 = OpAccessChain %88 %43 %48 +%91 = OpLoad %4 %90 +%93 = OpCompositeConstruct %92 %91 %91 %91 %91 +%94 = OpConvertSToF %23 %93 +%95 = OpMatrixTimesVector %23 %57 %94 +OpStore %49 %95 OpReturn OpFunctionEnd -%96 = OpFunction %2 None %51 -%95 = OpLabel -%94 = OpVariable %86 Function -OpBranch %97 +%98 = OpFunction %2 None %52 %97 = OpLabel -%99 = OpAccessChain %98 %33 %10 -%100 = OpAtomicLoad %4 %99 %12 %101 -%103 = OpAccessChain %98 %33 %10 -%102 = OpAtomicIAdd %4 %103 %12 %101 %17 -OpStore %94 %102 -%105 = OpAccessChain %98 %33 %10 -%104 = OpAtomicISub %4 %105 %12 %101 %17 -OpStore %94 %104 -%107 = OpAccessChain %98 %33 %10 -%106 = OpAtomicAnd %4 %107 %12 %101 %17 -OpStore %94 %106 -%109 = OpAccessChain %98 %33 %10 -%108 = OpAtomicOr %4 %109 %12 %101 %17 -OpStore %94 %108 -%111 = OpAccessChain %98 %33 %10 -%110 = OpAtomicXor %4 %111 %12 %101 %17 -OpStore %94 %110 -%113 = OpAccessChain %98 %33 %10 -%112 = OpAtomicSMin %4 %113 %12 %101 %17 -OpStore %94 %112 -%115 = OpAccessChain %98 %33 %10 -%114 = OpAtomicSMax %4 %115 %12 %101 %17 -OpStore %94 %114 -%117 = OpAccessChain %98 %33 %10 -%116 = OpAtomicExchange %4 %117 %12 %101 %17 -OpStore %94 %116 -%118 = OpAccessChain %98 %33 %10 -OpAtomicStore %118 %12 %101 %100 +%96 = OpVariable %88 Function +OpBranch %99 +%99 = OpLabel +%101 = OpAccessChain %100 %34 %10 +%102 = OpAtomicLoad %4 %101 %12 %103 +%105 = OpAccessChain %100 %34 %10 +%104 = OpAtomicIAdd %4 %105 %12 %103 %17 +OpStore %96 %104 +%107 = OpAccessChain %100 %34 %10 +%106 = OpAtomicISub %4 %107 %12 %103 %17 +OpStore %96 %106 +%109 = OpAccessChain %100 %34 %10 +%108 = OpAtomicAnd %4 %109 %12 %103 %17 +OpStore %96 %108 +%111 = OpAccessChain %100 %34 %10 +%110 = OpAtomicOr %4 %111 %12 %103 %17 +OpStore %96 %110 +%113 = OpAccessChain %100 %34 %10 +%112 = OpAtomicXor %4 %113 %12 %103 %17 +OpStore %96 %112 +%115 = OpAccessChain %100 %34 %10 +%114 = OpAtomicSMin %4 %115 %12 %103 %17 +OpStore %96 %114 +%117 = OpAccessChain %100 %34 %10 +%116 = OpAtomicSMax %4 %117 %12 %103 %17 +OpStore %96 %116 +%119 = OpAccessChain %100 %34 %10 +%118 = OpAtomicExchange %4 %119 %12 %103 %17 +OpStore %96 %118 +%120 = OpAccessChain %100 %34 %10 +OpAtomicStore %120 %12 %103 %102 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/246-collatz-comp.wgsl b/tests/out/wgsl/246-collatz-comp.wgsl index 45d08357cb..e983ee8dea 100644 --- a/tests/out/wgsl/246-collatz-comp.wgsl +++ b/tests/out/wgsl/246-collatz-comp.wgsl @@ -1,5 +1,5 @@ struct PrimeIndices { - indices: @stride(4) array; + indices: array; }; @group(0) @binding(0) diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 2e4ee4142a..28d29440f9 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -1,9 +1,13 @@ +struct AlignedWrapper { + value: i32; +}; + struct Bar { matrix: mat4x4; - matrix_array: @stride(16) array,2>; + matrix_array: array,2>; atom: atomic; - arr: @stride(8) array,2>; - data: @stride(8) array; + arr: array,2>; + data: array; }; @group(0) @binding(0) @@ -24,13 +28,13 @@ fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { let matrix = bar.matrix; let arr = bar.arr; let b = bar.matrix[3][0]; - let a = bar.data[(arrayLength((&bar.data)) - 2u)]; - let data_pointer = (&bar.data[0]); - let _e25 = read_from_private((&foo_1)); + let a = bar.data[(arrayLength((&bar.data)) - 2u)].value; + let data_pointer = (&bar.data[0].value); + let _e27 = read_from_private((&foo_1)); 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)); - bar.data[1] = 1; + bar.data[1].value = 1; c = array(a, i32(b), 3, 4, 5); c[(vi + 1u)] = 42; let value = c[vi]; diff --git a/tests/out/wgsl/bevy-pbr-frag.wgsl b/tests/out/wgsl/bevy-pbr-frag.wgsl index fa947b893e..b5913a6266 100644 --- a/tests/out/wgsl/bevy-pbr-frag.wgsl +++ b/tests/out/wgsl/bevy-pbr-frag.wgsl @@ -20,8 +20,8 @@ struct CameraPosition { struct Lights { AmbientColor: vec4; NumLights: vec4; - PointLights: @stride(48) array; - DirectionalLights: @stride(32) array; + PointLights: array; + DirectionalLights: array; }; struct StandardMaterial_base_color { diff --git a/tests/out/wgsl/boids.wgsl b/tests/out/wgsl/boids.wgsl index 7b1e9af1ea..b0f3180b29 100644 --- a/tests/out/wgsl/boids.wgsl +++ b/tests/out/wgsl/boids.wgsl @@ -14,7 +14,7 @@ struct SimParams { }; struct Particles { - particles: @stride(16) array; + particles: array; }; let NUM_PARTICLES: u32 = 1500u; diff --git a/tests/out/wgsl/collatz.wgsl b/tests/out/wgsl/collatz.wgsl index 054ecf3640..b4fb82c465 100644 --- a/tests/out/wgsl/collatz.wgsl +++ b/tests/out/wgsl/collatz.wgsl @@ -1,5 +1,5 @@ struct PrimeIndices { - data: @stride(4) array; + data: array; }; @group(0) @binding(0) diff --git a/tests/out/wgsl/constant-array-size-vert.wgsl b/tests/out/wgsl/constant-array-size-vert.wgsl index 9a767c8588..26e32eda7a 100644 --- a/tests/out/wgsl/constant-array-size-vert.wgsl +++ b/tests/out/wgsl/constant-array-size-vert.wgsl @@ -1,5 +1,5 @@ struct Data { - vecs: @stride(16) array,42u>; + vecs: array,42u>; }; @group(1) @binding(0) diff --git a/tests/out/wgsl/globals.wgsl b/tests/out/wgsl/globals.wgsl index 84705602af..6184f987c1 100644 --- a/tests/out/wgsl/globals.wgsl +++ b/tests/out/wgsl/globals.wgsl @@ -4,7 +4,7 @@ struct Foo { }; struct Dummy { - arr: @stride(8) array>; + arr: array>; }; let Foo_2: bool = true; diff --git a/tests/out/wgsl/pointers.wgsl b/tests/out/wgsl/pointers.wgsl index 0ac148e1ca..e06ecb0b15 100644 --- a/tests/out/wgsl/pointers.wgsl +++ b/tests/out/wgsl/pointers.wgsl @@ -1,5 +1,5 @@ struct DynamicArray { - arr: @stride(4) array; + arr: array; }; @group(0) @binding(0) diff --git a/tests/out/wgsl/shadow.wgsl b/tests/out/wgsl/shadow.wgsl index f6e9e50e66..d1e2ace3ee 100644 --- a/tests/out/wgsl/shadow.wgsl +++ b/tests/out/wgsl/shadow.wgsl @@ -9,7 +9,7 @@ struct Light { }; struct Lights { - data: @stride(96) array; + data: array; }; let c_ambient: vec3 = vec3(0.05000000074505806, 0.05000000074505806, 0.05000000074505806); diff --git a/tests/wgsl-errors.rs b/tests/wgsl-errors.rs index d6224ce8af..a04b50b804 100644 --- a/tests/wgsl-errors.rs +++ b/tests/wgsl-errors.rs @@ -369,22 +369,6 @@ fn unknown_conservative_depth() { ); } -#[test] -fn zero_array_stride() { - check( - r#" - type zero = @stride(0) array; - "#, - r#"error: array stride must not be zero - ┌─ wgsl:2:33 - │ -2 │ type zero = @stride(0) array; - │ ^ array stride must not be zero - -"#, - ); -} - #[test] fn struct_member_zero_size() { check( @@ -832,16 +816,6 @@ fn invalid_arrays() { }) } - check_validation_error! { - r#" - type Bad = @stride(2) array; - "#: - Err(naga::valid::ValidationError::Type { - error: naga::valid::TypeError::InsufficientArrayStride { stride: 2, base_size: 4 }, - .. - }) - } - check_validation_error! { "type Bad = array;", r#"