[naga wgsl-in] Automatic conversions for local var initializers.

This commit is contained in:
Jim Blandy 2023-11-25 14:07:07 -08:00 committed by Teodor Tanasoaia
parent 1676ee0dc0
commit f470103874
6 changed files with 518 additions and 62 deletions

View File

@ -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<f32> = vec2(1, 0);
var unit_x: vec2<f32> = 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

View File

@ -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 {

View File

@ -42,3 +42,79 @@ var<private> xafpaiai: array<i32, 2> = array(1, 2);
var<private> xafpaiaf: array<f32, 2> = array(1, 2.0);
var<private> xafpafai: array<f32, 2> = array(1.0, 2);
var<private> xafpafaf: array<f32, 2> = array(1.0, 2.0);
fn all_constant_arguments() {
var xvipaiai: vec2<i32> = vec2(42, 43);
var xvupaiai: vec2<u32> = vec2(44, 45);
var xvfpaiai: vec2<f32> = vec2(46, 47);
var xvupuai: vec2<u32> = vec2(42u, 43);
var xvupaiu: vec2<u32> = vec2(42, 43u);
var xvuuai: vec2<u32> = vec2<u32>(42u, 43);
var xvuaiu: vec2<u32> = vec2<u32>(42, 43u);
var xmfpaiaiaiai: mat2x2<f32> = mat2x2(1, 2, 3, 4);
var xmfpafaiaiai: mat2x2<f32> = mat2x2(1.0, 2, 3, 4);
var xmfpaiafaiai: mat2x2<f32> = mat2x2(1, 2.0, 3, 4);
var xmfpaiaiafai: mat2x2<f32> = mat2x2(1, 2, 3.0, 4);
var xmfpaiaiaiaf: mat2x2<f32> = mat2x2(1, 2, 3, 4.0);
var xmfp_faiaiai: mat2x2<f32> = mat2x2(1.0f, 2, 3, 4);
var xmfpai_faiai: mat2x2<f32> = mat2x2(1, 2.0f, 3, 4);
var xmfpaiai_fai: mat2x2<f32> = mat2x2(1, 2, 3.0f, 4);
var xmfpaiaiai_f: mat2x2<f32> = mat2x2(1, 2, 3, 4.0f);
var xvispai: vec2<i32> = vec2(1);
var xvfspaf: vec2<f32> = vec2(1.0);
var xvis_ai: vec2<i32> = vec2<i32>(1);
var xvus_ai: vec2<u32> = vec2<u32>(1);
var xvfs_ai: vec2<f32> = vec2<f32>(1);
var xvfs_af: vec2<f32> = vec2<f32>(1.0);
var xafafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xaf_faf: array<f32, 2> = array<f32, 2>(1.0f, 2.0);
var xafaf_f: array<f32, 2> = array<f32, 2>(1.0, 2.0f);
var xafaiai: array<f32, 2> = array<f32, 2>(1, 2);
var xai_iai: array<i32, 2> = array<i32, 2>(1i, 2);
var xaiai_i: array<i32, 2> = array<i32, 2>(1, 2i);
// Ideally these would infer the var type from the initializer,
// but we don't support that yet.
var xaipaiai: array<i32, 2> = array(1, 2);
var xafpaiai: array<f32, 2> = array(1, 2);
var xafpaiaf: array<f32, 2> = array(1, 2.0);
var xafpafai: array<f32, 2> = array(1.0, 2);
var xafpafaf: array<f32, 2> = array(1.0, 2.0);
}
fn mixed_constant_and_runtime_arguments() {
var u: u32;
var i: i32;
var f: f32;
var xvupuai: vec2<u32> = vec2(u, 43);
var xvupaiu: vec2<u32> = vec2(42, u);
var xvuuai: vec2<u32> = vec2<u32>(u, 43);
var xvuaiu: vec2<u32> = vec2<u32>(42, u);
var xmfp_faiaiai: mat2x2<f32> = mat2x2(f, 2, 3, 4);
var xmfpai_faiai: mat2x2<f32> = mat2x2(1, f, 3, 4);
var xmfpaiai_fai: mat2x2<f32> = mat2x2(1, 2, f, 4);
var xmfpaiaiai_f: mat2x2<f32> = mat2x2(1, 2, 3, f);
var xaf_faf: array<f32, 2> = array<f32, 2>(f, 2.0);
var xafaf_f: array<f32, 2> = array<f32, 2>(1.0, f);
var xaf_fai: array<f32, 2> = array<f32, 2>(f, 2);
var xafai_f: array<f32, 2> = array<f32, 2>(1, f);
var xai_iai: array<i32, 2> = array<i32, 2>(i, 2);
var xaiai_i: array<i32, 2> = array<i32, 2>(1, i);
var xafp_faf: array<f32, 2> = array(f, 2.0);
var xafpaf_f: array<f32, 2> = array(1.0, f);
var xafp_fai: array<f32, 2> = array(f, 2);
var xafpai_f: array<f32, 2> = array(1, f);
var xaip_iai: array<i32, 2> = array(i, 2);
var xaipai_i: array<i32, 2> = array(1, i);
}

View File

@ -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;
}

View File

@ -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
%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

View File

@ -1,25 +1,130 @@
var<private> xvipaiai: vec2<i32> = vec2<i32>(42, 43);
var<private> xvupaiai: vec2<u32> = vec2<u32>(44u, 45u);
var<private> xvfpaiai: vec2<f32> = vec2<f32>(46.0, 47.0);
var<private> xvupuai: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xvupaiu: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xvuuai: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xvuaiu: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xmfpaiaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpafaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiafaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiaiafai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiaiaiaf: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xvispai: vec2<i32> = vec2(1);
var<private> xvfspaf: vec2<f32> = vec2(1.0);
var<private> xvis_ai: vec2<i32> = vec2(1);
var<private> xvus_ai: vec2<u32> = vec2(1u);
var<private> xvfs_ai: vec2<f32> = vec2(1.0);
var<private> xvfs_af: vec2<f32> = vec2(1.0);
var<private> xafafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafaiai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpaiai: array<i32, 2> = array<i32, 2>(1, 2);
var<private> xafpaiaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpafai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xvipaiai_1: vec2<i32> = vec2<i32>(42, 43);
var<private> xvupaiai_1: vec2<u32> = vec2<u32>(44u, 45u);
var<private> xvfpaiai_1: vec2<f32> = vec2<f32>(46.0, 47.0);
var<private> xvupuai_2: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xvupaiu_2: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xvuuai_2: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xvuaiu_2: vec2<u32> = vec2<u32>(42u, 43u);
var<private> xmfpaiaiaiai_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpafaiaiai_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiafaiai_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiaiafai_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiaiaiaf_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xvispai_1: vec2<i32> = vec2(1);
var<private> xvfspaf_1: vec2<f32> = vec2(1.0);
var<private> xvis_ai_1: vec2<i32> = vec2(1);
var<private> xvus_ai_1: vec2<u32> = vec2(1u);
var<private> xvfs_ai_1: vec2<f32> = vec2(1.0);
var<private> xvfs_af_1: vec2<f32> = vec2(1.0);
var<private> xafafaf_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafaiai_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpaiai_1: array<i32, 2> = array<i32, 2>(1, 2);
var<private> xafpaiaf_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpafai_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpafaf_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
fn all_constant_arguments() {
var xvipaiai: vec2<i32> = vec2<i32>(42, 43);
var xvupaiai: vec2<u32> = vec2<u32>(44u, 45u);
var xvfpaiai: vec2<f32> = vec2<f32>(46.0, 47.0);
var xvupuai: vec2<u32> = vec2<u32>(42u, 43u);
var xvupaiu: vec2<u32> = vec2<u32>(42u, 43u);
var xvuuai: vec2<u32> = vec2<u32>(42u, 43u);
var xvuaiu: vec2<u32> = vec2<u32>(42u, 43u);
var xmfpaiaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpafaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpaiafaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpaiaiafai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpaiaiaiaf: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfp_faiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpai_faiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpaiai_fai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpaiaiai_f: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xvispai: vec2<i32> = vec2(1);
var xvfspaf: vec2<f32> = vec2(1.0);
var xvis_ai: vec2<i32> = vec2(1);
var xvus_ai: vec2<u32> = vec2(1u);
var xvfs_ai: vec2<f32> = vec2(1.0);
var xvfs_af: vec2<f32> = vec2(1.0);
var xafafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xaf_faf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafaf_f: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafaiai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xai_iai: array<i32, 2> = array<i32, 2>(1, 2);
var xaiai_i: array<i32, 2> = array<i32, 2>(1, 2);
var xaipaiai: array<i32, 2> = array<i32, 2>(1, 2);
var xafpaiai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafpaiaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafpafai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafpafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
}
fn mixed_constant_and_runtime_arguments() {
var u: u32;
var i: i32;
var f: f32;
var xvupuai_1: vec2<u32>;
var xvupaiu_1: vec2<u32>;
var xvuuai_1: vec2<u32>;
var xvuaiu_1: vec2<u32>;
var xmfp_faiaiai_1: mat2x2<f32>;
var xmfpai_faiai_1: mat2x2<f32>;
var xmfpaiai_fai_1: mat2x2<f32>;
var xmfpaiaiai_f_1: mat2x2<f32>;
var xaf_faf_1: array<f32, 2>;
var xafaf_f_1: array<f32, 2>;
var xaf_fai: array<f32, 2>;
var xafai_f: array<f32, 2>;
var xai_iai_1: array<i32, 2>;
var xaiai_i_1: array<i32, 2>;
var xafp_faf: array<f32, 2>;
var xafpaf_f: array<f32, 2>;
var xafp_fai: array<f32, 2>;
var xafpai_f: array<f32, 2>;
var xaip_iai: array<i32, 2>;
var xaipai_i: array<i32, 2>;
let _e3 = u;
xvupuai_1 = vec2<u32>(_e3, 43u);
let _e7 = u;
xvupaiu_1 = vec2<u32>(42u, _e7);
let _e11 = u;
xvuuai_1 = vec2<u32>(_e11, 43u);
let _e15 = u;
xvuaiu_1 = vec2<u32>(42u, _e15);
let _e19 = f;
xmfp_faiaiai_1 = mat2x2<f32>(vec2<f32>(_e19, 2.0), vec2<f32>(3.0, 4.0));
let _e27 = f;
xmfpai_faiai_1 = mat2x2<f32>(vec2<f32>(1.0, _e27), vec2<f32>(3.0, 4.0));
let _e35 = f;
xmfpaiai_fai_1 = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(_e35, 4.0));
let _e43 = f;
xmfpaiaiai_f_1 = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, _e43));
let _e51 = f;
xaf_faf_1 = array<f32, 2>(_e51, 2.0);
let _e55 = f;
xafaf_f_1 = array<f32, 2>(1.0, _e55);
let _e59 = f;
xaf_fai = array<f32, 2>(_e59, 2.0);
let _e63 = f;
xafai_f = array<f32, 2>(1.0, _e63);
let _e67 = i;
xai_iai_1 = array<i32, 2>(_e67, 2);
let _e71 = i;
xaiai_i_1 = array<i32, 2>(1, _e71);
let _e75 = f;
xafp_faf = array<f32, 2>(_e75, 2.0);
let _e79 = f;
xafpaf_f = array<f32, 2>(1.0, _e79);
let _e83 = f;
xafp_fai = array<f32, 2>(_e83, 2.0);
let _e87 = f;
xafpai_f = array<f32, 2>(1.0, _e87);
let _e91 = i;
xaip_iai = array<i32, 2>(_e91, 2);
let _e95 = i;
xaipai_i = array<i32, 2>(1, _e95);
return;
}