From 119b4efada475f95507f8f577bf1abfe3d529fd0 Mon Sep 17 00:00:00 2001 From: Andy Leiserson Date: Wed, 19 Nov 2025 17:06:49 -0800 Subject: [PATCH] [naga wgsl-in] Short-circuiting of && and || operators (#7339) Addresses parts of #4394 and #6302 --- CHANGELOG.md | 4 + naga/src/front/wgsl/lower/mod.rs | 296 ++++- naga/tests/in/wgsl/operators.wgsl | 6 + .../out/glsl/wgsl-operators.main.Compute.glsl | 57 +- naga/tests/out/hlsl/wgsl-operators.hlsl | 62 +- naga/tests/out/msl/wgsl-operators.msl | 61 +- naga/tests/out/spv/wgsl-operators.spvasm | 1109 +++++++++-------- naga/tests/out/wgsl/wgsl-operators.wgsl | 58 +- .../wgpu-gpu/subgroup_operations/shader.wgsl | 4 +- 9 files changed, 1092 insertions(+), 565 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 3b9c50c1b..a4b2b7880 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -152,6 +152,10 @@ By @SupaMaggie70Incorporated in [#8206](https://github.com/gfx-rs/wgpu/pull/8206 - Validate that buffers are unmapped in `write_buffer` calls. By @ErichDonGubler in [#8454](https://github.com/gfx-rs/wgpu/pull/8454). - Add WGSL parsing for mesh shaders. By @inner-daemons in [#8370](https://github.com/gfx-rs/wgpu/pull/8370). +#### naga + +- The `||` and `&&` operators now "short circuit", i.e., do not evaluate the RHS if the result can be determined from just the LHS. By @andyleiserson in [#7339](https://github.com/gfx-rs/wgpu/pull/7339). + #### DX12 - Align copies b/w textures and buffers via a single intermediate buffer per copy when `D3D12_FEATURE_DATA_D3D12_OPTIONS13.UnrestrictedBufferTextureCopyPitchSupported` is `false`. By @ErichDonGubler in [#7721](https://github.com/gfx-rs/wgpu/pull/7721). diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 33a1de6d5..c47941f0f 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -426,6 +426,13 @@ impl TypeContext for ExpressionContext<'_, '_, '_> { } impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { + const fn is_runtime(&self) -> bool { + match self.expr_type { + ExpressionContextType::Runtime(_) => true, + ExpressionContextType::Constant(_) | ExpressionContextType::Override => false, + } + } + #[allow(dead_code)] fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> { ExpressionContext { @@ -588,6 +595,16 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { } } + fn get(&self, handle: Handle) -> &crate::Expression { + match self.expr_type { + ExpressionContextType::Runtime(ref ctx) + | ExpressionContextType::Constant(Some(ref ctx)) => &ctx.function.expressions[handle], + ExpressionContextType::Constant(None) | ExpressionContextType::Override => { + &self.module.global_expressions[handle] + } + } + } + fn local( &mut self, local: &Handle, @@ -614,6 +631,52 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { } } + fn with_nested_runtime_expression_ctx<'a, F, T>( + &mut self, + span: Span, + f: F, + ) -> Result<'source, (T, crate::Block)> + where + for<'t> F: FnOnce(&mut ExpressionContext<'source, 't, 't>) -> Result<'source, T>, + { + let mut block = crate::Block::new(); + let rctx = match self.expr_type { + ExpressionContextType::Runtime(ref mut rctx) => Ok(rctx), + ExpressionContextType::Constant(_) | ExpressionContextType::Override => { + Err(Error::UnexpectedOperationInConstContext(span)) + } + }?; + + rctx.block + .extend(rctx.emitter.finish(&rctx.function.expressions)); + rctx.emitter.start(&rctx.function.expressions); + + let nested_rctx = LocalExpressionContext { + local_table: rctx.local_table, + function: rctx.function, + block: &mut block, + emitter: rctx.emitter, + typifier: rctx.typifier, + local_expression_kind_tracker: rctx.local_expression_kind_tracker, + }; + let mut nested_ctx = ExpressionContext { + expr_type: ExpressionContextType::Runtime(nested_rctx), + ast_expressions: self.ast_expressions, + types: self.types, + globals: self.globals, + module: self.module, + const_typifier: self.const_typifier, + layouter: self.layouter, + global_expression_kind_tracker: self.global_expression_kind_tracker, + }; + let ret = f(&mut nested_ctx)?; + + block.extend(rctx.emitter.finish(&rctx.function.expressions)); + rctx.emitter.start(&rctx.function.expressions); + + Ok((ret, block)) + } + fn gather_component( &mut self, expr: Handle, @@ -2375,6 +2438,130 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { expr.try_map(|handle| ctx.append_expression(handle, span)) } + /// Generate IR for the short-circuiting operators `&&` and `||`. + /// + /// `binary` has already lowered the LHS expression and resolved its type. + fn logical( + &mut self, + op: crate::BinaryOperator, + left: Handle, + right: Handle>, + span: Span, + ctx: &mut ExpressionContext<'source, '_, '_>, + ) -> Result<'source, Typed> { + debug_assert!( + op == crate::BinaryOperator::LogicalAnd || op == crate::BinaryOperator::LogicalOr + ); + + if ctx.is_runtime() { + // To simulate short-circuiting behavior, we want to generate IR + // like the following for `&&`. For `||`, the condition is `!_lhs` + // and the else value is `true`. + // + // var _e0: bool; + // if _lhs { + // _e0 = _rhs; + // } else { + // _e0 = false; + // } + + let (condition, else_val) = if op == crate::BinaryOperator::LogicalAnd { + let condition = left; + let else_val = ctx.append_expression( + crate::Expression::Literal(crate::Literal::Bool(false)), + span, + )?; + (condition, else_val) + } else { + let condition = ctx.append_expression( + crate::Expression::Unary { + op: crate::UnaryOperator::LogicalNot, + expr: left, + }, + span, + )?; + let else_val = ctx.append_expression( + crate::Expression::Literal(crate::Literal::Bool(true)), + span, + )?; + (condition, else_val) + }; + + let bool_ty = ctx.ensure_type_exists(crate::TypeInner::Scalar(crate::Scalar::BOOL)); + + let rctx = ctx.runtime_expression_ctx(span)?; + let result_var = rctx.function.local_variables.append( + crate::LocalVariable { + name: None, + ty: bool_ty, + init: None, + }, + span, + ); + let pointer = + ctx.append_expression(crate::Expression::LocalVariable(result_var), span)?; + + let (right, mut accept) = ctx.with_nested_runtime_expression_ctx(span, |ctx| { + let right = self.expression_for_abstract(right, ctx)?; + ctx.grow_types(right)?; + Ok(right) + })?; + + accept.push( + crate::Statement::Store { + pointer, + value: right, + }, + span, + ); + + let mut reject = crate::Block::with_capacity(1); + reject.push( + crate::Statement::Store { + pointer, + value: else_val, + }, + span, + ); + + let rctx = ctx.runtime_expression_ctx(span)?; + rctx.block.push( + crate::Statement::If { + condition, + accept, + reject, + }, + span, + ); + + Ok(Typed::Reference(crate::Expression::LocalVariable( + result_var, + ))) + } else { + let left_expr = ctx.get(left); + // Constant or override context in either function or module scope + let &crate::Expression::Literal(crate::Literal::Bool(left_val)) = left_expr else { + return Err(Box::new(Error::NotBool(span))); + }; + + if op == crate::BinaryOperator::LogicalAnd && !left_val + || op == crate::BinaryOperator::LogicalOr && left_val + { + // Short-circuit behavior: don't evaluate the RHS. Ideally we + // would do _some_ validity checks of the RHS here, but that's + // tricky, because the RHS is allowed to have things that aren't + // legal in const contexts. + + Ok(Typed::Plain(left_expr.clone())) + } else { + let right = self.expression_for_abstract(right, ctx)?; + ctx.grow_types(right)?; + + Ok(Typed::Plain(crate::Expression::Binary { op, left, right })) + } + } + } + fn binary( &mut self, op: ir::BinaryOperator, @@ -2383,57 +2570,74 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { span: Span, ctx: &mut ExpressionContext<'source, '_, '_>, ) -> Result<'source, Typed> { - // Load both operands. - let mut left = self.expression_for_abstract(left, ctx)?; - let mut right = self.expression_for_abstract(right, ctx)?; + if op == ir::BinaryOperator::LogicalAnd || op == ir::BinaryOperator::LogicalOr { + let left = self.expression_for_abstract(left, ctx)?; + ctx.grow_types(left)?; - // Convert `scalar op vector` to `vector op vector` by introducing - // `Splat` expressions. - ctx.binary_op_splat(op, &mut left, &mut right)?; - - // Apply automatic conversions. - match op { - ir::BinaryOperator::ShiftLeft | ir::BinaryOperator::ShiftRight => { - // Shift operators require the right operand to be `u32` or - // `vecN`. We can let the validator sort out vector length - // issues, but the right operand must be, or convert to, a u32 leaf - // scalar. - right = - ctx.try_automatic_conversion_for_leaf_scalar(right, ir::Scalar::U32, span)?; - - // Additionally, we must concretize the left operand if the right operand - // is not a const-expression. - // See https://www.w3.org/TR/WGSL/#overload-resolution-section. - // - // 2. Eliminate any candidate where one of its subexpressions resolves to - // an abstract type after feasible automatic conversions, but another of - // the candidate’s subexpressions is not a const-expression. - // - // We only have to explicitly do so for shifts as their operands may be - // of different types - for other binary ops this is achieved by finding - // the conversion consensus for both operands. - if !ctx.is_const(right) { - left = ctx.concretize(left)?; - } - } - - // All other operators follow the same pattern: reconcile the - // scalar leaf types. If there's no reconciliation possible, - // leave the expressions as they are: validation will report the - // problem. - _ => { - ctx.grow_types(left)?; + if !matches!( + resolve_inner!(ctx, left), + &ir::TypeInner::Scalar(ir::Scalar::BOOL) + ) { + // Pass it through as-is, will fail validation + let right = self.expression_for_abstract(right, ctx)?; ctx.grow_types(right)?; - if let Ok(consensus_scalar) = - ctx.automatic_conversion_consensus([left, right].iter()) - { - ctx.convert_to_leaf_scalar(&mut left, consensus_scalar)?; - ctx.convert_to_leaf_scalar(&mut right, consensus_scalar)?; + Ok(Typed::Plain(crate::Expression::Binary { op, left, right })) + } else { + self.logical(op, left, right, span, ctx) + } + } else { + // Load both operands. + let mut left = self.expression_for_abstract(left, ctx)?; + let mut right = self.expression_for_abstract(right, ctx)?; + + // Convert `scalar op vector` to `vector op vector` by introducing + // `Splat` expressions. + ctx.binary_op_splat(op, &mut left, &mut right)?; + + // Apply automatic conversions. + match op { + ir::BinaryOperator::ShiftLeft | ir::BinaryOperator::ShiftRight => { + // Shift operators require the right operand to be `u32` or + // `vecN`. We can let the validator sort out vector length + // issues, but the right operand must be, or convert to, a u32 leaf + // scalar. + right = + ctx.try_automatic_conversion_for_leaf_scalar(right, ir::Scalar::U32, span)?; + + // Additionally, we must concretize the left operand if the right operand + // is not a const-expression. + // See https://www.w3.org/TR/WGSL/#overload-resolution-section. + // + // 2. Eliminate any candidate where one of its subexpressions resolves to + // an abstract type after feasible automatic conversions, but another of + // the candidate’s subexpressions is not a const-expression. + // + // We only have to explicitly do so for shifts as their operands may be + // of different types - for other binary ops this is achieved by finding + // the conversion consensus for both operands. + if !ctx.is_const(right) { + left = ctx.concretize(left)?; + } + } + + // All other operators follow the same pattern: reconcile the + // scalar leaf types. If there's no reconciliation possible, + // leave the expressions as they are: validation will report the + // problem. + _ => { + ctx.grow_types(left)?; + ctx.grow_types(right)?; + if let Ok(consensus_scalar) = + ctx.automatic_conversion_consensus([left, right].iter()) + { + ctx.convert_to_leaf_scalar(&mut left, consensus_scalar)?; + ctx.convert_to_leaf_scalar(&mut right, consensus_scalar)?; + } } } - } - Ok(Typed::Plain(ir::Expression::Binary { op, left, right })) + Ok(Typed::Plain(ir::Expression::Binary { op, left, right })) + } } /// Generate Naga IR for call expressions and statements, and type diff --git a/naga/tests/in/wgsl/operators.wgsl b/naga/tests/in/wgsl/operators.wgsl index b3e47207b..8ecc63b02 100644 --- a/naga/tests/in/wgsl/operators.wgsl +++ b/naga/tests/in/wgsl/operators.wgsl @@ -40,6 +40,11 @@ fn bool_cast(x: vec3) -> vec3 { return vec3(y); } +fn p() -> bool { return true; } +fn q() -> bool { return false; } +fn r() -> bool { return true; } +fn s() -> bool { return false; } + fn logical() { let t = true; let f = false; @@ -55,6 +60,7 @@ fn logical() { let bitwise_or1 = vec3(t) | vec3(f); let bitwise_and0 = t & f; let bitwise_and1 = vec4(t) & vec4(f); + let short_circuit = (p() || q()) && (r() || s()); } fn arithmetic() { diff --git a/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl b/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl index cc66e08cc..7bf9518dd 100644 --- a/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl +++ b/naga/tests/out/glsl/wgsl-operators.main.Compute.glsl @@ -46,15 +46,68 @@ vec3 bool_cast(vec3 x) { return vec3(y); } +bool p() { + return true; +} + +bool q() { + return false; +} + +bool r() { + return true; +} + +bool s() { + return false; +} + void logical() { + bool local = false; + bool local_1 = false; + bool local_2 = false; + bool local_3 = false; + bool local_4 = false; bool neg0_ = !(true); bvec2 neg1_ = not(bvec2(true)); - bool or = (true || false); - bool and = (true && false); + if (!(true)) { + local = false; + } else { + local = true; + } + bool or = local; + if (true) { + local_1 = false; + } else { + local_1 = false; + } + bool and = local_1; bool bitwise_or0_ = (true || false); bvec3 bitwise_or1_ = bvec3(bvec3(true).x || bvec3(false).x, bvec3(true).y || bvec3(false).y, bvec3(true).z || bvec3(false).z); bool bitwise_and0_ = (true && false); bvec4 bitwise_and1_ = bvec4(bvec4(true).x && bvec4(false).x, bvec4(true).y && bvec4(false).y, bvec4(true).z && bvec4(false).z, bvec4(true).w && bvec4(false).w); + bool _e22 = p(); + if (!(_e22)) { + bool _e26 = q(); + local_2 = _e26; + } else { + local_2 = true; + } + bool _e28 = local_2; + if (_e28) { + bool _e31 = r(); + if (!(_e31)) { + bool _e35 = s(); + local_4 = _e35; + } else { + local_4 = true; + } + bool _e37 = local_4; + local_3 = _e37; + } else { + local_3 = false; + } + bool short_circuit = local_3; return; } diff --git a/naga/tests/out/hlsl/wgsl-operators.hlsl b/naga/tests/out/hlsl/wgsl-operators.hlsl index e584efe82..ff4f3a864 100644 --- a/naga/tests/out/hlsl/wgsl-operators.hlsl +++ b/naga/tests/out/hlsl/wgsl-operators.hlsl @@ -48,16 +48,74 @@ float3 bool_cast(float3 x) return float3(y); } +bool p() +{ + return true; +} + +bool q() +{ + return false; +} + +bool r() +{ + return true; +} + +bool s() +{ + return false; +} + void logical() { + bool local = (bool)0; + bool local_1 = (bool)0; + bool local_2 = (bool)0; + bool local_3 = (bool)0; + bool local_4 = (bool)0; + bool neg0_ = !(true); bool2 neg1_ = !((true).xx); - bool or_ = (true || false); - bool and_ = (true && false); + if (!(true)) { + local = false; + } else { + local = true; + } + bool or_ = local; + if (true) { + local_1 = false; + } else { + local_1 = false; + } + bool and_ = local_1; bool bitwise_or0_ = (true | false); bool3 bitwise_or1_ = ((true).xxx | (false).xxx); bool bitwise_and0_ = (true & false); bool4 bitwise_and1_ = ((true).xxxx & (false).xxxx); + const bool _e22 = p(); + if (!(_e22)) { + const bool _e26 = q(); + local_2 = _e26; + } else { + local_2 = true; + } + bool _e28 = local_2; + if (_e28) { + const bool _e31 = r(); + if (!(_e31)) { + const bool _e35 = s(); + local_4 = _e35; + } else { + local_4 = true; + } + bool _e37 = local_4; + local_3 = _e37; + } else { + local_3 = false; + } + bool short_circuit = local_3; return; } diff --git a/naga/tests/out/msl/wgsl-operators.msl b/naga/tests/out/msl/wgsl-operators.msl index 4f9fb2dfe..b45a3a9e2 100644 --- a/naga/tests/out/msl/wgsl-operators.msl +++ b/naga/tests/out/msl/wgsl-operators.msl @@ -56,16 +56,73 @@ metal::float3 bool_cast( return static_cast(y); } +bool p( +) { + return true; +} + +bool q( +) { + return false; +} + +bool r( +) { + return true; +} + +bool s( +) { + return false; +} + void logical( ) { + bool local = {}; + bool local_1 = {}; + bool local_2 = {}; + bool local_3 = {}; + bool local_4 = {}; bool neg0_ = !(true); metal::bool2 neg1_ = !(metal::bool2(true)); - bool or_ = true || false; - bool and_ = true && false; + if (!(true)) { + local = false; + } else { + local = true; + } + bool or_ = local; + if (true) { + local_1 = false; + } else { + local_1 = false; + } + bool and_ = local_1; bool bitwise_or0_ = true | false; metal::bool3 bitwise_or1_ = metal::bool3(true) | metal::bool3(false); bool bitwise_and0_ = true & false; metal::bool4 bitwise_and1_ = metal::bool4(true) & metal::bool4(false); + bool _e22 = p(); + if (!(_e22)) { + bool _e26 = q(); + local_2 = _e26; + } else { + local_2 = true; + } + bool _e28 = local_2; + if (_e28) { + bool _e31 = r(); + if (!(_e31)) { + bool _e35 = s(); + local_4 = _e35; + } else { + local_4 = true; + } + bool _e37 = local_4; + local_3 = _e37; + } else { + local_3 = false; + } + bool short_circuit = local_3; return; } diff --git a/naga/tests/out/spv/wgsl-operators.spvasm b/naga/tests/out/spv/wgsl-operators.spvasm index 9ffe1ae1e..fe417dce2 100644 --- a/naga/tests/out/spv/wgsl-operators.spvasm +++ b/naga/tests/out/spv/wgsl-operators.spvasm @@ -1,122 +1,129 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 533 +; Bound: 582 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %516 "main" %513 -OpExecutionMode %516 LocalSize 1 1 1 -OpDecorate %513 BuiltIn WorkgroupId +OpEntryPoint GLCompute %565 "main" %562 +OpExecutionMode %565 LocalSize 1 1 1 +OpDecorate %562 BuiltIn WorkgroupId %2 = OpTypeVoid %3 = OpTypeFloat 32 %4 = OpTypeVector %3 4 %5 = OpTypeInt 32 1 %6 = OpTypeVector %5 4 -%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 -%16 = OpConstantComposite %4 %15 %15 %15 %15 -%17 = OpConstant %3 0 -%18 = OpConstantComposite %4 %17 %17 %17 %17 -%19 = OpConstant %3 0.5 -%20 = OpConstantComposite %4 %19 %19 %19 %19 -%21 = OpConstant %5 1 -%22 = OpConstantComposite %6 %21 %21 %21 %21 -%25 = OpTypeFunction %4 -%26 = OpTypeBool -%27 = OpConstantTrue %26 +%7 = OpTypeBool +%8 = OpTypeVector %3 2 +%9 = OpTypeVector %3 3 +%11 = OpTypeInt 32 0 +%10 = OpTypeVector %11 3 +%12 = OpTypeMatrix %9 3 +%13 = OpTypeMatrix %9 4 +%14 = OpTypeMatrix %4 3 +%15 = OpTypeVector %5 3 +%16 = OpConstant %3 1 +%17 = OpConstantComposite %4 %16 %16 %16 %16 +%18 = OpConstant %3 0 +%19 = OpConstantComposite %4 %18 %18 %18 %18 +%20 = OpConstant %3 0.5 +%21 = OpConstantComposite %4 %20 %20 %20 %20 +%22 = OpConstant %5 1 +%23 = OpConstantComposite %6 %22 %22 %22 %22 +%26 = OpTypeFunction %4 +%27 = OpConstantTrue %7 %28 = OpConstant %5 0 %29 = OpConstant %3 0.1 %30 = OpConstantComposite %6 %28 %28 %28 %28 -%34 = OpTypeVector %26 4 +%34 = OpTypeVector %7 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 +%65 = OpConstantComposite %6 %22 %22 %22 %22 %72 = OpTypeFunction %4 %3 %5 %73 = OpConstant %3 2 -%74 = OpConstantComposite %7 %73 %73 +%74 = OpConstantComposite %8 %73 %73 %75 = OpConstant %3 4 -%76 = OpConstantComposite %7 %75 %75 +%76 = OpConstantComposite %8 %75 %75 %77 = OpConstant %3 8 -%78 = OpConstantComposite %7 %77 %77 +%78 = OpConstantComposite %8 %77 %77 %79 = OpConstant %5 2 %80 = OpConstantComposite %6 %79 %79 %79 %79 -%93 = OpTypeFunction %7 -%94 = OpConstantComposite %7 %15 %15 +%93 = OpTypeFunction %8 +%94 = OpConstantComposite %8 %16 %16 %95 = OpConstant %3 3 -%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 -%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 +%96 = OpConstantComposite %8 %95 %95 +%98 = OpTypePointer Function %8 +%110 = OpTypeFunction %9 %9 +%112 = OpTypeVector %7 3 +%113 = OpConstantComposite %9 %18 %18 %18 +%115 = OpConstantComposite %9 %16 %16 %16 +%119 = OpTypeFunction %7 +%123 = OpConstantFalse %7 +%133 = OpTypeFunction %2 +%134 = OpTypeVector %7 2 +%135 = OpConstantComposite %134 %27 %27 +%136 = OpConstantComposite %112 %27 %27 %27 +%137 = OpConstantComposite %112 %123 %123 %123 +%138 = OpConstantComposite %34 %27 %27 %27 %27 +%139 = OpConstantComposite %34 %123 %123 %123 %123 +%141 = OpTypePointer Function %7 +%142 = OpConstantNull %7 +%144 = OpConstantNull %7 +%146 = OpConstantNull %7 +%148 = OpConstantNull %7 +%150 = OpConstantNull %7 +%186 = OpTypeFunction %5 %5 %5 +%198 = OpTypeFunction %11 %11 %11 +%202 = OpConstant %11 0 +%204 = OpConstant %11 1 +%207 = OpTypeVector %5 2 +%209 = OpTypeFunction %207 %207 %207 +%213 = OpConstantComposite %207 %28 %28 +%215 = OpConstantComposite %207 %57 %57 +%216 = OpConstantComposite %207 %58 %58 +%221 = OpConstantComposite %207 %22 %22 +%225 = OpTypeFunction %10 %10 %10 +%229 = OpConstantComposite %10 %202 %202 %202 +%231 = OpConstantComposite %10 %204 %204 %204 +%270 = OpTypeVector %11 2 +%272 = OpTypeFunction %270 %270 %270 +%276 = OpConstantComposite %270 %202 %202 +%278 = OpConstantComposite %270 %204 %204 +%290 = OpConstant %11 2 +%291 = OpConstantComposite %207 %79 %79 +%292 = OpConstantComposite %10 %290 %290 %290 +%293 = OpConstantComposite %4 %73 %73 %73 %73 +%294 = OpConstantComposite %4 %16 %16 %16 %16 +%295 = OpConstantComposite %270 %290 %290 +%296 = OpConstantNull %12 +%297 = OpConstantNull %13 +%298 = OpConstantComposite %9 %73 %73 %73 +%299 = OpConstantNull %14 +%301 = OpTypePointer Function %5 +%302 = OpConstantNull %5 +%304 = OpConstantNull %5 +%470 = OpConstantNull %15 +%472 = OpConstantNull %5 +%474 = OpTypePointer Function %15 +%563 = OpTypePointer Input %10 +%562 = OpVariable %563 Input +%566 = OpConstantComposite %9 %16 %16 %16 +%25 = OpFunction %4 None %26 +%24 = OpLabel OpBranch %31 %31 = OpLabel -%32 = OpSelect %5 %27 %21 %28 +%32 = OpSelect %5 %27 %22 %28 %35 = OpCompositeConstruct %34 %27 %27 %27 %27 -%33 = OpSelect %4 %35 %16 %18 -%36 = OpExtInst %4 %1 FMix %18 %16 %20 +%33 = OpSelect %4 %35 %17 %19 +%36 = OpExtInst %4 %1 FMix %19 %17 %21 %38 = OpCompositeConstruct %4 %29 %29 %29 %29 -%37 = OpExtInst %4 %1 FMix %18 %16 %38 -%39 = OpBitcast %3 %21 -%40 = OpBitcast %4 %22 +%37 = OpExtInst %4 %1 FMix %19 %17 %38 +%39 = OpBitcast %3 %22 +%40 = OpBitcast %4 %23 %41 = OpCompositeConstruct %6 %32 %32 %32 %32 %42 = OpIAdd %6 %41 %30 %43 = OpConvertSToF %4 %42 @@ -147,10 +154,10 @@ OpFunctionEnd %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 +%82 = OpCompositeConstruct %8 %69 %69 +%83 = OpFAdd %8 %74 %82 +%84 = OpFSub %8 %83 %76 +%85 = OpFDiv %8 %84 %78 %86 = OpCompositeConstruct %6 %70 %70 %70 %70 %87 = OpFunctionCall %6 %50 %86 %80 %88 = OpVectorShuffle %4 %85 %85 0 1 0 1 @@ -158,460 +165,544 @@ OpBranch %81 %90 = OpFAdd %4 %88 %89 OpReturnValue %90 OpFunctionEnd -%92 = OpFunction %7 None %93 +%92 = OpFunction %8 None %93 %91 = OpLabel %97 = OpVariable %98 Function %74 OpBranch %99 %99 = OpLabel -%100 = OpLoad %7 %97 -%101 = OpFAdd %7 %100 %94 +%100 = OpLoad %8 %97 +%101 = OpFAdd %8 %100 %94 OpStore %97 %101 -%102 = OpLoad %7 %97 -%103 = OpFSub %7 %102 %96 +%102 = OpLoad %8 %97 +%103 = OpFSub %8 %102 %96 OpStore %97 %103 -%104 = OpLoad %7 %97 -%105 = OpFDiv %7 %104 %76 +%104 = OpLoad %8 %97 +%105 = OpFDiv %8 %104 %76 OpStore %97 %105 -%106 = OpLoad %7 %97 +%106 = OpLoad %8 %97 OpReturnValue %106 OpFunctionEnd -%109 = OpFunction %8 None %110 -%108 = OpFunctionParameter %8 +%109 = OpFunction %9 None %110 +%108 = OpFunctionParameter %9 %107 = OpLabel OpBranch %111 %111 = OpLabel %114 = OpFUnordNotEqual %112 %108 %113 -%116 = OpSelect %8 %114 %115 %113 +%116 = OpSelect %9 %114 %115 %113 OpReturnValue %116 OpFunctionEnd -%118 = OpFunction %2 None %119 +%118 = OpFunction %7 None %119 %117 = OpLabel +OpBranch %120 +%120 = OpLabel +OpReturnValue %27 +OpFunctionEnd +%122 = OpFunction %7 None %119 +%121 = OpLabel +OpBranch %124 +%124 = OpLabel +OpReturnValue %123 +OpFunctionEnd +%126 = OpFunction %7 None %119 +%125 = 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 +OpReturnValue %27 OpFunctionEnd -%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 +%129 = OpFunction %7 None %119 +%128 = OpLabel +OpBranch %130 +%130 = OpLabel +OpReturnValue %123 OpFunctionEnd -%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 -%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 -%175 = OpFunction %9 None %176 -%177 = OpFunctionParameter %9 -%178 = OpFunctionParameter %9 +%132 = OpFunction %2 None %133 +%131 = OpLabel +%149 = OpVariable %141 Function %150 +%143 = OpVariable %141 Function %144 +%147 = OpVariable %141 Function %148 +%140 = OpVariable %141 Function %142 +%145 = OpVariable %141 Function %146 +OpBranch %151 +%151 = OpLabel +%152 = OpLogicalNot %7 %27 +%153 = OpLogicalNot %134 %135 +%154 = OpLogicalNot %7 %27 +OpSelectionMerge %155 None +OpBranchConditional %154 %156 %157 +%156 = OpLabel +OpStore %140 %123 +OpBranch %155 +%157 = OpLabel +OpStore %140 %27 +OpBranch %155 +%155 = OpLabel +%158 = OpLoad %7 %140 +OpSelectionMerge %159 None +OpBranchConditional %27 %160 %161 +%160 = OpLabel +OpStore %143 %123 +OpBranch %159 +%161 = OpLabel +OpStore %143 %123 +OpBranch %159 +%159 = OpLabel +%162 = OpLoad %7 %143 +%163 = OpLogicalOr %7 %27 %123 +%164 = OpLogicalOr %112 %136 %137 +%165 = OpLogicalAnd %7 %27 %123 +%166 = OpLogicalAnd %34 %138 %139 +%167 = OpFunctionCall %7 %118 +%168 = OpLogicalNot %7 %167 +OpSelectionMerge %169 None +OpBranchConditional %168 %170 %171 +%170 = OpLabel +%172 = OpFunctionCall %7 %122 +OpStore %145 %172 +OpBranch %169 +%171 = OpLabel +OpStore %145 %27 +OpBranch %169 +%169 = OpLabel +%173 = OpLoad %7 %145 +OpSelectionMerge %174 None +OpBranchConditional %173 %175 %176 +%175 = OpLabel +%177 = OpFunctionCall %7 %126 +%178 = OpLogicalNot %7 %177 +OpSelectionMerge %179 None +OpBranchConditional %178 %180 %181 +%180 = OpLabel +%182 = OpFunctionCall %7 %129 +OpStore %149 %182 +OpBranch %179 +%181 = OpLabel +OpStore %149 %27 +OpBranch %179 %179 = OpLabel -%181 = OpIEqual %112 %178 %180 -%183 = OpSelect %9 %181 %182 %178 -%184 = OpUDiv %9 %177 %183 -OpReturnValue %184 +%183 = OpLoad %7 %149 +OpStore %147 %183 +OpBranch %174 +%176 = OpLabel +OpStore %147 %123 +OpBranch %174 +%174 = OpLabel +%184 = OpLoad %7 %147 +OpReturn OpFunctionEnd -%185 = OpFunction %5 None %137 -%186 = OpFunctionParameter %5 +%185 = OpFunction %5 None %186 %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 +%188 = OpFunctionParameter %5 +%189 = OpLabel +%190 = OpIEqual %7 %188 %28 +%191 = OpIEqual %7 %187 %57 +%192 = OpIEqual %7 %188 %58 +%193 = OpLogicalAnd %7 %191 %192 +%194 = OpLogicalOr %7 %190 %193 +%195 = OpSelect %5 %194 %22 %188 +%196 = OpSDiv %5 %187 %195 +OpReturnValue %196 OpFunctionEnd -%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 +%197 = OpFunction %11 None %198 +%199 = OpFunctionParameter %11 +%200 = OpFunctionParameter %11 +%201 = OpLabel +%203 = OpIEqual %7 %200 %202 +%205 = OpSelect %11 %203 %204 %200 +%206 = OpUDiv %11 %199 %205 +OpReturnValue %206 OpFunctionEnd -%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 +%208 = OpFunction %207 None %209 +%210 = OpFunctionParameter %207 +%211 = OpFunctionParameter %207 +%212 = OpLabel +%214 = OpIEqual %134 %211 %213 +%217 = OpIEqual %134 %210 %215 +%218 = OpIEqual %134 %211 %216 +%219 = OpLogicalAnd %134 %217 %218 +%220 = OpLogicalOr %134 %214 %219 +%222 = OpSelect %207 %220 %221 %211 +%223 = OpSDiv %207 %210 %222 +OpReturnValue %223 OpFunctionEnd -%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 +%224 = OpFunction %10 None %225 +%226 = OpFunctionParameter %10 +%227 = OpFunctionParameter %10 +%228 = OpLabel +%230 = OpIEqual %112 %227 %229 +%232 = OpSelect %10 %230 %231 %227 +%233 = OpUDiv %10 %226 %232 +OpReturnValue %233 OpFunctionEnd -%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 +%234 = OpFunction %5 None %186 +%235 = OpFunctionParameter %5 +%236 = OpFunctionParameter %5 +%237 = OpLabel +%238 = OpIEqual %7 %236 %28 +%239 = OpIEqual %7 %235 %57 +%240 = OpIEqual %7 %236 %58 +%241 = OpLogicalAnd %7 %239 %240 +%242 = OpLogicalOr %7 %238 %241 +%243 = OpSelect %5 %242 %22 %236 +%244 = OpSRem %5 %235 %243 +OpReturnValue %244 OpFunctionEnd -%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 +%245 = OpFunction %11 None %198 +%246 = OpFunctionParameter %11 +%247 = OpFunctionParameter %11 +%248 = OpLabel +%249 = OpIEqual %7 %247 %202 +%250 = OpSelect %11 %249 %204 %247 +%251 = OpUMod %11 %246 %250 +OpReturnValue %251 OpFunctionEnd -%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 +%252 = OpFunction %207 None %209 +%253 = OpFunctionParameter %207 +%254 = OpFunctionParameter %207 +%255 = OpLabel +%256 = OpIEqual %134 %254 %213 +%257 = OpIEqual %134 %253 %215 +%258 = OpIEqual %134 %254 %216 +%259 = OpLogicalAnd %134 %257 %258 +%260 = OpLogicalOr %134 %256 %259 +%261 = OpSelect %207 %260 %221 %254 +%262 = OpSRem %207 %253 %261 +OpReturnValue %262 +OpFunctionEnd +%263 = OpFunction %10 None %225 +%264 = OpFunctionParameter %10 +%265 = OpFunctionParameter %10 +%266 = OpLabel +%267 = OpIEqual %112 %265 %229 +%268 = OpSelect %10 %267 %231 %265 +%269 = OpUMod %10 %264 %268 +OpReturnValue %269 +OpFunctionEnd +%271 = OpFunction %270 None %272 +%273 = OpFunctionParameter %270 +%274 = OpFunctionParameter %270 +%275 = OpLabel +%277 = OpIEqual %134 %274 %276 +%279 = OpSelect %270 %277 %278 %274 +%280 = OpUDiv %270 %273 %279 +OpReturnValue %280 +OpFunctionEnd +%281 = OpFunction %270 None %272 +%282 = OpFunctionParameter %270 +%283 = OpFunctionParameter %270 +%284 = OpLabel +%285 = OpIEqual %134 %283 %276 +%286 = OpSelect %270 %285 %278 %283 +%287 = OpUMod %270 %282 %286 +OpReturnValue %287 +OpFunctionEnd +%289 = OpFunction %2 None %133 +%288 = OpLabel +%300 = OpVariable %301 Function %302 +%303 = OpVariable %301 Function %304 +OpBranch %305 +%305 = OpLabel +%306 = OpFNegate %3 %16 +%307 = OpSNegate %207 %221 +%308 = OpFNegate %8 %94 +%309 = OpIAdd %5 %79 %22 +%310 = OpIAdd %11 %290 %204 +%311 = OpFAdd %3 %73 %16 +%312 = OpIAdd %207 %291 %221 +%313 = OpIAdd %10 %292 %231 +%314 = OpFAdd %4 %293 %294 +%315 = OpISub %5 %79 %22 +%316 = OpISub %11 %290 %204 +%317 = OpFSub %3 %73 %16 +%318 = OpISub %207 %291 %221 +%319 = OpISub %10 %292 %231 +%320 = OpFSub %4 %293 %294 +%321 = OpIMul %5 %79 %22 +%322 = OpIMul %11 %290 %204 +%323 = OpFMul %3 %73 %16 +%324 = OpIMul %207 %291 %221 +%325 = OpIMul %10 %292 %231 +%326 = OpFMul %4 %293 %294 +%327 = OpFunctionCall %5 %185 %79 %22 +%328 = OpFunctionCall %11 %197 %290 %204 +%329 = OpFDiv %3 %73 %16 +%330 = OpFunctionCall %207 %208 %291 %221 +%331 = OpFunctionCall %10 %224 %292 %231 +%332 = OpFDiv %4 %293 %294 +%333 = OpFunctionCall %5 %234 %79 %22 +%334 = OpFunctionCall %11 %245 %290 %204 +%335 = OpFRem %3 %73 %16 +%336 = OpFunctionCall %207 %252 %291 %221 +%337 = OpFunctionCall %10 %263 %292 %231 +%338 = OpFRem %4 %293 %294 +OpBranch %339 +%339 = OpLabel +%341 = OpIAdd %207 %291 %221 +%342 = OpIAdd %207 %291 %221 +%343 = OpIAdd %270 %295 %278 +%344 = OpIAdd %270 %295 %278 +%345 = OpFAdd %8 %74 %94 +%346 = OpFAdd %8 %74 %94 +%347 = OpISub %207 %291 %221 +%348 = OpISub %207 %291 %221 +%349 = OpISub %270 %295 %278 +%350 = OpISub %270 %295 %278 +%351 = OpFSub %8 %74 %94 +%352 = OpFSub %8 %74 %94 +%354 = OpCompositeConstruct %207 %22 %22 +%353 = OpIMul %207 %291 %354 +%356 = OpCompositeConstruct %207 %79 %79 +%355 = OpIMul %207 %221 %356 +%358 = OpCompositeConstruct %270 %204 %204 +%357 = OpIMul %270 %295 %358 +%360 = OpCompositeConstruct %270 %290 %290 +%359 = OpIMul %270 %278 %360 +%361 = OpVectorTimesScalar %8 %74 %16 +%362 = OpVectorTimesScalar %8 %94 %73 +%363 = OpFunctionCall %207 %208 %291 %221 +%364 = OpFunctionCall %207 %208 %291 %221 +%365 = OpFunctionCall %270 %271 %295 %278 +%366 = OpFunctionCall %270 %271 %295 %278 +%367 = OpFDiv %8 %74 %94 +%368 = OpFDiv %8 %74 %94 +%369 = OpFunctionCall %207 %252 %291 %221 +%370 = OpFunctionCall %207 %252 %291 %221 +%371 = OpFunctionCall %270 %281 %295 %278 +%372 = OpFunctionCall %270 %281 %295 %278 +%373 = OpFRem %8 %74 %94 +%374 = OpFRem %8 %74 %94 +OpBranch %340 +%340 = OpLabel +%376 = OpCompositeExtract %9 %296 0 +%377 = OpCompositeExtract %9 %296 0 +%378 = OpFAdd %9 %376 %377 +%379 = OpCompositeExtract %9 %296 1 +%380 = OpCompositeExtract %9 %296 1 +%381 = OpFAdd %9 %379 %380 +%382 = OpCompositeExtract %9 %296 2 +%383 = OpCompositeExtract %9 %296 2 +%384 = OpFAdd %9 %382 %383 +%375 = OpCompositeConstruct %12 %378 %381 %384 +%386 = OpCompositeExtract %9 %296 0 +%387 = OpCompositeExtract %9 %296 0 +%388 = OpFSub %9 %386 %387 +%389 = OpCompositeExtract %9 %296 1 +%390 = OpCompositeExtract %9 %296 1 +%391 = OpFSub %9 %389 %390 +%392 = OpCompositeExtract %9 %296 2 +%393 = OpCompositeExtract %9 %296 2 +%394 = OpFSub %9 %392 %393 +%385 = OpCompositeConstruct %12 %388 %391 %394 +%395 = OpMatrixTimesScalar %12 %296 %16 +%396 = OpMatrixTimesScalar %12 %296 %73 +%397 = OpMatrixTimesVector %9 %297 %294 +%398 = OpVectorTimesMatrix %4 %298 %297 +%399 = OpMatrixTimesMatrix %12 %297 %299 +%400 = OpLoad %5 %300 +%401 = OpIAdd %5 %400 %57 +OpStore %303 %401 OpReturn OpFunctionEnd -%354 = OpFunction %2 None %119 -%353 = OpLabel -OpBranch %355 -%355 = OpLabel -%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 +%403 = OpFunction %2 None %133 +%402 = OpLabel +OpBranch %404 +%404 = OpLabel +%405 = OpNot %5 %22 +%406 = OpNot %11 %204 +%407 = OpNot %207 %221 +%408 = OpNot %10 %231 +%409 = OpBitwiseOr %5 %79 %22 +%410 = OpBitwiseOr %11 %290 %204 +%411 = OpBitwiseOr %207 %291 %221 +%412 = OpBitwiseOr %10 %292 %231 +%413 = OpBitwiseAnd %5 %79 %22 +%414 = OpBitwiseAnd %11 %290 %204 +%415 = OpBitwiseAnd %207 %291 %221 +%416 = OpBitwiseAnd %10 %292 %231 +%417 = OpBitwiseXor %5 %79 %22 +%418 = OpBitwiseXor %11 %290 %204 +%419 = OpBitwiseXor %207 %291 %221 +%420 = OpBitwiseXor %10 %292 %231 +%421 = OpShiftLeftLogical %5 %79 %204 +%422 = OpShiftLeftLogical %11 %290 %204 +%423 = OpShiftLeftLogical %207 %291 %278 +%424 = OpShiftLeftLogical %10 %292 %231 +%425 = OpShiftRightArithmetic %5 %79 %204 +%426 = OpShiftRightLogical %11 %290 %204 +%427 = OpShiftRightArithmetic %207 %291 %278 +%428 = OpShiftRightLogical %10 %292 %231 OpReturn OpFunctionEnd -%381 = OpFunction %2 None %119 -%380 = OpLabel -OpBranch %382 -%382 = OpLabel -%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 +%430 = OpFunction %2 None %133 +%429 = OpLabel +OpBranch %431 +%431 = OpLabel +%432 = OpIEqual %7 %79 %22 +%433 = OpIEqual %7 %290 %204 +%434 = OpFOrdEqual %7 %73 %16 +%435 = OpIEqual %134 %291 %221 +%436 = OpIEqual %112 %292 %231 +%437 = OpFOrdEqual %34 %293 %294 +%438 = OpINotEqual %7 %79 %22 +%439 = OpINotEqual %7 %290 %204 +%440 = OpFOrdNotEqual %7 %73 %16 +%441 = OpINotEqual %134 %291 %221 +%442 = OpINotEqual %112 %292 %231 +%443 = OpFOrdNotEqual %34 %293 %294 +%444 = OpSLessThan %7 %79 %22 +%445 = OpULessThan %7 %290 %204 +%446 = OpFOrdLessThan %7 %73 %16 +%447 = OpSLessThan %134 %291 %221 +%448 = OpULessThan %112 %292 %231 +%449 = OpFOrdLessThan %34 %293 %294 +%450 = OpSLessThanEqual %7 %79 %22 +%451 = OpULessThanEqual %7 %290 %204 +%452 = OpFOrdLessThanEqual %7 %73 %16 +%453 = OpSLessThanEqual %134 %291 %221 +%454 = OpULessThanEqual %112 %292 %231 +%455 = OpFOrdLessThanEqual %34 %293 %294 +%456 = OpSGreaterThan %7 %79 %22 +%457 = OpUGreaterThan %7 %290 %204 +%458 = OpFOrdGreaterThan %7 %73 %16 +%459 = OpSGreaterThan %134 %291 %221 +%460 = OpUGreaterThan %112 %292 %231 +%461 = OpFOrdGreaterThan %34 %293 %294 +%462 = OpSGreaterThanEqual %7 %79 %22 +%463 = OpUGreaterThanEqual %7 %290 %204 +%464 = OpFOrdGreaterThanEqual %7 %73 %16 +%465 = OpSGreaterThanEqual %134 %291 %221 +%466 = OpUGreaterThanEqual %112 %292 %231 +%467 = OpFOrdGreaterThanEqual %34 %293 %294 OpReturn OpFunctionEnd -%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 +%469 = OpFunction %2 None %133 +%468 = OpLabel +%471 = OpVariable %301 Function %472 +%473 = OpVariable %474 Function %470 +OpBranch %475 +%475 = OpLabel +OpStore %471 %22 +%476 = OpLoad %5 %471 +%477 = OpIAdd %5 %476 %22 +OpStore %471 %477 +%478 = OpLoad %5 %471 +%479 = OpISub %5 %478 %22 +OpStore %471 %479 +%480 = OpLoad %5 %471 +%481 = OpLoad %5 %471 +%482 = OpIMul %5 %481 %480 +OpStore %471 %482 +%483 = OpLoad %5 %471 +%484 = OpLoad %5 %471 +%485 = OpFunctionCall %5 %185 %484 %483 +OpStore %471 %485 +%486 = OpLoad %5 %471 +%487 = OpFunctionCall %5 %234 %486 %22 +OpStore %471 %487 +%488 = OpLoad %5 %471 +%489 = OpBitwiseAnd %5 %488 %28 +OpStore %471 %489 +%490 = OpLoad %5 %471 +%491 = OpBitwiseOr %5 %490 %28 +OpStore %471 %491 +%492 = OpLoad %5 %471 +%493 = OpBitwiseXor %5 %492 %28 +OpStore %471 %493 +%494 = OpLoad %5 %471 +%495 = OpShiftLeftLogical %5 %494 %290 +OpStore %471 %495 +%496 = OpLoad %5 %471 +%497 = OpShiftRightArithmetic %5 %496 %204 +OpStore %471 %497 +%498 = OpLoad %5 %471 +%499 = OpIAdd %5 %498 %22 +OpStore %471 %499 +%500 = OpLoad %5 %471 +%501 = OpISub %5 %500 %22 +OpStore %471 %501 +%502 = OpAccessChain %301 %473 %204 +%503 = OpLoad %5 %502 +%504 = OpIAdd %5 %503 %22 +%505 = OpAccessChain %301 %473 %204 +OpStore %505 %504 +%506 = OpAccessChain %301 %473 %204 +%507 = OpLoad %5 %506 +%508 = OpISub %5 %507 %22 +%509 = OpAccessChain %301 %473 %204 +OpStore %509 %508 OpReturn OpFunctionEnd -%462 = OpFunction %2 None %119 -%461 = OpLabel -OpBranch %463 -%463 = OpLabel -%464 = OpSNegate %5 %21 -%465 = OpSNegate %5 %21 -%466 = OpSNegate %5 %465 -%467 = OpSNegate %5 %21 -%468 = OpSNegate %5 %467 -%469 = OpSNegate %5 %21 -%470 = OpSNegate %5 %469 -%471 = OpSNegate %5 %21 -%472 = OpSNegate %5 %471 -%473 = OpSNegate %5 %472 -%474 = OpSNegate %5 %21 -%475 = OpSNegate %5 %474 -%476 = OpSNegate %5 %475 -%477 = OpSNegate %5 %476 -%478 = OpSNegate %5 %21 -%479 = OpSNegate %5 %478 -%480 = OpSNegate %5 %479 -%481 = OpSNegate %5 %480 -%482 = OpSNegate %5 %481 -%483 = OpSNegate %5 %21 -%484 = OpSNegate %5 %483 -%485 = OpSNegate %5 %484 -%486 = OpSNegate %5 %485 -%487 = OpSNegate %5 %486 -%488 = OpFNegate %3 %15 -%489 = OpFNegate %3 %15 -%490 = OpFNegate %3 %489 -%491 = OpFNegate %3 %15 -%492 = OpFNegate %3 %491 -%493 = OpFNegate %3 %15 -%494 = OpFNegate %3 %493 -%495 = OpFNegate %3 %15 -%496 = OpFNegate %3 %495 -%497 = OpFNegate %3 %496 -%498 = OpFNegate %3 %15 -%499 = OpFNegate %3 %498 -%500 = OpFNegate %3 %499 -%501 = OpFNegate %3 %500 -%502 = OpFNegate %3 %15 -%503 = OpFNegate %3 %502 -%504 = OpFNegate %3 %503 -%505 = OpFNegate %3 %504 -%506 = OpFNegate %3 %505 -%507 = OpFNegate %3 %15 -%508 = OpFNegate %3 %507 -%509 = OpFNegate %3 %508 -%510 = OpFNegate %3 %509 -%511 = OpFNegate %3 %510 -OpReturn -OpFunctionEnd -%516 = OpFunction %2 None %119 +%511 = OpFunction %2 None %133 +%510 = OpLabel +OpBranch %512 %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 +%513 = OpSNegate %5 %22 +%514 = OpSNegate %5 %22 +%515 = OpSNegate %5 %514 +%516 = OpSNegate %5 %22 +%517 = OpSNegate %5 %516 +%518 = OpSNegate %5 %22 +%519 = OpSNegate %5 %518 +%520 = OpSNegate %5 %22 +%521 = OpSNegate %5 %520 +%522 = OpSNegate %5 %521 +%523 = OpSNegate %5 %22 +%524 = OpSNegate %5 %523 +%525 = OpSNegate %5 %524 +%526 = OpSNegate %5 %525 +%527 = OpSNegate %5 %22 +%528 = OpSNegate %5 %527 +%529 = OpSNegate %5 %528 +%530 = OpSNegate %5 %529 +%531 = OpSNegate %5 %530 +%532 = OpSNegate %5 %22 +%533 = OpSNegate %5 %532 +%534 = OpSNegate %5 %533 +%535 = OpSNegate %5 %534 +%536 = OpSNegate %5 %535 +%537 = OpFNegate %3 %16 +%538 = OpFNegate %3 %16 +%539 = OpFNegate %3 %538 +%540 = OpFNegate %3 %16 +%541 = OpFNegate %3 %540 +%542 = OpFNegate %3 %16 +%543 = OpFNegate %3 %542 +%544 = OpFNegate %3 %16 +%545 = OpFNegate %3 %544 +%546 = OpFNegate %3 %545 +%547 = OpFNegate %3 %16 +%548 = OpFNegate %3 %547 +%549 = OpFNegate %3 %548 +%550 = OpFNegate %3 %549 +%551 = OpFNegate %3 %16 +%552 = OpFNegate %3 %551 +%553 = OpFNegate %3 %552 +%554 = OpFNegate %3 %553 +%555 = OpFNegate %3 %554 +%556 = OpFNegate %3 %16 +%557 = OpFNegate %3 %556 +%558 = OpFNegate %3 %557 +%559 = OpFNegate %3 %558 +%560 = OpFNegate %3 %559 +OpReturn +OpFunctionEnd +%565 = OpFunction %2 None %133 +%561 = OpLabel +%564 = OpLoad %10 %562 +OpBranch %567 +%567 = OpLabel +%568 = OpFunctionCall %4 %25 +%569 = OpCompositeExtract %11 %564 0 +%570 = OpConvertUToF %3 %569 +%571 = OpCompositeExtract %11 %564 1 +%572 = OpBitcast %5 %571 +%573 = OpFunctionCall %4 %71 %570 %572 +%574 = OpFunctionCall %8 %92 +%575 = OpFunctionCall %9 %109 %566 +%576 = OpFunctionCall %2 %132 +%577 = OpFunctionCall %2 %289 +%578 = OpFunctionCall %2 %403 +%579 = OpFunctionCall %2 %430 +%580 = OpFunctionCall %2 %469 +%581 = OpFunctionCall %2 %511 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 e0664e4e1..c46524bbd 100644 --- a/naga/tests/out/wgsl/wgsl-operators.wgsl +++ b/naga/tests/out/wgsl/wgsl-operators.wgsl @@ -39,15 +39,69 @@ fn bool_cast(x: vec3) -> vec3 { return vec3(y); } +fn p() -> bool { + return true; +} + +fn q() -> bool { + return false; +} + +fn r() -> bool { + return true; +} + +fn s() -> bool { + return false; +} + fn logical() { + var local: bool; + var local_1: bool; + var local_2: bool; + var local_3: bool; + var local_4: bool; + let neg0_ = !(true); let neg1_ = !(vec2(true)); - let or = (true || false); - let and = (true && false); + if !(true) { + local = false; + } else { + local = true; + } + let or = local; + if true { + local_1 = false; + } else { + local_1 = false; + } + let and = local_1; let bitwise_or0_ = (true | false); let bitwise_or1_ = (vec3(true) | vec3(false)); let bitwise_and0_ = (true & false); let bitwise_and1_ = (vec4(true) & vec4(false)); + let _e22 = p(); + if !(_e22) { + let _e26 = q(); + local_2 = _e26; + } else { + local_2 = true; + } + let _e28 = local_2; + if _e28 { + let _e31 = r(); + if !(_e31) { + let _e35 = s(); + local_4 = _e35; + } else { + local_4 = true; + } + let _e37 = local_4; + local_3 = _e37; + } else { + local_3 = false; + } + let short_circuit = local_3; return; } diff --git a/tests/tests/wgpu-gpu/subgroup_operations/shader.wgsl b/tests/tests/wgpu-gpu/subgroup_operations/shader.wgsl index 454f35ea9..f536ad4bf 100644 --- a/tests/tests/wgpu-gpu/subgroup_operations/shader.wgsl +++ b/tests/tests/wgpu-gpu/subgroup_operations/shader.wgsl @@ -107,8 +107,8 @@ fn main( add_result_to_mask(&passed, 21u, subgroupBroadcast(subgroup_invocation_id, 1u) == 1u); add_result_to_mask(&passed, 22u, subgroupShuffle(subgroup_invocation_id, subgroup_invocation_id) == subgroup_invocation_id); add_result_to_mask(&passed, 23u, subgroupShuffle(subgroup_invocation_id, subgroup_size - 1u - subgroup_invocation_id) == subgroup_size - 1u - subgroup_invocation_id); - add_result_to_mask(&passed, 24u, subgroup_invocation_id == subgroup_size - 1u || subgroupShuffleDown(subgroup_invocation_id, 1u) == subgroup_invocation_id + 1u); - add_result_to_mask(&passed, 25u, subgroup_invocation_id == 0u || subgroupShuffleUp(subgroup_invocation_id, 1u) == subgroup_invocation_id - 1u); + add_result_to_mask(&passed, 24u, subgroupShuffleDown(subgroup_invocation_id, 1u) == subgroup_invocation_id + 1u || subgroup_invocation_id == subgroup_size - 1u); + add_result_to_mask(&passed, 25u, subgroupShuffleUp(subgroup_invocation_id, 1u) == subgroup_invocation_id - 1u || subgroup_invocation_id == 0u); add_result_to_mask(&passed, 26u, subgroupShuffleXor(subgroup_invocation_id, subgroup_size - 1u) == (subgroup_invocation_id ^ (subgroup_size - 1u))); // Mac/Apple will fail this test.