From 21324b8feaecd5808ea32997dc4a64cbef8640ab Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Mon, 27 Sep 2021 18:49:28 -0400 Subject: [PATCH] Update WGSL grammar for pointer access. (#1312) * Update WGSL grammar for pointer access. Comes with a small test, which revealed a number of issues in the backends. * Validate pointer arguments to functions to only have function/private/workgroup classes. Comes with a small test. Also, "pointer-access.spv" test is temporarily disabled. --- src/back/glsl/mod.rs | 15 +- src/back/hlsl/writer.rs | 11 +- src/back/msl/writer.rs | 13 +- src/back/wgsl/writer.rs | 20 +- src/front/wgsl/lexer.rs | 24 +- src/front/wgsl/mod.rs | 17 +- src/front/wgsl/tests.rs | 4 +- src/valid/function.rs | 19 ++ tests/in/access.wgsl | 8 + tests/in/image.wgsl | 2 +- tests/out/glsl/access.atomics.Compute.glsl | 5 + tests/out/glsl/access.foo.Vertex.glsl | 6 + tests/out/hlsl/access.hlsl | 7 + tests/out/msl/access.msl | 14 +- tests/out/spv/access.spvasm | 296 +++++++++++---------- tests/out/wgsl/access.wgsl | 7 + tests/out/wgsl/image.wgsl | 4 +- tests/snapshots.rs | 17 +- tests/wgsl-errors.rs | 20 +- 19 files changed, 312 insertions(+), 197 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 4cb694ccd8..9b171fd9e9 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -715,7 +715,9 @@ impl<'a, W: Write> Writer<'a, W> { TypeInner::Pointer { .. } | TypeInner::Struct { .. } | TypeInner::Image { .. } - | TypeInner::Sampler { .. } => unreachable!(), + | TypeInner::Sampler { .. } => { + return Err(Error::Custom(format!("Unable to write type {:?}", inner))) + } } Ok(()) @@ -1332,7 +1334,14 @@ impl<'a, W: Write> Writer<'a, W> { // This is where we can generate intermediate constants for some expression types. Statement::Emit(ref range) => { for handle in range.clone() { - let expr_name = if let Some(name) = ctx.named_expressions.get(&handle) { + let info = &ctx.info[handle]; + let ptr_class = info.ty.inner_with(&self.module.types).pointer_class(); + let expr_name = if ptr_class.is_some() { + // GLSL can't save a pointer-valued expression in a variable, + // but we shouldn't ever need to: they should never be named expressions, + // and none of the expression types flagged by bake_ref_count can be pointer-valued. + None + } else if let Some(name) = ctx.named_expressions.get(&handle) { // Front end provides names for all variables at the start of writing. // But we write them to step by step. We need to recache them // Otherwise, we could accidentally write variable name instead of full expression. @@ -1340,7 +1349,7 @@ impl<'a, W: Write> Writer<'a, W> { Some(self.namer.call_unique(name)) } else { let min_ref_count = ctx.expressions[handle].bake_ref_count(); - if min_ref_count <= ctx.info[handle].ref_count { + if min_ref_count <= info.ref_count { Some(format!("{}{}", super::BAKE_PREFIX, handle.index())) } else { None diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index d48daf1aae..339523f40a 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1057,7 +1057,14 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { match *stmt { Statement::Emit(ref range) => { for handle in range.clone() { - let expr_name = if let Some(name) = func_ctx.named_expressions.get(&handle) { + let info = &func_ctx.info[handle]; + let ptr_class = info.ty.inner_with(&module.types).pointer_class(); + let expr_name = if ptr_class.is_some() { + // HLSL can't save a pointer-valued expression in a variable, + // but we shouldn't ever need to: they should never be named expressions, + // and none of the expression types flagged by bake_ref_count can be pointer-valued. + None + } else if let Some(name) = func_ctx.named_expressions.get(&handle) { // Front end provides names for all variables at the start of writing. // But we write them to step by step. We need to recache them // Otherwise, we could accidentally write variable name instead of full expression. @@ -1065,7 +1072,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Some(self.namer.call_unique(name)) } else { let min_ref_count = func_ctx.expressions[handle].bake_ref_count(); - if min_ref_count <= func_ctx.info[handle].ref_count { + if min_ref_count <= info.ref_count { Some(format!("_expr{}", handle.index())) } else { None diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 8f0368908b..78d56aa0c7 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1355,7 +1355,7 @@ impl Writer { )?; } TypeResolution::Value(ref other) => { - log::error!("Type {:?} isn't a known local", other); + log::warn!("Type {:?} isn't a known local", other); //TEMP! return Err(Error::FeatureNotImplemented("weird local type".to_string())); } } @@ -1383,7 +1383,14 @@ impl Writer { match *statement { crate::Statement::Emit(ref range) => { for handle in range.clone() { - let expr_name = if let Some(name) = + let info = &context.expression.info[handle]; + let ptr_class = info + .ty + .inner_with(&context.expression.module.types) + .pointer_class(); + let expr_name = if ptr_class.is_some() { + None // don't bake pointer expressions (just yet) + } else if let Some(name) = context.expression.function.named_expressions.get(&handle) { // Front end provides names for all variables at the start of writing. @@ -1394,7 +1401,7 @@ impl Writer { } else { let min_ref_count = context.expression.function.expressions[handle].bake_ref_count(); - if min_ref_count <= context.expression.info[handle].ref_count { + if min_ref_count <= info.ref_count { Some(format!("{}{}", back::BAKE_PREFIX, handle.index())) } else { None diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index c95aa33755..c108d90e8f 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -525,10 +525,13 @@ impl Writer { "storage_", "", storage_format_str(format), - if access.contains(crate::StorageAccess::STORE) { - ",write" + if access.contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE) + { + ",read_write" + } else if access.contains(crate::StorageAccess::LOAD) { + ",read" } else { - "" + ",write" }, ), }; @@ -639,6 +642,7 @@ impl Writer { inner ))); } + write!(self.out, ">")?; } _ => { return Err(Error::Unimplemented(format!( @@ -666,6 +670,7 @@ impl Writer { match *stmt { Statement::Emit(ref range) => { for handle in range.clone() { + let info = &func_ctx.info[handle]; let expr_name = if let Some(name) = func_ctx.named_expressions.get(&handle) { // Front end provides names for all variables at the start of writing. // But we write them to step by step. We need to recache them @@ -682,8 +687,7 @@ impl Writer { | Expression::ImageSample { .. } => true, _ => false, }; - if min_ref_count <= func_ctx.info[handle].ref_count || required_baking_expr - { + if min_ref_count <= info.ref_count || required_baking_expr { // If expression contains unsupported builtin we should skip it if let Expression::Load { pointer } = func_ctx.expressions[handle] { if let Expression::AccessIndex { base, index } = @@ -809,8 +813,8 @@ impl Writer { } let func_name = &self.names[&NameKey::Function(function)]; write!(self.out, "{}(", func_name)?; - for (index, argument) in arguments.iter().enumerate() { - self.write_expr(module, *argument, func_ctx)?; + for (index, &argument) in arguments.iter().enumerate() { + self.write_expr(module, argument, func_ctx)?; // Only write a comma if isn't the last element if index != arguments.len().saturating_sub(1) { // The leading space is for readability only @@ -1199,14 +1203,12 @@ impl Writer { self.write_expr(module, right, func_ctx)?; write!(self.out, ")")?; } - // TODO: copy-paste from glsl-out Expression::Access { base, index } => { self.write_expr_with_indirection(module, base, func_ctx, indirection)?; write!(self.out, "[")?; self.write_expr(module, index, func_ctx)?; write!(self.out, "]")? } - // TODO: copy-paste from glsl-out Expression::AccessIndex { base, index } => { let base_ty_res = &func_ctx.info[base].ty; let mut resolved = base_ty_res.inner_with(&module.types); diff --git a/src/front/wgsl/lexer.rs b/src/front/wgsl/lexer.rs index 54a2660ea6..91f2a5ead2 100644 --- a/src/front/wgsl/lexer.rs +++ b/src/front/wgsl/lexer.rs @@ -558,24 +558,24 @@ impl<'a> Lexer<'a> { Ok(pair) } - // TODO relocate storage texture specifics + pub(super) fn next_storage_access(&mut self) -> Result> { + let (ident, span) = self.next_ident_with_span()?; + match ident { + "read" => Ok(crate::StorageAccess::LOAD), + "write" => Ok(crate::StorageAccess::STORE), + "read_write" => Ok(crate::StorageAccess::LOAD | crate::StorageAccess::STORE), + _ => Err(Error::UnknownAccess(span)), + } + } + pub(super) fn next_format_generic( &mut self, ) -> Result<(crate::StorageFormat, crate::StorageAccess), Error<'a>> { self.expect(Token::Paren('<'))?; let (ident, ident_span) = self.next_ident_with_span()?; let format = conv::map_storage_format(ident, ident_span)?; - let access = if self.skip(Token::Separator(',')) { - let (raw, span) = self.next_ident_with_span()?; - match raw { - "read" => crate::StorageAccess::LOAD, - "write" => crate::StorageAccess::STORE, - "read_write" => crate::StorageAccess::all(), - _ => return Err(Error::UnknownAccess(span)), - } - } else { - crate::StorageAccess::LOAD - }; + self.expect(Token::Separator(','))?; + let access = self.next_storage_access()?; self.expect(Token::Paren('>'))?; Ok((format, access)) } diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index 95b0e174c4..800851a744 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -2587,13 +2587,7 @@ impl Parser { class = Some(match class_str { "storage" => { let access = if lexer.skip(Token::Separator(',')) { - let (ident, span) = lexer.next_ident_with_span()?; - match ident { - "read" => crate::StorageAccess::LOAD, - "write" => crate::StorageAccess::STORE, - "read_write" => crate::StorageAccess::all(), - _ => return Err(Error::UnknownAccess(span)), - } + lexer.next_storage_access()? } else { // defaulting to `read` crate::StorageAccess::LOAD @@ -2836,9 +2830,16 @@ impl Parser { "ptr" => { lexer.expect_generic_paren('<')?; let (ident, span) = lexer.next_ident_with_span()?; - let class = conv::map_storage_class(ident, span)?; + let mut class = conv::map_storage_class(ident, span)?; lexer.expect(Token::Separator(','))?; let (base, _access) = self.parse_type_decl(lexer, None, type_arena, const_arena)?; + if let crate::StorageClass::Storage { ref mut access } = class { + *access = if lexer.skip(Token::Separator(',')) { + lexer.next_storage_access()? + } else { + crate::StorageAccess::LOAD + }; + } lexer.expect_generic_paren('>')?; crate::TypeInner::Pointer { base, class } } diff --git a/src/front/wgsl/tests.rs b/src/front/wgsl/tests.rs index 432b39f466..da1670c145 100644 --- a/src/front/wgsl/tests.rs +++ b/src/front/wgsl/tests.rs @@ -92,7 +92,7 @@ fn parse_types() { parse_str("var t: texture_cube_array;").unwrap(); parse_str("var t: texture_multisampled_2d;").unwrap(); parse_str("var t: texture_storage_1d;").unwrap(); - parse_str("var t: texture_storage_3d;").unwrap(); + parse_str("var t: texture_storage_3d;").unwrap(); } #[test] @@ -305,7 +305,7 @@ fn parse_texture_load() { .unwrap(); parse_str( " - var t: texture_storage_1d_array; + var t: texture_storage_1d_array; fn foo() { let r: vec4 = textureLoad(t, 10, 2); } diff --git a/src/valid/function.rs b/src/valid/function.rs index 792b600e24..23d4c00743 100644 --- a/src/valid/function.rs +++ b/src/valid/function.rs @@ -76,6 +76,12 @@ pub enum FunctionError { }, #[error("Argument '{name}' at index {index} has a type that can't be passed into functions.")] InvalidArgumentType { index: usize, name: String }, + #[error("Argument '{name}' at index {index} is a pointer of class {class:?}, which can't be passed into functions.")] + InvalidArgumentPointerClass { + index: usize, + name: String, + class: crate::StorageClass, + }, #[error("There are instructions after `return`/`break`/`continue`")] InstructionsAfterReturn, #[error("The `break` is used outside of a `loop` or `switch` context")] @@ -696,6 +702,19 @@ impl super::Validator { name: argument.name.clone().unwrap_or_default(), }); } + match module.types[argument.ty].inner.pointer_class() { + Some(crate::StorageClass::Private) + | Some(crate::StorageClass::Function) + | Some(crate::StorageClass::WorkGroup) + | None => {} + Some(other) => { + return Err(FunctionError::InvalidArgumentPointerClass { + index, + name: argument.name.clone().unwrap_or_default(), + class: other, + }) + } + } } self.valid_expression_set.clear(); diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index 8ab506c276..f4af84913d 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -11,6 +11,10 @@ struct Bar { [[group(0), binding(0)]] var bar: Bar; +fn read_from_private(foo: ptr) -> f32 { + return *foo; +} + [[stage(vertex)]] fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { var foo: f32 = 0.0; @@ -25,6 +29,10 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { let b = bar.matrix[index].x; let a = bar.data[arrayLength(&bar.data) - 2u]; + // test pointer types + let pointer: ptr = &bar.data[0]; + 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)); diff --git a/tests/in/image.wgsl b/tests/in/image.wgsl index 74378d4fe9..c4621a407f 100644 --- a/tests/in/image.wgsl +++ b/tests/in/image.wgsl @@ -5,7 +5,7 @@ var image_multisampled_src: texture_multisampled_2d; [[group(0), binding(4)]] var image_depth_multisampled_src: texture_depth_multisampled_2d; [[group(0), binding(1)]] -var image_storage_src: texture_storage_2d; +var image_storage_src: texture_storage_2d; [[group(0), binding(5)]] var image_array_src: texture_2d_array; [[group(0), binding(6)]] diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index 2af1e78275..958cd5a6ac 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -13,6 +13,11 @@ buffer Bar_block_0Cs { } _group_0_binding_0; +float read_from_private(inout float foo2) { + float _e2 = foo2; + return _e2; +} + void main() { int tmp = 0; int value = _group_0_binding_0.atom; diff --git a/tests/out/glsl/access.foo.Vertex.glsl b/tests/out/glsl/access.foo.Vertex.glsl index 9cf0b27aff..97248f5465 100644 --- a/tests/out/glsl/access.foo.Vertex.glsl +++ b/tests/out/glsl/access.foo.Vertex.glsl @@ -11,6 +11,11 @@ buffer Bar_block_0Vs { } _group_0_binding_0; +float read_from_private(inout float foo2) { + float _e2 = foo2; + return _e2; +} + void main() { uint vi = uint(gl_VertexID); float foo1 = 0.0; @@ -21,6 +26,7 @@ void main() { uvec2 arr[2] = _group_0_binding_0.arr; float b = _group_0_binding_0.matrix[3][0]; int a = _group_0_binding_0.data[(uint(_group_0_binding_0.data.length()) - 2u)]; + float _e25 = read_from_private(foo1); _group_0_binding_0.matrix[1][2] = 1.0; _group_0_binding_0.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); _group_0_binding_0.arr = uvec2[2](uvec2(0u), uvec2(1u)); diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index e9af3b1101..fe8c701ca6 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -1,6 +1,12 @@ RWByteAddressBuffer bar : register(u0); +float read_from_private(inout float foo2) +{ + float _expr2 = foo2; + return _expr2; +} + uint NagaBufferLengthRW(RWByteAddressBuffer buffer) { uint ret; @@ -19,6 +25,7 @@ float4 foo(uint vi : SV_VertexID) : SV_Position uint2 arr[2] = {asuint(bar.Load2(72+0)), asuint(bar.Load2(72+8))}; float b = asfloat(bar.Load(0+48+0)); int a = asint(bar.Load((((NagaBufferLengthRW(bar) - 88) / 4) - 2u)*4+88)); + const float _e25 = read_from_private(foo1); 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)); diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index dd9622a60d..44134f29d8 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -17,10 +17,17 @@ struct Bar { type3 arr; type5 data; }; -struct type9 { +struct type11 { int inner[5]; }; +float read_from_private( + thread float& foo2 +) { + float _e2 = foo2; + return _e2; +} + struct fooInput { }; struct fooOutput { @@ -32,17 +39,18 @@ vertex fooOutput foo( , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { float foo1 = 0.0; - type9 c; + type11 c; float baz = foo1; foo1 = 1.0; metal::float4x4 matrix = bar.matrix; type3 arr = bar.arr; float b = bar.matrix[3].x; int a = bar.data[(1 + (_buffer_sizes.size0 - 88 - 4) / 4) - 2u]; + float _e25 = read_from_private(foo1); 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] = type3 {metal::uint2(0u), metal::uint2(1u)}.inner[_i]; - for(int _i=0; _i<5; ++_i) c.inner[_i] = type9 {a, static_cast(b), 3, 4, 5}.inner[_i]; + for(int _i=0; _i<5; ++_i) c.inner[_i] = type11 {a, static_cast(b), 3, 4, 5}.inner[_i]; c.inner[vi + 1u] = 42; int value = c.inner[vi]; return fooOutput { matrix * static_cast(metal::int4(value)) }; diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index f143347e5f..3bdea548bd 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,41 +1,42 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 106 +; Bound: 114 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %39 "foo" %34 %37 -OpEntryPoint GLCompute %83 "atomics" -OpExecutionMode %83 LocalSize 1 1 1 +OpEntryPoint Vertex %47 "foo" %42 %45 +OpEntryPoint GLCompute %91 "atomics" +OpExecutionMode %91 LocalSize 1 1 1 OpSource GLSL 450 -OpMemberName %25 0 "matrix" -OpMemberName %25 1 "atom" -OpMemberName %25 2 "arr" -OpMemberName %25 3 "data" -OpName %25 "Bar" -OpName %27 "bar" -OpName %29 "foo" -OpName %31 "c" -OpName %34 "vi" -OpName %39 "foo" -OpName %81 "tmp" -OpName %83 "atomics" -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 72 -OpMemberDecorate %25 3 Offset 88 -OpDecorate %26 ArrayStride 4 -OpDecorate %27 DescriptorSet 0 -OpDecorate %27 Binding 0 -OpDecorate %34 BuiltIn VertexIndex -OpDecorate %37 BuiltIn Position +OpMemberName %26 0 "matrix" +OpMemberName %26 1 "atom" +OpMemberName %26 2 "arr" +OpMemberName %26 3 "data" +OpName %26 "Bar" +OpName %30 "bar" +OpName %34 "read_from_private" +OpName %38 "foo" +OpName %39 "c" +OpName %42 "vi" +OpName %47 "foo" +OpName %89 "tmp" +OpName %91 "atomics" +OpDecorate %24 ArrayStride 8 +OpDecorate %25 ArrayStride 4 +OpDecorate %26 Block +OpMemberDecorate %26 0 Offset 0 +OpMemberDecorate %26 0 ColMajor +OpMemberDecorate %26 0 MatrixStride 16 +OpMemberDecorate %26 1 Offset 64 +OpMemberDecorate %26 2 Offset 72 +OpMemberDecorate %26 3 Offset 88 +OpDecorate %29 ArrayStride 4 +OpDecorate %30 DescriptorSet 0 +OpDecorate %30 Binding 0 +OpDecorate %42 BuiltIn VertexIndex +OpDecorate %45 BuiltIn Position %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 @@ -45,120 +46,131 @@ OpDecorate %37 BuiltIn Position %9 = OpTypeInt 32 0 %8 = OpConstant %9 3 %10 = OpConstant %9 2 -%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 %4 %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 -%50 = OpTypePointer StorageBuffer %6 -%53 = OpTypePointer StorageBuffer %24 -%56 = OpTypePointer StorageBuffer %4 -%73 = OpTypePointer Function %4 -%77 = OpTypeVector %4 4 -%85 = OpTypePointer StorageBuffer %4 -%88 = OpConstant %9 64 -%39 = OpFunction %2 None %40 -%33 = OpLabel -%29 = OpVariable %30 Function %5 -%31 = OpVariable %32 Function -%36 = OpLoad %9 %34 -OpBranch %41 +%11 = OpConstant %4 0 +%12 = OpConstant %4 1 +%13 = OpConstant %6 2.0 +%14 = OpConstant %6 3.0 +%15 = OpConstant %9 0 +%16 = OpConstant %9 1 +%17 = OpConstant %4 5 +%18 = OpConstant %4 3 +%19 = OpConstant %4 4 +%20 = OpConstant %4 42 +%22 = OpTypeVector %6 4 +%21 = OpTypeMatrix %22 4 +%23 = OpTypeVector %9 2 +%24 = OpTypeArray %23 %3 +%25 = OpTypeRuntimeArray %4 +%26 = OpTypeStruct %21 %4 %24 %25 +%27 = OpTypePointer Function %6 +%28 = OpTypePointer StorageBuffer %4 +%29 = OpTypeArray %4 %17 +%31 = OpTypePointer StorageBuffer %26 +%30 = OpVariable %31 StorageBuffer +%35 = OpTypeFunction %6 %27 +%40 = OpTypePointer Function %29 +%43 = OpTypePointer Input %9 +%42 = OpVariable %43 Input +%46 = OpTypePointer Output %22 +%45 = OpVariable %46 Output +%48 = OpTypeFunction %2 +%51 = OpTypePointer StorageBuffer %21 +%54 = OpTypePointer StorageBuffer %24 +%57 = OpTypePointer StorageBuffer %22 +%58 = OpTypePointer StorageBuffer %6 +%61 = OpTypePointer StorageBuffer %25 +%81 = OpTypePointer Function %4 +%85 = OpTypeVector %4 4 +%93 = OpTypePointer StorageBuffer %4 +%96 = OpConstant %9 64 +%34 = OpFunction %6 None %35 +%33 = OpFunctionParameter %27 +%32 = OpLabel +OpBranch %36 +%36 = OpLabel +%37 = OpLoad %6 %33 +OpReturnValue %37 +OpFunctionEnd +%47 = OpFunction %2 None %48 %41 = OpLabel -%42 = OpLoad %6 %29 -OpStore %29 %7 -%44 = OpAccessChain %43 %27 %14 -%45 = OpLoad %20 %44 -%47 = OpAccessChain %46 %27 %10 -%48 = OpLoad %23 %47 -%51 = OpAccessChain %50 %27 %14 %8 %14 -%52 = OpLoad %6 %51 -%54 = OpArrayLength %9 %27 3 -%55 = OpISub %9 %54 %10 -%57 = OpAccessChain %56 %27 %8 %55 -%58 = OpLoad %4 %57 -%59 = OpAccessChain %50 %27 %14 %15 %10 -OpStore %59 %7 -%60 = OpCompositeConstruct %21 %5 %5 %5 %5 -%61 = OpCompositeConstruct %21 %7 %7 %7 %7 -%62 = OpCompositeConstruct %21 %12 %12 %12 %12 -%63 = OpCompositeConstruct %21 %13 %13 %13 %13 -%64 = OpCompositeConstruct %20 %60 %61 %62 %63 -%65 = OpAccessChain %43 %27 %14 -OpStore %65 %64 -%66 = OpCompositeConstruct %22 %14 %14 -%67 = OpCompositeConstruct %22 %15 %15 -%68 = OpCompositeConstruct %23 %66 %67 -%69 = OpAccessChain %46 %27 %10 -OpStore %69 %68 -%70 = OpConvertFToS %4 %52 -%71 = OpCompositeConstruct %26 %58 %70 %17 %18 %16 -OpStore %31 %71 -%72 = OpIAdd %9 %36 %15 -%74 = OpAccessChain %73 %31 %72 -OpStore %74 %19 -%75 = OpAccessChain %73 %31 %36 -%76 = OpLoad %4 %75 -%78 = OpCompositeConstruct %77 %76 %76 %76 %76 -%79 = OpConvertSToF %21 %78 -%80 = OpMatrixTimesVector %21 %45 %79 -OpStore %37 %80 +%38 = OpVariable %27 Function %5 +%39 = OpVariable %40 Function +%44 = OpLoad %9 %42 +OpBranch %49 +%49 = OpLabel +%50 = OpLoad %6 %38 +OpStore %38 %7 +%52 = OpAccessChain %51 %30 %15 +%53 = OpLoad %21 %52 +%55 = OpAccessChain %54 %30 %10 +%56 = OpLoad %24 %55 +%59 = OpAccessChain %58 %30 %15 %8 %15 +%60 = OpLoad %6 %59 +%62 = OpArrayLength %9 %30 3 +%63 = OpISub %9 %62 %10 +%64 = OpAccessChain %28 %30 %8 %63 +%65 = OpLoad %4 %64 +%66 = OpFunctionCall %6 %34 %38 +%67 = OpAccessChain %58 %30 %15 %16 %10 +OpStore %67 %7 +%68 = OpCompositeConstruct %22 %5 %5 %5 %5 +%69 = OpCompositeConstruct %22 %7 %7 %7 %7 +%70 = OpCompositeConstruct %22 %13 %13 %13 %13 +%71 = OpCompositeConstruct %22 %14 %14 %14 %14 +%72 = OpCompositeConstruct %21 %68 %69 %70 %71 +%73 = OpAccessChain %51 %30 %15 +OpStore %73 %72 +%74 = OpCompositeConstruct %23 %15 %15 +%75 = OpCompositeConstruct %23 %16 %16 +%76 = OpCompositeConstruct %24 %74 %75 +%77 = OpAccessChain %54 %30 %10 +OpStore %77 %76 +%78 = OpConvertFToS %4 %60 +%79 = OpCompositeConstruct %29 %65 %78 %18 %19 %17 +OpStore %39 %79 +%80 = OpIAdd %9 %44 %16 +%82 = OpAccessChain %81 %39 %80 +OpStore %82 %20 +%83 = OpAccessChain %81 %39 %44 +%84 = OpLoad %4 %83 +%86 = OpCompositeConstruct %85 %84 %84 %84 %84 +%87 = OpConvertSToF %22 %86 +%88 = OpMatrixTimesVector %22 %53 %87 +OpStore %45 %88 OpReturn OpFunctionEnd -%83 = OpFunction %2 None %40 -%82 = OpLabel -%81 = OpVariable %73 Function -OpBranch %84 -%84 = OpLabel -%86 = OpAccessChain %85 %27 %15 -%87 = OpAtomicLoad %4 %86 %11 %88 -%90 = OpAccessChain %85 %27 %15 -%89 = OpAtomicIAdd %4 %90 %11 %88 %16 -OpStore %81 %89 -%92 = OpAccessChain %85 %27 %15 -%91 = OpAtomicISub %4 %92 %11 %88 %16 -OpStore %81 %91 -%94 = OpAccessChain %85 %27 %15 -%93 = OpAtomicAnd %4 %94 %11 %88 %16 -OpStore %81 %93 -%96 = OpAccessChain %85 %27 %15 -%95 = OpAtomicOr %4 %96 %11 %88 %16 -OpStore %81 %95 -%98 = OpAccessChain %85 %27 %15 -%97 = OpAtomicXor %4 %98 %11 %88 %16 -OpStore %81 %97 -%100 = OpAccessChain %85 %27 %15 -%99 = OpAtomicSMin %4 %100 %11 %88 %16 -OpStore %81 %99 -%102 = OpAccessChain %85 %27 %15 -%101 = OpAtomicSMax %4 %102 %11 %88 %16 -OpStore %81 %101 -%104 = OpAccessChain %85 %27 %15 -%103 = OpAtomicExchange %4 %104 %11 %88 %16 -OpStore %81 %103 -%105 = OpAccessChain %85 %27 %15 -OpAtomicStore %105 %11 %88 %87 +%91 = OpFunction %2 None %48 +%90 = OpLabel +%89 = OpVariable %81 Function +OpBranch %92 +%92 = OpLabel +%94 = OpAccessChain %93 %30 %16 +%95 = OpAtomicLoad %4 %94 %12 %96 +%98 = OpAccessChain %93 %30 %16 +%97 = OpAtomicIAdd %4 %98 %12 %96 %17 +OpStore %89 %97 +%100 = OpAccessChain %93 %30 %16 +%99 = OpAtomicISub %4 %100 %12 %96 %17 +OpStore %89 %99 +%102 = OpAccessChain %93 %30 %16 +%101 = OpAtomicAnd %4 %102 %12 %96 %17 +OpStore %89 %101 +%104 = OpAccessChain %93 %30 %16 +%103 = OpAtomicOr %4 %104 %12 %96 %17 +OpStore %89 %103 +%106 = OpAccessChain %93 %30 %16 +%105 = OpAtomicXor %4 %106 %12 %96 %17 +OpStore %89 %105 +%108 = OpAccessChain %93 %30 %16 +%107 = OpAtomicSMin %4 %108 %12 %96 %17 +OpStore %89 %107 +%110 = OpAccessChain %93 %30 %16 +%109 = OpAtomicSMax %4 %110 %12 %96 %17 +OpStore %89 %109 +%112 = OpAccessChain %93 %30 %16 +%111 = OpAtomicExchange %4 %112 %12 %96 %17 +OpStore %89 %111 +%113 = OpAccessChain %93 %30 %16 +OpAtomicStore %113 %12 %96 %95 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 1439604190..11ef27c132 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -9,6 +9,11 @@ struct Bar { [[group(0), binding(0)]] var bar: Bar; +fn read_from_private(foo2: ptr) -> f32 { + let e2: f32 = (*foo2); + return e2; +} + [[stage(vertex)]] fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { var foo1: f32 = 0.0; @@ -20,6 +25,8 @@ fn foo([[builtin(vertex_index)]] vi: u32) -> [[builtin(position)]] vec4 { let arr: array,2> = bar.arr; let b: f32 = bar.matrix[3][0]; let a: i32 = bar.data[(arrayLength((&bar.data)) - 2u)]; + let pointer1: ptr = (&bar.data[0]); + let e25: f32 = read_from_private((&foo1)); 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)); diff --git a/tests/out/wgsl/image.wgsl b/tests/out/wgsl/image.wgsl index 64e3cdae82..504a630bba 100644 --- a/tests/out/wgsl/image.wgsl +++ b/tests/out/wgsl/image.wgsl @@ -5,11 +5,11 @@ var image_multisampled_src: texture_multisampled_2d; [[group(0), binding(4)]] var image_depth_multisampled_src: texture_depth_multisampled_2d; [[group(0), binding(1)]] -var image_storage_src: texture_storage_2d; +var image_storage_src: texture_storage_2d; [[group(0), binding(5)]] var image_array_src: texture_2d_array; [[group(0), binding(6)]] -var image_dup_src: texture_storage_1d; +var image_dup_src: texture_storage_1d; [[group(0), binding(2)]] var image_dst: texture_storage_1d; [[group(0), binding(0)]] diff --git a/tests/snapshots.rs b/tests/snapshots.rs index b34aadf842..35ed0cf131 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -267,7 +267,8 @@ fn write_output_msl( allow_point_size: true, }; - let (string, tr_info) = msl::write_string(module, info, options, &pipeline_options).unwrap(); + let (string, tr_info) = + msl::write_string(module, info, options, &pipeline_options).expect("Metal write failed"); for (ep, result) in module.entry_points.iter().zip(tr_info.entry_point_names) { if let Err(error) = result { @@ -308,9 +309,9 @@ fn write_output_glsl( }; let mut buffer = String::new(); - let mut writer = - glsl::Writer::new(&mut buffer, module, info, options, &pipeline_options).unwrap(); - writer.write().unwrap(); + let mut writer = glsl::Writer::new(&mut buffer, module, info, options, &pipeline_options) + .expect("GLSL init failed"); + writer.write().expect("GLSL write failed"); fs::write( destination.join(format!("glsl/{}.{}.{:?}.glsl", file_name, ep_name, stage)), @@ -344,7 +345,7 @@ fn write_output_hlsl( let mut buffer = String::new(); let mut writer = hlsl::Writer::new(&mut buffer, options); - let reflection_info = writer.write(module, info).unwrap(); + let reflection_info = writer.write(module, info).expect("HLSL write failed"); fs::write(destination.join(format!("hlsl/{}.hlsl", file_name)), buffer).unwrap(); @@ -417,7 +418,7 @@ fn write_output_wgsl( ) { use naga::back::wgsl; - let string = wgsl::write_string(module, info).unwrap(); + let string = wgsl::write_string(module, info).expect("WGSL write failed"); fs::write(destination.join(format!("wgsl/{}.wgsl", file_name)), string).unwrap(); } @@ -572,8 +573,8 @@ fn convert_spv_inverse_hyperbolic_trig_functions() { } #[cfg(all(feature = "spv-in", feature = "spv-out"))] -#[test] -fn convert_spv_pointer_access() { +//#[test] //TODO: https://github.com/gfx-rs/naga/issues/1432 +fn _convert_spv_pointer_access() { convert_spv("pointer-access", true, Targets::SPIRV); } diff --git a/tests/wgsl-errors.rs b/tests/wgsl-errors.rs index f1b3fc55b0..8ccfdfa0ed 100644 --- a/tests/wgsl-errors.rs +++ b/tests/wgsl-errors.rs @@ -677,14 +677,30 @@ fn invalid_functions() { if function_name == "unacceptable_unsized" && argument_name == "arg" } - // A *valid* way to pass an unsized value. check_validation_error! { " struct Unsized { data: array; }; - fn acceptable_ptr_to_unsized(okay: ptr) { } + fn acceptable_pointer_to_unsized(arg: ptr) { } ": Ok(_) } + + check_validation_error! { + " + struct Unsized { data: array; }; + fn unacceptable_uniform_class(arg: ptr) { } + ": + Err(naga::valid::ValidationError::Function { + name: function_name, + error: naga::valid::FunctionError::InvalidArgumentPointerClass { + index: 0, + name: argument_name, + class: naga::StorageClass::Uniform, + }, + .. + }) + if function_name == "unacceptable_uniform_class" && argument_name == "arg" + } } #[test]