From f470103874272520f2bbaa4bf64e142fcef12635 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Sat, 25 Nov 2023 14:07:07 -0800 Subject: [PATCH] [naga wgsl-in] Automatic conversions for local `var` initializers. --- CHANGELOG.md | 7 +- naga/src/front/wgsl/lower/mod.rs | 70 ++++---- naga/tests/in/abstract-types-var.wgsl | 76 +++++++++ naga/tests/out/msl/abstract-types-var.msl | 105 ++++++++++++ naga/tests/out/spv/abstract-types-var.spvasm | 169 ++++++++++++++++++- naga/tests/out/wgsl/abstract-types-var.wgsl | 153 ++++++++++++++--- 6 files changed, 518 insertions(+), 62 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 1a4fe90827..2b402c86dc 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -104,7 +104,7 @@ Passing an owned value `window` to `Surface` will return a `Surface<'static>`. S - Introduce a new `Scalar` struct type for use in Naga's IR, and update all frontend, middle, and backend code appropriately. By @jimblandy in [#4673](https://github.com/gfx-rs/wgpu/pull/4673). - Add more metal keywords. By @fornwall in [#4707](https://github.com/gfx-rs/wgpu/pull/4707). -- Add partial support for WGSL abstract types (@jimblandy in [#4743](https://github.com/gfx-rs/wgpu/pull/4743)). +- Add partial support for WGSL abstract types (@jimblandy in [#4743](https://github.com/gfx-rs/wgpu/pull/4743), [#4755](https://github.com/gfx-rs/wgpu/pull/4755)). Abstract types make numeric literals easier to use, by automatically converting literals and other constant expressions @@ -121,9 +121,10 @@ Passing an owned value `window` to `Surface` will return a `Surface<'static>`. S Even though the literals are abstract integers, Naga recognizes that it is safe and necessary to convert them to `f32` values in order to build the vector. You can also use abstract values as - initializers for global constants, like this: + initializers for global constants and global and local variables, + like this: - const unit_x: vec2 = vec2(1, 0); + var unit_x: vec2 = vec2(1, 0); The literals `1` and `0` are abstract integers, and the expression `vec2(1, 0)` is an abstract vector. However, Naga recognizes that diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 27ed142286..6486e6cf6b 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1162,45 +1162,49 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { return Ok(()); } ast::LocalDecl::Var(ref v) => { - let mut emitter = Emitter::default(); - emitter.start(&ctx.function.expressions); - - let initializer = match v.init { - Some(init) => Some( - self.expression(init, &mut ctx.as_expression(block, &mut emitter))?, - ), - None => None, - }; - let explicit_ty = - v.ty.map(|ty| self.resolve_ast_type(ty, &mut ctx.as_global())) + v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_global())) .transpose()?; - let ty = match (explicit_ty, initializer) { - (Some(explicit), Some(initializer)) => { - let mut ctx = ctx.as_expression(block, &mut emitter); - let initializer_ty = resolve_inner!(ctx, initializer); - if !ctx.module.types[explicit] - .inner - .equivalent(initializer_ty, &ctx.module.types) - { - let gctx = &ctx.module.to_ctx(); - return Err(Error::InitializationTypeMismatch { + let mut emitter = Emitter::default(); + emitter.start(&ctx.function.expressions); + let mut ectx = ctx.as_expression(block, &mut emitter); + + let ty; + let initializer; + match (v.init, explicit_ty) { + (Some(init), Some(explicit_ty)) => { + let init = self.expression_for_abstract(init, &mut ectx)?; + let ty_res = crate::proc::TypeResolution::Handle(explicit_ty); + let init = ectx + .try_automatic_conversions(init, &ty_res, v.name.span) + .map_err(|error| match error { + Error::AutoConversion { + dest_span: _, + dest_type, + source_span: _, + source_type, + } => Error::InitializationTypeMismatch { name: v.name.span, - expected: explicit.to_wgsl(gctx), - got: initializer_ty.to_wgsl(gctx), - }); - } - explicit + expected: dest_type, + got: source_type, + }, + other => other, + })?; + ty = explicit_ty; + initializer = Some(init); } - (Some(explicit), None) => explicit, - (None, Some(initializer)) => ctx - .as_expression(block, &mut emitter) - .register_type(initializer)?, - (None, None) => { - return Err(Error::MissingType(v.name.span)); + (Some(init), None) => { + let concretized = self.expression(init, &mut ectx)?; + ty = ectx.register_type(concretized)?; + initializer = Some(concretized); } - }; + (None, Some(explicit_ty)) => { + ty = explicit_ty; + initializer = None; + } + (None, None) => return Err(Error::MissingType(v.name.span)), + } let (const_initializer, initializer) = { match initializer { diff --git a/naga/tests/in/abstract-types-var.wgsl b/naga/tests/in/abstract-types-var.wgsl index 7eeb69065c..a733888530 100644 --- a/naga/tests/in/abstract-types-var.wgsl +++ b/naga/tests/in/abstract-types-var.wgsl @@ -42,3 +42,79 @@ var xafpaiai: array = array(1, 2); var xafpaiaf: array = array(1, 2.0); var xafpafai: array = array(1.0, 2); var xafpafaf: array = array(1.0, 2.0); + +fn all_constant_arguments() { + var xvipaiai: vec2 = vec2(42, 43); + var xvupaiai: vec2 = vec2(44, 45); + var xvfpaiai: vec2 = vec2(46, 47); + + var xvupuai: vec2 = vec2(42u, 43); + var xvupaiu: vec2 = vec2(42, 43u); + + var xvuuai: vec2 = vec2(42u, 43); + var xvuaiu: vec2 = vec2(42, 43u); + + var xmfpaiaiaiai: mat2x2 = mat2x2(1, 2, 3, 4); + var xmfpafaiaiai: mat2x2 = mat2x2(1.0, 2, 3, 4); + var xmfpaiafaiai: mat2x2 = mat2x2(1, 2.0, 3, 4); + var xmfpaiaiafai: mat2x2 = mat2x2(1, 2, 3.0, 4); + var xmfpaiaiaiaf: mat2x2 = mat2x2(1, 2, 3, 4.0); + + var xmfp_faiaiai: mat2x2 = mat2x2(1.0f, 2, 3, 4); + var xmfpai_faiai: mat2x2 = mat2x2(1, 2.0f, 3, 4); + var xmfpaiai_fai: mat2x2 = mat2x2(1, 2, 3.0f, 4); + var xmfpaiaiai_f: mat2x2 = mat2x2(1, 2, 3, 4.0f); + + var xvispai: vec2 = vec2(1); + var xvfspaf: vec2 = vec2(1.0); + var xvis_ai: vec2 = vec2(1); + var xvus_ai: vec2 = vec2(1); + var xvfs_ai: vec2 = vec2(1); + var xvfs_af: vec2 = vec2(1.0); + + var xafafaf: array = array(1.0, 2.0); + var xaf_faf: array = array(1.0f, 2.0); + var xafaf_f: array = array(1.0, 2.0f); + var xafaiai: array = array(1, 2); + var xai_iai: array = array(1i, 2); + var xaiai_i: array = array(1, 2i); + + // Ideally these would infer the var type from the initializer, + // but we don't support that yet. + var xaipaiai: array = array(1, 2); + var xafpaiai: array = array(1, 2); + var xafpaiaf: array = array(1, 2.0); + var xafpafai: array = array(1.0, 2); + var xafpafaf: array = array(1.0, 2.0); +} + +fn mixed_constant_and_runtime_arguments() { + var u: u32; + var i: i32; + var f: f32; + + var xvupuai: vec2 = vec2(u, 43); + var xvupaiu: vec2 = vec2(42, u); + + var xvuuai: vec2 = vec2(u, 43); + var xvuaiu: vec2 = vec2(42, u); + + var xmfp_faiaiai: mat2x2 = mat2x2(f, 2, 3, 4); + var xmfpai_faiai: mat2x2 = mat2x2(1, f, 3, 4); + var xmfpaiai_fai: mat2x2 = mat2x2(1, 2, f, 4); + var xmfpaiaiai_f: mat2x2 = mat2x2(1, 2, 3, f); + + var xaf_faf: array = array(f, 2.0); + var xafaf_f: array = array(1.0, f); + var xaf_fai: array = array(f, 2); + var xafai_f: array = array(1, f); + var xai_iai: array = array(i, 2); + var xaiai_i: array = array(1, i); + + var xafp_faf: array = array(f, 2.0); + var xafpaf_f: array = array(1.0, f); + var xafp_fai: array = array(f, 2); + var xafpai_f: array = array(1, f); + var xaip_iai: array = array(i, 2); + var xaipai_i: array = array(1, i); +} diff --git a/naga/tests/out/msl/abstract-types-var.msl b/naga/tests/out/msl/abstract-types-var.msl index 63f1cdfa75..45096f8672 100644 --- a/naga/tests/out/msl/abstract-types-var.msl +++ b/naga/tests/out/msl/abstract-types-var.msl @@ -10,3 +10,108 @@ struct type_5 { struct type_7 { int inner[2]; }; + +void all_constant_arguments( +) { + metal::int2 xvipaiai = metal::int2(42, 43); + metal::uint2 xvupaiai = metal::uint2(44u, 45u); + metal::float2 xvfpaiai = metal::float2(46.0, 47.0); + metal::uint2 xvupuai = metal::uint2(42u, 43u); + metal::uint2 xvupaiu = metal::uint2(42u, 43u); + metal::uint2 xvuuai = metal::uint2(42u, 43u); + metal::uint2 xvuaiu = metal::uint2(42u, 43u); + metal::float2x2 xmfpaiaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfpafaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfpaiafaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfpaiaiafai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfpaiaiaiaf = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfp_faiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfpai_faiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfpaiai_fai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::float2x2 xmfpaiaiai_f = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0)); + metal::int2 xvispai = metal::int2(1); + metal::float2 xvfspaf = metal::float2(1.0); + metal::int2 xvis_ai = metal::int2(1); + metal::uint2 xvus_ai = metal::uint2(1u); + metal::float2 xvfs_ai = metal::float2(1.0); + metal::float2 xvfs_af = metal::float2(1.0); + type_5 xafafaf = type_5 {1.0, 2.0}; + type_5 xaf_faf = type_5 {1.0, 2.0}; + type_5 xafaf_f = type_5 {1.0, 2.0}; + type_5 xafaiai = type_5 {1.0, 2.0}; + type_7 xai_iai = type_7 {1, 2}; + type_7 xaiai_i = type_7 {1, 2}; + type_7 xaipaiai = type_7 {1, 2}; + type_5 xafpaiai = type_5 {1.0, 2.0}; + type_5 xafpaiaf = type_5 {1.0, 2.0}; + type_5 xafpafai = type_5 {1.0, 2.0}; + type_5 xafpafaf = type_5 {1.0, 2.0}; +} + +void mixed_constant_and_runtime_arguments( +) { + uint u = {}; + int i = {}; + float f = {}; + metal::uint2 xvupuai_1 = {}; + metal::uint2 xvupaiu_1 = {}; + metal::uint2 xvuuai_1 = {}; + metal::uint2 xvuaiu_1 = {}; + metal::float2x2 xmfp_faiaiai_1 = {}; + metal::float2x2 xmfpai_faiai_1 = {}; + metal::float2x2 xmfpaiai_fai_1 = {}; + metal::float2x2 xmfpaiaiai_f_1 = {}; + type_5 xaf_faf_1 = {}; + type_5 xafaf_f_1 = {}; + type_5 xaf_fai = {}; + type_5 xafai_f = {}; + type_7 xai_iai_1 = {}; + type_7 xaiai_i_1 = {}; + type_5 xafp_faf = {}; + type_5 xafpaf_f = {}; + type_5 xafp_fai = {}; + type_5 xafpai_f = {}; + type_7 xaip_iai = {}; + type_7 xaipai_i = {}; + uint _e3 = u; + xvupuai_1 = metal::uint2(_e3, 43u); + uint _e7 = u; + xvupaiu_1 = metal::uint2(42u, _e7); + uint _e11 = u; + xvuuai_1 = metal::uint2(_e11, 43u); + uint _e15 = u; + xvuaiu_1 = metal::uint2(42u, _e15); + float _e19 = f; + xmfp_faiaiai_1 = metal::float2x2(metal::float2(_e19, 2.0), metal::float2(3.0, 4.0)); + float _e27 = f; + xmfpai_faiai_1 = metal::float2x2(metal::float2(1.0, _e27), metal::float2(3.0, 4.0)); + float _e35 = f; + xmfpaiai_fai_1 = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(_e35, 4.0)); + float _e43 = f; + xmfpaiaiai_f_1 = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, _e43)); + float _e51 = f; + xaf_faf_1 = type_5 {_e51, 2.0}; + float _e55 = f; + xafaf_f_1 = type_5 {1.0, _e55}; + float _e59 = f; + xaf_fai = type_5 {_e59, 2.0}; + float _e63 = f; + xafai_f = type_5 {1.0, _e63}; + int _e67 = i; + xai_iai_1 = type_7 {_e67, 2}; + int _e71 = i; + xaiai_i_1 = type_7 {1, _e71}; + float _e75 = f; + xafp_faf = type_5 {_e75, 2.0}; + float _e79 = f; + xafpaf_f = type_5 {1.0, _e79}; + float _e83 = f; + xafp_fai = type_5 {_e83, 2.0}; + float _e87 = f; + xafpai_f = type_5 {1.0, _e87}; + int _e91 = i; + xaip_iai = type_7 {_e91, 2}; + int _e95 = i; + xaipai_i = type_7 {1, _e95}; + return; +} diff --git a/naga/tests/out/spv/abstract-types-var.spvasm b/naga/tests/out/spv/abstract-types-var.spvasm index ee9b4c1c72..1b4b0664b4 100644 --- a/naga/tests/out/spv/abstract-types-var.spvasm +++ b/naga/tests/out/spv/abstract-types-var.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 70 +; Bound: 209 OpCapability Shader OpCapability Linkage %1 = OpExtInstImport "GLSL.std.450" @@ -75,4 +75,169 @@ OpDecorate %12 ArrayStride 4 %65 = OpVariable %66 Private %39 %67 = OpVariable %63 Private %37 %68 = OpVariable %63 Private %37 -%69 = OpVariable %63 Private %37 \ No newline at end of file +%69 = OpVariable %63 Private %37 +%72 = OpTypeFunction %2 +%74 = OpTypePointer Function %3 +%76 = OpTypePointer Function %5 +%78 = OpTypePointer Function %7 +%84 = OpTypePointer Function %9 +%100 = OpTypePointer Function %10 +%105 = OpTypePointer Function %12 +%116 = OpTypePointer Function %6 +%117 = OpConstantNull %6 +%119 = OpTypePointer Function %4 +%120 = OpConstantNull %4 +%122 = OpTypePointer Function %8 +%123 = OpConstantNull %8 +%125 = OpConstantNull %5 +%127 = OpConstantNull %5 +%129 = OpConstantNull %5 +%131 = OpConstantNull %5 +%133 = OpConstantNull %9 +%135 = OpConstantNull %9 +%137 = OpConstantNull %9 +%139 = OpConstantNull %9 +%141 = OpConstantNull %10 +%143 = OpConstantNull %10 +%145 = OpConstantNull %10 +%147 = OpConstantNull %10 +%149 = OpConstantNull %12 +%151 = OpConstantNull %12 +%153 = OpConstantNull %10 +%155 = OpConstantNull %10 +%157 = OpConstantNull %10 +%159 = OpConstantNull %10 +%161 = OpConstantNull %12 +%163 = OpConstantNull %12 +%71 = OpFunction %2 None %72 +%70 = OpLabel +%109 = OpVariable %100 Function %37 +%106 = OpVariable %105 Function %39 +%102 = OpVariable %100 Function %37 +%98 = OpVariable %78 Function %34 +%95 = OpVariable %74 Function %33 +%92 = OpVariable %84 Function %31 +%89 = OpVariable %84 Function %31 +%86 = OpVariable %84 Function %31 +%82 = OpVariable %76 Function %24 +%79 = OpVariable %76 Function %24 +%73 = OpVariable %74 Function %15 +%110 = OpVariable %100 Function %37 +%107 = OpVariable %105 Function %39 +%103 = OpVariable %100 Function %37 +%99 = OpVariable %100 Function %37 +%96 = OpVariable %76 Function %36 +%93 = OpVariable %74 Function %33 +%90 = OpVariable %84 Function %31 +%87 = OpVariable %84 Function %31 +%83 = OpVariable %84 Function %31 +%80 = OpVariable %76 Function %24 +%75 = OpVariable %76 Function %18 +%111 = OpVariable %100 Function %37 +%108 = OpVariable %100 Function %37 +%104 = OpVariable %105 Function %39 +%101 = OpVariable %100 Function %37 +%97 = OpVariable %78 Function %34 +%94 = OpVariable %78 Function %34 +%91 = OpVariable %84 Function %31 +%88 = OpVariable %84 Function %31 +%85 = OpVariable %84 Function %31 +%81 = OpVariable %76 Function %24 +%77 = OpVariable %78 Function %21 +OpBranch %112 +%112 = OpLabel +OpReturn +OpFunctionEnd +%114 = OpFunction %2 None %72 +%113 = OpLabel +%162 = OpVariable %105 Function %163 +%156 = OpVariable %100 Function %157 +%150 = OpVariable %105 Function %151 +%144 = OpVariable %100 Function %145 +%138 = OpVariable %84 Function %139 +%132 = OpVariable %84 Function %133 +%126 = OpVariable %76 Function %127 +%118 = OpVariable %119 Function %120 +%160 = OpVariable %105 Function %161 +%154 = OpVariable %100 Function %155 +%148 = OpVariable %105 Function %149 +%142 = OpVariable %100 Function %143 +%136 = OpVariable %84 Function %137 +%130 = OpVariable %76 Function %131 +%124 = OpVariable %76 Function %125 +%115 = OpVariable %116 Function %117 +%158 = OpVariable %100 Function %159 +%152 = OpVariable %100 Function %153 +%146 = OpVariable %100 Function %147 +%140 = OpVariable %100 Function %141 +%134 = OpVariable %84 Function %135 +%128 = OpVariable %76 Function %129 +%121 = OpVariable %122 Function %123 +OpBranch %164 +%164 = OpLabel +%165 = OpLoad %6 %115 +%166 = OpCompositeConstruct %5 %165 %23 +OpStore %124 %166 +%167 = OpLoad %6 %115 +%168 = OpCompositeConstruct %5 %22 %167 +OpStore %126 %168 +%169 = OpLoad %6 %115 +%170 = OpCompositeConstruct %5 %169 %23 +OpStore %128 %170 +%171 = OpLoad %6 %115 +%172 = OpCompositeConstruct %5 %22 %171 +OpStore %130 %172 +%173 = OpLoad %8 %121 +%174 = OpCompositeConstruct %7 %173 %26 +%175 = OpCompositeConstruct %9 %174 %30 +OpStore %132 %175 +%176 = OpLoad %8 %121 +%177 = OpCompositeConstruct %7 %25 %176 +%178 = OpCompositeConstruct %9 %177 %30 +OpStore %134 %178 +%179 = OpLoad %8 %121 +%180 = OpCompositeConstruct %7 %179 %29 +%181 = OpCompositeConstruct %9 %27 %180 +OpStore %136 %181 +%182 = OpLoad %8 %121 +%183 = OpCompositeConstruct %7 %28 %182 +%184 = OpCompositeConstruct %9 %27 %183 +OpStore %138 %184 +%185 = OpLoad %8 %121 +%186 = OpCompositeConstruct %10 %185 %26 +OpStore %140 %186 +%187 = OpLoad %8 %121 +%188 = OpCompositeConstruct %10 %25 %187 +OpStore %142 %188 +%189 = OpLoad %8 %121 +%190 = OpCompositeConstruct %10 %189 %26 +OpStore %144 %190 +%191 = OpLoad %8 %121 +%192 = OpCompositeConstruct %10 %25 %191 +OpStore %146 %192 +%193 = OpLoad %4 %118 +%194 = OpCompositeConstruct %12 %193 %38 +OpStore %148 %194 +%195 = OpLoad %4 %118 +%196 = OpCompositeConstruct %12 %32 %195 +OpStore %150 %196 +%197 = OpLoad %8 %121 +%198 = OpCompositeConstruct %10 %197 %26 +OpStore %152 %198 +%199 = OpLoad %8 %121 +%200 = OpCompositeConstruct %10 %25 %199 +OpStore %154 %200 +%201 = OpLoad %8 %121 +%202 = OpCompositeConstruct %10 %201 %26 +OpStore %156 %202 +%203 = OpLoad %8 %121 +%204 = OpCompositeConstruct %10 %25 %203 +OpStore %158 %204 +%205 = OpLoad %4 %118 +%206 = OpCompositeConstruct %12 %205 %38 +OpStore %160 %206 +%207 = OpLoad %4 %118 +%208 = OpCompositeConstruct %12 %32 %207 +OpStore %162 %208 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/abstract-types-var.wgsl b/naga/tests/out/wgsl/abstract-types-var.wgsl index cf05b83cde..ae95ff2e82 100644 --- a/naga/tests/out/wgsl/abstract-types-var.wgsl +++ b/naga/tests/out/wgsl/abstract-types-var.wgsl @@ -1,25 +1,130 @@ -var xvipaiai: vec2 = vec2(42, 43); -var xvupaiai: vec2 = vec2(44u, 45u); -var xvfpaiai: vec2 = vec2(46.0, 47.0); -var xvupuai: vec2 = vec2(42u, 43u); -var xvupaiu: vec2 = vec2(42u, 43u); -var xvuuai: vec2 = vec2(42u, 43u); -var xvuaiu: vec2 = vec2(42u, 43u); -var xmfpaiaiaiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); -var xmfpafaiaiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); -var xmfpaiafaiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); -var xmfpaiaiafai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); -var xmfpaiaiaiaf: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); -var xvispai: vec2 = vec2(1); -var xvfspaf: vec2 = vec2(1.0); -var xvis_ai: vec2 = vec2(1); -var xvus_ai: vec2 = vec2(1u); -var xvfs_ai: vec2 = vec2(1.0); -var xvfs_af: vec2 = vec2(1.0); -var xafafaf: array = array(1.0, 2.0); -var xafaiai: array = array(1.0, 2.0); -var xafpaiai: array = array(1, 2); -var xafpaiaf: array = array(1.0, 2.0); -var xafpafai: array = array(1.0, 2.0); -var xafpafaf: array = array(1.0, 2.0); +var xvipaiai_1: vec2 = vec2(42, 43); +var xvupaiai_1: vec2 = vec2(44u, 45u); +var xvfpaiai_1: vec2 = vec2(46.0, 47.0); +var xvupuai_2: vec2 = vec2(42u, 43u); +var xvupaiu_2: vec2 = vec2(42u, 43u); +var xvuuai_2: vec2 = vec2(42u, 43u); +var xvuaiu_2: vec2 = vec2(42u, 43u); +var xmfpaiaiaiai_1: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); +var xmfpafaiaiai_1: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); +var xmfpaiafaiai_1: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); +var xmfpaiaiafai_1: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); +var xmfpaiaiaiaf_1: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); +var xvispai_1: vec2 = vec2(1); +var xvfspaf_1: vec2 = vec2(1.0); +var xvis_ai_1: vec2 = vec2(1); +var xvus_ai_1: vec2 = vec2(1u); +var xvfs_ai_1: vec2 = vec2(1.0); +var xvfs_af_1: vec2 = vec2(1.0); +var xafafaf_1: array = array(1.0, 2.0); +var xafaiai_1: array = array(1.0, 2.0); +var xafpaiai_1: array = array(1, 2); +var xafpaiaf_1: array = array(1.0, 2.0); +var xafpafai_1: array = array(1.0, 2.0); +var xafpafaf_1: array = array(1.0, 2.0); + +fn all_constant_arguments() { + var xvipaiai: vec2 = vec2(42, 43); + var xvupaiai: vec2 = vec2(44u, 45u); + var xvfpaiai: vec2 = vec2(46.0, 47.0); + var xvupuai: vec2 = vec2(42u, 43u); + var xvupaiu: vec2 = vec2(42u, 43u); + var xvuuai: vec2 = vec2(42u, 43u); + var xvuaiu: vec2 = vec2(42u, 43u); + var xmfpaiaiaiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfpafaiaiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfpaiafaiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfpaiaiafai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfpaiaiaiaf: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfp_faiaiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfpai_faiai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfpaiai_fai: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xmfpaiaiai_f: mat2x2 = mat2x2(vec2(1.0, 2.0), vec2(3.0, 4.0)); + var xvispai: vec2 = vec2(1); + var xvfspaf: vec2 = vec2(1.0); + var xvis_ai: vec2 = vec2(1); + var xvus_ai: vec2 = vec2(1u); + var xvfs_ai: vec2 = vec2(1.0); + var xvfs_af: vec2 = vec2(1.0); + var xafafaf: array = array(1.0, 2.0); + var xaf_faf: array = array(1.0, 2.0); + var xafaf_f: array = array(1.0, 2.0); + var xafaiai: array = array(1.0, 2.0); + var xai_iai: array = array(1, 2); + var xaiai_i: array = array(1, 2); + var xaipaiai: array = array(1, 2); + var xafpaiai: array = array(1.0, 2.0); + var xafpaiaf: array = array(1.0, 2.0); + var xafpafai: array = array(1.0, 2.0); + var xafpafaf: array = array(1.0, 2.0); + +} + +fn mixed_constant_and_runtime_arguments() { + var u: u32; + var i: i32; + var f: f32; + var xvupuai_1: vec2; + var xvupaiu_1: vec2; + var xvuuai_1: vec2; + var xvuaiu_1: vec2; + var xmfp_faiaiai_1: mat2x2; + var xmfpai_faiai_1: mat2x2; + var xmfpaiai_fai_1: mat2x2; + var xmfpaiaiai_f_1: mat2x2; + var xaf_faf_1: array; + var xafaf_f_1: array; + var xaf_fai: array; + var xafai_f: array; + var xai_iai_1: array; + var xaiai_i_1: array; + var xafp_faf: array; + var xafpaf_f: array; + var xafp_fai: array; + var xafpai_f: array; + var xaip_iai: array; + var xaipai_i: array; + + let _e3 = u; + xvupuai_1 = vec2(_e3, 43u); + let _e7 = u; + xvupaiu_1 = vec2(42u, _e7); + let _e11 = u; + xvuuai_1 = vec2(_e11, 43u); + let _e15 = u; + xvuaiu_1 = vec2(42u, _e15); + let _e19 = f; + xmfp_faiaiai_1 = mat2x2(vec2(_e19, 2.0), vec2(3.0, 4.0)); + let _e27 = f; + xmfpai_faiai_1 = mat2x2(vec2(1.0, _e27), vec2(3.0, 4.0)); + let _e35 = f; + xmfpaiai_fai_1 = mat2x2(vec2(1.0, 2.0), vec2(_e35, 4.0)); + let _e43 = f; + xmfpaiaiai_f_1 = mat2x2(vec2(1.0, 2.0), vec2(3.0, _e43)); + let _e51 = f; + xaf_faf_1 = array(_e51, 2.0); + let _e55 = f; + xafaf_f_1 = array(1.0, _e55); + let _e59 = f; + xaf_fai = array(_e59, 2.0); + let _e63 = f; + xafai_f = array(1.0, _e63); + let _e67 = i; + xai_iai_1 = array(_e67, 2); + let _e71 = i; + xaiai_i_1 = array(1, _e71); + let _e75 = f; + xafp_faf = array(_e75, 2.0); + let _e79 = f; + xafpaf_f = array(1.0, _e79); + let _e83 = f; + xafp_fai = array(_e83, 2.0); + let _e87 = f; + xafpai_f = array(1.0, _e87); + let _e91 = i; + xaip_iai = array(_e91, 2); + let _e95 = i; + xaipai_i = array(1, _e95); + return; +}