diff --git a/src/valid/function.rs b/src/valid/function.rs index 06aa27c84b..f792e2d045 100644 --- a/src/valid/function.rs +++ b/src/valid/function.rs @@ -993,12 +993,7 @@ impl super::Validator { #[cfg(feature = "validate")] for (index, argument) in fun.arguments.iter().enumerate() { match module.types[argument.ty].inner.pointer_space() { - Some( - crate::AddressSpace::Private - | crate::AddressSpace::Function - | crate::AddressSpace::WorkGroup, - ) - | None => {} + Some(crate::AddressSpace::Private | crate::AddressSpace::Function) | None => {} Some(other) => { return Err(FunctionError::InvalidArgumentPointerSpace { index, diff --git a/src/valid/type.rs b/src/valid/type.rs index f8ceb463c6..9589d676f5 100644 --- a/src/valid/type.rs +++ b/src/valid/type.rs @@ -164,12 +164,14 @@ fn check_member_layout( /// `TypeFlags::empty()`. /// /// Pointers passed as arguments to user-defined functions must be in the -/// `Function`, `Private`, or `Workgroup` storage space. +/// `Function` or `Private` address space. const fn ptr_space_argument_flag(space: crate::AddressSpace) -> TypeFlags { use crate::AddressSpace as As; match space { - As::Function | As::Private | As::WorkGroup => TypeFlags::ARGUMENT, - As::Uniform | As::Storage { .. } | As::Handle | As::PushConstant => TypeFlags::empty(), + As::Function | As::Private => TypeFlags::ARGUMENT, + As::Uniform | As::Storage { .. } | As::Handle | As::PushConstant | As::WorkGroup => { + TypeFlags::empty() + } } } @@ -316,7 +318,7 @@ impl super::Validator { return Err(TypeError::InvalidPointerBase(base)); } - // Runtime-sized values can only live in the `Storage` storage + // Runtime-sized values can only live in the `Storage` address // space, so it's useless to have a pointer to such a type in // any other space. // @@ -336,7 +338,7 @@ impl super::Validator { } } - // `Validator::validate_function` actually checks the storage + // `Validator::validate_function` actually checks the address // space of pointer arguments explicitly before checking the // `ARGUMENT` flag, to give better error messages. But it seems // best to set `ARGUMENT` accurately anyway. @@ -364,7 +366,7 @@ impl super::Validator { // `InvalidPointerBase` or `InvalidPointerToUnsized`. self.check_width(kind, width)?; - // `Validator::validate_function` actually checks the storage + // `Validator::validate_function` actually checks the address // space of pointer arguments explicitly before checking the // `ARGUMENT` flag, to give better error messages. But it seems // best to set `ARGUMENT` accurately anyway. diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index c9097f913b..956a694aaa 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -151,9 +151,7 @@ fn foo_frag() -> @location(0) vec4 { return vec4(0.0); } -var val: u32; - -fn assign_through_ptr_fn(p: ptr) { +fn assign_through_ptr_fn(p: ptr) { *p = 42u; } @@ -163,8 +161,9 @@ fn assign_array_through_ptr_fn(foo: ptr, 2>>) { @compute @workgroup_size(1) fn assign_through_ptr() { - var arr = array, 2>(vec4(6.0), vec4(7.0)); - + var val = 33u; assign_through_ptr_fn(&val); + + var arr = array, 2>(vec4(6.0), vec4(7.0)); assign_array_through_ptr_fn(&arr); } diff --git a/tests/out/analysis/access.info.ron b/tests/out/analysis/access.info.ron index 4a020e48c4..9151a25749 100644 --- a/tests/out/analysis/access.info.ron +++ b/tests/out/analysis/access.info.ron @@ -46,7 +46,6 @@ ("READ"), (""), (""), - (""), ], expressions: [ ( @@ -1144,7 +1143,6 @@ (""), (""), ("READ"), - (""), ], expressions: [ ( @@ -2414,7 +2412,6 @@ (""), (""), (""), - (""), ], expressions: [ ( @@ -2454,7 +2451,6 @@ (""), (""), (""), - (""), ], expressions: [ ( @@ -2503,7 +2499,6 @@ (""), (""), (""), - (""), ], expressions: [ ( @@ -2546,7 +2541,6 @@ (""), (""), (""), - (""), ], expressions: [ ( @@ -2638,7 +2632,6 @@ ("READ"), ("READ"), ("READ"), - (""), ], expressions: [ ( @@ -3302,7 +3295,6 @@ (""), ("WRITE"), (""), - (""), ], expressions: [ ( @@ -3736,9 +3728,32 @@ (""), (""), (""), - ("READ"), ], expressions: [ + ( + uniformity: ( + non_uniform_result: None, + requirements: (""), + ), + ref_count: 1, + assignable_global: None, + ty: Value(Scalar( + kind: Uint, + width: 4, + )), + ), + ( + uniformity: ( + non_uniform_result: Some(2), + requirements: (""), + ), + ref_count: 2, + assignable_global: None, + ty: Value(Pointer( + base: 1, + space: Function, + )), + ), ( uniformity: ( non_uniform_result: None, @@ -3800,7 +3815,7 @@ ), ( uniformity: ( - non_uniform_result: Some(6), + non_uniform_result: Some(8), requirements: (""), ), ref_count: 2, @@ -3810,18 +3825,6 @@ space: Function, )), ), - ( - uniformity: ( - non_uniform_result: None, - requirements: (""), - ), - ref_count: 1, - assignable_global: Some(6), - ty: Value(Pointer( - base: 1, - space: WorkGroup, - )), - ), ], sampling: [], dual_source_blending: false, diff --git a/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/tests/out/glsl/access.assign_through_ptr.Compute.glsl index ab5bd9a3fb..ad0ebb5401 100644 --- a/tests/out/glsl/access.assign_through_ptr.Compute.glsl +++ b/tests/out/glsl/access.assign_through_ptr.Compute.glsl @@ -19,8 +19,6 @@ struct Baz { struct MatCx2InArray { mat4x2 am[2]; }; -shared uint val; - float read_from_private(inout float foo_1) { float _e1 = foo_1; @@ -42,14 +40,11 @@ void assign_array_through_ptr_fn(inout vec4 foo_2[2]) { } void main() { - if (gl_LocalInvocationID == uvec3(0u)) { - val = 0u; - } - memoryBarrierShared(); - barrier(); + uint val = 0u; vec4 arr[2] = vec4[2](vec4(0.0), vec4(0.0)); - arr = vec4[2](vec4(6.0), vec4(7.0)); + val = 33u; assign_through_ptr_fn(val); + arr = vec4[2](vec4(6.0), vec4(7.0)); assign_array_through_ptr_fn(arr); return; } diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 465643bb44..74c12efb2d 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -81,7 +81,6 @@ RWByteAddressBuffer bar : register(u0); cbuffer baz : register(b1) { Baz baz; } RWByteAddressBuffer qux : register(u2); cbuffer nested_mat_cx2_ : register(b3) { MatCx2InArray nested_mat_cx2_; } -groupshared uint val; Baz ConstructBaz(float3x2 arg0) { Baz ret = (Baz)0; @@ -293,16 +292,14 @@ float4 foo_frag() : SV_Target0 } [numthreads(1, 1, 1)] -void assign_through_ptr(uint3 __local_invocation_id : SV_GroupThreadID) +void assign_through_ptr() { - if (all(__local_invocation_id == uint3(0u, 0u, 0u))) { - val = (uint)0; - } - GroupMemoryBarrierWithGroupSync(); + uint val = (uint)0; float4 arr[2] = (float4[2])0; - arr = Constructarray2_float4_((6.0).xxxx, (7.0).xxxx); + val = 33u; assign_through_ptr_fn(val); + arr = Constructarray2_float4_((6.0).xxxx, (7.0).xxxx); assign_array_through_ptr_fn(arr); return; } diff --git a/tests/out/ir/access.compact.ron b/tests/out/ir/access.compact.ron index b725765d46..019f178b01 100644 --- a/tests/out/ir/access.compact.ron +++ b/tests/out/ir/access.compact.ron @@ -279,7 +279,7 @@ name: None, inner: Pointer( base: 1, - space: WorkGroup, + space: Function, ), ), ( @@ -356,13 +356,6 @@ ty: 20, init: None, ), - ( - name: Some("val"), - space: WorkGroup, - binding: None, - ty: 1, - init: None, - ), ], const_expressions: [ Literal(U32(0)), @@ -2157,6 +2150,11 @@ arguments: [], result: None, local_variables: [ + ( + name: Some("val"), + ty: 1, + init: None, + ), ( name: Some("arr"), ty: 28, @@ -2164,51 +2162,56 @@ ), ], expressions: [ + Literal(U32(33)), + LocalVariable(1), Literal(F32(6.0)), Splat( size: Quad, - value: 1, + value: 3, ), Literal(F32(7.0)), Splat( size: Quad, - value: 3, + value: 5, ), Compose( ty: 28, components: [ - 2, 4, + 6, ], ), - LocalVariable(1), - GlobalVariable(6), + LocalVariable(2), ], named_expressions: {}, body: [ - Emit(( - start: 1, - end: 2, - )), - Emit(( - start: 3, - end: 5, - )), Store( - pointer: 6, - value: 5, + pointer: 2, + value: 1, ), Call( function: 5, arguments: [ - 7, + 2, ], result: None, ), + Emit(( + start: 3, + end: 4, + )), + Emit(( + start: 5, + end: 7, + )), + Store( + pointer: 8, + value: 7, + ), Call( function: 6, arguments: [ - 6, + 8, ], result: None, ), diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index a200e75a91..829801ad7c 100644 --- a/tests/out/ir/access.ron +++ b/tests/out/ir/access.ron @@ -312,7 +312,7 @@ name: None, inner: Pointer( base: 1, - space: WorkGroup, + space: Function, ), ), ( @@ -389,13 +389,6 @@ ty: 21, init: None, ), - ( - name: Some("val"), - space: WorkGroup, - binding: None, - ty: 1, - init: None, - ), ], const_expressions: [ Literal(U32(0)), @@ -2244,6 +2237,11 @@ arguments: [], result: None, local_variables: [ + ( + name: Some("val"), + ty: 1, + init: None, + ), ( name: Some("arr"), ty: 32, @@ -2251,51 +2249,56 @@ ), ], expressions: [ + Literal(U32(33)), + LocalVariable(1), Literal(F32(6.0)), Splat( size: Quad, - value: 1, + value: 3, ), Literal(F32(7.0)), Splat( size: Quad, - value: 3, + value: 5, ), Compose( ty: 32, components: [ - 2, 4, + 6, ], ), - LocalVariable(1), - GlobalVariable(6), + LocalVariable(2), ], named_expressions: {}, body: [ - Emit(( - start: 1, - end: 2, - )), - Emit(( - start: 3, - end: 5, - )), Store( - pointer: 6, - value: 5, + pointer: 2, + value: 1, ), Call( function: 5, arguments: [ - 7, + 2, ], result: None, ), + Emit(( + start: 3, + end: 4, + )), + Emit(( + start: 5, + end: 7, + )), + Store( + pointer: 8, + value: 7, + ), Call( function: 6, arguments: [ - 6, + 8, ], result: None, ), diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index c35139a8cf..2cc637a864 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -150,7 +150,7 @@ float test_arr_as_arg( } void assign_through_ptr_fn( - threadgroup uint& p + thread uint& p ) { p = 42u; return; @@ -215,16 +215,12 @@ fragment foo_fragOutput foo_frag( kernel void assign_through_ptr( - metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]] -, threadgroup uint& val ) { - if (metal::all(__local_invocation_id == metal::uint3(0u))) { - val = {}; - } - metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint val = {}; type_22 arr = {}; - arr = type_22 {metal::float4(6.0), metal::float4(7.0)}; + val = 33u; assign_through_ptr_fn(val); + arr = type_22 {metal::float4(6.0), metal::float4(7.0)}; assign_array_through_ptr_fn(arr); return; } diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index d5b6524ee7..b7815af872 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,16 +1,16 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 322 +; Bound: 311 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %232 "foo_vert" %227 %230 -OpEntryPoint Fragment %282 "foo_frag" %281 -OpEntryPoint GLCompute %302 "assign_through_ptr" %305 -OpExecutionMode %282 OriginUpperLeft -OpExecutionMode %302 LocalSize 1 1 1 +OpEntryPoint Vertex %231 "foo_vert" %226 %229 +OpEntryPoint Fragment %281 "foo_frag" %280 +OpEntryPoint GLCompute %303 "assign_through_ptr" +OpExecutionMode %281 OriginUpperLeft +OpExecutionMode %303 LocalSize 1 1 1 OpMemberName %6 0 "a" OpMemberName %6 1 "b" OpMemberName %6 2 "c" @@ -33,28 +33,28 @@ OpName %42 "bar" OpName %44 "baz" OpName %47 "qux" OpName %50 "nested_mat_cx2" -OpName %53 "val" -OpName %54 "idx" -OpName %57 "t" -OpName %61 "test_matrix_within_struct_accesses" -OpName %132 "idx" -OpName %133 "t" -OpName %137 "test_matrix_within_array_within_struct_accesses" -OpName %195 "foo" -OpName %196 "read_from_private" -OpName %201 "a" -OpName %202 "test_arr_as_arg" -OpName %208 "p" -OpName %209 "assign_through_ptr_fn" -OpName %214 "foo" -OpName %215 "assign_array_through_ptr_fn" -OpName %221 "foo" -OpName %223 "c2" -OpName %227 "vi" -OpName %232 "foo_vert" -OpName %282 "foo_frag" -OpName %299 "arr" -OpName %302 "assign_through_ptr" +OpName %53 "idx" +OpName %56 "t" +OpName %60 "test_matrix_within_struct_accesses" +OpName %131 "idx" +OpName %132 "t" +OpName %136 "test_matrix_within_array_within_struct_accesses" +OpName %194 "foo" +OpName %195 "read_from_private" +OpName %200 "a" +OpName %201 "test_arr_as_arg" +OpName %207 "p" +OpName %208 "assign_through_ptr_fn" +OpName %213 "foo" +OpName %214 "assign_array_through_ptr_fn" +OpName %220 "foo" +OpName %222 "c2" +OpName %226 "vi" +OpName %231 "foo_vert" +OpName %281 "foo_frag" +OpName %298 "val" +OpName %300 "arr" +OpName %303 "assign_through_ptr" OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %6 1 Offset 16 OpMemberDecorate %6 2 Offset 28 @@ -99,10 +99,9 @@ OpDecorate %50 DescriptorSet 0 OpDecorate %50 Binding 3 OpDecorate %51 Block OpMemberDecorate %51 0 Offset 0 -OpDecorate %227 BuiltIn VertexIndex -OpDecorate %230 BuiltIn Position -OpDecorate %281 Location 0 -OpDecorate %305 BuiltIn LocalInvocationId +OpDecorate %226 BuiltIn VertexIndex +OpDecorate %229 BuiltIn Position +OpDecorate %280 Location 0 %2 = OpTypeVoid %3 = OpTypeInt 32 0 %4 = OpTypeVector %3 3 @@ -134,7 +133,7 @@ OpDecorate %305 BuiltIn LocalInvocationId %29 = OpTypeArray %28 %30 %31 = OpTypeVector %10 4 %32 = OpTypeArray %5 %30 -%33 = OpTypePointer Workgroup %3 +%33 = OpTypePointer Function %3 %34 = OpTypeArray %31 %14 %35 = OpTypePointer Function %34 %36 = OpConstant %3 0 @@ -154,342 +153,326 @@ OpDecorate %305 BuiltIn LocalInvocationId %51 = OpTypeStruct %26 %52 = OpTypePointer Uniform %51 %50 = OpVariable %52 Uniform -%53 = OpVariable %33 Workgroup -%55 = OpTypePointer Function %5 -%56 = OpConstantNull %5 -%58 = OpTypePointer Function %22 -%59 = OpConstantNull %22 -%62 = OpTypeFunction %2 -%63 = OpTypePointer Uniform %22 -%65 = OpConstant %5 1 -%66 = OpConstant %10 1.0 -%67 = OpConstant %10 2.0 -%68 = OpConstant %10 3.0 -%69 = OpConstant %10 6.0 -%70 = OpConstant %10 5.0 -%71 = OpConstant %10 4.0 -%72 = OpConstant %10 9.0 -%73 = OpConstant %10 90.0 -%74 = OpConstant %10 10.0 -%75 = OpConstant %10 20.0 -%76 = OpConstant %10 30.0 -%77 = OpConstant %10 40.0 -%81 = OpTypePointer Uniform %21 -%84 = OpTypePointer Uniform %12 -%90 = OpTypePointer Uniform %10 -%91 = OpConstant %3 1 -%111 = OpTypePointer Function %21 -%117 = OpTypePointer Function %12 -%123 = OpTypePointer Function %10 -%134 = OpTypePointer Function %26 -%135 = OpConstantNull %26 -%138 = OpTypePointer Uniform %26 -%140 = OpConstantNull %25 -%141 = OpConstant %10 8.0 -%142 = OpConstant %10 7.0 -%146 = OpTypePointer Uniform %25 -%149 = OpTypePointer Uniform %24 -%172 = OpTypePointer Function %25 -%174 = OpTypePointer Function %24 -%197 = OpTypeFunction %10 %27 -%203 = OpTypeFunction %10 %29 -%210 = OpTypeFunction %2 %33 -%211 = OpConstant %3 42 -%216 = OpTypeFunction %2 %35 -%222 = OpConstantNull %10 -%224 = OpTypePointer Function %32 -%225 = OpConstantNull %32 -%228 = OpTypePointer Input %3 -%227 = OpVariable %228 Input -%231 = OpTypePointer Output %31 -%230 = OpVariable %231 Output -%234 = OpTypePointer StorageBuffer %23 -%237 = OpConstant %10 0.0 -%238 = OpConstant %3 3 -%239 = OpConstant %5 3 -%240 = OpConstant %5 4 -%241 = OpConstant %5 5 -%242 = OpConstant %5 42 -%243 = OpConstantNull %29 -%248 = OpTypePointer StorageBuffer %8 -%251 = OpTypePointer StorageBuffer %18 -%252 = OpConstant %3 4 -%255 = OpTypePointer StorageBuffer %9 -%256 = OpTypePointer StorageBuffer %10 -%259 = OpTypePointer StorageBuffer %19 -%262 = OpTypePointer StorageBuffer %7 -%263 = OpTypePointer StorageBuffer %5 -%275 = OpTypeVector %5 4 -%281 = OpVariable %231 Output -%284 = OpConstantNull %23 -%300 = OpConstantNull %34 -%304 = OpConstantNull %3 -%306 = OpTypePointer Input %4 -%305 = OpVariable %306 Input -%308 = OpConstantNull %4 -%310 = OpTypeBool -%309 = OpTypeVector %310 3 -%315 = OpConstant %3 264 -%61 = OpFunction %2 None %62 -%60 = OpLabel -%54 = OpVariable %55 Function %56 -%57 = OpVariable %58 Function %59 -%64 = OpAccessChain %63 %44 %36 -OpBranch %78 -%78 = OpLabel -OpStore %54 %65 -%79 = OpLoad %5 %54 -%80 = OpISub %5 %79 %65 -OpStore %54 %80 -%82 = OpAccessChain %81 %64 %36 -%83 = OpLoad %21 %82 -%85 = OpAccessChain %84 %64 %36 %36 -%86 = OpLoad %12 %85 -%87 = OpLoad %5 %54 -%88 = OpAccessChain %84 %64 %36 %87 -%89 = OpLoad %12 %88 -%92 = OpAccessChain %90 %64 %36 %36 %91 -%93 = OpLoad %10 %92 -%94 = OpLoad %5 %54 -%95 = OpAccessChain %90 %64 %36 %36 %94 -%96 = OpLoad %10 %95 -%97 = OpLoad %5 %54 -%98 = OpAccessChain %90 %64 %36 %97 %91 -%99 = OpLoad %10 %98 -%100 = OpLoad %5 %54 -%101 = OpLoad %5 %54 -%102 = OpAccessChain %90 %64 %36 %100 %101 -%103 = OpLoad %10 %102 +%54 = OpTypePointer Function %5 +%55 = OpConstantNull %5 +%57 = OpTypePointer Function %22 +%58 = OpConstantNull %22 +%61 = OpTypeFunction %2 +%62 = OpTypePointer Uniform %22 +%64 = OpConstant %5 1 +%65 = OpConstant %10 1.0 +%66 = OpConstant %10 2.0 +%67 = OpConstant %10 3.0 +%68 = OpConstant %10 6.0 +%69 = OpConstant %10 5.0 +%70 = OpConstant %10 4.0 +%71 = OpConstant %10 9.0 +%72 = OpConstant %10 90.0 +%73 = OpConstant %10 10.0 +%74 = OpConstant %10 20.0 +%75 = OpConstant %10 30.0 +%76 = OpConstant %10 40.0 +%80 = OpTypePointer Uniform %21 +%83 = OpTypePointer Uniform %12 +%89 = OpTypePointer Uniform %10 +%90 = OpConstant %3 1 +%110 = OpTypePointer Function %21 +%116 = OpTypePointer Function %12 +%122 = OpTypePointer Function %10 +%133 = OpTypePointer Function %26 +%134 = OpConstantNull %26 +%137 = OpTypePointer Uniform %26 +%139 = OpConstantNull %25 +%140 = OpConstant %10 8.0 +%141 = OpConstant %10 7.0 +%145 = OpTypePointer Uniform %25 +%148 = OpTypePointer Uniform %24 +%171 = OpTypePointer Function %25 +%173 = OpTypePointer Function %24 +%196 = OpTypeFunction %10 %27 +%202 = OpTypeFunction %10 %29 +%209 = OpTypeFunction %2 %33 +%210 = OpConstant %3 42 +%215 = OpTypeFunction %2 %35 +%221 = OpConstantNull %10 +%223 = OpTypePointer Function %32 +%224 = OpConstantNull %32 +%227 = OpTypePointer Input %3 +%226 = OpVariable %227 Input +%230 = OpTypePointer Output %31 +%229 = OpVariable %230 Output +%233 = OpTypePointer StorageBuffer %23 +%236 = OpConstant %10 0.0 +%237 = OpConstant %3 3 +%238 = OpConstant %5 3 +%239 = OpConstant %5 4 +%240 = OpConstant %5 5 +%241 = OpConstant %5 42 +%242 = OpConstantNull %29 +%247 = OpTypePointer StorageBuffer %8 +%250 = OpTypePointer StorageBuffer %18 +%251 = OpConstant %3 4 +%254 = OpTypePointer StorageBuffer %9 +%255 = OpTypePointer StorageBuffer %10 +%258 = OpTypePointer StorageBuffer %19 +%261 = OpTypePointer StorageBuffer %7 +%262 = OpTypePointer StorageBuffer %5 +%274 = OpTypeVector %5 4 +%280 = OpVariable %230 Output +%283 = OpConstantNull %23 +%299 = OpConstantNull %3 +%301 = OpConstantNull %34 +%304 = OpConstant %3 33 +%60 = OpFunction %2 None %61 +%59 = OpLabel +%53 = OpVariable %54 Function %55 +%56 = OpVariable %57 Function %58 +%63 = OpAccessChain %62 %44 %36 +OpBranch %77 +%77 = OpLabel +OpStore %53 %64 +%78 = OpLoad %5 %53 +%79 = OpISub %5 %78 %64 +OpStore %53 %79 +%81 = OpAccessChain %80 %63 %36 +%82 = OpLoad %21 %81 +%84 = OpAccessChain %83 %63 %36 %36 +%85 = OpLoad %12 %84 +%86 = OpLoad %5 %53 +%87 = OpAccessChain %83 %63 %36 %86 +%88 = OpLoad %12 %87 +%91 = OpAccessChain %89 %63 %36 %36 %90 +%92 = OpLoad %10 %91 +%93 = OpLoad %5 %53 +%94 = OpAccessChain %89 %63 %36 %36 %93 +%95 = OpLoad %10 %94 +%96 = OpLoad %5 %53 +%97 = OpAccessChain %89 %63 %36 %96 %90 +%98 = OpLoad %10 %97 +%99 = OpLoad %5 %53 +%100 = OpLoad %5 %53 +%101 = OpAccessChain %89 %63 %36 %99 %100 +%102 = OpLoad %10 %101 +%103 = OpCompositeConstruct %12 %65 %65 %104 = OpCompositeConstruct %12 %66 %66 %105 = OpCompositeConstruct %12 %67 %67 -%106 = OpCompositeConstruct %12 %68 %68 -%107 = OpCompositeConstruct %21 %104 %105 %106 -%108 = OpCompositeConstruct %22 %107 -OpStore %57 %108 -%109 = OpLoad %5 %54 -%110 = OpIAdd %5 %109 %65 -OpStore %54 %110 +%106 = OpCompositeConstruct %21 %103 %104 %105 +%107 = OpCompositeConstruct %22 %106 +OpStore %56 %107 +%108 = OpLoad %5 %53 +%109 = OpIAdd %5 %108 %64 +OpStore %53 %109 +%111 = OpCompositeConstruct %12 %68 %68 %112 = OpCompositeConstruct %12 %69 %69 %113 = OpCompositeConstruct %12 %70 %70 -%114 = OpCompositeConstruct %12 %71 %71 -%115 = OpCompositeConstruct %21 %112 %113 %114 -%116 = OpAccessChain %111 %57 %36 -OpStore %116 %115 -%118 = OpCompositeConstruct %12 %72 %72 -%119 = OpAccessChain %117 %57 %36 %36 -OpStore %119 %118 -%120 = OpLoad %5 %54 -%121 = OpCompositeConstruct %12 %73 %73 -%122 = OpAccessChain %117 %57 %36 %120 -OpStore %122 %121 -%124 = OpAccessChain %123 %57 %36 %36 %91 -OpStore %124 %74 -%125 = OpLoad %5 %54 -%126 = OpAccessChain %123 %57 %36 %36 %125 -OpStore %126 %75 -%127 = OpLoad %5 %54 -%128 = OpAccessChain %123 %57 %36 %127 %91 -OpStore %128 %76 -%129 = OpLoad %5 %54 -%130 = OpLoad %5 %54 -%131 = OpAccessChain %123 %57 %36 %129 %130 -OpStore %131 %77 +%114 = OpCompositeConstruct %21 %111 %112 %113 +%115 = OpAccessChain %110 %56 %36 +OpStore %115 %114 +%117 = OpCompositeConstruct %12 %71 %71 +%118 = OpAccessChain %116 %56 %36 %36 +OpStore %118 %117 +%119 = OpLoad %5 %53 +%120 = OpCompositeConstruct %12 %72 %72 +%121 = OpAccessChain %116 %56 %36 %119 +OpStore %121 %120 +%123 = OpAccessChain %122 %56 %36 %36 %90 +OpStore %123 %73 +%124 = OpLoad %5 %53 +%125 = OpAccessChain %122 %56 %36 %36 %124 +OpStore %125 %74 +%126 = OpLoad %5 %53 +%127 = OpAccessChain %122 %56 %36 %126 %90 +OpStore %127 %75 +%128 = OpLoad %5 %53 +%129 = OpLoad %5 %53 +%130 = OpAccessChain %122 %56 %36 %128 %129 +OpStore %130 %76 OpReturn OpFunctionEnd -%137 = OpFunction %2 None %62 -%136 = OpLabel -%132 = OpVariable %55 Function %56 -%133 = OpVariable %134 Function %135 -%139 = OpAccessChain %138 %50 %36 -OpBranch %143 -%143 = OpLabel -OpStore %132 %65 -%144 = OpLoad %5 %132 -%145 = OpISub %5 %144 %65 -OpStore %132 %145 -%147 = OpAccessChain %146 %139 %36 -%148 = OpLoad %25 %147 -%150 = OpAccessChain %149 %139 %36 %36 -%151 = OpLoad %24 %150 -%152 = OpAccessChain %84 %139 %36 %36 %36 -%153 = OpLoad %12 %152 -%154 = OpLoad %5 %132 -%155 = OpAccessChain %84 %139 %36 %36 %154 -%156 = OpLoad %12 %155 -%157 = OpAccessChain %90 %139 %36 %36 %36 %91 -%158 = OpLoad %10 %157 -%159 = OpLoad %5 %132 -%160 = OpAccessChain %90 %139 %36 %36 %36 %159 -%161 = OpLoad %10 %160 -%162 = OpLoad %5 %132 -%163 = OpAccessChain %90 %139 %36 %36 %162 %91 -%164 = OpLoad %10 %163 -%165 = OpLoad %5 %132 -%166 = OpLoad %5 %132 -%167 = OpAccessChain %90 %139 %36 %36 %165 %166 -%168 = OpLoad %10 %167 -%169 = OpCompositeConstruct %26 %140 -OpStore %133 %169 -%170 = OpLoad %5 %132 -%171 = OpIAdd %5 %170 %65 -OpStore %132 %171 -%173 = OpAccessChain %172 %133 %36 -OpStore %173 %140 +%136 = OpFunction %2 None %61 +%135 = OpLabel +%131 = OpVariable %54 Function %55 +%132 = OpVariable %133 Function %134 +%138 = OpAccessChain %137 %50 %36 +OpBranch %142 +%142 = OpLabel +OpStore %131 %64 +%143 = OpLoad %5 %131 +%144 = OpISub %5 %143 %64 +OpStore %131 %144 +%146 = OpAccessChain %145 %138 %36 +%147 = OpLoad %25 %146 +%149 = OpAccessChain %148 %138 %36 %36 +%150 = OpLoad %24 %149 +%151 = OpAccessChain %83 %138 %36 %36 %36 +%152 = OpLoad %12 %151 +%153 = OpLoad %5 %131 +%154 = OpAccessChain %83 %138 %36 %36 %153 +%155 = OpLoad %12 %154 +%156 = OpAccessChain %89 %138 %36 %36 %36 %90 +%157 = OpLoad %10 %156 +%158 = OpLoad %5 %131 +%159 = OpAccessChain %89 %138 %36 %36 %36 %158 +%160 = OpLoad %10 %159 +%161 = OpLoad %5 %131 +%162 = OpAccessChain %89 %138 %36 %36 %161 %90 +%163 = OpLoad %10 %162 +%164 = OpLoad %5 %131 +%165 = OpLoad %5 %131 +%166 = OpAccessChain %89 %138 %36 %36 %164 %165 +%167 = OpLoad %10 %166 +%168 = OpCompositeConstruct %26 %139 +OpStore %132 %168 +%169 = OpLoad %5 %131 +%170 = OpIAdd %5 %169 %64 +OpStore %131 %170 +%172 = OpAccessChain %171 %132 %36 +OpStore %172 %139 +%174 = OpCompositeConstruct %12 %140 %140 %175 = OpCompositeConstruct %12 %141 %141 -%176 = OpCompositeConstruct %12 %142 %142 +%176 = OpCompositeConstruct %12 %68 %68 %177 = OpCompositeConstruct %12 %69 %69 -%178 = OpCompositeConstruct %12 %70 %70 -%179 = OpCompositeConstruct %24 %175 %176 %177 %178 -%180 = OpAccessChain %174 %133 %36 %36 -OpStore %180 %179 -%181 = OpCompositeConstruct %12 %72 %72 -%182 = OpAccessChain %117 %133 %36 %36 %36 -OpStore %182 %181 -%183 = OpLoad %5 %132 -%184 = OpCompositeConstruct %12 %73 %73 -%185 = OpAccessChain %117 %133 %36 %36 %183 -OpStore %185 %184 -%186 = OpAccessChain %123 %133 %36 %36 %36 %91 -OpStore %186 %74 -%187 = OpLoad %5 %132 -%188 = OpAccessChain %123 %133 %36 %36 %36 %187 -OpStore %188 %75 -%189 = OpLoad %5 %132 -%190 = OpAccessChain %123 %133 %36 %36 %189 %91 -OpStore %190 %76 -%191 = OpLoad %5 %132 -%192 = OpLoad %5 %132 -%193 = OpAccessChain %123 %133 %36 %36 %191 %192 -OpStore %193 %77 +%178 = OpCompositeConstruct %24 %174 %175 %176 %177 +%179 = OpAccessChain %173 %132 %36 %36 +OpStore %179 %178 +%180 = OpCompositeConstruct %12 %71 %71 +%181 = OpAccessChain %116 %132 %36 %36 %36 +OpStore %181 %180 +%182 = OpLoad %5 %131 +%183 = OpCompositeConstruct %12 %72 %72 +%184 = OpAccessChain %116 %132 %36 %36 %182 +OpStore %184 %183 +%185 = OpAccessChain %122 %132 %36 %36 %36 %90 +OpStore %185 %73 +%186 = OpLoad %5 %131 +%187 = OpAccessChain %122 %132 %36 %36 %36 %186 +OpStore %187 %74 +%188 = OpLoad %5 %131 +%189 = OpAccessChain %122 %132 %36 %36 %188 %90 +OpStore %189 %75 +%190 = OpLoad %5 %131 +%191 = OpLoad %5 %131 +%192 = OpAccessChain %122 %132 %36 %36 %190 %191 +OpStore %192 %76 OpReturn OpFunctionEnd -%196 = OpFunction %10 None %197 -%195 = OpFunctionParameter %27 -%194 = OpLabel -OpBranch %198 -%198 = OpLabel -%199 = OpLoad %10 %195 -OpReturnValue %199 +%195 = OpFunction %10 None %196 +%194 = OpFunctionParameter %27 +%193 = OpLabel +OpBranch %197 +%197 = OpLabel +%198 = OpLoad %10 %194 +OpReturnValue %198 OpFunctionEnd -%202 = OpFunction %10 None %203 -%201 = OpFunctionParameter %29 -%200 = OpLabel -OpBranch %204 -%204 = OpLabel -%205 = OpCompositeExtract %28 %201 4 -%206 = OpCompositeExtract %10 %205 9 -OpReturnValue %206 +%201 = OpFunction %10 None %202 +%200 = OpFunctionParameter %29 +%199 = OpLabel +OpBranch %203 +%203 = OpLabel +%204 = OpCompositeExtract %28 %200 4 +%205 = OpCompositeExtract %10 %204 9 +OpReturnValue %205 OpFunctionEnd -%209 = OpFunction %2 None %210 -%208 = OpFunctionParameter %33 -%207 = OpLabel -OpBranch %212 -%212 = OpLabel -OpStore %208 %211 +%208 = OpFunction %2 None %209 +%207 = OpFunctionParameter %33 +%206 = OpLabel +OpBranch %211 +%211 = OpLabel +OpStore %207 %210 OpReturn OpFunctionEnd -%215 = OpFunction %2 None %216 -%214 = OpFunctionParameter %35 -%213 = OpLabel -OpBranch %217 -%217 = OpLabel +%214 = OpFunction %2 None %215 +%213 = OpFunctionParameter %35 +%212 = OpLabel +OpBranch %216 +%216 = OpLabel +%217 = OpCompositeConstruct %31 %65 %65 %65 %65 %218 = OpCompositeConstruct %31 %66 %66 %66 %66 -%219 = OpCompositeConstruct %31 %67 %67 %67 %67 -%220 = OpCompositeConstruct %34 %218 %219 -OpStore %214 %220 +%219 = OpCompositeConstruct %34 %217 %218 +OpStore %213 %219 OpReturn OpFunctionEnd -%232 = OpFunction %2 None %62 -%226 = OpLabel -%221 = OpVariable %27 Function %222 -%223 = OpVariable %224 Function %225 -%229 = OpLoad %3 %227 -%233 = OpAccessChain %63 %44 %36 -%235 = OpAccessChain %234 %47 %36 -%236 = OpAccessChain %138 %50 %36 -OpBranch %244 -%244 = OpLabel -OpStore %221 %237 -%245 = OpLoad %10 %221 -OpStore %221 %66 -%246 = OpFunctionCall %2 %61 -%247 = OpFunctionCall %2 %137 -%249 = OpAccessChain %248 %42 %36 -%250 = OpLoad %8 %249 -%253 = OpAccessChain %251 %42 %252 -%254 = OpLoad %18 %253 -%257 = OpAccessChain %256 %42 %36 %238 %36 -%258 = OpLoad %10 %257 -%260 = OpArrayLength %3 %42 5 -%261 = OpISub %3 %260 %14 -%264 = OpAccessChain %263 %42 %30 %261 %36 -%265 = OpLoad %5 %264 -%266 = OpLoad %23 %235 -%267 = OpFunctionCall %10 %196 %221 -%268 = OpConvertFToS %5 %258 -%269 = OpCompositeConstruct %32 %265 %268 %239 %240 %241 -OpStore %223 %269 -%270 = OpIAdd %3 %229 %91 -%271 = OpAccessChain %55 %223 %270 -OpStore %271 %242 -%272 = OpAccessChain %55 %223 %229 -%273 = OpLoad %5 %272 -%274 = OpFunctionCall %10 %202 %243 -%276 = OpCompositeConstruct %275 %273 %273 %273 %273 -%277 = OpConvertSToF %31 %276 -%278 = OpMatrixTimesVector %9 %250 %277 -%279 = OpCompositeConstruct %31 %278 %67 -OpStore %230 %279 +%231 = OpFunction %2 None %61 +%225 = OpLabel +%220 = OpVariable %27 Function %221 +%222 = OpVariable %223 Function %224 +%228 = OpLoad %3 %226 +%232 = OpAccessChain %62 %44 %36 +%234 = OpAccessChain %233 %47 %36 +%235 = OpAccessChain %137 %50 %36 +OpBranch %243 +%243 = OpLabel +OpStore %220 %236 +%244 = OpLoad %10 %220 +OpStore %220 %65 +%245 = OpFunctionCall %2 %60 +%246 = OpFunctionCall %2 %136 +%248 = OpAccessChain %247 %42 %36 +%249 = OpLoad %8 %248 +%252 = OpAccessChain %250 %42 %251 +%253 = OpLoad %18 %252 +%256 = OpAccessChain %255 %42 %36 %237 %36 +%257 = OpLoad %10 %256 +%259 = OpArrayLength %3 %42 5 +%260 = OpISub %3 %259 %14 +%263 = OpAccessChain %262 %42 %30 %260 %36 +%264 = OpLoad %5 %263 +%265 = OpLoad %23 %234 +%266 = OpFunctionCall %10 %195 %220 +%267 = OpConvertFToS %5 %257 +%268 = OpCompositeConstruct %32 %264 %267 %238 %239 %240 +OpStore %222 %268 +%269 = OpIAdd %3 %228 %90 +%270 = OpAccessChain %54 %222 %269 +OpStore %270 %241 +%271 = OpAccessChain %54 %222 %228 +%272 = OpLoad %5 %271 +%273 = OpFunctionCall %10 %201 %242 +%275 = OpCompositeConstruct %274 %272 %272 %272 %272 +%276 = OpConvertSToF %31 %275 +%277 = OpMatrixTimesVector %9 %249 %276 +%278 = OpCompositeConstruct %31 %277 %66 +OpStore %229 %278 OpReturn OpFunctionEnd -%282 = OpFunction %2 None %62 -%280 = OpLabel -%283 = OpAccessChain %234 %47 %36 -OpBranch %285 -%285 = OpLabel -%286 = OpAccessChain %256 %42 %36 %91 %14 -OpStore %286 %66 -%287 = OpCompositeConstruct %9 %237 %237 %237 +%281 = OpFunction %2 None %61 +%279 = OpLabel +%282 = OpAccessChain %233 %47 %36 +OpBranch %284 +%284 = OpLabel +%285 = OpAccessChain %255 %42 %36 %90 %14 +OpStore %285 %65 +%286 = OpCompositeConstruct %9 %236 %236 %236 +%287 = OpCompositeConstruct %9 %65 %65 %65 %288 = OpCompositeConstruct %9 %66 %66 %66 %289 = OpCompositeConstruct %9 %67 %67 %67 -%290 = OpCompositeConstruct %9 %68 %68 %68 -%291 = OpCompositeConstruct %8 %287 %288 %289 %290 -%292 = OpAccessChain %248 %42 %36 -OpStore %292 %291 -%293 = OpCompositeConstruct %17 %36 %36 -%294 = OpCompositeConstruct %17 %91 %91 -%295 = OpCompositeConstruct %18 %293 %294 -%296 = OpAccessChain %251 %42 %252 -OpStore %296 %295 -%297 = OpAccessChain %263 %42 %30 %91 %36 -OpStore %297 %65 -OpStore %283 %284 -%298 = OpCompositeConstruct %31 %237 %237 %237 %237 -OpStore %281 %298 +%290 = OpCompositeConstruct %8 %286 %287 %288 %289 +%291 = OpAccessChain %247 %42 %36 +OpStore %291 %290 +%292 = OpCompositeConstruct %17 %36 %36 +%293 = OpCompositeConstruct %17 %90 %90 +%294 = OpCompositeConstruct %18 %292 %293 +%295 = OpAccessChain %250 %42 %251 +OpStore %295 %294 +%296 = OpAccessChain %262 %42 %30 %90 %36 +OpStore %296 %64 +OpStore %282 %283 +%297 = OpCompositeConstruct %31 %236 %236 %236 %236 +OpStore %280 %297 OpReturn OpFunctionEnd -%302 = OpFunction %2 None %62 -%301 = OpLabel -%299 = OpVariable %35 Function %300 -OpBranch %303 -%303 = OpLabel -%307 = OpLoad %4 %305 -%311 = OpIEqual %309 %307 %308 -%312 = OpAll %310 %311 -OpSelectionMerge %313 None -OpBranchConditional %312 %314 %313 -%314 = OpLabel -OpStore %53 %304 -OpBranch %313 -%313 = OpLabel -OpControlBarrier %14 %14 %315 -OpBranch %316 -%316 = OpLabel -%317 = OpCompositeConstruct %31 %69 %69 %69 %69 -%318 = OpCompositeConstruct %31 %142 %142 %142 %142 -%319 = OpCompositeConstruct %34 %317 %318 -OpStore %299 %319 -%320 = OpFunctionCall %2 %209 %53 -%321 = OpFunctionCall %2 %215 %299 +%303 = OpFunction %2 None %61 +%302 = OpLabel +%298 = OpVariable %33 Function %299 +%300 = OpVariable %35 Function %301 +OpBranch %305 +%305 = OpLabel +OpStore %298 %304 +%306 = OpFunctionCall %2 %208 %298 +%307 = OpCompositeConstruct %31 %68 %68 %68 %68 +%308 = OpCompositeConstruct %31 %141 %141 %141 %141 +%309 = OpCompositeConstruct %34 %307 %308 +OpStore %300 %309 +%310 = OpFunctionCall %2 %214 %300 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index c01c4247c5..3f60dfd839 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -34,7 +34,6 @@ var baz: Baz; var qux: vec2; @group(0) @binding(3) var nested_mat_cx2_: MatCx2InArray; -var val: u32; fn test_matrix_within_struct_accesses() { var idx: i32; @@ -121,7 +120,7 @@ fn test_arr_as_arg(a: array, 5>) -> f32 { return a[4][9]; } -fn assign_through_ptr_fn(p: ptr) { +fn assign_through_ptr_fn(p: ptr) { (*p) = 42u; return; } @@ -167,10 +166,12 @@ fn foo_frag() -> @location(0) vec4 { @compute @workgroup_size(1, 1, 1) fn assign_through_ptr() { + var val: u32; var arr: array, 2>; - arr = array, 2>(vec4(6.0), vec4(7.0)); + val = 33u; assign_through_ptr_fn((&val)); + arr = array, 2>(vec4(6.0), vec4(7.0)); assign_array_through_ptr_fn((&arr)); return; } diff --git a/tests/wgsl-errors.rs b/tests/wgsl-errors.rs index 48318f2804..7f5ab3a821 100644 --- a/tests/wgsl-errors.rs +++ b/tests/wgsl-errors.rs @@ -1064,7 +1064,7 @@ fn invalid_functions() { }) } - // Pointers of these storage classes cannot be passed as arguments. + // Pointers of these address spaces cannot be passed as arguments. check_validation! { "fn unacceptable_ptr_space(arg: ptr>) { }": Err(naga::valid::ValidationError::Function { @@ -1078,7 +1078,6 @@ fn invalid_functions() { }) if function_name == "unacceptable_ptr_space" && argument_name == "arg" } - check_validation! { "fn unacceptable_ptr_space(arg: ptr) { }": Err(naga::valid::ValidationError::Function { @@ -1092,6 +1091,19 @@ fn invalid_functions() { }) if function_name == "unacceptable_ptr_space" && argument_name == "arg" } + check_validation! { + "fn unacceptable_ptr_space(arg: ptr) { }": + Err(naga::valid::ValidationError::Function { + name: function_name, + source: naga::valid::FunctionError::InvalidArgumentPointerSpace { + index: 0, + name: argument_name, + space: naga::AddressSpace::WorkGroup, + }, + .. + }) + if function_name == "unacceptable_ptr_space" && argument_name == "arg" + } check_validation! { "