From dd2d53814de01ac163936cb2209dc33593e99de2 Mon Sep 17 00:00:00 2001 From: Jamie Nicol Date: Thu, 6 Feb 2025 10:09:01 +0000 Subject: [PATCH] [naga wgsl-in] Do not eagerly concretize global const declarations of abstract types Instead allow the const to be converted and each time it is const evaluated as part of another expression. This allows an abstract const to be used as a different type depending on the context. A consequence of this is that abstract types may now find their way to the validation stage, which we don't want. We therefore additionally now ensure that the compact pass removes global constants of abstract types. This will have no *functional* effect on shaders generated by the backends, as the expressions belonging to the abstract consts in the IR will not actually be used, as any usage in the input shader will have been const-evaluated away. Certain unused const declarations will now be removed, however, as can be seen by the effect on the snapshot outputs. --- CHANGELOG.md | 1 + naga/src/compact/mod.rs | 21 ++- naga/src/front/wgsl/lower/mod.rs | 27 ++-- .../out/glsl/constructors.main.Compute.glsl | 2 - naga/tests/out/hlsl/constructors.hlsl | 14 +- naga/tests/out/ir/const_assert.compact.ron | 28 +--- naga/tests/out/ir/const_assert.ron | 28 +--- naga/tests/out/ir/local-const.compact.ron | 59 ++++---- naga/tests/out/ir/local-const.ron | 59 ++++---- naga/tests/out/msl/abstract-types-const.msl | 8 -- .../out/msl/abstract-types-operators.msl | 1 - naga/tests/out/msl/constructors.msl | 10 +- .../tests/out/spv/abstract-types-const.spvasm | 18 +-- naga/tests/out/spv/constructors.spvasm | 127 +++++++++--------- naga/tests/out/wgsl/abstract-types-const.wgsl | 8 -- .../out/wgsl/abstract-types-operators.wgsl | 1 - naga/tests/out/wgsl/const_assert.wgsl | 3 - naga/tests/out/wgsl/constructors.wgsl | 2 - naga/tests/out/wgsl/local-const.wgsl | 4 +- 19 files changed, 178 insertions(+), 243 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 2779fab13a..160f99a356 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -182,6 +182,7 @@ By @brodycj in [#6924](https://github.com/gfx-rs/wgpu/pull/6924). - Error if structs have two fields with the same name. By @SparkyPotato in [#7088](https://github.com/gfx-rs/wgpu/pull/7088). - Forward '--keep-coordinate-space' flag to GLSL backend in naga-cli. By @cloone8 in [#7206](https://github.com/gfx-rs/wgpu/pull/7206). - Allow template lists to have a trailing comma. By @KentSlaney in [#7142](https://github.com/gfx-rs/wgpu/pull/7142). +- Allow WGSL const declarations to have abstract types. By @jamienicol in [#7055](https://github.com/gfx-rs/wgpu/pull/7055). #### General diff --git a/naga/src/compact/mod.rs b/naga/src/compact/mod.rs index c9a7c93977..25835984eb 100644 --- a/naga/src/compact/mod.rs +++ b/naga/src/compact/mod.rs @@ -110,14 +110,27 @@ pub fn compact(module: &mut crate::Module) { log::trace!("tracing special types"); module_tracer.trace_special_types(&module.special_types); - // We treat all named constants as used by definition. + // We treat all named constants as used by definition, unless they have an + // abstract type as we do not want those reaching the validator. log::trace!("tracing named constants"); for (handle, constant) in module.constants.iter() { if constant.name.is_some() { log::trace!("tracing constant {:?}", constant.name.as_ref().unwrap()); - module_tracer.constants_used.insert(handle); - module_tracer.types_used.insert(constant.ty); - module_tracer.global_expressions_used.insert(constant.init); + // If the type is an array (of an array, etc) then we must check whether the + // type of the innermost array's base type is abstract. + let mut ty = constant.ty; + while let crate::TypeInner::Array { base, .. } = module.types[ty].inner { + ty = base; + } + if !module.types[ty] + .inner + .scalar() + .is_some_and(|s| s.is_abstract()) + { + module_tracer.constants_used.insert(handle); + module_tracer.types_used.insert(constant.ty); + module_tracer.global_expressions_used.insert(constant.init); + } } } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 70900e0db3..6a4851850b 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1072,8 +1072,13 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const())) .transpose()?; - let (ty, initializer) = - self.type_and_init(v.name, v.init, explicit_ty, &mut ctx.as_override())?; + let (ty, initializer) = self.type_and_init( + v.name, + v.init, + explicit_ty, + false, + &mut ctx.as_override(), + )?; let binding = if let Some(ref binding) = v.binding { Some(crate::ResourceBinding { @@ -1106,7 +1111,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { .transpose()?; let (ty, init) = - self.type_and_init(c.name, Some(c.init), explicit_ty, &mut ectx)?; + self.type_and_init(c.name, Some(c.init), explicit_ty, true, &mut ectx)?; let init = init.expect("Global const must have init"); let handle = ctx.module.constants.append( @@ -1128,7 +1133,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let mut ectx = ctx.as_override(); - let (ty, init) = self.type_and_init(o.name, o.init, explicit_ty, &mut ectx)?; + let (ty, init) = + self.type_and_init(o.name, o.init, explicit_ty, false, &mut ectx)?; let id = o.id.map(|id| self.const_u32(id, &mut ctx.as_const())) @@ -1201,6 +1207,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { name: ast::Ident<'source>, init: Option>>, explicit_ty: Option>, + allow_abstract: bool, ectx: &mut ExpressionContext<'source, '_, '_>, ) -> Result<(Handle, Option>), Error<'source>> { let ty; @@ -1234,9 +1241,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { initializer = Some(init); } (Some(init), None) => { - let concretized = self.expression(init, ectx)?; - ty = ectx.register_type(concretized)?; - initializer = Some(concretized); + let mut init = self.expression_for_abstract(init, ectx)?; + if !allow_abstract { + init = ectx.concretize(init)?; + } + ty = ectx.register_type(init)?; + initializer = Some(init); } (None, Some(explicit_ty)) => { ty = explicit_ty; @@ -1481,7 +1491,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let mut ectx = ctx.as_expression(block, &mut emitter); let (ty, initializer) = - self.type_and_init(v.name, v.init, explicit_ty, &mut ectx)?; + self.type_and_init(v.name, v.init, explicit_ty, false, &mut ectx)?; let (const_initializer, initializer) = { match initializer { @@ -1544,6 +1554,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { c.name, Some(c.init), explicit_ty, + false, &mut ectx.as_const(), )?; let init = init.expect("Local const must have init"); diff --git a/naga/tests/out/glsl/constructors.main.Compute.glsl b/naga/tests/out/glsl/constructors.main.Compute.glsl index cff178f200..ef31b76d9b 100644 --- a/naga/tests/out/glsl/constructors.main.Compute.glsl +++ b/naga/tests/out/glsl/constructors.main.Compute.glsl @@ -9,7 +9,6 @@ struct Foo { vec4 a; int b; }; -const vec3 const2_ = vec3(0.0, 1.0, 2.0); const mat2x2 const3_ = mat2x2(vec2(0.0, 1.0), vec2(2.0, 3.0)); const mat2x2 const4_[1] = mat2x2[1](mat2x2(vec2(0.0, 1.0), vec2(2.0, 3.0))); const bool cz0_ = false; @@ -20,7 +19,6 @@ const uvec2 cz4_ = uvec2(0u); const mat2x2 cz5_ = mat2x2(0.0); const Foo cz6_[3] = Foo[3](Foo(vec4(0.0), 0), Foo(vec4(0.0), 0), Foo(vec4(0.0), 0)); const Foo cz7_ = Foo(vec4(0.0), 0); -const int cp3_[4] = int[4](0, 1, 2, 3); void main() { diff --git a/naga/tests/out/hlsl/constructors.hlsl b/naga/tests/out/hlsl/constructors.hlsl index 5cc3702630..e21548b7a1 100644 --- a/naga/tests/out/hlsl/constructors.hlsl +++ b/naga/tests/out/hlsl/constructors.hlsl @@ -12,12 +12,6 @@ ret_Constructarray1_float2x2_ Constructarray1_float2x2_(float2x2 arg0) { return ret; } -typedef int ret_Constructarray4_int_[4]; -ret_Constructarray4_int_ Constructarray4_int_(int arg0, int arg1, int arg2, int arg3) { - int ret[4] = { arg0, arg1, arg2, arg3 }; - return ret; -} - bool ZeroValuebool() { return (bool)0; } @@ -51,7 +45,6 @@ Foo ZeroValueFoo() { return (Foo)0; } -static const float3 const2_ = float3(0.0, 1.0, 2.0); static const float2x2 const3_ = float2x2(float2(0.0, 1.0), float2(2.0, 3.0)); static const float2x2 const4_[1] = Constructarray1_float2x2_(float2x2(float2(0.0, 1.0), float2(2.0, 3.0))); static const bool cz0_ = ZeroValuebool(); @@ -62,7 +55,6 @@ static const uint2 cz4_ = ZeroValueuint2(); static const float2x2 cz5_ = ZeroValuefloat2x2(); static const Foo cz6_[3] = ZeroValuearray3_Foo_(); static const Foo cz7_ = ZeroValueFoo(); -static const int cp3_[4] = Constructarray4_int_(int(0), int(1), int(2), int(3)); Foo ConstructFoo(float4 arg0, int arg1) { Foo ret = (Foo)0; @@ -71,6 +63,12 @@ Foo ConstructFoo(float4 arg0, int arg1) { return ret; } +typedef int ret_Constructarray4_int_[4]; +ret_Constructarray4_int_ Constructarray4_int_(int arg0, int arg1, int arg2, int arg3) { + int ret[4] = { arg0, arg1, arg2, arg3 }; + return ret; +} + float2x3 ZeroValuefloat2x3() { return (float2x3)0; } diff --git a/naga/tests/out/ir/const_assert.compact.ron b/naga/tests/out/ir/const_assert.compact.ron index 03c540b601..15fa331fb6 100644 --- a/naga/tests/out/ir/const_assert.compact.ron +++ b/naga/tests/out/ir/const_assert.compact.ron @@ -1,36 +1,14 @@ ( - types: [ - ( - name: None, - inner: Scalar(( - kind: Sint, - width: 4, - )), - ), - ], + types: [], special_types: ( ray_desc: None, ray_intersection: None, predeclared_types: {}, ), - constants: [ - ( - name: Some("x"), - ty: 0, - init: 0, - ), - ( - name: Some("y"), - ty: 0, - init: 1, - ), - ], + constants: [], overrides: [], global_variables: [], - global_expressions: [ - Literal(I32(1)), - Literal(I32(2)), - ], + global_expressions: [], functions: [ ( name: Some("foo"), diff --git a/naga/tests/out/ir/const_assert.ron b/naga/tests/out/ir/const_assert.ron index 03c540b601..15fa331fb6 100644 --- a/naga/tests/out/ir/const_assert.ron +++ b/naga/tests/out/ir/const_assert.ron @@ -1,36 +1,14 @@ ( - types: [ - ( - name: None, - inner: Scalar(( - kind: Sint, - width: 4, - )), - ), - ], + types: [], special_types: ( ray_desc: None, ray_intersection: None, predeclared_types: {}, ), - constants: [ - ( - name: Some("x"), - ty: 0, - init: 0, - ), - ( - name: Some("y"), - ty: 0, - init: 1, - ), - ], + constants: [], overrides: [], global_variables: [], - global_expressions: [ - Literal(I32(1)), - Literal(I32(2)), - ], + global_expressions: [], functions: [ ( name: Some("foo"), diff --git a/naga/tests/out/ir/local-const.compact.ron b/naga/tests/out/ir/local-const.compact.ron index fc3dc81e71..53fa73819a 100644 --- a/naga/tests/out/ir/local-const.compact.ron +++ b/naga/tests/out/ir/local-const.compact.ron @@ -39,52 +39,27 @@ ), constants: [ ( - name: Some("ga"), + name: Some("gb"), ty: 0, init: 0, ), - ( - name: Some("gb"), - ty: 0, - init: 1, - ), ( name: Some("gc"), ty: 1, - init: 2, + init: 1, ), ( name: Some("gd"), ty: 2, - init: 3, - ), - ( - name: Some("ge"), - ty: 3, - init: 4, - ), - ( - name: Some("gf"), - ty: 2, - init: 5, + init: 2, ), ], overrides: [], global_variables: [], global_expressions: [ - Literal(I32(4)), Literal(I32(4)), Literal(U32(4)), Literal(F32(4.0)), - Compose( - ty: 3, - components: [ - 0, - 0, - 0, - ], - ), - Literal(F32(2.0)), ], functions: [ ( @@ -106,12 +81,22 @@ ], ), Literal(F32(2.0)), + Literal(I32(4)), Constant(0), Constant(1), Constant(2), - Constant(3), - Constant(4), - Constant(5), + Literal(I32(4)), + Literal(I32(4)), + Literal(I32(4)), + Compose( + ty: 3, + components: [ + 10, + 11, + 12, + ], + ), + Literal(F32(2.0)), ], named_expressions: { 0: "a", @@ -124,14 +109,22 @@ 7: "bg", 8: "cg", 9: "dg", - 10: "eg", - 11: "fg", + 13: "eg", + 14: "fg", }, body: [ Emit(( start: 4, end: 5, )), + Emit(( + start: 0, + end: 0, + )), + Emit(( + start: 13, + end: 14, + )), Return( value: None, ), diff --git a/naga/tests/out/ir/local-const.ron b/naga/tests/out/ir/local-const.ron index fc3dc81e71..53fa73819a 100644 --- a/naga/tests/out/ir/local-const.ron +++ b/naga/tests/out/ir/local-const.ron @@ -39,52 +39,27 @@ ), constants: [ ( - name: Some("ga"), + name: Some("gb"), ty: 0, init: 0, ), - ( - name: Some("gb"), - ty: 0, - init: 1, - ), ( name: Some("gc"), ty: 1, - init: 2, + init: 1, ), ( name: Some("gd"), ty: 2, - init: 3, - ), - ( - name: Some("ge"), - ty: 3, - init: 4, - ), - ( - name: Some("gf"), - ty: 2, - init: 5, + init: 2, ), ], overrides: [], global_variables: [], global_expressions: [ - Literal(I32(4)), Literal(I32(4)), Literal(U32(4)), Literal(F32(4.0)), - Compose( - ty: 3, - components: [ - 0, - 0, - 0, - ], - ), - Literal(F32(2.0)), ], functions: [ ( @@ -106,12 +81,22 @@ ], ), Literal(F32(2.0)), + Literal(I32(4)), Constant(0), Constant(1), Constant(2), - Constant(3), - Constant(4), - Constant(5), + Literal(I32(4)), + Literal(I32(4)), + Literal(I32(4)), + Compose( + ty: 3, + components: [ + 10, + 11, + 12, + ], + ), + Literal(F32(2.0)), ], named_expressions: { 0: "a", @@ -124,14 +109,22 @@ 7: "bg", 8: "cg", 9: "dg", - 10: "eg", - 11: "fg", + 13: "eg", + 14: "fg", }, body: [ Emit(( start: 4, end: 5, )), + Emit(( + start: 0, + end: 0, + )), + Emit(( + start: 13, + end: 14, + )), Return( value: None, ), diff --git a/naga/tests/out/msl/abstract-types-const.msl b/naga/tests/out/msl/abstract-types-const.msl index 16a3e35b9a..35efedb44a 100644 --- a/naga/tests/out/msl/abstract-types-const.msl +++ b/naga/tests/out/msl/abstract-types-const.msl @@ -23,20 +23,12 @@ constant metal::float2x2 xmfpafaiaiai = metal::float2x2(metal::float2(1.0, 2.0), constant metal::float2x2 xmfpaiafaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); constant metal::float2x2 xmfpaiaiafai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); constant metal::float2x2 xmfpaiaiaiaf = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); -constant metal::float2x2 imfpaiaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); -constant metal::float2x2 imfpafaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); -constant metal::float2x2 imfpafafafaf = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); -constant metal::int2 ivispai = metal::int2(1); -constant metal::float2 ivfspaf = metal::float2(1.0); constant metal::int2 ivis_ai = metal::int2(1); constant metal::uint2 ivus_ai = metal::uint2(1u); constant metal::float2 ivfs_ai = metal::float2(1.0); constant metal::float2 ivfs_af = metal::float2(1.0); constant type_7 iafafaf = type_7 {1.0, 2.0}; constant type_7 iafaiai = type_7 {1.0, 2.0}; -constant type_7 iafpafaf = type_7 {1.0, 2.0}; -constant type_7 iafpaiaf = type_7 {1.0, 2.0}; -constant type_7 iafpafai = type_7 {1.0, 2.0}; constant type_7 xafpafaf = type_7 {1.0, 2.0}; constant S s_f_i_u = S {1.0, 1, 1u}; constant S s_f_iai = S {1.0, 1, 1u}; diff --git a/naga/tests/out/msl/abstract-types-operators.msl b/naga/tests/out/msl/abstract-types-operators.msl index 6bb3ba4380..a7a972fdb5 100644 --- a/naga/tests/out/msl/abstract-types-operators.msl +++ b/naga/tests/out/msl/abstract-types-operators.msl @@ -37,7 +37,6 @@ constant int shr_iai_u = 0; constant uint shr_uaiai = 0u; constant uint shr_uai_u = 0u; constant int wgpu_4492_ = -2147483648; -constant int wgpu_4492_2_ = -2147483648; void runtime_values( ) { diff --git a/naga/tests/out/msl/constructors.msl b/naga/tests/out/msl/constructors.msl index 1084d8e8aa..3e2c096227 100644 --- a/naga/tests/out/msl/constructors.msl +++ b/naga/tests/out/msl/constructors.msl @@ -8,27 +8,25 @@ struct Foo { metal::float4 a; int b; }; -struct type_6 { +struct type_5 { metal::float2x2 inner[1]; }; -struct type_10 { +struct type_9 { Foo inner[3]; }; struct type_11 { int inner[4]; }; -constant metal::float3 const2_ = metal::float3(0.0, 1.0, 2.0); constant metal::float2x2 const3_ = metal::float2x2(metal::float2(0.0, 1.0), metal::float2(2.0, 3.0)); -constant type_6 const4_ = type_6 {metal::float2x2(metal::float2(0.0, 1.0), metal::float2(2.0, 3.0))}; +constant type_5 const4_ = type_5 {metal::float2x2(metal::float2(0.0, 1.0), metal::float2(2.0, 3.0))}; constant bool cz0_ = bool {}; constant int cz1_ = int {}; constant uint cz2_ = uint {}; constant float cz3_ = float {}; constant metal::uint2 cz4_ = metal::uint2 {}; constant metal::float2x2 cz5_ = metal::float2x2 {}; -constant type_10 cz6_ = type_10 {}; +constant type_9 cz6_ = type_9 {}; constant Foo cz7_ = Foo {}; -constant type_11 cp3_ = type_11 {0, 1, 2, 3}; kernel void main_( ) { diff --git a/naga/tests/out/spv/abstract-types-const.spvasm b/naga/tests/out/spv/abstract-types-const.spvasm index edebb600ac..af8cfb0283 100644 --- a/naga/tests/out/spv/abstract-types-const.spvasm +++ b/naga/tests/out/spv/abstract-types-const.spvasm @@ -16,11 +16,11 @@ OpMemberDecorate %12 2 Offset 8 %5 = OpTypeFloat 32 %6 = OpTypeVector %5 2 %7 = OpTypeMatrix %6 2 -%9 = OpTypeInt 32 1 -%8 = OpTypeVector %9 2 +%8 = OpTypeInt 32 1 +%9 = OpTypeVector %8 2 %11 = OpConstant %3 2 %10 = OpTypeArray %5 %11 -%12 = OpTypeStruct %5 %9 %3 +%12 = OpTypeStruct %5 %8 %3 %13 = OpTypeVector %5 3 %14 = OpConstant %3 42 %15 = OpConstant %3 43 @@ -35,12 +35,12 @@ OpMemberDecorate %12 2 Offset 8 %24 = OpConstant %5 4.0 %25 = OpConstantComposite %6 %23 %24 %26 = OpConstantComposite %7 %22 %25 -%27 = OpConstant %9 1 -%28 = OpConstantComposite %8 %27 %27 -%29 = OpConstantComposite %6 %20 %20 -%30 = OpConstant %3 1 -%31 = OpConstantComposite %4 %30 %30 +%27 = OpConstant %8 1 +%28 = OpConstantComposite %9 %27 %27 +%29 = OpConstant %3 1 +%30 = OpConstantComposite %4 %29 %29 +%31 = OpConstantComposite %6 %20 %20 %32 = OpConstantComposite %10 %20 %21 -%33 = OpConstantComposite %12 %20 %27 %30 +%33 = OpConstantComposite %12 %20 %27 %29 %34 = OpConstantComposite %13 %20 %21 %23 %35 = OpConstantComposite %6 %21 %23 \ No newline at end of file diff --git a/naga/tests/out/spv/constructors.spvasm b/naga/tests/out/spv/constructors.spvasm index ec83d04822..4aa6a6f663 100644 --- a/naga/tests/out/spv/constructors.spvasm +++ b/naga/tests/out/spv/constructors.spvasm @@ -1,84 +1,83 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 68 +; Bound: 67 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %44 "main" -OpExecutionMode %44 LocalSize 1 1 1 +OpEntryPoint GLCompute %38 "main" +OpExecutionMode %38 LocalSize 1 1 1 OpMemberDecorate %6 0 Offset 0 OpMemberDecorate %6 1 Offset 16 -OpDecorate %10 ArrayStride 16 -OpDecorate %15 ArrayStride 32 +OpDecorate %9 ArrayStride 16 +OpDecorate %14 ArrayStride 32 OpDecorate %17 ArrayStride 4 %2 = OpTypeVoid %3 = OpTypeFloat 32 %4 = OpTypeVector %3 4 %5 = OpTypeInt 32 1 %6 = OpTypeStruct %4 %5 -%7 = OpTypeVector %3 3 -%9 = OpTypeVector %3 2 -%8 = OpTypeMatrix %9 2 -%12 = OpTypeInt 32 0 -%11 = OpConstant %12 1 -%10 = OpTypeArray %8 %11 -%13 = OpTypeBool -%14 = OpTypeVector %12 2 -%16 = OpConstant %12 3 -%15 = OpTypeArray %6 %16 -%18 = OpConstant %12 4 +%8 = OpTypeVector %3 2 +%7 = OpTypeMatrix %8 2 +%11 = OpTypeInt 32 0 +%10 = OpConstant %11 1 +%9 = OpTypeArray %7 %10 +%12 = OpTypeBool +%13 = OpTypeVector %11 2 +%15 = OpConstant %11 3 +%14 = OpTypeArray %6 %15 +%16 = OpTypeMatrix %4 4 +%18 = OpConstant %11 4 %17 = OpTypeArray %5 %18 -%19 = OpTypeMatrix %4 4 -%20 = OpTypeMatrix %7 2 +%20 = OpTypeVector %3 3 +%19 = OpTypeMatrix %20 2 %21 = OpConstant %3 0.0 %22 = OpConstant %3 1.0 %23 = OpConstant %3 2.0 -%24 = OpConstantComposite %7 %21 %22 %23 -%25 = OpConstant %3 3.0 -%26 = OpConstantComposite %9 %21 %22 -%27 = OpConstantComposite %9 %23 %25 -%28 = OpConstantComposite %8 %26 %27 -%29 = OpConstantComposite %10 %28 -%30 = OpConstantNull %13 -%31 = OpConstantNull %5 -%32 = OpConstantNull %12 -%33 = OpConstantNull %3 -%34 = OpConstantNull %14 -%35 = OpConstantNull %8 -%36 = OpConstantNull %15 -%37 = OpConstantNull %6 -%38 = OpConstant %5 0 -%39 = OpConstant %5 1 -%40 = OpConstant %5 2 -%41 = OpConstant %5 3 -%42 = OpConstantComposite %17 %38 %39 %40 %41 -%45 = OpTypeFunction %2 -%46 = OpConstantComposite %4 %22 %22 %22 %22 -%47 = OpConstantComposite %6 %46 %39 -%48 = OpConstantComposite %9 %22 %21 -%49 = OpConstantComposite %8 %48 %26 -%50 = OpConstantComposite %4 %22 %21 %21 %21 -%51 = OpConstantComposite %4 %21 %22 %21 %21 -%52 = OpConstantComposite %4 %21 %21 %22 %21 -%53 = OpConstantComposite %4 %21 %21 %21 %22 -%54 = OpConstantComposite %19 %50 %51 %52 %53 -%55 = OpConstant %12 0 -%56 = OpConstantComposite %14 %55 %55 -%57 = OpConstantComposite %9 %21 %21 -%58 = OpConstantComposite %8 %57 %57 -%59 = OpConstantFalse %13 -%60 = OpConstantComposite %14 %55 %55 -%61 = OpConstantComposite %7 %21 %21 %21 -%62 = OpConstantComposite %20 %61 %61 -%63 = OpConstantNull %20 -%65 = OpTypePointer Function %6 -%66 = OpConstantNull %6 -%44 = OpFunction %2 None %45 -%43 = OpLabel -%64 = OpVariable %65 Function %66 -OpBranch %67 -%67 = OpLabel -OpStore %64 %47 +%24 = OpConstant %3 3.0 +%25 = OpConstantComposite %8 %21 %22 +%26 = OpConstantComposite %8 %23 %24 +%27 = OpConstantComposite %7 %25 %26 +%28 = OpConstantComposite %9 %27 +%29 = OpConstantNull %12 +%30 = OpConstantNull %5 +%31 = OpConstantNull %11 +%32 = OpConstantNull %3 +%33 = OpConstantNull %13 +%34 = OpConstantNull %7 +%35 = OpConstantNull %14 +%36 = OpConstantNull %6 +%39 = OpTypeFunction %2 +%40 = OpConstantComposite %4 %22 %22 %22 %22 +%41 = OpConstant %5 1 +%42 = OpConstantComposite %6 %40 %41 +%43 = OpConstantComposite %8 %22 %21 +%44 = OpConstantComposite %7 %43 %25 +%45 = OpConstantComposite %4 %22 %21 %21 %21 +%46 = OpConstantComposite %4 %21 %22 %21 %21 +%47 = OpConstantComposite %4 %21 %21 %22 %21 +%48 = OpConstantComposite %4 %21 %21 %21 %22 +%49 = OpConstantComposite %16 %45 %46 %47 %48 +%50 = OpConstant %11 0 +%51 = OpConstantComposite %13 %50 %50 +%52 = OpConstantComposite %8 %21 %21 +%53 = OpConstantComposite %7 %52 %52 +%54 = OpConstant %5 0 +%55 = OpConstant %5 2 +%56 = OpConstant %5 3 +%57 = OpConstantComposite %17 %54 %41 %55 %56 +%58 = OpConstantFalse %12 +%59 = OpConstantComposite %13 %50 %50 +%60 = OpConstantComposite %20 %21 %21 %21 +%61 = OpConstantComposite %19 %60 %60 +%62 = OpConstantNull %19 +%64 = OpTypePointer Function %6 +%65 = OpConstantNull %6 +%38 = OpFunction %2 None %39 +%37 = OpLabel +%63 = OpVariable %64 Function %65 +OpBranch %66 +%66 = OpLabel +OpStore %63 %42 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/abstract-types-const.wgsl b/naga/tests/out/wgsl/abstract-types-const.wgsl index 195247bbbc..649740f8c8 100644 --- a/naga/tests/out/wgsl/abstract-types-const.wgsl +++ b/naga/tests/out/wgsl/abstract-types-const.wgsl @@ -15,20 +15,12 @@ const xmfpafaiaiai: mat2x2 = mat2x2(vec2(1f, 2f), vec2(3f, 4 const xmfpaiafaiai: mat2x2 = mat2x2(vec2(1f, 2f), vec2(3f, 4f)); const xmfpaiaiafai: mat2x2 = mat2x2(vec2(1f, 2f), vec2(3f, 4f)); const xmfpaiaiaiaf: mat2x2 = mat2x2(vec2(1f, 2f), vec2(3f, 4f)); -const imfpaiaiaiai: mat2x2 = mat2x2(vec2(1f, 2f), vec2(3f, 4f)); -const imfpafaiaiai: mat2x2 = mat2x2(vec2(1f, 2f), vec2(3f, 4f)); -const imfpafafafaf: mat2x2 = mat2x2(vec2(1f, 2f), vec2(3f, 4f)); -const ivispai: vec2 = vec2(1i); -const ivfspaf: vec2 = vec2(1f); const ivis_ai: vec2 = vec2(1i); const ivus_ai: vec2 = vec2(1u); const ivfs_ai: vec2 = vec2(1f); const ivfs_af: vec2 = vec2(1f); const iafafaf: array = array(1f, 2f); const iafaiai: array = array(1f, 2f); -const iafpafaf: array = array(1f, 2f); -const iafpaiaf: array = array(1f, 2f); -const iafpafai: array = array(1f, 2f); const xafpafaf: array = array(1f, 2f); const s_f_i_u: S = S(1f, 1i, 1u); const s_f_iai: S = S(1f, 1i, 1u); diff --git a/naga/tests/out/wgsl/abstract-types-operators.wgsl b/naga/tests/out/wgsl/abstract-types-operators.wgsl index 6b8ffceb14..d86477de0b 100644 --- a/naga/tests/out/wgsl/abstract-types-operators.wgsl +++ b/naga/tests/out/wgsl/abstract-types-operators.wgsl @@ -28,7 +28,6 @@ const shr_iai_u: i32 = 0i; const shr_uaiai: u32 = 0u; const shr_uai_u: u32 = 0u; const wgpu_4492_: i32 = i32(-2147483648); -const wgpu_4492_2_: i32 = i32(-2147483648); var a: array; diff --git a/naga/tests/out/wgsl/const_assert.wgsl b/naga/tests/out/wgsl/const_assert.wgsl index fe1c1c02f4..6271a6de8f 100644 --- a/naga/tests/out/wgsl/const_assert.wgsl +++ b/naga/tests/out/wgsl/const_assert.wgsl @@ -1,6 +1,3 @@ -const x: i32 = 1i; -const y: i32 = 2i; - fn foo() { return; } diff --git a/naga/tests/out/wgsl/constructors.wgsl b/naga/tests/out/wgsl/constructors.wgsl index 622903d7eb..f053e1f6ef 100644 --- a/naga/tests/out/wgsl/constructors.wgsl +++ b/naga/tests/out/wgsl/constructors.wgsl @@ -3,7 +3,6 @@ struct Foo { b: i32, } -const const2_: vec3 = vec3(0f, 1f, 2f); const const3_: mat2x2 = mat2x2(vec2(0f, 1f), vec2(2f, 3f)); const const4_: array, 1> = array, 1>(mat2x2(vec2(0f, 1f), vec2(2f, 3f))); const cz0_: bool = bool(); @@ -14,7 +13,6 @@ const cz4_: vec2 = vec2(); const cz5_: mat2x2 = mat2x2(); const cz6_: array = array(); const cz7_: Foo = Foo(); -const cp3_: array = array(0i, 1i, 2i, 3i); @compute @workgroup_size(1, 1, 1) fn main() { diff --git a/naga/tests/out/wgsl/local-const.wgsl b/naga/tests/out/wgsl/local-const.wgsl index 06e52a9277..535a1bfb14 100644 --- a/naga/tests/out/wgsl/local-const.wgsl +++ b/naga/tests/out/wgsl/local-const.wgsl @@ -1,12 +1,10 @@ -const ga: i32 = 4i; const gb: i32 = 4i; const gc: u32 = 4u; const gd: f32 = 4f; -const ge: vec3 = vec3(4i, 4i, 4i); -const gf: f32 = 2f; fn const_in_fn() { const e = vec3(4i, 4i, 4i); + const eg = vec3(4i, 4i, 4i); return; }