From db3c35db90172e2bde5e15e11564260be0090da5 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Mon, 14 Apr 2025 14:49:08 -0400 Subject: [PATCH] fix(naga): properly impl. auto. type conv. for `select` --- CHANGELOG.md | 1 + cts_runner/test.lst | 3 + naga/src/front/wgsl/error.rs | 28 + naga/src/front/wgsl/lower/mod.rs | 61 +- naga/src/proc/constant_evaluator.rs | 147 ++- naga/tests/in/wgsl/select.wgsl | 41 + naga/tests/naga/validation.rs | 50 +- naga/tests/naga/wgsl_errors.rs | 97 +- .../out/glsl/wgsl-operators.main.Compute.glsl | 2 +- .../out/glsl/wgsl-select.main.Compute.glsl | 17 + naga/tests/out/hlsl/wgsl-operators.hlsl | 2 +- naga/tests/out/hlsl/wgsl-select.hlsl | 11 + naga/tests/out/hlsl/wgsl-select.ron | 12 + naga/tests/out/msl/wgsl-operators.msl | 2 +- naga/tests/out/msl/wgsl-select.msl | 16 + naga/tests/out/spv/wgsl-operators.spvasm | 1098 ++++++++--------- naga/tests/out/spv/wgsl-select.spvasm | 47 + naga/tests/out/wgsl/wgsl-operators.wgsl | 2 +- naga/tests/out/wgsl/wgsl-select.wgsl | 10 + 19 files changed, 1026 insertions(+), 621 deletions(-) create mode 100644 naga/tests/in/wgsl/select.wgsl create mode 100644 naga/tests/out/glsl/wgsl-select.main.Compute.glsl create mode 100644 naga/tests/out/hlsl/wgsl-select.hlsl create mode 100644 naga/tests/out/hlsl/wgsl-select.ron create mode 100644 naga/tests/out/msl/wgsl-select.msl create mode 100644 naga/tests/out/spv/wgsl-select.spvasm create mode 100644 naga/tests/out/wgsl/wgsl-select.wgsl diff --git a/CHANGELOG.md b/CHANGELOG.md index ce305a515..46051e53e 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -77,6 +77,7 @@ Naga now infers the correct binding layout when a resource appears only in an as - Properly apply WGSL's automatic conversions to the arguments to texture sampling functions. By @jimblandy in [#7548](https://github.com/gfx-rs/wgpu/pull/7548). - Properly evaluate `abs(most negative abstract int)`. By @jimblandy in [#7507](https://github.com/gfx-rs/wgpu/pull/7507). - Generate vectorized code for `[un]pack4x{I,U}8[Clamp]` on SPIR-V and MSL 2.1+. By @robamler in [#7664](https://github.com/gfx-rs/wgpu/pull/7664). +- Fix typing for `select`, which had issues particularly with a lack of automatic type conversion. By @ErichDonGubler in [#7572](https://github.com/gfx-rs/wgpu/pull/7572). #### DX12 diff --git a/cts_runner/test.lst b/cts_runner/test.lst index 8adf85fcd..e91fcd25d 100644 --- a/cts_runner/test.lst +++ b/cts_runner/test.lst @@ -18,6 +18,9 @@ webgpu:api,operation,rendering,depth:* webgpu:api,operation,rendering,draw:* webgpu:api,operation,shader_module,compilation_info:* webgpu:api,operation,uncapturederror:iff_uncaptured:* +//FAIL: webgpu:shader,execution,expression,call,builtin,select:* +// - Fails with `const`/abstract int cases on all platforms because of . +// - Fails with `vec3` & `f16` cases on macOS because of . //FAIL: webgpu:api,operation,uncapturederror:onuncapturederror_order_wrt_addEventListener // There are also two unimplemented SKIPs in uncapturederror not enumerated here. webgpu:api,validation,encoding,queries,general:occlusion_query,query_type:* diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs index c4249b017..1123caff6 100644 --- a/naga/src/front/wgsl/error.rs +++ b/naga/src/front/wgsl/error.rs @@ -399,6 +399,16 @@ pub(crate) enum Error<'a> { on_what: DiagnosticAttributeNotSupportedPosition, spans: Vec, }, + SelectUnexpectedArgumentType { + arg_span: Span, + arg_type: String, + }, + SelectRejectAndAcceptHaveNoCommonType { + reject_span: Span, + reject_type: String, + accept_span: Span, + accept_type: String, + }, } impl From for Error<'_> { @@ -1342,6 +1352,24 @@ impl<'a> Error<'a> { ], } } + Error::SelectUnexpectedArgumentType { arg_span, ref arg_type } => ParseError { + message: "unexpected argument type for `select` call".into(), + labels: vec![(arg_span, format!("this value of type {arg_type}").into())], + notes: vec!["expected a scalar or a `vecN` of scalars".into()], + }, + Error::SelectRejectAndAcceptHaveNoCommonType { + reject_span, + ref reject_type, + accept_span, + ref accept_type, + } => ParseError { + message: "type mismatch for reject and accept values in `select` call".into(), + labels: vec![ + (reject_span, format!("reject value of type {reject_type}").into()), + (accept_span, format!("accept value of type {accept_type}").into()), + ], + notes: vec![], + }, } } } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 4bf800181..ebf19cc55 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1,6 +1,7 @@ use alloc::{ borrow::ToOwned, boxed::Box, + format, string::{String, ToString}, vec::Vec, }; @@ -2541,12 +2542,68 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { "select" => { let mut args = ctx.prepare_args(arguments, 3, span); - let reject = self.expression(args.next()?, ctx)?; - let accept = self.expression(args.next()?, ctx)?; + let reject_orig = args.next()?; + let accept_orig = args.next()?; + let mut values = [ + self.expression_for_abstract(reject_orig, ctx)?, + self.expression_for_abstract(accept_orig, ctx)?, + ]; let condition = self.expression(args.next()?, ctx)?; args.finish()?; + let diagnostic_details = + |ctx: &ExpressionContext<'_, '_, '_>, + ty_res: &proc::TypeResolution, + orig_expr| { + ( + ctx.ast_expressions.get_span(orig_expr), + format!("`{}`", ctx.as_diagnostic_display(ty_res)), + ) + }; + for (&value, orig_value) in + values.iter().zip([reject_orig, accept_orig]) + { + let value_ty_res = resolve!(ctx, value); + if value_ty_res + .inner_with(&ctx.module.types) + .vector_size_and_scalar() + .is_none() + { + let (arg_span, arg_type) = + diagnostic_details(ctx, value_ty_res, orig_value); + return Err(Box::new(Error::SelectUnexpectedArgumentType { + arg_span, + arg_type, + })); + } + } + let mut consensus_scalar = ctx + .automatic_conversion_consensus(&values) + .map_err(|_idx| { + let [reject, accept] = values; + let [(reject_span, reject_type), (accept_span, accept_type)] = + [(reject_orig, reject), (accept_orig, accept)].map( + |(orig_expr, expr)| { + let ty_res = &ctx.typifier()[expr]; + diagnostic_details(ctx, ty_res, orig_expr) + }, + ); + Error::SelectRejectAndAcceptHaveNoCommonType { + reject_span, + reject_type, + accept_span, + accept_type, + } + })?; + if !ctx.is_const(condition) { + consensus_scalar = consensus_scalar.concretize(); + } + + ctx.convert_slice_to_common_leaf_scalar(&mut values, consensus_scalar)?; + + let [reject, accept] = values; + ir::Expression::Select { reject, accept, diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index bf0aea303..4dba552f0 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -563,6 +563,27 @@ pub enum ConstantEvaluatorError { RuntimeExpr, #[error("Unexpected override-expression")] OverrideExpr, + #[error("Expected boolean expression for condition argument of `select`, got something else")] + SelectScalarConditionNotABool, + #[error( + "Expected vectors of the same size for reject and accept args., got {:?} and {:?}", + reject, + accept + )] + SelectVecRejectAcceptSizeMismatch { + reject: crate::VectorSize, + accept: crate::VectorSize, + }, + #[error("Expected boolean vector for condition arg., got something else")] + SelectConditionNotAVecBool, + #[error( + "Expected same number of vector components between condition, accept, and reject args., got something else", + )] + SelectConditionVecSizeMismatch, + #[error( + "Expected reject and accept args. to be scalars of vectors of the same type, got something else", + )] + SelectAcceptRejectTypeMismatch, } impl<'a> ConstantEvaluator<'a> { @@ -904,9 +925,19 @@ impl<'a> ConstantEvaluator<'a> { )), } } - Expression::Select { .. } => Err(ConstantEvaluatorError::NotImplemented( - "select built-in function".into(), - )), + Expression::Select { + reject, + accept, + condition, + } => { + let mut arg = |expr| self.check_and_get(expr); + + let reject = arg(reject)?; + let accept = arg(accept)?; + let condition = arg(condition)?; + + self.select(reject, accept, condition, span) + } Expression::Relational { fun, argument } => { let argument = self.check_and_get(argument)?; self.relational(fun, argument, span) @@ -2501,6 +2532,116 @@ impl<'a> ConstantEvaluator<'a> { Ok(resolution) } + + fn select( + &mut self, + reject: Handle, + accept: Handle, + condition: Handle, + span: Span, + ) -> Result, ConstantEvaluatorError> { + let mut arg = |arg| self.eval_zero_value_and_splat(arg, span); + + let reject = arg(reject)?; + let accept = arg(accept)?; + let condition = arg(condition)?; + + let select_single_component = + |this: &mut Self, reject_scalar, reject, accept, condition| { + let accept = this.cast(accept, reject_scalar, span)?; + if condition { + Ok(accept) + } else { + Ok(reject) + } + }; + + match (&self.expressions[reject], &self.expressions[accept]) { + (&Expression::Literal(reject_lit), &Expression::Literal(_accept_lit)) => { + let reject_scalar = reject_lit.scalar(); + let &Expression::Literal(Literal::Bool(condition)) = &self.expressions[condition] + else { + return Err(ConstantEvaluatorError::SelectScalarConditionNotABool); + }; + select_single_component(self, reject_scalar, reject, accept, condition) + } + ( + &Expression::Compose { + ty: reject_ty, + components: ref reject_components, + }, + &Expression::Compose { + ty: accept_ty, + components: ref accept_components, + }, + ) => { + let ty_deets = |ty| { + let (size, scalar) = self.types[ty].inner.vector_size_and_scalar().unwrap(); + (size.unwrap(), scalar) + }; + + let expected_vec_size = { + let [(reject_vec_size, _), (accept_vec_size, _)] = + [reject_ty, accept_ty].map(ty_deets); + + if reject_vec_size != accept_vec_size { + return Err(ConstantEvaluatorError::SelectVecRejectAcceptSizeMismatch { + reject: reject_vec_size, + accept: accept_vec_size, + }); + } + reject_vec_size + }; + + let condition_components = match self.expressions[condition] { + Expression::Literal(Literal::Bool(condition)) => { + vec![condition; (expected_vec_size as u8).into()] + } + Expression::Compose { + ty: condition_ty, + components: ref condition_components, + } => { + let (condition_vec_size, condition_scalar) = ty_deets(condition_ty); + if condition_scalar.kind != ScalarKind::Bool { + return Err(ConstantEvaluatorError::SelectConditionNotAVecBool); + } + if condition_vec_size != expected_vec_size { + return Err(ConstantEvaluatorError::SelectConditionVecSizeMismatch); + } + condition_components + .iter() + .copied() + .map(|component| match &self.expressions[component] { + &Expression::Literal(Literal::Bool(condition)) => condition, + _ => unreachable!(), + }) + .collect() + } + + _ => return Err(ConstantEvaluatorError::SelectConditionNotAVecBool), + }; + + let evaluated = Expression::Compose { + ty: reject_ty, + components: reject_components + .clone() + .into_iter() + .zip(accept_components.clone().into_iter()) + .zip(condition_components.into_iter()) + .map(|((reject, accept), condition)| { + let reject_scalar = match &self.expressions[reject] { + &Expression::Literal(lit) => lit.scalar(), + _ => unreachable!(), + }; + select_single_component(self, reject_scalar, reject, accept, condition) + }) + .collect::>()?, + }; + self.register_evaluated_expr(evaluated, span) + } + _ => Err(ConstantEvaluatorError::SelectAcceptRejectTypeMismatch), + } + } } fn first_trailing_bit(concrete_int: ConcreteInt<1>) -> ConcreteInt<1> { diff --git a/naga/tests/in/wgsl/select.wgsl b/naga/tests/in/wgsl/select.wgsl new file mode 100644 index 000000000..748e8635f --- /dev/null +++ b/naga/tests/in/wgsl/select.wgsl @@ -0,0 +1,41 @@ +const_assert select(0xdeadbeef, 42f, false) == 0xdeadbeef; +const_assert select(0xdeadbeefu, 42, false) == 0xdeadbeefu; +const_assert select(0xdeadi, 42, false) == 0xdeadi; + +const_assert select(42f, 0xdeadbeef, true) == 0xdeadbeef; +const_assert select(42, 0xdeadbeefu, true) == 0xdeadbeefu; +const_assert select(42, 0xdeadi, true) == 0xdeadi; + +const_assert select(42f, 9001, true) == 9001; +const_assert select(42f, 9001, true) == 9001f; +const_assert select(42, 9001i, true) == 9001; +const_assert select(42, 9001u, true) == 9001; + +const_assert select(9001, 42f, false) == 9001; +const_assert select(9001, 42f, false) == 9001f; +const_assert select(9001i, 42, false) == 9001; +const_assert select(9001u, 42, false) == 9001; + +const_assert !select(false, true, false); +const_assert select(false, true, true); +const_assert select(true, false, false); +const_assert !select(true, false, true); + +const_assert all(select(vec2(2f), vec2(), true) == vec2(0)); +const_assert all(select(vec2(1), vec2(2f), false) == vec2(1)); +const_assert all(select(vec2(1), vec2(2f), false) == vec2(1)); +const_assert all(select(vec2(1), vec2(2f), vec2(false, false)) == vec2(1)); +const_assert all(select(vec2(1), vec2(2f), vec2(true)) == vec2(2)); +const_assert all(select(vec2(1), vec2(2f), vec2(true)) == vec2(2)); +const_assert all(select(vec2(1), vec2(2f), vec2(true, false)) == vec2(2, 1)); + +const_assert all(select(vec3(1), vec3(2f), vec3(true)) == vec3(2)); +const_assert all(select(vec4(1), vec4(2f), vec4(true)) == vec4(2)); + +@compute @workgroup_size(1, 1) +fn main() { + _ = select(1, 2f, false); + + var x0 = vec2(1, 2); + var i1: vec2 = select(vec2(1., 0.), vec2(0., 1.), (x0.x < x0.y)); +} diff --git a/naga/tests/naga/validation.rs b/naga/tests/naga/validation.rs index 5b06bd3bc..099a0f79a 100644 --- a/naga/tests/naga/validation.rs +++ b/naga/tests/naga/validation.rs @@ -648,9 +648,8 @@ fn binding_arrays_cannot_hold_scalars() { #[cfg(feature = "wgsl-in")] #[test] fn validation_error_messages() { - let cases = [ - ( - r#"@group(0) @binding(0) var my_sampler: sampler; + let cases = [( + r#"@group(0) @binding(0) var my_sampler: sampler; fn foo(tex: texture_2d) -> vec4 { return textureSampleLevel(tex, my_sampler, vec2f(0, 0), 0.0); @@ -660,7 +659,7 @@ fn validation_error_messages() { foo(); } "#, - "\ + "\ error: Function [1] 'main' is invalid ┌─ wgsl:7:17 │ \n7 │ ╭ fn main() { @@ -671,48 +670,7 @@ error: Function [1] 'main' is invalid = Requires 1 arguments, but 0 are provided ", - ), - ( - "\ -@compute @workgroup_size(1, 1) -fn main() { - // Bad: `9001` isn't a `bool`. - _ = select(1, 2, 9001); -} -", - "\ -error: Entry point main at Compute is invalid - ┌─ wgsl:4:9 - │ -4 │ _ = select(1, 2, 9001); - │ ^^^^^^ naga::ir::Expression [3] - │ - = Expression [3] is invalid - = Expected selection condition to be a boolean value, got Scalar(Scalar { kind: Sint, width: 4 }) - -", - ), - ( - "\ -@compute @workgroup_size(1, 1) -fn main() { - // Bad: `bool` and abstract int args. don't match. - _ = select(true, 1, false); -} -", - "\ -error: Entry point main at Compute is invalid - ┌─ wgsl:4:9 - │ -4 │ _ = select(true, 1, false); - │ ^^^^^^ naga::ir::Expression [3] - │ - = Expression [3] is invalid - = Expected selection argument types to match, but reject value of type Scalar(Scalar { kind: Bool, width: 1 }) does not match accept value of value Scalar(Scalar { kind: Sint, width: 4 }) - -", - ), - ]; + )]; for (source, expected_err) in cases { let module = naga::front::wgsl::parse_str(source).unwrap(); diff --git a/naga/tests/naga/wgsl_errors.rs b/naga/tests/naga/wgsl_errors.rs index a02f88441..1784715a3 100644 --- a/naga/tests/naga/wgsl_errors.rs +++ b/naga/tests/naga/wgsl_errors.rs @@ -2012,8 +2012,9 @@ fn invalid_runtime_sized_arrays() { #[test] fn select() { - check_validation! { - " + let snapshots = [ + ( + " fn select_pointers(which: bool) -> i32 { var x: i32 = 1; var y: i32 = 2; @@ -2021,7 +2022,19 @@ fn select() { return *p; } ", - " + "\ +error: unexpected argument type for `select` call + ┌─ wgsl:5:28 + │ +5 │ let p = select(&x, &y, which); + │ ^^ this value of type `ptr` + │ + = note: expected a scalar or a `vecN` of scalars + +", + ), + ( + " fn select_arrays(which: bool) -> i32 { var x: array; var y: array; @@ -2029,7 +2042,19 @@ fn select() { return s[0]; } ", - " + "\ +error: unexpected argument type for `select` call + ┌─ wgsl:5:28 + │ +5 │ let s = select(x, y, which); + │ ^ this value of type `array` + │ + = note: expected a scalar or a `vecN` of scalars + +", + ), + ( + " struct S { member: i32 } fn select_structs(which: bool) -> S { var x: S = S(1); @@ -2037,18 +2062,58 @@ fn select() { let s = select(x, y, which); return s; } - ": - Err( - naga::valid::ValidationError::Function { - name, - source: naga::valid::FunctionError::Expression { - source: naga::valid::ExpressionError::SelectConditionNotABool { .. }, - .. - }, - .. - }, - ) - if name.starts_with("select_") + ", + "\ +error: unexpected argument type for `select` call + ┌─ wgsl:6:28 + │ +6 │ let s = select(x, y, which); + │ ^ this value of type `S` + │ + = note: expected a scalar or a `vecN` of scalars + +", + ), + ( + " + @compute @workgroup_size(1, 1) + fn main() { + // Bad: `9001` isn't a `bool`. + _ = select(1, 2, 9001); + } + ", + "\ +error: Expected boolean expression for condition argument of `select`, got something else + ┌─ wgsl:5:17 + │ +5 │ _ = select(1, 2, 9001); + │ ^^^^^^ see msg + +", + ), + ( + " + @compute @workgroup_size(1, 1) + fn main() { + // Bad: `bool` and abstract int args. don't match. + _ = select(true, 1, false); + } + ", + "\ +error: type mismatch for reject and accept values in `select` call + ┌─ wgsl:5:24 + │ +5 │ _ = select(true, 1, false); + │ ^^^^ ^ accept value of type `{AbstractInt}` + │ │ + │ reject value of type `bool` + +", + ), + ]; + + for (input, snapshot) in snapshots { + check(input, snapshot); } } diff --git a/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl b/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl index aadbf7231..cc66e08cc 100644 --- a/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl +++ b/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl @@ -14,7 +14,7 @@ const ivec4 v_i32_one = ivec4(1, 1, 1, 1); vec4 builtins() { int s1_ = (true ? 1 : 0); vec4 s2_ = (true ? v_f32_one : v_f32_zero); - vec4 s3_ = mix(v_f32_one, v_f32_zero, bvec4(false, false, false, false)); + vec4 s3_ = vec4(1.0, 1.0, 1.0, 1.0); vec4 m1_ = mix(v_f32_zero, v_f32_one, v_f32_half); vec4 m2_ = mix(v_f32_zero, v_f32_one, 0.1); float b1_ = intBitsToFloat(1); diff --git a/naga/tests/out/glsl/wgsl-select.main.Compute.glsl b/naga/tests/out/glsl/wgsl-select.main.Compute.glsl new file mode 100644 index 000000000..a5de0917b --- /dev/null +++ b/naga/tests/out/glsl/wgsl-select.main.Compute.glsl @@ -0,0 +1,17 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + + +void main() { + ivec2 x0_ = ivec2(1, 2); + vec2 i1_ = vec2(0.0); + int _e12 = x0_.x; + int _e14 = x0_.y; + i1_ = ((_e12 < _e14) ? vec2(0.0, 1.0) : vec2(1.0, 0.0)); + return; +} + diff --git a/naga/tests/out/hlsl/wgsl-operators.hlsl b/naga/tests/out/hlsl/wgsl-operators.hlsl index d76518b73..e584efe82 100644 --- a/naga/tests/out/hlsl/wgsl-operators.hlsl +++ b/naga/tests/out/hlsl/wgsl-operators.hlsl @@ -7,7 +7,7 @@ float4 builtins() { int s1_ = (true ? int(1) : int(0)); float4 s2_ = (true ? v_f32_one : v_f32_zero); - float4 s3_ = (bool4(false, false, false, false) ? v_f32_zero : v_f32_one); + float4 s3_ = float4(1.0, 1.0, 1.0, 1.0); float4 m1_ = lerp(v_f32_zero, v_f32_one, v_f32_half); float4 m2_ = lerp(v_f32_zero, v_f32_one, 0.1); float b1_ = asfloat(int(1)); diff --git a/naga/tests/out/hlsl/wgsl-select.hlsl b/naga/tests/out/hlsl/wgsl-select.hlsl new file mode 100644 index 000000000..d644e903f --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-select.hlsl @@ -0,0 +1,11 @@ +[numthreads(1, 1, 1)] +void main() +{ + int2 x0_ = int2(int(1), int(2)); + float2 i1_ = (float2)0; + + int _e12 = x0_.x; + int _e14 = x0_.y; + i1_ = ((_e12 < _e14) ? float2(0.0, 1.0) : float2(1.0, 0.0)); + return; +} diff --git a/naga/tests/out/hlsl/wgsl-select.ron b/naga/tests/out/hlsl/wgsl-select.ron new file mode 100644 index 000000000..a07b03300 --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-select.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/msl/wgsl-operators.msl b/naga/tests/out/msl/wgsl-operators.msl index d43996c80..4f9fb2dfe 100644 --- a/naga/tests/out/msl/wgsl-operators.msl +++ b/naga/tests/out/msl/wgsl-operators.msl @@ -13,7 +13,7 @@ metal::float4 builtins( ) { int s1_ = true ? 1 : 0; metal::float4 s2_ = true ? v_f32_one : v_f32_zero; - metal::float4 s3_ = metal::select(v_f32_one, v_f32_zero, metal::bool4(false, false, false, false)); + metal::float4 s3_ = metal::float4(1.0, 1.0, 1.0, 1.0); metal::float4 m1_ = metal::mix(v_f32_zero, v_f32_one, v_f32_half); metal::float4 m2_ = metal::mix(v_f32_zero, v_f32_one, 0.1); float b1_ = as_type(1); diff --git a/naga/tests/out/msl/wgsl-select.msl b/naga/tests/out/msl/wgsl-select.msl new file mode 100644 index 000000000..d193acda7 --- /dev/null +++ b/naga/tests/out/msl/wgsl-select.msl @@ -0,0 +1,16 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + + +kernel void main_( +) { + metal::int2 x0_ = metal::int2(1, 2); + metal::float2 i1_ = {}; + int _e12 = x0_.x; + int _e14 = x0_.y; + i1_ = (_e12 < _e14) ? metal::float2(0.0, 1.0) : metal::float2(1.0, 0.0); + return; +} diff --git a/naga/tests/out/spv/wgsl-operators.spvasm b/naga/tests/out/spv/wgsl-operators.spvasm index cdcb905df..cf24cc810 100644 --- a/naga/tests/out/spv/wgsl-operators.spvasm +++ b/naga/tests/out/spv/wgsl-operators.spvasm @@ -1,619 +1,617 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 535 +; Bound: 533 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %518 "main" %515 -OpExecutionMode %518 LocalSize 1 1 1 -OpDecorate %515 BuiltIn WorkgroupId +OpEntryPoint GLCompute %516 "main" %513 +OpExecutionMode %516 LocalSize 1 1 1 +OpDecorate %513 BuiltIn WorkgroupId %2 = OpTypeVoid %3 = OpTypeFloat 32 %4 = OpTypeVector %3 4 %5 = OpTypeInt 32 1 %6 = OpTypeVector %5 4 -%8 = OpTypeBool -%7 = OpTypeVector %8 4 -%9 = OpTypeVector %3 2 -%10 = OpTypeVector %3 3 -%12 = OpTypeInt 32 0 -%11 = OpTypeVector %12 3 -%13 = OpTypeMatrix %10 3 -%14 = OpTypeMatrix %10 4 -%15 = OpTypeMatrix %4 3 -%16 = OpTypeVector %5 3 -%17 = OpConstant %3 1.0 +%7 = OpTypeVector %3 2 +%8 = OpTypeVector %3 3 +%10 = OpTypeInt 32 0 +%9 = OpTypeVector %10 3 +%11 = OpTypeMatrix %8 3 +%12 = OpTypeMatrix %8 4 +%13 = OpTypeMatrix %4 3 +%14 = OpTypeVector %5 3 +%15 = OpConstant %3 1.0 +%16 = OpConstantComposite %4 %15 %15 %15 %15 +%17 = OpConstant %3 0.0 %18 = OpConstantComposite %4 %17 %17 %17 %17 -%19 = OpConstant %3 0.0 +%19 = OpConstant %3 0.5 %20 = OpConstantComposite %4 %19 %19 %19 %19 -%21 = OpConstant %3 0.5 -%22 = OpConstantComposite %4 %21 %21 %21 %21 -%23 = OpConstant %5 1 -%24 = OpConstantComposite %6 %23 %23 %23 %23 -%27 = OpTypeFunction %4 -%28 = OpConstantTrue %8 -%29 = OpConstant %5 0 -%30 = OpConstantFalse %8 -%31 = OpConstantComposite %7 %30 %30 %30 %30 -%32 = OpConstant %3 0.1 -%33 = OpConstantComposite %6 %29 %29 %29 %29 -%54 = OpTypeFunction %6 %6 %6 -%58 = OpConstantComposite %6 %29 %29 %29 %29 -%60 = OpConstant %5 -2147483648 -%61 = OpConstant %5 -1 -%62 = OpConstantComposite %6 %60 %60 %60 %60 -%63 = OpConstantComposite %6 %61 %61 %61 %61 -%68 = OpConstantComposite %6 %23 %23 %23 %23 -%75 = OpTypeFunction %4 %3 %5 -%76 = OpConstant %3 2.0 -%77 = OpConstantComposite %9 %76 %76 -%78 = OpConstant %3 4.0 -%79 = OpConstantComposite %9 %78 %78 -%80 = OpConstant %3 8.0 -%81 = OpConstantComposite %9 %80 %80 -%82 = OpConstant %5 2 -%83 = OpConstantComposite %6 %82 %82 %82 %82 -%96 = OpTypeFunction %9 -%97 = OpConstantComposite %9 %17 %17 -%98 = OpConstant %3 3.0 -%99 = OpConstantComposite %9 %98 %98 -%101 = OpTypePointer Function %9 -%113 = OpTypeFunction %10 %10 -%115 = OpTypeVector %8 3 -%116 = OpConstantComposite %10 %19 %19 %19 -%118 = OpConstantComposite %10 %17 %17 %17 -%122 = OpTypeFunction %2 -%123 = OpTypeVector %8 2 -%124 = OpConstantComposite %123 %28 %28 -%125 = OpConstantComposite %115 %28 %28 %28 -%126 = OpConstantComposite %115 %30 %30 %30 -%127 = OpConstantComposite %7 %28 %28 %28 %28 -%128 = OpConstantComposite %7 %30 %30 %30 %30 -%139 = OpTypeFunction %5 %5 %5 -%151 = OpTypeFunction %12 %12 %12 -%155 = OpConstant %12 0 -%157 = OpConstant %12 1 -%160 = OpTypeVector %5 2 -%162 = OpTypeFunction %160 %160 %160 -%166 = OpConstantComposite %160 %29 %29 -%168 = OpConstantComposite %160 %60 %60 -%169 = OpConstantComposite %160 %61 %61 -%174 = OpConstantComposite %160 %23 %23 -%178 = OpTypeFunction %11 %11 %11 -%182 = OpConstantComposite %11 %155 %155 %155 -%184 = OpConstantComposite %11 %157 %157 %157 -%223 = OpTypeVector %12 2 -%225 = OpTypeFunction %223 %223 %223 -%229 = OpConstantComposite %223 %155 %155 -%231 = OpConstantComposite %223 %157 %157 -%243 = OpConstant %12 2 -%244 = OpConstantComposite %160 %82 %82 -%245 = OpConstantComposite %11 %243 %243 %243 -%246 = OpConstantComposite %4 %76 %76 %76 %76 -%247 = OpConstantComposite %4 %17 %17 %17 %17 -%248 = OpConstantComposite %223 %243 %243 -%249 = OpConstantNull %13 -%250 = OpConstantNull %14 -%251 = OpConstantComposite %10 %76 %76 %76 -%252 = OpConstantNull %15 -%254 = OpTypePointer Function %5 +%21 = OpConstant %5 1 +%22 = OpConstantComposite %6 %21 %21 %21 %21 +%25 = OpTypeFunction %4 +%26 = OpTypeBool +%27 = OpConstantTrue %26 +%28 = OpConstant %5 0 +%29 = OpConstant %3 0.1 +%30 = OpConstantComposite %6 %28 %28 %28 %28 +%34 = OpTypeVector %26 4 +%51 = OpTypeFunction %6 %6 %6 +%55 = OpConstantComposite %6 %28 %28 %28 %28 +%57 = OpConstant %5 -2147483648 +%58 = OpConstant %5 -1 +%59 = OpConstantComposite %6 %57 %57 %57 %57 +%60 = OpConstantComposite %6 %58 %58 %58 %58 +%65 = OpConstantComposite %6 %21 %21 %21 %21 +%72 = OpTypeFunction %4 %3 %5 +%73 = OpConstant %3 2.0 +%74 = OpConstantComposite %7 %73 %73 +%75 = OpConstant %3 4.0 +%76 = OpConstantComposite %7 %75 %75 +%77 = OpConstant %3 8.0 +%78 = OpConstantComposite %7 %77 %77 +%79 = OpConstant %5 2 +%80 = OpConstantComposite %6 %79 %79 %79 %79 +%93 = OpTypeFunction %7 +%94 = OpConstantComposite %7 %15 %15 +%95 = OpConstant %3 3.0 +%96 = OpConstantComposite %7 %95 %95 +%98 = OpTypePointer Function %7 +%110 = OpTypeFunction %8 %8 +%112 = OpTypeVector %26 3 +%113 = OpConstantComposite %8 %17 %17 %17 +%115 = OpConstantComposite %8 %15 %15 %15 +%119 = OpTypeFunction %2 +%120 = OpConstantFalse %26 +%121 = OpTypeVector %26 2 +%122 = OpConstantComposite %121 %27 %27 +%123 = OpConstantComposite %112 %27 %27 %27 +%124 = OpConstantComposite %112 %120 %120 %120 +%125 = OpConstantComposite %34 %27 %27 %27 %27 +%126 = OpConstantComposite %34 %120 %120 %120 %120 +%137 = OpTypeFunction %5 %5 %5 +%149 = OpTypeFunction %10 %10 %10 +%153 = OpConstant %10 0 +%155 = OpConstant %10 1 +%158 = OpTypeVector %5 2 +%160 = OpTypeFunction %158 %158 %158 +%164 = OpConstantComposite %158 %28 %28 +%166 = OpConstantComposite %158 %57 %57 +%167 = OpConstantComposite %158 %58 %58 +%172 = OpConstantComposite %158 %21 %21 +%176 = OpTypeFunction %9 %9 %9 +%180 = OpConstantComposite %9 %153 %153 %153 +%182 = OpConstantComposite %9 %155 %155 %155 +%221 = OpTypeVector %10 2 +%223 = OpTypeFunction %221 %221 %221 +%227 = OpConstantComposite %221 %153 %153 +%229 = OpConstantComposite %221 %155 %155 +%241 = OpConstant %10 2 +%242 = OpConstantComposite %158 %79 %79 +%243 = OpConstantComposite %9 %241 %241 %241 +%244 = OpConstantComposite %4 %73 %73 %73 %73 +%245 = OpConstantComposite %4 %15 %15 %15 %15 +%246 = OpConstantComposite %221 %241 %241 +%247 = OpConstantNull %11 +%248 = OpConstantNull %12 +%249 = OpConstantComposite %8 %73 %73 %73 +%250 = OpConstantNull %13 +%252 = OpTypePointer Function %5 +%253 = OpConstantNull %5 %255 = OpConstantNull %5 -%257 = OpConstantNull %5 -%423 = OpConstantNull %16 -%425 = OpConstantNull %5 -%427 = OpTypePointer Function %16 -%516 = OpTypePointer Input %11 -%515 = OpVariable %516 Input -%519 = OpConstantComposite %10 %17 %17 %17 -%26 = OpFunction %4 None %27 -%25 = OpLabel -OpBranch %34 -%34 = OpLabel -%35 = OpSelect %5 %28 %23 %29 -%37 = OpCompositeConstruct %7 %28 %28 %28 %28 -%36 = OpSelect %4 %37 %18 %20 -%38 = OpSelect %4 %31 %20 %18 -%39 = OpExtInst %4 %1 FMix %20 %18 %22 -%41 = OpCompositeConstruct %4 %32 %32 %32 %32 -%40 = OpExtInst %4 %1 FMix %20 %18 %41 -%42 = OpBitcast %3 %23 -%43 = OpBitcast %4 %24 -%44 = OpCompositeConstruct %6 %35 %35 %35 %35 -%45 = OpIAdd %6 %44 %33 -%46 = OpConvertSToF %4 %45 -%47 = OpFAdd %4 %46 %36 -%48 = OpFAdd %4 %47 %39 +%421 = OpConstantNull %14 +%423 = OpConstantNull %5 +%425 = OpTypePointer Function %14 +%514 = OpTypePointer Input %9 +%513 = OpVariable %514 Input +%517 = OpConstantComposite %8 %15 %15 %15 +%24 = OpFunction %4 None %25 +%23 = OpLabel +OpBranch %31 +%31 = OpLabel +%32 = OpSelect %5 %27 %21 %28 +%35 = OpCompositeConstruct %34 %27 %27 %27 %27 +%33 = OpSelect %4 %35 %16 %18 +%36 = OpExtInst %4 %1 FMix %18 %16 %20 +%38 = OpCompositeConstruct %4 %29 %29 %29 %29 +%37 = OpExtInst %4 %1 FMix %18 %16 %38 +%39 = OpBitcast %3 %21 +%40 = OpBitcast %4 %22 +%41 = OpCompositeConstruct %6 %32 %32 %32 %32 +%42 = OpIAdd %6 %41 %30 +%43 = OpConvertSToF %4 %42 +%44 = OpFAdd %4 %43 %33 +%45 = OpFAdd %4 %44 %36 +%46 = OpFAdd %4 %45 %37 +%47 = OpCompositeConstruct %4 %39 %39 %39 %39 +%48 = OpFAdd %4 %46 %47 %49 = OpFAdd %4 %48 %40 -%50 = OpCompositeConstruct %4 %42 %42 %42 %42 -%51 = OpFAdd %4 %49 %50 -%52 = OpFAdd %4 %51 %43 -OpReturnValue %52 +OpReturnValue %49 OpFunctionEnd -%53 = OpFunction %6 None %54 -%55 = OpFunctionParameter %6 -%56 = OpFunctionParameter %6 -%57 = OpLabel -%59 = OpIEqual %7 %56 %58 -%64 = OpIEqual %7 %55 %62 -%65 = OpIEqual %7 %56 %63 -%66 = OpLogicalAnd %7 %64 %65 -%67 = OpLogicalOr %7 %59 %66 -%69 = OpSelect %6 %67 %68 %56 -%70 = OpSRem %6 %55 %69 -OpReturnValue %70 +%50 = OpFunction %6 None %51 +%52 = OpFunctionParameter %6 +%53 = OpFunctionParameter %6 +%54 = OpLabel +%56 = OpIEqual %34 %53 %55 +%61 = OpIEqual %34 %52 %59 +%62 = OpIEqual %34 %53 %60 +%63 = OpLogicalAnd %34 %61 %62 +%64 = OpLogicalOr %34 %56 %63 +%66 = OpSelect %6 %64 %65 %53 +%67 = OpSRem %6 %52 %66 +OpReturnValue %67 OpFunctionEnd -%74 = OpFunction %4 None %75 -%72 = OpFunctionParameter %3 -%73 = OpFunctionParameter %5 -%71 = OpLabel -OpBranch %84 -%84 = OpLabel -%85 = OpCompositeConstruct %9 %72 %72 -%86 = OpFAdd %9 %77 %85 -%87 = OpFSub %9 %86 %79 -%88 = OpFDiv %9 %87 %81 -%89 = OpCompositeConstruct %6 %73 %73 %73 %73 -%90 = OpFunctionCall %6 %53 %89 %83 -%91 = OpVectorShuffle %4 %88 %88 0 1 0 1 -%92 = OpConvertSToF %4 %90 -%93 = OpFAdd %4 %91 %92 -OpReturnValue %93 +%71 = OpFunction %4 None %72 +%69 = OpFunctionParameter %3 +%70 = OpFunctionParameter %5 +%68 = OpLabel +OpBranch %81 +%81 = OpLabel +%82 = OpCompositeConstruct %7 %69 %69 +%83 = OpFAdd %7 %74 %82 +%84 = OpFSub %7 %83 %76 +%85 = OpFDiv %7 %84 %78 +%86 = OpCompositeConstruct %6 %70 %70 %70 %70 +%87 = OpFunctionCall %6 %50 %86 %80 +%88 = OpVectorShuffle %4 %85 %85 0 1 0 1 +%89 = OpConvertSToF %4 %87 +%90 = OpFAdd %4 %88 %89 +OpReturnValue %90 OpFunctionEnd -%95 = OpFunction %9 None %96 -%94 = OpLabel -%100 = OpVariable %101 Function %77 -OpBranch %102 -%102 = OpLabel -%103 = OpLoad %9 %100 -%104 = OpFAdd %9 %103 %97 -OpStore %100 %104 -%105 = OpLoad %9 %100 -%106 = OpFSub %9 %105 %99 -OpStore %100 %106 -%107 = OpLoad %9 %100 -%108 = OpFDiv %9 %107 %79 -OpStore %100 %108 -%109 = OpLoad %9 %100 -OpReturnValue %109 +%92 = OpFunction %7 None %93 +%91 = OpLabel +%97 = OpVariable %98 Function %74 +OpBranch %99 +%99 = OpLabel +%100 = OpLoad %7 %97 +%101 = OpFAdd %7 %100 %94 +OpStore %97 %101 +%102 = OpLoad %7 %97 +%103 = OpFSub %7 %102 %96 +OpStore %97 %103 +%104 = OpLoad %7 %97 +%105 = OpFDiv %7 %104 %76 +OpStore %97 %105 +%106 = OpLoad %7 %97 +OpReturnValue %106 OpFunctionEnd -%112 = OpFunction %10 None %113 -%111 = OpFunctionParameter %10 -%110 = OpLabel -OpBranch %114 -%114 = OpLabel -%117 = OpFUnordNotEqual %115 %111 %116 -%119 = OpSelect %10 %117 %118 %116 -OpReturnValue %119 +%109 = OpFunction %8 None %110 +%108 = OpFunctionParameter %8 +%107 = OpLabel +OpBranch %111 +%111 = OpLabel +%114 = OpFUnordNotEqual %112 %108 %113 +%116 = OpSelect %8 %114 %115 %113 +OpReturnValue %116 OpFunctionEnd -%121 = OpFunction %2 None %122 -%120 = OpLabel -OpBranch %129 -%129 = OpLabel -%130 = OpLogicalNot %8 %28 -%131 = OpLogicalNot %123 %124 -%132 = OpLogicalOr %8 %28 %30 -%133 = OpLogicalAnd %8 %28 %30 -%134 = OpLogicalOr %8 %28 %30 -%135 = OpLogicalOr %115 %125 %126 -%136 = OpLogicalAnd %8 %28 %30 -%137 = OpLogicalAnd %7 %127 %128 +%118 = OpFunction %2 None %119 +%117 = OpLabel +OpBranch %127 +%127 = OpLabel +%128 = OpLogicalNot %26 %27 +%129 = OpLogicalNot %121 %122 +%130 = OpLogicalOr %26 %27 %120 +%131 = OpLogicalAnd %26 %27 %120 +%132 = OpLogicalOr %26 %27 %120 +%133 = OpLogicalOr %112 %123 %124 +%134 = OpLogicalAnd %26 %27 %120 +%135 = OpLogicalAnd %34 %125 %126 OpReturn OpFunctionEnd -%138 = OpFunction %5 None %139 -%140 = OpFunctionParameter %5 -%141 = OpFunctionParameter %5 -%142 = OpLabel -%143 = OpIEqual %8 %141 %29 -%144 = OpIEqual %8 %140 %60 -%145 = OpIEqual %8 %141 %61 -%146 = OpLogicalAnd %8 %144 %145 -%147 = OpLogicalOr %8 %143 %146 -%148 = OpSelect %5 %147 %23 %141 -%149 = OpSDiv %5 %140 %148 -OpReturnValue %149 +%136 = OpFunction %5 None %137 +%138 = OpFunctionParameter %5 +%139 = OpFunctionParameter %5 +%140 = OpLabel +%141 = OpIEqual %26 %139 %28 +%142 = OpIEqual %26 %138 %57 +%143 = OpIEqual %26 %139 %58 +%144 = OpLogicalAnd %26 %142 %143 +%145 = OpLogicalOr %26 %141 %144 +%146 = OpSelect %5 %145 %21 %139 +%147 = OpSDiv %5 %138 %146 +OpReturnValue %147 OpFunctionEnd -%150 = OpFunction %12 None %151 -%152 = OpFunctionParameter %12 -%153 = OpFunctionParameter %12 -%154 = OpLabel -%156 = OpIEqual %8 %153 %155 -%158 = OpSelect %12 %156 %157 %153 -%159 = OpUDiv %12 %152 %158 -OpReturnValue %159 +%148 = OpFunction %10 None %149 +%150 = OpFunctionParameter %10 +%151 = OpFunctionParameter %10 +%152 = OpLabel +%154 = OpIEqual %26 %151 %153 +%156 = OpSelect %10 %154 %155 %151 +%157 = OpUDiv %10 %150 %156 +OpReturnValue %157 OpFunctionEnd -%161 = OpFunction %160 None %162 -%163 = OpFunctionParameter %160 -%164 = OpFunctionParameter %160 -%165 = OpLabel -%167 = OpIEqual %123 %164 %166 -%170 = OpIEqual %123 %163 %168 -%171 = OpIEqual %123 %164 %169 -%172 = OpLogicalAnd %123 %170 %171 -%173 = OpLogicalOr %123 %167 %172 -%175 = OpSelect %160 %173 %174 %164 -%176 = OpSDiv %160 %163 %175 -OpReturnValue %176 +%159 = OpFunction %158 None %160 +%161 = OpFunctionParameter %158 +%162 = OpFunctionParameter %158 +%163 = OpLabel +%165 = OpIEqual %121 %162 %164 +%168 = OpIEqual %121 %161 %166 +%169 = OpIEqual %121 %162 %167 +%170 = OpLogicalAnd %121 %168 %169 +%171 = OpLogicalOr %121 %165 %170 +%173 = OpSelect %158 %171 %172 %162 +%174 = OpSDiv %158 %161 %173 +OpReturnValue %174 OpFunctionEnd -%177 = OpFunction %11 None %178 -%179 = OpFunctionParameter %11 -%180 = OpFunctionParameter %11 -%181 = OpLabel -%183 = OpIEqual %115 %180 %182 -%185 = OpSelect %11 %183 %184 %180 -%186 = OpUDiv %11 %179 %185 -OpReturnValue %186 +%175 = OpFunction %9 None %176 +%177 = OpFunctionParameter %9 +%178 = OpFunctionParameter %9 +%179 = OpLabel +%181 = OpIEqual %112 %178 %180 +%183 = OpSelect %9 %181 %182 %178 +%184 = OpUDiv %9 %177 %183 +OpReturnValue %184 OpFunctionEnd -%187 = OpFunction %5 None %139 -%188 = OpFunctionParameter %5 -%189 = OpFunctionParameter %5 -%190 = OpLabel -%191 = OpIEqual %8 %189 %29 -%192 = OpIEqual %8 %188 %60 -%193 = OpIEqual %8 %189 %61 -%194 = OpLogicalAnd %8 %192 %193 -%195 = OpLogicalOr %8 %191 %194 -%196 = OpSelect %5 %195 %23 %189 -%197 = OpSRem %5 %188 %196 -OpReturnValue %197 +%185 = OpFunction %5 None %137 +%186 = OpFunctionParameter %5 +%187 = OpFunctionParameter %5 +%188 = OpLabel +%189 = OpIEqual %26 %187 %28 +%190 = OpIEqual %26 %186 %57 +%191 = OpIEqual %26 %187 %58 +%192 = OpLogicalAnd %26 %190 %191 +%193 = OpLogicalOr %26 %189 %192 +%194 = OpSelect %5 %193 %21 %187 +%195 = OpSRem %5 %186 %194 +OpReturnValue %195 OpFunctionEnd -%198 = OpFunction %12 None %151 -%199 = OpFunctionParameter %12 -%200 = OpFunctionParameter %12 -%201 = OpLabel -%202 = OpIEqual %8 %200 %155 -%203 = OpSelect %12 %202 %157 %200 -%204 = OpUMod %12 %199 %203 -OpReturnValue %204 +%196 = OpFunction %10 None %149 +%197 = OpFunctionParameter %10 +%198 = OpFunctionParameter %10 +%199 = OpLabel +%200 = OpIEqual %26 %198 %153 +%201 = OpSelect %10 %200 %155 %198 +%202 = OpUMod %10 %197 %201 +OpReturnValue %202 OpFunctionEnd -%205 = OpFunction %160 None %162 -%206 = OpFunctionParameter %160 -%207 = OpFunctionParameter %160 -%208 = OpLabel -%209 = OpIEqual %123 %207 %166 -%210 = OpIEqual %123 %206 %168 -%211 = OpIEqual %123 %207 %169 -%212 = OpLogicalAnd %123 %210 %211 -%213 = OpLogicalOr %123 %209 %212 -%214 = OpSelect %160 %213 %174 %207 -%215 = OpSRem %160 %206 %214 -OpReturnValue %215 +%203 = OpFunction %158 None %160 +%204 = OpFunctionParameter %158 +%205 = OpFunctionParameter %158 +%206 = OpLabel +%207 = OpIEqual %121 %205 %164 +%208 = OpIEqual %121 %204 %166 +%209 = OpIEqual %121 %205 %167 +%210 = OpLogicalAnd %121 %208 %209 +%211 = OpLogicalOr %121 %207 %210 +%212 = OpSelect %158 %211 %172 %205 +%213 = OpSRem %158 %204 %212 +OpReturnValue %213 OpFunctionEnd -%216 = OpFunction %11 None %178 -%217 = OpFunctionParameter %11 -%218 = OpFunctionParameter %11 -%219 = OpLabel -%220 = OpIEqual %115 %218 %182 -%221 = OpSelect %11 %220 %184 %218 -%222 = OpUMod %11 %217 %221 -OpReturnValue %222 +%214 = OpFunction %9 None %176 +%215 = OpFunctionParameter %9 +%216 = OpFunctionParameter %9 +%217 = OpLabel +%218 = OpIEqual %112 %216 %180 +%219 = OpSelect %9 %218 %182 %216 +%220 = OpUMod %9 %215 %219 +OpReturnValue %220 OpFunctionEnd -%224 = OpFunction %223 None %225 -%226 = OpFunctionParameter %223 -%227 = OpFunctionParameter %223 -%228 = OpLabel -%230 = OpIEqual %123 %227 %229 -%232 = OpSelect %223 %230 %231 %227 -%233 = OpUDiv %223 %226 %232 -OpReturnValue %233 +%222 = OpFunction %221 None %223 +%224 = OpFunctionParameter %221 +%225 = OpFunctionParameter %221 +%226 = OpLabel +%228 = OpIEqual %121 %225 %227 +%230 = OpSelect %221 %228 %229 %225 +%231 = OpUDiv %221 %224 %230 +OpReturnValue %231 OpFunctionEnd -%234 = OpFunction %223 None %225 -%235 = OpFunctionParameter %223 -%236 = OpFunctionParameter %223 -%237 = OpLabel -%238 = OpIEqual %123 %236 %229 -%239 = OpSelect %223 %238 %231 %236 -%240 = OpUMod %223 %235 %239 -OpReturnValue %240 +%232 = OpFunction %221 None %223 +%233 = OpFunctionParameter %221 +%234 = OpFunctionParameter %221 +%235 = OpLabel +%236 = OpIEqual %121 %234 %227 +%237 = OpSelect %221 %236 %229 %234 +%238 = OpUMod %221 %233 %237 +OpReturnValue %238 OpFunctionEnd -%242 = OpFunction %2 None %122 -%241 = OpLabel -%253 = OpVariable %254 Function %255 -%256 = OpVariable %254 Function %257 -OpBranch %258 -%258 = OpLabel -%259 = OpFNegate %3 %17 -%260 = OpSNegate %160 %174 -%261 = OpFNegate %9 %97 -%262 = OpIAdd %5 %82 %23 -%263 = OpIAdd %12 %243 %157 -%264 = OpFAdd %3 %76 %17 -%265 = OpIAdd %160 %244 %174 -%266 = OpIAdd %11 %245 %184 -%267 = OpFAdd %4 %246 %247 -%268 = OpISub %5 %82 %23 -%269 = OpISub %12 %243 %157 -%270 = OpFSub %3 %76 %17 -%271 = OpISub %160 %244 %174 -%272 = OpISub %11 %245 %184 -%273 = OpFSub %4 %246 %247 -%274 = OpIMul %5 %82 %23 -%275 = OpIMul %12 %243 %157 -%276 = OpFMul %3 %76 %17 -%277 = OpIMul %160 %244 %174 -%278 = OpIMul %11 %245 %184 -%279 = OpFMul %4 %246 %247 -%280 = OpFunctionCall %5 %138 %82 %23 -%281 = OpFunctionCall %12 %150 %243 %157 -%282 = OpFDiv %3 %76 %17 -%283 = OpFunctionCall %160 %161 %244 %174 -%284 = OpFunctionCall %11 %177 %245 %184 -%285 = OpFDiv %4 %246 %247 -%286 = OpFunctionCall %5 %187 %82 %23 -%287 = OpFunctionCall %12 %198 %243 %157 -%288 = OpFRem %3 %76 %17 -%289 = OpFunctionCall %160 %205 %244 %174 -%290 = OpFunctionCall %11 %216 %245 %184 -%291 = OpFRem %4 %246 %247 -OpBranch %292 -%292 = OpLabel -%294 = OpIAdd %160 %244 %174 -%295 = OpIAdd %160 %244 %174 -%296 = OpIAdd %223 %248 %231 -%297 = OpIAdd %223 %248 %231 -%298 = OpFAdd %9 %77 %97 -%299 = OpFAdd %9 %77 %97 -%300 = OpISub %160 %244 %174 -%301 = OpISub %160 %244 %174 -%302 = OpISub %223 %248 %231 -%303 = OpISub %223 %248 %231 -%304 = OpFSub %9 %77 %97 -%305 = OpFSub %9 %77 %97 -%307 = OpCompositeConstruct %160 %23 %23 -%306 = OpIMul %160 %244 %307 -%309 = OpCompositeConstruct %160 %82 %82 -%308 = OpIMul %160 %174 %309 -%311 = OpCompositeConstruct %223 %157 %157 -%310 = OpIMul %223 %248 %311 -%313 = OpCompositeConstruct %223 %243 %243 -%312 = OpIMul %223 %231 %313 -%314 = OpVectorTimesScalar %9 %77 %17 -%315 = OpVectorTimesScalar %9 %97 %76 -%316 = OpFunctionCall %160 %161 %244 %174 -%317 = OpFunctionCall %160 %161 %244 %174 -%318 = OpFunctionCall %223 %224 %248 %231 -%319 = OpFunctionCall %223 %224 %248 %231 -%320 = OpFDiv %9 %77 %97 -%321 = OpFDiv %9 %77 %97 -%322 = OpFunctionCall %160 %205 %244 %174 -%323 = OpFunctionCall %160 %205 %244 %174 -%324 = OpFunctionCall %223 %234 %248 %231 -%325 = OpFunctionCall %223 %234 %248 %231 -%326 = OpFRem %9 %77 %97 -%327 = OpFRem %9 %77 %97 -OpBranch %293 -%293 = OpLabel -%329 = OpCompositeExtract %10 %249 0 -%330 = OpCompositeExtract %10 %249 0 -%331 = OpFAdd %10 %329 %330 -%332 = OpCompositeExtract %10 %249 1 -%333 = OpCompositeExtract %10 %249 1 -%334 = OpFAdd %10 %332 %333 -%335 = OpCompositeExtract %10 %249 2 -%336 = OpCompositeExtract %10 %249 2 -%337 = OpFAdd %10 %335 %336 -%328 = OpCompositeConstruct %13 %331 %334 %337 -%339 = OpCompositeExtract %10 %249 0 -%340 = OpCompositeExtract %10 %249 0 -%341 = OpFSub %10 %339 %340 -%342 = OpCompositeExtract %10 %249 1 -%343 = OpCompositeExtract %10 %249 1 -%344 = OpFSub %10 %342 %343 -%345 = OpCompositeExtract %10 %249 2 -%346 = OpCompositeExtract %10 %249 2 -%347 = OpFSub %10 %345 %346 -%338 = OpCompositeConstruct %13 %341 %344 %347 -%348 = OpMatrixTimesScalar %13 %249 %17 -%349 = OpMatrixTimesScalar %13 %249 %76 -%350 = OpMatrixTimesVector %10 %250 %247 -%351 = OpVectorTimesMatrix %4 %251 %250 -%352 = OpMatrixTimesMatrix %13 %250 %252 -%353 = OpLoad %5 %253 -%354 = OpIAdd %5 %353 %60 -OpStore %256 %354 +%240 = OpFunction %2 None %119 +%239 = OpLabel +%251 = OpVariable %252 Function %253 +%254 = OpVariable %252 Function %255 +OpBranch %256 +%256 = OpLabel +%257 = OpFNegate %3 %15 +%258 = OpSNegate %158 %172 +%259 = OpFNegate %7 %94 +%260 = OpIAdd %5 %79 %21 +%261 = OpIAdd %10 %241 %155 +%262 = OpFAdd %3 %73 %15 +%263 = OpIAdd %158 %242 %172 +%264 = OpIAdd %9 %243 %182 +%265 = OpFAdd %4 %244 %245 +%266 = OpISub %5 %79 %21 +%267 = OpISub %10 %241 %155 +%268 = OpFSub %3 %73 %15 +%269 = OpISub %158 %242 %172 +%270 = OpISub %9 %243 %182 +%271 = OpFSub %4 %244 %245 +%272 = OpIMul %5 %79 %21 +%273 = OpIMul %10 %241 %155 +%274 = OpFMul %3 %73 %15 +%275 = OpIMul %158 %242 %172 +%276 = OpIMul %9 %243 %182 +%277 = OpFMul %4 %244 %245 +%278 = OpFunctionCall %5 %136 %79 %21 +%279 = OpFunctionCall %10 %148 %241 %155 +%280 = OpFDiv %3 %73 %15 +%281 = OpFunctionCall %158 %159 %242 %172 +%282 = OpFunctionCall %9 %175 %243 %182 +%283 = OpFDiv %4 %244 %245 +%284 = OpFunctionCall %5 %185 %79 %21 +%285 = OpFunctionCall %10 %196 %241 %155 +%286 = OpFRem %3 %73 %15 +%287 = OpFunctionCall %158 %203 %242 %172 +%288 = OpFunctionCall %9 %214 %243 %182 +%289 = OpFRem %4 %244 %245 +OpBranch %290 +%290 = OpLabel +%292 = OpIAdd %158 %242 %172 +%293 = OpIAdd %158 %242 %172 +%294 = OpIAdd %221 %246 %229 +%295 = OpIAdd %221 %246 %229 +%296 = OpFAdd %7 %74 %94 +%297 = OpFAdd %7 %74 %94 +%298 = OpISub %158 %242 %172 +%299 = OpISub %158 %242 %172 +%300 = OpISub %221 %246 %229 +%301 = OpISub %221 %246 %229 +%302 = OpFSub %7 %74 %94 +%303 = OpFSub %7 %74 %94 +%305 = OpCompositeConstruct %158 %21 %21 +%304 = OpIMul %158 %242 %305 +%307 = OpCompositeConstruct %158 %79 %79 +%306 = OpIMul %158 %172 %307 +%309 = OpCompositeConstruct %221 %155 %155 +%308 = OpIMul %221 %246 %309 +%311 = OpCompositeConstruct %221 %241 %241 +%310 = OpIMul %221 %229 %311 +%312 = OpVectorTimesScalar %7 %74 %15 +%313 = OpVectorTimesScalar %7 %94 %73 +%314 = OpFunctionCall %158 %159 %242 %172 +%315 = OpFunctionCall %158 %159 %242 %172 +%316 = OpFunctionCall %221 %222 %246 %229 +%317 = OpFunctionCall %221 %222 %246 %229 +%318 = OpFDiv %7 %74 %94 +%319 = OpFDiv %7 %74 %94 +%320 = OpFunctionCall %158 %203 %242 %172 +%321 = OpFunctionCall %158 %203 %242 %172 +%322 = OpFunctionCall %221 %232 %246 %229 +%323 = OpFunctionCall %221 %232 %246 %229 +%324 = OpFRem %7 %74 %94 +%325 = OpFRem %7 %74 %94 +OpBranch %291 +%291 = OpLabel +%327 = OpCompositeExtract %8 %247 0 +%328 = OpCompositeExtract %8 %247 0 +%329 = OpFAdd %8 %327 %328 +%330 = OpCompositeExtract %8 %247 1 +%331 = OpCompositeExtract %8 %247 1 +%332 = OpFAdd %8 %330 %331 +%333 = OpCompositeExtract %8 %247 2 +%334 = OpCompositeExtract %8 %247 2 +%335 = OpFAdd %8 %333 %334 +%326 = OpCompositeConstruct %11 %329 %332 %335 +%337 = OpCompositeExtract %8 %247 0 +%338 = OpCompositeExtract %8 %247 0 +%339 = OpFSub %8 %337 %338 +%340 = OpCompositeExtract %8 %247 1 +%341 = OpCompositeExtract %8 %247 1 +%342 = OpFSub %8 %340 %341 +%343 = OpCompositeExtract %8 %247 2 +%344 = OpCompositeExtract %8 %247 2 +%345 = OpFSub %8 %343 %344 +%336 = OpCompositeConstruct %11 %339 %342 %345 +%346 = OpMatrixTimesScalar %11 %247 %15 +%347 = OpMatrixTimesScalar %11 %247 %73 +%348 = OpMatrixTimesVector %8 %248 %245 +%349 = OpVectorTimesMatrix %4 %249 %248 +%350 = OpMatrixTimesMatrix %11 %248 %250 +%351 = OpLoad %5 %251 +%352 = OpIAdd %5 %351 %57 +OpStore %254 %352 OpReturn OpFunctionEnd -%356 = OpFunction %2 None %122 +%354 = OpFunction %2 None %119 +%353 = OpLabel +OpBranch %355 %355 = OpLabel -OpBranch %357 -%357 = OpLabel -%358 = OpNot %5 %23 -%359 = OpNot %12 %157 -%360 = OpNot %160 %174 -%361 = OpNot %11 %184 -%362 = OpBitwiseOr %5 %82 %23 -%363 = OpBitwiseOr %12 %243 %157 -%364 = OpBitwiseOr %160 %244 %174 -%365 = OpBitwiseOr %11 %245 %184 -%366 = OpBitwiseAnd %5 %82 %23 -%367 = OpBitwiseAnd %12 %243 %157 -%368 = OpBitwiseAnd %160 %244 %174 -%369 = OpBitwiseAnd %11 %245 %184 -%370 = OpBitwiseXor %5 %82 %23 -%371 = OpBitwiseXor %12 %243 %157 -%372 = OpBitwiseXor %160 %244 %174 -%373 = OpBitwiseXor %11 %245 %184 -%374 = OpShiftLeftLogical %5 %82 %157 -%375 = OpShiftLeftLogical %12 %243 %157 -%376 = OpShiftLeftLogical %160 %244 %231 -%377 = OpShiftLeftLogical %11 %245 %184 -%378 = OpShiftRightArithmetic %5 %82 %157 -%379 = OpShiftRightLogical %12 %243 %157 -%380 = OpShiftRightArithmetic %160 %244 %231 -%381 = OpShiftRightLogical %11 %245 %184 +%356 = OpNot %5 %21 +%357 = OpNot %10 %155 +%358 = OpNot %158 %172 +%359 = OpNot %9 %182 +%360 = OpBitwiseOr %5 %79 %21 +%361 = OpBitwiseOr %10 %241 %155 +%362 = OpBitwiseOr %158 %242 %172 +%363 = OpBitwiseOr %9 %243 %182 +%364 = OpBitwiseAnd %5 %79 %21 +%365 = OpBitwiseAnd %10 %241 %155 +%366 = OpBitwiseAnd %158 %242 %172 +%367 = OpBitwiseAnd %9 %243 %182 +%368 = OpBitwiseXor %5 %79 %21 +%369 = OpBitwiseXor %10 %241 %155 +%370 = OpBitwiseXor %158 %242 %172 +%371 = OpBitwiseXor %9 %243 %182 +%372 = OpShiftLeftLogical %5 %79 %155 +%373 = OpShiftLeftLogical %10 %241 %155 +%374 = OpShiftLeftLogical %158 %242 %229 +%375 = OpShiftLeftLogical %9 %243 %182 +%376 = OpShiftRightArithmetic %5 %79 %155 +%377 = OpShiftRightLogical %10 %241 %155 +%378 = OpShiftRightArithmetic %158 %242 %229 +%379 = OpShiftRightLogical %9 %243 %182 OpReturn OpFunctionEnd -%383 = OpFunction %2 None %122 +%381 = OpFunction %2 None %119 +%380 = OpLabel +OpBranch %382 %382 = OpLabel -OpBranch %384 -%384 = OpLabel -%385 = OpIEqual %8 %82 %23 -%386 = OpIEqual %8 %243 %157 -%387 = OpFOrdEqual %8 %76 %17 -%388 = OpIEqual %123 %244 %174 -%389 = OpIEqual %115 %245 %184 -%390 = OpFOrdEqual %7 %246 %247 -%391 = OpINotEqual %8 %82 %23 -%392 = OpINotEqual %8 %243 %157 -%393 = OpFOrdNotEqual %8 %76 %17 -%394 = OpINotEqual %123 %244 %174 -%395 = OpINotEqual %115 %245 %184 -%396 = OpFOrdNotEqual %7 %246 %247 -%397 = OpSLessThan %8 %82 %23 -%398 = OpULessThan %8 %243 %157 -%399 = OpFOrdLessThan %8 %76 %17 -%400 = OpSLessThan %123 %244 %174 -%401 = OpULessThan %115 %245 %184 -%402 = OpFOrdLessThan %7 %246 %247 -%403 = OpSLessThanEqual %8 %82 %23 -%404 = OpULessThanEqual %8 %243 %157 -%405 = OpFOrdLessThanEqual %8 %76 %17 -%406 = OpSLessThanEqual %123 %244 %174 -%407 = OpULessThanEqual %115 %245 %184 -%408 = OpFOrdLessThanEqual %7 %246 %247 -%409 = OpSGreaterThan %8 %82 %23 -%410 = OpUGreaterThan %8 %243 %157 -%411 = OpFOrdGreaterThan %8 %76 %17 -%412 = OpSGreaterThan %123 %244 %174 -%413 = OpUGreaterThan %115 %245 %184 -%414 = OpFOrdGreaterThan %7 %246 %247 -%415 = OpSGreaterThanEqual %8 %82 %23 -%416 = OpUGreaterThanEqual %8 %243 %157 -%417 = OpFOrdGreaterThanEqual %8 %76 %17 -%418 = OpSGreaterThanEqual %123 %244 %174 -%419 = OpUGreaterThanEqual %115 %245 %184 -%420 = OpFOrdGreaterThanEqual %7 %246 %247 +%383 = OpIEqual %26 %79 %21 +%384 = OpIEqual %26 %241 %155 +%385 = OpFOrdEqual %26 %73 %15 +%386 = OpIEqual %121 %242 %172 +%387 = OpIEqual %112 %243 %182 +%388 = OpFOrdEqual %34 %244 %245 +%389 = OpINotEqual %26 %79 %21 +%390 = OpINotEqual %26 %241 %155 +%391 = OpFOrdNotEqual %26 %73 %15 +%392 = OpINotEqual %121 %242 %172 +%393 = OpINotEqual %112 %243 %182 +%394 = OpFOrdNotEqual %34 %244 %245 +%395 = OpSLessThan %26 %79 %21 +%396 = OpULessThan %26 %241 %155 +%397 = OpFOrdLessThan %26 %73 %15 +%398 = OpSLessThan %121 %242 %172 +%399 = OpULessThan %112 %243 %182 +%400 = OpFOrdLessThan %34 %244 %245 +%401 = OpSLessThanEqual %26 %79 %21 +%402 = OpULessThanEqual %26 %241 %155 +%403 = OpFOrdLessThanEqual %26 %73 %15 +%404 = OpSLessThanEqual %121 %242 %172 +%405 = OpULessThanEqual %112 %243 %182 +%406 = OpFOrdLessThanEqual %34 %244 %245 +%407 = OpSGreaterThan %26 %79 %21 +%408 = OpUGreaterThan %26 %241 %155 +%409 = OpFOrdGreaterThan %26 %73 %15 +%410 = OpSGreaterThan %121 %242 %172 +%411 = OpUGreaterThan %112 %243 %182 +%412 = OpFOrdGreaterThan %34 %244 %245 +%413 = OpSGreaterThanEqual %26 %79 %21 +%414 = OpUGreaterThanEqual %26 %241 %155 +%415 = OpFOrdGreaterThanEqual %26 %73 %15 +%416 = OpSGreaterThanEqual %121 %242 %172 +%417 = OpUGreaterThanEqual %112 %243 %182 +%418 = OpFOrdGreaterThanEqual %34 %244 %245 OpReturn OpFunctionEnd -%422 = OpFunction %2 None %122 -%421 = OpLabel -%424 = OpVariable %254 Function %425 -%426 = OpVariable %427 Function %423 -OpBranch %428 -%428 = OpLabel -OpStore %424 %23 -%429 = OpLoad %5 %424 -%430 = OpIAdd %5 %429 %23 -OpStore %424 %430 -%431 = OpLoad %5 %424 -%432 = OpISub %5 %431 %23 -OpStore %424 %432 -%433 = OpLoad %5 %424 -%434 = OpLoad %5 %424 -%435 = OpIMul %5 %434 %433 -OpStore %424 %435 -%436 = OpLoad %5 %424 -%437 = OpLoad %5 %424 -%438 = OpFunctionCall %5 %138 %437 %436 -OpStore %424 %438 -%439 = OpLoad %5 %424 -%440 = OpFunctionCall %5 %187 %439 %23 -OpStore %424 %440 -%441 = OpLoad %5 %424 -%442 = OpBitwiseAnd %5 %441 %29 -OpStore %424 %442 -%443 = OpLoad %5 %424 -%444 = OpBitwiseOr %5 %443 %29 -OpStore %424 %444 -%445 = OpLoad %5 %424 -%446 = OpBitwiseXor %5 %445 %29 -OpStore %424 %446 -%447 = OpLoad %5 %424 -%448 = OpShiftLeftLogical %5 %447 %243 -OpStore %424 %448 -%449 = OpLoad %5 %424 -%450 = OpShiftRightArithmetic %5 %449 %157 -OpStore %424 %450 -%451 = OpLoad %5 %424 -%452 = OpIAdd %5 %451 %23 -OpStore %424 %452 -%453 = OpLoad %5 %424 -%454 = OpISub %5 %453 %23 -OpStore %424 %454 -%455 = OpAccessChain %254 %426 %157 -%456 = OpLoad %5 %455 -%457 = OpIAdd %5 %456 %23 -%458 = OpAccessChain %254 %426 %157 -OpStore %458 %457 -%459 = OpAccessChain %254 %426 %157 -%460 = OpLoad %5 %459 -%461 = OpISub %5 %460 %23 -%462 = OpAccessChain %254 %426 %157 -OpStore %462 %461 +%420 = OpFunction %2 None %119 +%419 = OpLabel +%422 = OpVariable %252 Function %423 +%424 = OpVariable %425 Function %421 +OpBranch %426 +%426 = OpLabel +OpStore %422 %21 +%427 = OpLoad %5 %422 +%428 = OpIAdd %5 %427 %21 +OpStore %422 %428 +%429 = OpLoad %5 %422 +%430 = OpISub %5 %429 %21 +OpStore %422 %430 +%431 = OpLoad %5 %422 +%432 = OpLoad %5 %422 +%433 = OpIMul %5 %432 %431 +OpStore %422 %433 +%434 = OpLoad %5 %422 +%435 = OpLoad %5 %422 +%436 = OpFunctionCall %5 %136 %435 %434 +OpStore %422 %436 +%437 = OpLoad %5 %422 +%438 = OpFunctionCall %5 %185 %437 %21 +OpStore %422 %438 +%439 = OpLoad %5 %422 +%440 = OpBitwiseAnd %5 %439 %28 +OpStore %422 %440 +%441 = OpLoad %5 %422 +%442 = OpBitwiseOr %5 %441 %28 +OpStore %422 %442 +%443 = OpLoad %5 %422 +%444 = OpBitwiseXor %5 %443 %28 +OpStore %422 %444 +%445 = OpLoad %5 %422 +%446 = OpShiftLeftLogical %5 %445 %241 +OpStore %422 %446 +%447 = OpLoad %5 %422 +%448 = OpShiftRightArithmetic %5 %447 %155 +OpStore %422 %448 +%449 = OpLoad %5 %422 +%450 = OpIAdd %5 %449 %21 +OpStore %422 %450 +%451 = OpLoad %5 %422 +%452 = OpISub %5 %451 %21 +OpStore %422 %452 +%453 = OpAccessChain %252 %424 %155 +%454 = OpLoad %5 %453 +%455 = OpIAdd %5 %454 %21 +%456 = OpAccessChain %252 %424 %155 +OpStore %456 %455 +%457 = OpAccessChain %252 %424 %155 +%458 = OpLoad %5 %457 +%459 = OpISub %5 %458 %21 +%460 = OpAccessChain %252 %424 %155 +OpStore %460 %459 OpReturn OpFunctionEnd -%464 = OpFunction %2 None %122 +%462 = OpFunction %2 None %119 +%461 = OpLabel +OpBranch %463 %463 = OpLabel -OpBranch %465 -%465 = OpLabel -%466 = OpSNegate %5 %23 -%467 = OpSNegate %5 %23 +%464 = OpSNegate %5 %21 +%465 = OpSNegate %5 %21 +%466 = OpSNegate %5 %465 +%467 = OpSNegate %5 %21 %468 = OpSNegate %5 %467 -%469 = OpSNegate %5 %23 +%469 = OpSNegate %5 %21 %470 = OpSNegate %5 %469 -%471 = OpSNegate %5 %23 +%471 = OpSNegate %5 %21 %472 = OpSNegate %5 %471 -%473 = OpSNegate %5 %23 -%474 = OpSNegate %5 %473 +%473 = OpSNegate %5 %472 +%474 = OpSNegate %5 %21 %475 = OpSNegate %5 %474 -%476 = OpSNegate %5 %23 +%476 = OpSNegate %5 %475 %477 = OpSNegate %5 %476 -%478 = OpSNegate %5 %477 +%478 = OpSNegate %5 %21 %479 = OpSNegate %5 %478 -%480 = OpSNegate %5 %23 +%480 = OpSNegate %5 %479 %481 = OpSNegate %5 %480 %482 = OpSNegate %5 %481 -%483 = OpSNegate %5 %482 +%483 = OpSNegate %5 %21 %484 = OpSNegate %5 %483 -%485 = OpSNegate %5 %23 +%485 = OpSNegate %5 %484 %486 = OpSNegate %5 %485 %487 = OpSNegate %5 %486 -%488 = OpSNegate %5 %487 -%489 = OpSNegate %5 %488 -%490 = OpFNegate %3 %17 -%491 = OpFNegate %3 %17 +%488 = OpFNegate %3 %15 +%489 = OpFNegate %3 %15 +%490 = OpFNegate %3 %489 +%491 = OpFNegate %3 %15 %492 = OpFNegate %3 %491 -%493 = OpFNegate %3 %17 +%493 = OpFNegate %3 %15 %494 = OpFNegate %3 %493 -%495 = OpFNegate %3 %17 +%495 = OpFNegate %3 %15 %496 = OpFNegate %3 %495 -%497 = OpFNegate %3 %17 -%498 = OpFNegate %3 %497 +%497 = OpFNegate %3 %496 +%498 = OpFNegate %3 %15 %499 = OpFNegate %3 %498 -%500 = OpFNegate %3 %17 +%500 = OpFNegate %3 %499 %501 = OpFNegate %3 %500 -%502 = OpFNegate %3 %501 +%502 = OpFNegate %3 %15 %503 = OpFNegate %3 %502 -%504 = OpFNegate %3 %17 +%504 = OpFNegate %3 %503 %505 = OpFNegate %3 %504 %506 = OpFNegate %3 %505 -%507 = OpFNegate %3 %506 +%507 = OpFNegate %3 %15 %508 = OpFNegate %3 %507 -%509 = OpFNegate %3 %17 +%509 = OpFNegate %3 %508 %510 = OpFNegate %3 %509 %511 = OpFNegate %3 %510 -%512 = OpFNegate %3 %511 -%513 = OpFNegate %3 %512 OpReturn OpFunctionEnd -%518 = OpFunction %2 None %122 -%514 = OpLabel -%517 = OpLoad %11 %515 -OpBranch %520 -%520 = OpLabel -%521 = OpFunctionCall %4 %26 -%522 = OpCompositeExtract %12 %517 0 -%523 = OpConvertUToF %3 %522 -%524 = OpCompositeExtract %12 %517 1 -%525 = OpBitcast %5 %524 -%526 = OpFunctionCall %4 %74 %523 %525 -%527 = OpFunctionCall %9 %95 -%528 = OpFunctionCall %10 %112 %519 -%529 = OpFunctionCall %2 %121 -%530 = OpFunctionCall %2 %242 -%531 = OpFunctionCall %2 %356 -%532 = OpFunctionCall %2 %383 -%533 = OpFunctionCall %2 %422 -%534 = OpFunctionCall %2 %464 +%516 = OpFunction %2 None %119 +%512 = OpLabel +%515 = OpLoad %9 %513 +OpBranch %518 +%518 = OpLabel +%519 = OpFunctionCall %4 %24 +%520 = OpCompositeExtract %10 %515 0 +%521 = OpConvertUToF %3 %520 +%522 = OpCompositeExtract %10 %515 1 +%523 = OpBitcast %5 %522 +%524 = OpFunctionCall %4 %71 %521 %523 +%525 = OpFunctionCall %7 %92 +%526 = OpFunctionCall %8 %109 %517 +%527 = OpFunctionCall %2 %118 +%528 = OpFunctionCall %2 %240 +%529 = OpFunctionCall %2 %354 +%530 = OpFunctionCall %2 %381 +%531 = OpFunctionCall %2 %420 +%532 = OpFunctionCall %2 %462 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/wgsl-select.spvasm b/naga/tests/out/spv/wgsl-select.spvasm new file mode 100644 index 000000000..d0011c464 --- /dev/null +++ b/naga/tests/out/spv/wgsl-select.spvasm @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 36 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %8 "main" +OpExecutionMode %8 LocalSize 1 1 1 +%2 = OpTypeVoid +%4 = OpTypeFloat 32 +%3 = OpTypeVector %4 2 +%6 = OpTypeInt 32 1 +%5 = OpTypeVector %6 2 +%9 = OpTypeFunction %2 +%10 = OpConstant %4 1.0 +%11 = OpConstant %6 1 +%12 = OpConstant %6 2 +%13 = OpConstantComposite %5 %11 %12 +%14 = OpConstant %4 0.0 +%15 = OpConstantComposite %3 %10 %14 +%16 = OpConstantComposite %3 %14 %10 +%18 = OpTypePointer Function %5 +%20 = OpTypePointer Function %3 +%21 = OpConstantNull %3 +%23 = OpTypePointer Function %6 +%25 = OpTypeInt 32 0 +%24 = OpConstant %25 0 +%28 = OpConstant %25 1 +%31 = OpTypeBool +%34 = OpTypeVector %31 2 +%8 = OpFunction %2 None %9 +%7 = OpLabel +%17 = OpVariable %18 Function %13 +%19 = OpVariable %20 Function %21 +OpBranch %22 +%22 = OpLabel +%26 = OpAccessChain %23 %17 %24 +%27 = OpLoad %6 %26 +%29 = OpAccessChain %23 %17 %28 +%30 = OpLoad %6 %29 +%32 = OpSLessThan %31 %27 %30 +%35 = OpCompositeConstruct %34 %32 %32 +%33 = OpSelect %3 %35 %16 %15 +OpStore %19 %33 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/wgsl-operators.wgsl b/naga/tests/out/wgsl/wgsl-operators.wgsl index 0506880a6..e0664e4e1 100644 --- a/naga/tests/out/wgsl/wgsl-operators.wgsl +++ b/naga/tests/out/wgsl/wgsl-operators.wgsl @@ -6,7 +6,7 @@ const v_i32_one: vec4 = vec4(1i, 1i, 1i, 1i); fn builtins() -> vec4 { let s1_ = select(0i, 1i, true); let s2_ = select(v_f32_zero, v_f32_one, true); - let s3_ = select(v_f32_one, v_f32_zero, vec4(false, false, false, false)); + let s3_ = vec4(1f, 1f, 1f, 1f); let m1_ = mix(v_f32_zero, v_f32_one, v_f32_half); let m2_ = mix(v_f32_zero, v_f32_one, 0.1f); let b1_ = bitcast(1i); diff --git a/naga/tests/out/wgsl/wgsl-select.wgsl b/naga/tests/out/wgsl/wgsl-select.wgsl new file mode 100644 index 000000000..b580dad3f --- /dev/null +++ b/naga/tests/out/wgsl/wgsl-select.wgsl @@ -0,0 +1,10 @@ +@compute @workgroup_size(1, 1, 1) +fn main() { + var x0_: vec2 = vec2(1i, 2i); + var i1_: vec2; + + let _e12 = x0_.x; + let _e14 = x0_.y; + i1_ = select(vec2(1f, 0f), vec2(0f, 1f), (_e12 < _e14)); + return; +}