diff --git a/CHANGELOG.md b/CHANGELOG.md index 4c33eef76b..39054efb0b 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -68,7 +68,8 @@ This feature allowed you to call `global_id` on any wgpu opaque handle to get a #### Naga -- Naga's WGSL front and back ends now have experimental support for 64-bit floating-point literals: `1.0lf` denotes an `f64` value. There has been experimental support for an `f64` type for a while, but until now there was no syntax for writing literals with that type. As before, Naga module validation rejects `f64` values unless `naga::valid::Capabilities::FLOAT64` is requested. By @jimblandy in [#4747](https://github.com/gfx-rs/wgpu/pull/4747). +- Naga'sn WGSL front and back ends now have experimental support for 64-bit floating-point literals: `1.0lf` denotes an `f64` value. There has been experimental support for an `f64` type for a while, but until now there was no syntax for writing literals with that type. As before, Naga module validation rejects `f64` values unless `naga::valid::Capabilities::FLOAT64` is requested. By @jimblandy in [#4747](https://github.com/gfx-rs/wgpu/pull/4747). +- Naga constant evaluation can now process binary operators whose operands are both vectors. By @jimblandy in [#4861](https://github.com/gfx-rs/wgpu/pull/4861). ### Changes diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs index f2db433e81..fffa28817b 100644 --- a/naga/src/front/wgsl/error.rs +++ b/naga/src/front/wgsl/error.rs @@ -257,6 +257,12 @@ pub enum Error<'a> { source_span: Span, source_type: String, }, + ConcretizationFailed { + expr_span: Span, + expr_type: String, + scalar: String, + inner: ConstantEvaluatorError, + }, } impl<'a> Error<'a> { @@ -731,7 +737,19 @@ impl<'a> Error<'a> { ) ], notes: vec![], - } + }, + Error::ConcretizationFailed { expr_span, ref expr_type, ref scalar, ref inner } => ParseError { + message: format!("failed to convert expression to a concrete type: {}", inner), + labels: vec![ + ( + expr_span, + format!("this expression has type {}", expr_type).into(), + ) + ], + notes: vec![ + format!("the expression should have been converted to have {} scalar type", scalar), + ] + }, } } } diff --git a/naga/src/front/wgsl/lower/conversion.rs b/naga/src/front/wgsl/lower/conversion.rs index 0819c13a66..01acf8b9d3 100644 --- a/naga/src/front/wgsl/lower/conversion.rs +++ b/naga/src/front/wgsl/lower/conversion.rs @@ -174,11 +174,23 @@ impl<'source, 'temp, 'out> super::ExpressionContext<'source, 'temp, 'out> { if let Some(scalar) = inner.automatically_convertible_scalar(&self.module.types) { let concretized = scalar.concretize(); if concretized != scalar { - let span = self.get_expression_span(expr); + assert!(scalar.is_abstract()); + let expr_span = self.get_expression_span(expr); expr = self .as_const_evaluator() - .cast_array(expr, concretized, span) - .map_err(|err| super::Error::ConstantEvaluatorError(err, span))?; + .cast_array(expr, concretized, expr_span) + .map_err(|err| { + // A `TypeResolution` includes the type's full name, if + // it has one. Also, avoid holding the borrow of `inner` + // across the call to `cast_array`. + let expr_type = &self.typifier()[expr]; + super::Error::ConcretizationFailed { + expr_span, + expr_type: expr_type.to_wgsl(&self.module.to_ctx()), + scalar: concretized.to_wgsl(), + inner: err, + } + })?; } } diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 826cb1da0d..5c71bc167d 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1361,12 +1361,107 @@ impl<'a> ConstantEvaluator<'a> { } Expression::Compose { ty, components } } + ( + &Expression::Compose { + components: ref left_components, + ty: left_ty, + }, + &Expression::Compose { + components: ref right_components, + ty: right_ty, + }, + ) => { + // We have to make a copy of the component lists, because the + // call to `binary_op_vector` needs `&mut self`, but `self` owns + // the component lists. + let left_flattened = crate::proc::flatten_compose( + left_ty, + left_components, + self.expressions, + self.types, + ); + let right_flattened = crate::proc::flatten_compose( + right_ty, + right_components, + self.expressions, + self.types, + ); + + // `flatten_compose` doesn't return an `ExactSizeIterator`, so + // make a reasonable guess of the capacity we'll need. + let mut flattened = Vec::with_capacity(left_components.len()); + flattened.extend(left_flattened.zip(right_flattened)); + + match (&self.types[left_ty].inner, &self.types[right_ty].inner) { + ( + &TypeInner::Vector { + size: left_size, .. + }, + &TypeInner::Vector { + size: right_size, .. + }, + ) if left_size == right_size => { + self.binary_op_vector(op, left_size, &flattened, left_ty, span)? + } + _ => return Err(ConstantEvaluatorError::InvalidBinaryOpArgs), + } + } _ => return Err(ConstantEvaluatorError::InvalidBinaryOpArgs), }; self.register_evaluated_expr(expr, span) } + fn binary_op_vector( + &mut self, + op: BinaryOperator, + size: crate::VectorSize, + components: &[(Handle, Handle)], + left_ty: Handle, + span: Span, + ) -> Result { + let ty = match op { + // Relational operators produce vectors of booleans. + BinaryOperator::Equal + | BinaryOperator::NotEqual + | BinaryOperator::Less + | BinaryOperator::LessEqual + | BinaryOperator::Greater + | BinaryOperator::GreaterEqual => self.types.insert( + Type { + name: None, + inner: TypeInner::Vector { + size, + scalar: crate::Scalar::BOOL, + }, + }, + span, + ), + + // Other operators produce the same type as their left + // operand. + BinaryOperator::Add + | BinaryOperator::Subtract + | BinaryOperator::Multiply + | BinaryOperator::Divide + | BinaryOperator::Modulo + | BinaryOperator::And + | BinaryOperator::ExclusiveOr + | BinaryOperator::InclusiveOr + | BinaryOperator::LogicalAnd + | BinaryOperator::LogicalOr + | BinaryOperator::ShiftLeft + | BinaryOperator::ShiftRight => left_ty, + }; + + let components = components + .iter() + .map(|&(left, right)| self.binary_op(op, left, right, span)) + .collect::, _>>()?; + + Ok(Expression::Compose { ty, components }) + } + /// Deep copy `expr` from `expressions` into `self.expressions`. /// /// Return the root of the new copy. diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index e375bb1af3..deddc7033d 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -115,6 +115,16 @@ impl super::Scalar { width: crate::ABSTRACT_WIDTH, }; + pub const fn is_abstract(self) -> bool { + match self.kind { + crate::ScalarKind::AbstractInt | crate::ScalarKind::AbstractFloat => true, + crate::ScalarKind::Sint + | crate::ScalarKind::Uint + | crate::ScalarKind::Float + | crate::ScalarKind::Bool => false, + } + } + /// Construct a float `Scalar` with the given width. /// /// This is especially common when dealing with diff --git a/naga/src/valid/compose.rs b/naga/src/valid/compose.rs index bbb0d61d2a..c21e98c6f2 100644 --- a/naga/src/valid/compose.rs +++ b/naga/src/valid/compose.rs @@ -32,7 +32,12 @@ pub fn validate_compose( scalar: comp_scalar, } if comp_scalar == scalar => comp_size as u32, ref other => { - log::error!("Vector component[{}] type {:?}", index, other); + log::error!( + "Vector component[{}] type {:?}, building {:?}", + index, + other, + scalar + ); return Err(ComposeError::ComponentType { index: index as u32, }); diff --git a/naga/tests/in/const-exprs.wgsl b/naga/tests/in/const-exprs.wgsl index 9465d7a515..ee9304ce45 100644 --- a/naga/tests/in/const-exprs.wgsl +++ b/naga/tests/in/const-exprs.wgsl @@ -84,3 +84,6 @@ fn map_texture_kind(texture_kind: i32) -> u32 { fn compose_of_splat() { var x = vec4f(vec3f(1.0), 2.0).wzyx; } + +const add_vec = vec2(1.0f) + vec2(3.0f, 4.0f); +const compare_vec = vec2(3.0f) == vec2(3.0f, 4.0f); diff --git a/naga/tests/in/operators.wgsl b/naga/tests/in/operators.wgsl index 40d1ce8ead..293b84ab11 100644 --- a/naga/tests/in/operators.wgsl +++ b/naga/tests/in/operators.wgsl @@ -21,9 +21,9 @@ fn builtins() -> vec4 { return vec4(vec4(s1) + v_i32_zero) + s2 + m1 + m2 + b1 + b2; } -fn splat() -> vec4 { - let a = (1.0 + vec2(2.0) - 3.0) / 4.0; - let b = vec4(5) % 2; +fn splat(m: f32, n: i32) -> vec4 { + let a = (2.0 + vec2(m) - 4.0) / 8.0; + let b = vec4(n) % 2; return a.xyxy + vec4(b); } @@ -280,9 +280,9 @@ fn assignment() { } @compute @workgroup_size(1) -fn main() { +fn main(@builtin(workgroup_id) id: vec3) { builtins(); - splat(); + splat(f32(id.x), i32(id.y)); bool_cast(v_f32_one.xyz); logical(); diff --git a/naga/tests/out/glsl/const-exprs.main.Compute.glsl b/naga/tests/out/glsl/const-exprs.main.Compute.glsl index 85a7d01773..b095345de9 100644 --- a/naga/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/naga/tests/out/glsl/const-exprs.main.Compute.glsl @@ -17,6 +17,8 @@ const vec4 DIV = vec4(0.44444445, 0.0, 0.0, 0.0); const int TEXTURE_KIND_REGULAR = 0; const int TEXTURE_KIND_WARP = 1; const int TEXTURE_KIND_SKY = 2; +const vec2 add_vec = vec2(4.0, 5.0); +const bvec2 compare_vec = bvec2(true, false); void swizzle_of_compose() { diff --git a/naga/tests/out/glsl/operators.main.Compute.glsl b/naga/tests/out/glsl/operators.main.Compute.glsl index 83dd56f8a4..006bce205e 100644 --- a/naga/tests/out/glsl/operators.main.Compute.glsl +++ b/naga/tests/out/glsl/operators.main.Compute.glsl @@ -23,9 +23,9 @@ vec4 builtins() { return (((((vec4((ivec4(s1_) + v_i32_zero)) + s2_) + m1_) + m2_) + vec4(b1_)) + b2_); } -vec4 splat() { - vec2 a_2 = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0)); - ivec4 b = (ivec4(5) % ivec4(2)); +vec4 splat(float m, int n) { + vec2 a_2 = (((vec2(2.0) + vec2(m)) - vec2(4.0)) / vec2(8.0)); + ivec4 b = (ivec4(n) % ivec4(2)); return (a_2.xyxy + vec4(b)); } @@ -247,9 +247,10 @@ void negation_avoids_prefix_decrement() { } void main() { - vec4 _e0 = builtins(); - vec4 _e1 = splat(); - vec3 _e6 = bool_cast(vec3(1.0, 1.0, 1.0)); + uvec3 id = gl_WorkGroupID; + vec4 _e1 = builtins(); + vec4 _e6 = splat(float(id.x), int(id.y)); + vec3 _e11 = bool_cast(vec3(1.0, 1.0, 1.0)); logical(); arithmetic(); bit(); diff --git a/naga/tests/out/hlsl/const-exprs.hlsl b/naga/tests/out/hlsl/const-exprs.hlsl index b9a2d1c905..4cc2491ce8 100644 --- a/naga/tests/out/hlsl/const-exprs.hlsl +++ b/naga/tests/out/hlsl/const-exprs.hlsl @@ -10,6 +10,8 @@ static const float4 DIV = float4(0.44444445, 0.0, 0.0, 0.0); static const int TEXTURE_KIND_REGULAR = 0; static const int TEXTURE_KIND_WARP = 1; static const int TEXTURE_KIND_SKY = 2; +static const float2 add_vec = float2(4.0, 5.0); +static const bool2 compare_vec = bool2(true, false); void swizzle_of_compose() { diff --git a/naga/tests/out/hlsl/operators.hlsl b/naga/tests/out/hlsl/operators.hlsl index 6d18d07ed6..58ec5a170d 100644 --- a/naga/tests/out/hlsl/operators.hlsl +++ b/naga/tests/out/hlsl/operators.hlsl @@ -16,10 +16,10 @@ float4 builtins() return (((((float4(((s1_).xxxx + v_i32_zero)) + s2_) + m1_) + m2_) + (b1_).xxxx) + b2_); } -float4 splat() +float4 splat(float m, int n) { - float2 a_2 = ((((1.0).xx + (2.0).xx) - (3.0).xx) / (4.0).xx); - int4 b = ((5).xxxx % (2).xxxx); + float2 a_2 = ((((2.0).xx + (m).xx) - (4.0).xx) / (8.0).xx); + int4 b = ((n).xxxx % (2).xxxx); return (a_2.xyxy + float4(b)); } @@ -251,11 +251,11 @@ void negation_avoids_prefix_decrement() } [numthreads(1, 1, 1)] -void main() +void main(uint3 id : SV_GroupID) { - const float4 _e0 = builtins(); - const float4 _e1 = splat(); - const float3 _e6 = bool_cast(float3(1.0, 1.0, 1.0)); + const float4 _e1 = builtins(); + const float4 _e6 = splat(float(id.x), int(id.y)); + const float3 _e11 = bool_cast(float3(1.0, 1.0, 1.0)); logical(); arithmetic(); bit(); diff --git a/naga/tests/out/msl/const-exprs.msl b/naga/tests/out/msl/const-exprs.msl index 33e281f37c..7798ae62b3 100644 --- a/naga/tests/out/msl/const-exprs.msl +++ b/naga/tests/out/msl/const-exprs.msl @@ -16,6 +16,8 @@ constant metal::float4 DIV = metal::float4(0.44444445, 0.0, 0.0, 0.0); constant int TEXTURE_KIND_REGULAR = 0; constant int TEXTURE_KIND_WARP = 1; constant int TEXTURE_KIND_SKY = 2; +constant metal::float2 add_vec = metal::float2(4.0, 5.0); +constant metal::bool2 compare_vec = metal::bool2(true, false); void swizzle_of_compose( ) { diff --git a/naga/tests/out/msl/operators.msl b/naga/tests/out/msl/operators.msl index 960a87f01c..85fba28c33 100644 --- a/naga/tests/out/msl/operators.msl +++ b/naga/tests/out/msl/operators.msl @@ -23,9 +23,11 @@ metal::float4 builtins( } metal::float4 splat( + float m, + int n ) { - metal::float2 a_2 = ((metal::float2(1.0) + metal::float2(2.0)) - metal::float2(3.0)) / metal::float2(4.0); - metal::int4 b = metal::int4(5) % metal::int4(2); + metal::float2 a_2 = ((metal::float2(2.0) + metal::float2(m)) - metal::float2(4.0)) / metal::float2(8.0); + metal::int4 b = metal::int4(n) % metal::int4(2); return a_2.xyxy + static_cast(b); } @@ -255,11 +257,14 @@ void negation_avoids_prefix_decrement( int p7_ = -(-(-(-(-(1))))); } +struct main_Input { +}; kernel void main_( + metal::uint3 id [[threadgroup_position_in_grid]] ) { - metal::float4 _e0 = builtins(); - metal::float4 _e1 = splat(); - metal::float3 _e6 = bool_cast(metal::float3(1.0, 1.0, 1.0)); + metal::float4 _e1 = builtins(); + metal::float4 _e6 = splat(static_cast(id.x), static_cast(id.y)); + metal::float3 _e11 = bool_cast(metal::float3(1.0, 1.0, 1.0)); logical(); arithmetic(); bit(); diff --git a/naga/tests/out/spv/const-exprs.spvasm b/naga/tests/out/spv/const-exprs.spvasm index 09a30da111..22fef53749 100644 --- a/naga/tests/out/spv/const-exprs.spvasm +++ b/naga/tests/out/spv/const-exprs.spvasm @@ -1,143 +1,152 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 100 +; Bound: 109 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %91 "main" -OpExecutionMode %91 LocalSize 2 3 1 +OpEntryPoint GLCompute %100 "main" +OpExecutionMode %100 LocalSize 2 3 1 %2 = OpTypeVoid %3 = OpTypeInt 32 0 %4 = OpTypeInt 32 1 %5 = OpTypeVector %4 4 %7 = OpTypeFloat 32 %6 = OpTypeVector %7 4 -%8 = OpConstant %3 2 -%9 = OpConstant %4 3 -%10 = OpConstant %4 4 -%11 = OpConstant %4 8 -%12 = OpConstant %7 3.141 -%13 = OpConstant %7 6.282 -%14 = OpConstant %7 0.44444445 -%15 = OpConstant %7 0.0 -%16 = OpConstantComposite %6 %14 %15 %15 %15 -%17 = OpConstant %4 0 -%18 = OpConstant %4 1 -%19 = OpConstant %4 2 -%22 = OpTypeFunction %2 -%23 = OpConstantComposite %5 %10 %9 %19 %18 -%25 = OpTypePointer Function %5 -%30 = OpTypePointer Function %4 -%34 = OpConstant %4 6 -%39 = OpConstant %4 30 -%40 = OpConstant %4 70 -%43 = OpConstantNull %4 -%45 = OpConstantNull %4 -%48 = OpConstantNull %5 -%59 = OpConstant %4 -4 -%60 = OpConstantComposite %5 %59 %59 %59 %59 -%69 = OpConstant %7 1.0 -%70 = OpConstant %7 2.0 -%71 = OpConstantComposite %6 %70 %69 %69 %69 -%73 = OpTypePointer Function %6 -%78 = OpTypeFunction %3 %4 -%79 = OpConstant %3 10 -%80 = OpConstant %3 20 -%81 = OpConstant %3 30 -%82 = OpConstant %3 0 -%89 = OpConstantNull %3 -%21 = OpFunction %2 None %22 -%20 = OpLabel -%24 = OpVariable %25 Function %23 -OpBranch %26 -%26 = OpLabel +%8 = OpTypeVector %7 2 +%10 = OpTypeBool +%9 = OpTypeVector %10 2 +%11 = OpConstant %3 2 +%12 = OpConstant %4 3 +%13 = OpConstant %4 4 +%14 = OpConstant %4 8 +%15 = OpConstant %7 3.141 +%16 = OpConstant %7 6.282 +%17 = OpConstant %7 0.44444445 +%18 = OpConstant %7 0.0 +%19 = OpConstantComposite %6 %17 %18 %18 %18 +%20 = OpConstant %4 0 +%21 = OpConstant %4 1 +%22 = OpConstant %4 2 +%23 = OpConstant %7 4.0 +%24 = OpConstant %7 5.0 +%25 = OpConstantComposite %8 %23 %24 +%26 = OpConstantTrue %10 +%27 = OpConstantFalse %10 +%28 = OpConstantComposite %9 %26 %27 +%31 = OpTypeFunction %2 +%32 = OpConstantComposite %5 %13 %12 %22 %21 +%34 = OpTypePointer Function %5 +%39 = OpTypePointer Function %4 +%43 = OpConstant %4 6 +%48 = OpConstant %4 30 +%49 = OpConstant %4 70 +%52 = OpConstantNull %4 +%54 = OpConstantNull %4 +%57 = OpConstantNull %5 +%68 = OpConstant %4 -4 +%69 = OpConstantComposite %5 %68 %68 %68 %68 +%78 = OpConstant %7 1.0 +%79 = OpConstant %7 2.0 +%80 = OpConstantComposite %6 %79 %78 %78 %78 +%82 = OpTypePointer Function %6 +%87 = OpTypeFunction %3 %4 +%88 = OpConstant %3 10 +%89 = OpConstant %3 20 +%90 = OpConstant %3 30 +%91 = OpConstant %3 0 +%98 = OpConstantNull %3 +%30 = OpFunction %2 None %31 +%29 = OpLabel +%33 = OpVariable %34 Function %32 +OpBranch %35 +%35 = OpLabel OpReturn OpFunctionEnd -%28 = OpFunction %2 None %22 -%27 = OpLabel -%29 = OpVariable %30 Function %19 -OpBranch %31 -%31 = OpLabel -OpReturn -OpFunctionEnd -%33 = OpFunction %2 None %22 -%32 = OpLabel -%35 = OpVariable %30 Function %34 -OpBranch %36 +%37 = OpFunction %2 None %31 %36 = OpLabel +%38 = OpVariable %39 Function %22 +OpBranch %40 +%40 = OpLabel OpReturn OpFunctionEnd -%38 = OpFunction %2 None %22 -%37 = OpLabel -%47 = OpVariable %25 Function %48 -%42 = OpVariable %30 Function %43 -%46 = OpVariable %30 Function %40 -%41 = OpVariable %30 Function %39 -%44 = OpVariable %30 Function %45 -OpBranch %49 -%49 = OpLabel -%50 = OpLoad %4 %41 -OpStore %42 %50 -%51 = OpLoad %4 %42 -OpStore %44 %51 -%52 = OpLoad %4 %41 -%53 = OpLoad %4 %42 -%54 = OpLoad %4 %44 -%55 = OpLoad %4 %46 -%56 = OpCompositeConstruct %5 %52 %53 %54 %55 -OpStore %47 %56 +%42 = OpFunction %2 None %31 +%41 = OpLabel +%44 = OpVariable %39 Function %43 +OpBranch %45 +%45 = OpLabel OpReturn OpFunctionEnd -%58 = OpFunction %2 None %22 -%57 = OpLabel -%61 = OpVariable %25 Function %60 -OpBranch %62 -%62 = OpLabel +%47 = OpFunction %2 None %31 +%46 = OpLabel +%56 = OpVariable %34 Function %57 +%51 = OpVariable %39 Function %52 +%55 = OpVariable %39 Function %49 +%50 = OpVariable %39 Function %48 +%53 = OpVariable %39 Function %54 +OpBranch %58 +%58 = OpLabel +%59 = OpLoad %4 %50 +OpStore %51 %59 +%60 = OpLoad %4 %51 +OpStore %53 %60 +%61 = OpLoad %4 %50 +%62 = OpLoad %4 %51 +%63 = OpLoad %4 %53 +%64 = OpLoad %4 %55 +%65 = OpCompositeConstruct %5 %61 %62 %63 %64 +OpStore %56 %65 OpReturn OpFunctionEnd -%64 = OpFunction %2 None %22 -%63 = OpLabel -%65 = OpVariable %25 Function %60 -OpBranch %66 +%67 = OpFunction %2 None %31 %66 = OpLabel +%70 = OpVariable %34 Function %69 +OpBranch %71 +%71 = OpLabel OpReturn OpFunctionEnd -%68 = OpFunction %2 None %22 -%67 = OpLabel -%72 = OpVariable %73 Function %71 -OpBranch %74 -%74 = OpLabel +%73 = OpFunction %2 None %31 +%72 = OpLabel +%74 = OpVariable %34 Function %69 +OpBranch %75 +%75 = OpLabel OpReturn OpFunctionEnd -%77 = OpFunction %3 None %78 -%76 = OpFunctionParameter %4 -%75 = OpLabel +%77 = OpFunction %2 None %31 +%76 = OpLabel +%81 = OpVariable %82 Function %80 OpBranch %83 %83 = OpLabel -OpSelectionMerge %84 None -OpSwitch %76 %88 0 %85 1 %86 2 %87 -%85 = OpLabel -OpReturnValue %79 -%86 = OpLabel -OpReturnValue %80 -%87 = OpLabel -OpReturnValue %81 -%88 = OpLabel -OpReturnValue %82 -%84 = OpLabel -OpReturnValue %89 +OpReturn OpFunctionEnd -%91 = OpFunction %2 None %22 -%90 = OpLabel +%86 = OpFunction %3 None %87 +%85 = OpFunctionParameter %4 +%84 = OpLabel OpBranch %92 %92 = OpLabel -%93 = OpFunctionCall %2 %21 -%94 = OpFunctionCall %2 %28 -%95 = OpFunctionCall %2 %33 -%96 = OpFunctionCall %2 %38 -%97 = OpFunctionCall %2 %58 -%98 = OpFunctionCall %2 %64 -%99 = OpFunctionCall %2 %68 +OpSelectionMerge %93 None +OpSwitch %85 %97 0 %94 1 %95 2 %96 +%94 = OpLabel +OpReturnValue %88 +%95 = OpLabel +OpReturnValue %89 +%96 = OpLabel +OpReturnValue %90 +%97 = OpLabel +OpReturnValue %91 +%93 = OpLabel +OpReturnValue %98 +OpFunctionEnd +%100 = OpFunction %2 None %31 +%99 = OpLabel +OpBranch %101 +%101 = OpLabel +%102 = OpFunctionCall %2 %30 +%103 = OpFunctionCall %2 %37 +%104 = OpFunctionCall %2 %42 +%105 = OpFunctionCall %2 %47 +%106 = OpFunctionCall %2 %67 +%107 = OpFunctionCall %2 %73 +%108 = OpFunctionCall %2 %77 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/operators.spvasm b/naga/tests/out/spv/operators.spvasm index 89bc4ccf28..974623bc70 100644 --- a/naga/tests/out/spv/operators.spvasm +++ b/naga/tests/out/spv/operators.spvasm @@ -1,12 +1,13 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 377 +; Bound: 389 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %366 "main" -OpExecutionMode %366 LocalSize 1 1 1 +OpEntryPoint GLCompute %374 "main" %371 +OpExecutionMode %374 LocalSize 1 1 1 +OpDecorate %371 BuiltIn WorkgroupId %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpTypeVector %4 4 @@ -20,418 +21,430 @@ OpExecutionMode %366 LocalSize 1 1 1 %12 = OpTypeMatrix %10 4 %13 = OpTypeMatrix %3 3 %14 = OpTypeVector %6 3 -%15 = OpConstant %4 1.0 -%16 = OpConstantComposite %3 %15 %15 %15 %15 -%17 = OpConstant %4 0.0 +%16 = OpTypeInt 32 0 +%15 = OpTypeVector %16 3 +%17 = OpConstant %4 1.0 %18 = OpConstantComposite %3 %17 %17 %17 %17 -%19 = OpConstant %4 0.5 +%19 = OpConstant %4 0.0 %20 = OpConstantComposite %3 %19 %19 %19 %19 -%21 = OpConstant %6 1 -%22 = OpConstantComposite %5 %21 %21 %21 %21 -%25 = OpTypeFunction %3 -%26 = OpConstantTrue %8 -%27 = OpConstant %6 0 -%28 = OpConstantFalse %8 -%29 = OpConstantComposite %7 %28 %28 %28 %28 -%30 = OpConstant %4 0.1 -%31 = OpConstantComposite %5 %27 %27 %27 %27 -%53 = OpConstant %4 2.0 -%54 = OpConstantComposite %9 %53 %53 -%55 = OpConstantComposite %9 %15 %15 -%56 = OpConstant %4 3.0 -%57 = OpConstantComposite %9 %56 %56 -%58 = OpConstant %4 4.0 +%21 = OpConstant %4 0.5 +%22 = OpConstantComposite %3 %21 %21 %21 %21 +%23 = OpConstant %6 1 +%24 = OpConstantComposite %5 %23 %23 %23 %23 +%27 = OpTypeFunction %3 +%28 = OpConstantTrue %8 +%29 = OpConstant %6 0 +%30 = OpConstantFalse %8 +%31 = OpConstantComposite %7 %30 %30 %30 %30 +%32 = OpConstant %4 0.1 +%33 = OpConstantComposite %5 %29 %29 %29 %29 +%57 = OpTypeFunction %3 %4 %6 +%58 = OpConstant %4 2.0 %59 = OpConstantComposite %9 %58 %58 -%60 = OpConstant %6 5 -%61 = OpConstantComposite %5 %60 %60 %60 %60 -%62 = OpConstant %6 2 -%63 = OpConstantComposite %5 %62 %62 %62 %62 -%74 = OpTypeFunction %9 -%76 = OpTypePointer Function %9 -%88 = OpTypeFunction %10 %10 -%90 = OpTypeVector %8 3 -%91 = OpConstantComposite %10 %17 %17 %17 -%93 = OpConstantComposite %10 %15 %15 %15 -%97 = OpTypeFunction %2 -%98 = OpTypeVector %8 2 -%99 = OpConstantComposite %98 %26 %26 -%100 = OpConstantComposite %90 %26 %26 %26 -%101 = OpConstantComposite %90 %28 %28 %28 -%102 = OpConstantComposite %7 %26 %26 %26 %26 -%103 = OpConstantComposite %7 %28 %28 %28 %28 -%115 = OpTypeInt 32 0 -%116 = OpConstant %115 1 -%117 = OpConstant %115 2 -%118 = OpTypeVector %6 2 -%119 = OpConstantComposite %118 %21 %21 -%120 = OpConstantComposite %118 %62 %62 -%121 = OpTypeVector %115 3 -%122 = OpConstantComposite %121 %117 %117 %117 -%123 = OpConstantComposite %121 %116 %116 %116 -%124 = OpConstantComposite %3 %53 %53 %53 %53 -%125 = OpConstantComposite %3 %15 %15 %15 %15 -%126 = OpTypeVector %115 2 -%127 = OpConstantComposite %126 %117 %117 -%128 = OpConstantComposite %126 %116 %116 -%129 = OpConstantNull %11 -%130 = OpConstantNull %12 -%131 = OpConstantComposite %10 %53 %53 %53 -%132 = OpConstantNull %13 -%296 = OpConstantNull %14 -%298 = OpTypePointer Function %6 -%299 = OpConstantNull %6 -%301 = OpTypePointer Function %14 -%329 = OpTypePointer Function %6 -%367 = OpConstantComposite %10 %15 %15 %15 -%24 = OpFunction %3 None %25 -%23 = OpLabel -OpBranch %32 -%32 = OpLabel -%33 = OpSelect %6 %26 %21 %27 -%35 = OpCompositeConstruct %7 %26 %26 %26 %26 -%34 = OpSelect %3 %35 %16 %18 -%36 = OpSelect %3 %29 %18 %16 -%37 = OpExtInst %3 %1 FMix %18 %16 %20 -%39 = OpCompositeConstruct %3 %30 %30 %30 %30 -%38 = OpExtInst %3 %1 FMix %18 %16 %39 -%40 = OpBitcast %4 %21 -%41 = OpBitcast %3 %22 -%42 = OpCompositeConstruct %5 %33 %33 %33 %33 -%43 = OpIAdd %5 %42 %31 -%44 = OpConvertSToF %3 %43 -%45 = OpFAdd %3 %44 %34 -%46 = OpFAdd %3 %45 %37 -%47 = OpFAdd %3 %46 %38 -%48 = OpCompositeConstruct %3 %40 %40 %40 %40 -%49 = OpFAdd %3 %47 %48 -%50 = OpFAdd %3 %49 %41 -OpReturnValue %50 +%60 = OpConstant %4 4.0 +%61 = OpConstantComposite %9 %60 %60 +%62 = OpConstant %4 8.0 +%63 = OpConstantComposite %9 %62 %62 +%64 = OpConstant %6 2 +%65 = OpConstantComposite %5 %64 %64 %64 %64 +%78 = OpTypeFunction %9 +%79 = OpConstantComposite %9 %17 %17 +%80 = OpConstant %4 3.0 +%81 = OpConstantComposite %9 %80 %80 +%83 = OpTypePointer Function %9 +%95 = OpTypeFunction %10 %10 +%97 = OpTypeVector %8 3 +%98 = OpConstantComposite %10 %19 %19 %19 +%100 = OpConstantComposite %10 %17 %17 %17 +%104 = OpTypeFunction %2 +%105 = OpTypeVector %8 2 +%106 = OpConstantComposite %105 %28 %28 +%107 = OpConstantComposite %97 %28 %28 %28 +%108 = OpConstantComposite %97 %30 %30 %30 +%109 = OpConstantComposite %7 %28 %28 %28 %28 +%110 = OpConstantComposite %7 %30 %30 %30 %30 +%122 = OpConstant %16 1 +%123 = OpConstant %16 2 +%124 = OpTypeVector %6 2 +%125 = OpConstantComposite %124 %23 %23 +%126 = OpConstantComposite %124 %64 %64 +%127 = OpConstantComposite %15 %123 %123 %123 +%128 = OpConstantComposite %15 %122 %122 %122 +%129 = OpConstantComposite %3 %58 %58 %58 %58 +%130 = OpConstantComposite %3 %17 %17 %17 %17 +%131 = OpTypeVector %16 2 +%132 = OpConstantComposite %131 %123 %123 +%133 = OpConstantComposite %131 %122 %122 +%134 = OpConstantNull %11 +%135 = OpConstantNull %12 +%136 = OpConstantComposite %10 %58 %58 %58 +%137 = OpConstantNull %13 +%301 = OpConstantNull %14 +%303 = OpTypePointer Function %6 +%304 = OpConstantNull %6 +%306 = OpTypePointer Function %14 +%334 = OpTypePointer Function %6 +%372 = OpTypePointer Input %15 +%371 = OpVariable %372 Input +%375 = OpConstantComposite %10 %17 %17 %17 +%26 = OpFunction %3 None %27 +%25 = OpLabel +OpBranch %34 +%34 = OpLabel +%35 = OpSelect %6 %28 %23 %29 +%37 = OpCompositeConstruct %7 %28 %28 %28 %28 +%36 = OpSelect %3 %37 %18 %20 +%38 = OpSelect %3 %31 %20 %18 +%39 = OpExtInst %3 %1 FMix %20 %18 %22 +%41 = OpCompositeConstruct %3 %32 %32 %32 %32 +%40 = OpExtInst %3 %1 FMix %20 %18 %41 +%42 = OpBitcast %4 %23 +%43 = OpBitcast %3 %24 +%44 = OpCompositeConstruct %5 %35 %35 %35 %35 +%45 = OpIAdd %5 %44 %33 +%46 = OpConvertSToF %3 %45 +%47 = OpFAdd %3 %46 %36 +%48 = OpFAdd %3 %47 %39 +%49 = OpFAdd %3 %48 %40 +%50 = OpCompositeConstruct %3 %42 %42 %42 %42 +%51 = OpFAdd %3 %49 %50 +%52 = OpFAdd %3 %51 %43 +OpReturnValue %52 OpFunctionEnd -%52 = OpFunction %3 None %25 -%51 = OpLabel -OpBranch %64 -%64 = OpLabel -%65 = OpFAdd %9 %55 %54 -%66 = OpFSub %9 %65 %57 -%67 = OpFDiv %9 %66 %59 -%68 = OpSRem %5 %61 %63 -%69 = OpVectorShuffle %3 %67 %67 0 1 0 1 -%70 = OpConvertSToF %3 %68 -%71 = OpFAdd %3 %69 %70 -OpReturnValue %71 +%56 = OpFunction %3 None %57 +%54 = OpFunctionParameter %4 +%55 = OpFunctionParameter %6 +%53 = OpLabel +OpBranch %66 +%66 = OpLabel +%67 = OpCompositeConstruct %9 %54 %54 +%68 = OpFAdd %9 %59 %67 +%69 = OpFSub %9 %68 %61 +%70 = OpFDiv %9 %69 %63 +%71 = OpCompositeConstruct %5 %55 %55 %55 %55 +%72 = OpSRem %5 %71 %65 +%73 = OpVectorShuffle %3 %70 %70 0 1 0 1 +%74 = OpConvertSToF %3 %72 +%75 = OpFAdd %3 %73 %74 +OpReturnValue %75 OpFunctionEnd -%73 = OpFunction %9 None %74 -%72 = OpLabel -%75 = OpVariable %76 Function %54 -OpBranch %77 -%77 = OpLabel -%78 = OpLoad %9 %75 -%79 = OpFAdd %9 %78 %55 -OpStore %75 %79 -%80 = OpLoad %9 %75 -%81 = OpFSub %9 %80 %57 -OpStore %75 %81 -%82 = OpLoad %9 %75 -%83 = OpFDiv %9 %82 %59 -OpStore %75 %83 -%84 = OpLoad %9 %75 -OpReturnValue %84 +%77 = OpFunction %9 None %78 +%76 = OpLabel +%82 = OpVariable %83 Function %59 +OpBranch %84 +%84 = OpLabel +%85 = OpLoad %9 %82 +%86 = OpFAdd %9 %85 %79 +OpStore %82 %86 +%87 = OpLoad %9 %82 +%88 = OpFSub %9 %87 %81 +OpStore %82 %88 +%89 = OpLoad %9 %82 +%90 = OpFDiv %9 %89 %61 +OpStore %82 %90 +%91 = OpLoad %9 %82 +OpReturnValue %91 OpFunctionEnd -%87 = OpFunction %10 None %88 -%86 = OpFunctionParameter %10 -%85 = OpLabel -OpBranch %89 -%89 = OpLabel -%92 = OpFUnordNotEqual %90 %86 %91 -%94 = OpSelect %10 %92 %93 %91 -OpReturnValue %94 +%94 = OpFunction %10 None %95 +%93 = OpFunctionParameter %10 +%92 = OpLabel +OpBranch %96 +%96 = OpLabel +%99 = OpFUnordNotEqual %97 %93 %98 +%101 = OpSelect %10 %99 %100 %98 +OpReturnValue %101 OpFunctionEnd -%96 = OpFunction %2 None %97 -%95 = OpLabel -OpBranch %104 -%104 = OpLabel -%105 = OpLogicalNot %8 %26 -%106 = OpLogicalNot %98 %99 -%107 = OpLogicalOr %8 %26 %28 -%108 = OpLogicalAnd %8 %26 %28 -%109 = OpLogicalOr %8 %26 %28 -%110 = OpLogicalOr %90 %100 %101 -%111 = OpLogicalAnd %8 %26 %28 -%112 = OpLogicalAnd %7 %102 %103 +%103 = OpFunction %2 None %104 +%102 = OpLabel +OpBranch %111 +%111 = OpLabel +%112 = OpLogicalNot %8 %28 +%113 = OpLogicalNot %105 %106 +%114 = OpLogicalOr %8 %28 %30 +%115 = OpLogicalAnd %8 %28 %30 +%116 = OpLogicalOr %8 %28 %30 +%117 = OpLogicalOr %97 %107 %108 +%118 = OpLogicalAnd %8 %28 %30 +%119 = OpLogicalAnd %7 %109 %110 OpReturn OpFunctionEnd -%114 = OpFunction %2 None %97 -%113 = OpLabel -OpBranch %133 -%133 = OpLabel -%134 = OpFNegate %4 %15 -%135 = OpSNegate %118 %119 -%136 = OpFNegate %9 %55 -%137 = OpIAdd %6 %62 %21 -%138 = OpIAdd %115 %117 %116 -%139 = OpFAdd %4 %53 %15 -%140 = OpIAdd %118 %120 %119 -%141 = OpIAdd %121 %122 %123 -%142 = OpFAdd %3 %124 %125 -%143 = OpISub %6 %62 %21 -%144 = OpISub %115 %117 %116 -%145 = OpFSub %4 %53 %15 -%146 = OpISub %118 %120 %119 -%147 = OpISub %121 %122 %123 -%148 = OpFSub %3 %124 %125 -%149 = OpIMul %6 %62 %21 -%150 = OpIMul %115 %117 %116 -%151 = OpFMul %4 %53 %15 -%152 = OpIMul %118 %120 %119 -%153 = OpIMul %121 %122 %123 -%154 = OpFMul %3 %124 %125 -%155 = OpSDiv %6 %62 %21 -%156 = OpUDiv %115 %117 %116 -%157 = OpFDiv %4 %53 %15 -%158 = OpSDiv %118 %120 %119 -%159 = OpUDiv %121 %122 %123 -%160 = OpFDiv %3 %124 %125 -%161 = OpSRem %6 %62 %21 -%162 = OpUMod %115 %117 %116 -%163 = OpFRem %4 %53 %15 -%164 = OpSRem %118 %120 %119 -%165 = OpUMod %121 %122 %123 -%166 = OpFRem %3 %124 %125 -OpBranch %167 -%167 = OpLabel -%169 = OpIAdd %118 %120 %119 -%170 = OpIAdd %118 %120 %119 -%171 = OpIAdd %126 %127 %128 -%172 = OpIAdd %126 %127 %128 -%173 = OpFAdd %9 %54 %55 -%174 = OpFAdd %9 %54 %55 -%175 = OpISub %118 %120 %119 -%176 = OpISub %118 %120 %119 -%177 = OpISub %126 %127 %128 -%178 = OpISub %126 %127 %128 -%179 = OpFSub %9 %54 %55 -%180 = OpFSub %9 %54 %55 -%182 = OpCompositeConstruct %118 %21 %21 -%181 = OpIMul %118 %120 %182 -%184 = OpCompositeConstruct %118 %62 %62 -%183 = OpIMul %118 %119 %184 -%186 = OpCompositeConstruct %126 %116 %116 -%185 = OpIMul %126 %127 %186 -%188 = OpCompositeConstruct %126 %117 %117 -%187 = OpIMul %126 %128 %188 -%189 = OpVectorTimesScalar %9 %54 %15 -%190 = OpVectorTimesScalar %9 %55 %53 -%191 = OpSDiv %118 %120 %119 -%192 = OpSDiv %118 %120 %119 -%193 = OpUDiv %126 %127 %128 -%194 = OpUDiv %126 %127 %128 -%195 = OpFDiv %9 %54 %55 -%196 = OpFDiv %9 %54 %55 -%197 = OpSRem %118 %120 %119 -%198 = OpSRem %118 %120 %119 -%199 = OpUMod %126 %127 %128 -%200 = OpUMod %126 %127 %128 -%201 = OpFRem %9 %54 %55 -%202 = OpFRem %9 %54 %55 -OpBranch %168 -%168 = OpLabel -%204 = OpCompositeExtract %10 %129 0 -%205 = OpCompositeExtract %10 %129 0 -%206 = OpFAdd %10 %204 %205 -%207 = OpCompositeExtract %10 %129 1 -%208 = OpCompositeExtract %10 %129 1 -%209 = OpFAdd %10 %207 %208 -%210 = OpCompositeExtract %10 %129 2 -%211 = OpCompositeExtract %10 %129 2 -%212 = OpFAdd %10 %210 %211 -%203 = OpCompositeConstruct %11 %206 %209 %212 -%214 = OpCompositeExtract %10 %129 0 -%215 = OpCompositeExtract %10 %129 0 -%216 = OpFSub %10 %214 %215 -%217 = OpCompositeExtract %10 %129 1 -%218 = OpCompositeExtract %10 %129 1 -%219 = OpFSub %10 %217 %218 -%220 = OpCompositeExtract %10 %129 2 -%221 = OpCompositeExtract %10 %129 2 -%222 = OpFSub %10 %220 %221 -%213 = OpCompositeConstruct %11 %216 %219 %222 -%223 = OpMatrixTimesScalar %11 %129 %15 -%224 = OpMatrixTimesScalar %11 %129 %53 -%225 = OpMatrixTimesVector %10 %130 %125 -%226 = OpVectorTimesMatrix %3 %131 %130 -%227 = OpMatrixTimesMatrix %11 %130 %132 +%121 = OpFunction %2 None %104 +%120 = OpLabel +OpBranch %138 +%138 = OpLabel +%139 = OpFNegate %4 %17 +%140 = OpSNegate %124 %125 +%141 = OpFNegate %9 %79 +%142 = OpIAdd %6 %64 %23 +%143 = OpIAdd %16 %123 %122 +%144 = OpFAdd %4 %58 %17 +%145 = OpIAdd %124 %126 %125 +%146 = OpIAdd %15 %127 %128 +%147 = OpFAdd %3 %129 %130 +%148 = OpISub %6 %64 %23 +%149 = OpISub %16 %123 %122 +%150 = OpFSub %4 %58 %17 +%151 = OpISub %124 %126 %125 +%152 = OpISub %15 %127 %128 +%153 = OpFSub %3 %129 %130 +%154 = OpIMul %6 %64 %23 +%155 = OpIMul %16 %123 %122 +%156 = OpFMul %4 %58 %17 +%157 = OpIMul %124 %126 %125 +%158 = OpIMul %15 %127 %128 +%159 = OpFMul %3 %129 %130 +%160 = OpSDiv %6 %64 %23 +%161 = OpUDiv %16 %123 %122 +%162 = OpFDiv %4 %58 %17 +%163 = OpSDiv %124 %126 %125 +%164 = OpUDiv %15 %127 %128 +%165 = OpFDiv %3 %129 %130 +%166 = OpSRem %6 %64 %23 +%167 = OpUMod %16 %123 %122 +%168 = OpFRem %4 %58 %17 +%169 = OpSRem %124 %126 %125 +%170 = OpUMod %15 %127 %128 +%171 = OpFRem %3 %129 %130 +OpBranch %172 +%172 = OpLabel +%174 = OpIAdd %124 %126 %125 +%175 = OpIAdd %124 %126 %125 +%176 = OpIAdd %131 %132 %133 +%177 = OpIAdd %131 %132 %133 +%178 = OpFAdd %9 %59 %79 +%179 = OpFAdd %9 %59 %79 +%180 = OpISub %124 %126 %125 +%181 = OpISub %124 %126 %125 +%182 = OpISub %131 %132 %133 +%183 = OpISub %131 %132 %133 +%184 = OpFSub %9 %59 %79 +%185 = OpFSub %9 %59 %79 +%187 = OpCompositeConstruct %124 %23 %23 +%186 = OpIMul %124 %126 %187 +%189 = OpCompositeConstruct %124 %64 %64 +%188 = OpIMul %124 %125 %189 +%191 = OpCompositeConstruct %131 %122 %122 +%190 = OpIMul %131 %132 %191 +%193 = OpCompositeConstruct %131 %123 %123 +%192 = OpIMul %131 %133 %193 +%194 = OpVectorTimesScalar %9 %59 %17 +%195 = OpVectorTimesScalar %9 %79 %58 +%196 = OpSDiv %124 %126 %125 +%197 = OpSDiv %124 %126 %125 +%198 = OpUDiv %131 %132 %133 +%199 = OpUDiv %131 %132 %133 +%200 = OpFDiv %9 %59 %79 +%201 = OpFDiv %9 %59 %79 +%202 = OpSRem %124 %126 %125 +%203 = OpSRem %124 %126 %125 +%204 = OpUMod %131 %132 %133 +%205 = OpUMod %131 %132 %133 +%206 = OpFRem %9 %59 %79 +%207 = OpFRem %9 %59 %79 +OpBranch %173 +%173 = OpLabel +%209 = OpCompositeExtract %10 %134 0 +%210 = OpCompositeExtract %10 %134 0 +%211 = OpFAdd %10 %209 %210 +%212 = OpCompositeExtract %10 %134 1 +%213 = OpCompositeExtract %10 %134 1 +%214 = OpFAdd %10 %212 %213 +%215 = OpCompositeExtract %10 %134 2 +%216 = OpCompositeExtract %10 %134 2 +%217 = OpFAdd %10 %215 %216 +%208 = OpCompositeConstruct %11 %211 %214 %217 +%219 = OpCompositeExtract %10 %134 0 +%220 = OpCompositeExtract %10 %134 0 +%221 = OpFSub %10 %219 %220 +%222 = OpCompositeExtract %10 %134 1 +%223 = OpCompositeExtract %10 %134 1 +%224 = OpFSub %10 %222 %223 +%225 = OpCompositeExtract %10 %134 2 +%226 = OpCompositeExtract %10 %134 2 +%227 = OpFSub %10 %225 %226 +%218 = OpCompositeConstruct %11 %221 %224 %227 +%228 = OpMatrixTimesScalar %11 %134 %17 +%229 = OpMatrixTimesScalar %11 %134 %58 +%230 = OpMatrixTimesVector %10 %135 %130 +%231 = OpVectorTimesMatrix %3 %136 %135 +%232 = OpMatrixTimesMatrix %11 %135 %137 OpReturn OpFunctionEnd -%229 = OpFunction %2 None %97 -%228 = OpLabel -OpBranch %230 -%230 = OpLabel -%231 = OpNot %6 %21 -%232 = OpNot %115 %116 -%233 = OpNot %118 %119 -%234 = OpNot %121 %123 -%235 = OpBitwiseOr %6 %62 %21 -%236 = OpBitwiseOr %115 %117 %116 -%237 = OpBitwiseOr %118 %120 %119 -%238 = OpBitwiseOr %121 %122 %123 -%239 = OpBitwiseAnd %6 %62 %21 -%240 = OpBitwiseAnd %115 %117 %116 -%241 = OpBitwiseAnd %118 %120 %119 -%242 = OpBitwiseAnd %121 %122 %123 -%243 = OpBitwiseXor %6 %62 %21 -%244 = OpBitwiseXor %115 %117 %116 -%245 = OpBitwiseXor %118 %120 %119 -%246 = OpBitwiseXor %121 %122 %123 -%247 = OpShiftLeftLogical %6 %62 %116 -%248 = OpShiftLeftLogical %115 %117 %116 -%249 = OpShiftLeftLogical %118 %120 %128 -%250 = OpShiftLeftLogical %121 %122 %123 -%251 = OpShiftRightArithmetic %6 %62 %116 -%252 = OpShiftRightLogical %115 %117 %116 -%253 = OpShiftRightArithmetic %118 %120 %128 -%254 = OpShiftRightLogical %121 %122 %123 +%234 = OpFunction %2 None %104 +%233 = OpLabel +OpBranch %235 +%235 = OpLabel +%236 = OpNot %6 %23 +%237 = OpNot %16 %122 +%238 = OpNot %124 %125 +%239 = OpNot %15 %128 +%240 = OpBitwiseOr %6 %64 %23 +%241 = OpBitwiseOr %16 %123 %122 +%242 = OpBitwiseOr %124 %126 %125 +%243 = OpBitwiseOr %15 %127 %128 +%244 = OpBitwiseAnd %6 %64 %23 +%245 = OpBitwiseAnd %16 %123 %122 +%246 = OpBitwiseAnd %124 %126 %125 +%247 = OpBitwiseAnd %15 %127 %128 +%248 = OpBitwiseXor %6 %64 %23 +%249 = OpBitwiseXor %16 %123 %122 +%250 = OpBitwiseXor %124 %126 %125 +%251 = OpBitwiseXor %15 %127 %128 +%252 = OpShiftLeftLogical %6 %64 %122 +%253 = OpShiftLeftLogical %16 %123 %122 +%254 = OpShiftLeftLogical %124 %126 %133 +%255 = OpShiftLeftLogical %15 %127 %128 +%256 = OpShiftRightArithmetic %6 %64 %122 +%257 = OpShiftRightLogical %16 %123 %122 +%258 = OpShiftRightArithmetic %124 %126 %133 +%259 = OpShiftRightLogical %15 %127 %128 OpReturn OpFunctionEnd -%256 = OpFunction %2 None %97 -%255 = OpLabel -OpBranch %257 -%257 = OpLabel -%258 = OpIEqual %8 %62 %21 -%259 = OpIEqual %8 %117 %116 -%260 = OpFOrdEqual %8 %53 %15 -%261 = OpIEqual %98 %120 %119 -%262 = OpIEqual %90 %122 %123 -%263 = OpFOrdEqual %7 %124 %125 -%264 = OpINotEqual %8 %62 %21 -%265 = OpINotEqual %8 %117 %116 -%266 = OpFOrdNotEqual %8 %53 %15 -%267 = OpINotEqual %98 %120 %119 -%268 = OpINotEqual %90 %122 %123 -%269 = OpFOrdNotEqual %7 %124 %125 -%270 = OpSLessThan %8 %62 %21 -%271 = OpULessThan %8 %117 %116 -%272 = OpFOrdLessThan %8 %53 %15 -%273 = OpSLessThan %98 %120 %119 -%274 = OpULessThan %90 %122 %123 -%275 = OpFOrdLessThan %7 %124 %125 -%276 = OpSLessThanEqual %8 %62 %21 -%277 = OpULessThanEqual %8 %117 %116 -%278 = OpFOrdLessThanEqual %8 %53 %15 -%279 = OpSLessThanEqual %98 %120 %119 -%280 = OpULessThanEqual %90 %122 %123 -%281 = OpFOrdLessThanEqual %7 %124 %125 -%282 = OpSGreaterThan %8 %62 %21 -%283 = OpUGreaterThan %8 %117 %116 -%284 = OpFOrdGreaterThan %8 %53 %15 -%285 = OpSGreaterThan %98 %120 %119 -%286 = OpUGreaterThan %90 %122 %123 -%287 = OpFOrdGreaterThan %7 %124 %125 -%288 = OpSGreaterThanEqual %8 %62 %21 -%289 = OpUGreaterThanEqual %8 %117 %116 -%290 = OpFOrdGreaterThanEqual %8 %53 %15 -%291 = OpSGreaterThanEqual %98 %120 %119 -%292 = OpUGreaterThanEqual %90 %122 %123 -%293 = OpFOrdGreaterThanEqual %7 %124 %125 +%261 = OpFunction %2 None %104 +%260 = OpLabel +OpBranch %262 +%262 = OpLabel +%263 = OpIEqual %8 %64 %23 +%264 = OpIEqual %8 %123 %122 +%265 = OpFOrdEqual %8 %58 %17 +%266 = OpIEqual %105 %126 %125 +%267 = OpIEqual %97 %127 %128 +%268 = OpFOrdEqual %7 %129 %130 +%269 = OpINotEqual %8 %64 %23 +%270 = OpINotEqual %8 %123 %122 +%271 = OpFOrdNotEqual %8 %58 %17 +%272 = OpINotEqual %105 %126 %125 +%273 = OpINotEqual %97 %127 %128 +%274 = OpFOrdNotEqual %7 %129 %130 +%275 = OpSLessThan %8 %64 %23 +%276 = OpULessThan %8 %123 %122 +%277 = OpFOrdLessThan %8 %58 %17 +%278 = OpSLessThan %105 %126 %125 +%279 = OpULessThan %97 %127 %128 +%280 = OpFOrdLessThan %7 %129 %130 +%281 = OpSLessThanEqual %8 %64 %23 +%282 = OpULessThanEqual %8 %123 %122 +%283 = OpFOrdLessThanEqual %8 %58 %17 +%284 = OpSLessThanEqual %105 %126 %125 +%285 = OpULessThanEqual %97 %127 %128 +%286 = OpFOrdLessThanEqual %7 %129 %130 +%287 = OpSGreaterThan %8 %64 %23 +%288 = OpUGreaterThan %8 %123 %122 +%289 = OpFOrdGreaterThan %8 %58 %17 +%290 = OpSGreaterThan %105 %126 %125 +%291 = OpUGreaterThan %97 %127 %128 +%292 = OpFOrdGreaterThan %7 %129 %130 +%293 = OpSGreaterThanEqual %8 %64 %23 +%294 = OpUGreaterThanEqual %8 %123 %122 +%295 = OpFOrdGreaterThanEqual %8 %58 %17 +%296 = OpSGreaterThanEqual %105 %126 %125 +%297 = OpUGreaterThanEqual %97 %127 %128 +%298 = OpFOrdGreaterThanEqual %7 %129 %130 OpReturn OpFunctionEnd -%295 = OpFunction %2 None %97 -%294 = OpLabel -%297 = OpVariable %298 Function %299 -%300 = OpVariable %301 Function %296 -OpBranch %302 -%302 = OpLabel -OpStore %297 %21 -%303 = OpLoad %6 %297 -%304 = OpIAdd %6 %303 %21 -OpStore %297 %304 -%305 = OpLoad %6 %297 -%306 = OpISub %6 %305 %21 -OpStore %297 %306 -%307 = OpLoad %6 %297 -%308 = OpLoad %6 %297 -%309 = OpIMul %6 %308 %307 -OpStore %297 %309 -%310 = OpLoad %6 %297 -%311 = OpLoad %6 %297 -%312 = OpSDiv %6 %311 %310 -OpStore %297 %312 -%313 = OpLoad %6 %297 -%314 = OpSRem %6 %313 %21 -OpStore %297 %314 -%315 = OpLoad %6 %297 -%316 = OpBitwiseAnd %6 %315 %27 -OpStore %297 %316 -%317 = OpLoad %6 %297 -%318 = OpBitwiseOr %6 %317 %27 -OpStore %297 %318 -%319 = OpLoad %6 %297 -%320 = OpBitwiseXor %6 %319 %27 -OpStore %297 %320 -%321 = OpLoad %6 %297 -%322 = OpShiftLeftLogical %6 %321 %117 -OpStore %297 %322 -%323 = OpLoad %6 %297 -%324 = OpShiftRightArithmetic %6 %323 %116 -OpStore %297 %324 -%325 = OpLoad %6 %297 -%326 = OpIAdd %6 %325 %21 -OpStore %297 %326 -%327 = OpLoad %6 %297 -%328 = OpISub %6 %327 %21 -OpStore %297 %328 -%330 = OpAccessChain %329 %300 %21 -%331 = OpLoad %6 %330 -%332 = OpIAdd %6 %331 %21 -%333 = OpAccessChain %329 %300 %21 -OpStore %333 %332 -%334 = OpAccessChain %329 %300 %21 -%335 = OpLoad %6 %334 -%336 = OpISub %6 %335 %21 -%337 = OpAccessChain %329 %300 %21 -OpStore %337 %336 +%300 = OpFunction %2 None %104 +%299 = OpLabel +%302 = OpVariable %303 Function %304 +%305 = OpVariable %306 Function %301 +OpBranch %307 +%307 = OpLabel +OpStore %302 %23 +%308 = OpLoad %6 %302 +%309 = OpIAdd %6 %308 %23 +OpStore %302 %309 +%310 = OpLoad %6 %302 +%311 = OpISub %6 %310 %23 +OpStore %302 %311 +%312 = OpLoad %6 %302 +%313 = OpLoad %6 %302 +%314 = OpIMul %6 %313 %312 +OpStore %302 %314 +%315 = OpLoad %6 %302 +%316 = OpLoad %6 %302 +%317 = OpSDiv %6 %316 %315 +OpStore %302 %317 +%318 = OpLoad %6 %302 +%319 = OpSRem %6 %318 %23 +OpStore %302 %319 +%320 = OpLoad %6 %302 +%321 = OpBitwiseAnd %6 %320 %29 +OpStore %302 %321 +%322 = OpLoad %6 %302 +%323 = OpBitwiseOr %6 %322 %29 +OpStore %302 %323 +%324 = OpLoad %6 %302 +%325 = OpBitwiseXor %6 %324 %29 +OpStore %302 %325 +%326 = OpLoad %6 %302 +%327 = OpShiftLeftLogical %6 %326 %123 +OpStore %302 %327 +%328 = OpLoad %6 %302 +%329 = OpShiftRightArithmetic %6 %328 %122 +OpStore %302 %329 +%330 = OpLoad %6 %302 +%331 = OpIAdd %6 %330 %23 +OpStore %302 %331 +%332 = OpLoad %6 %302 +%333 = OpISub %6 %332 %23 +OpStore %302 %333 +%335 = OpAccessChain %334 %305 %23 +%336 = OpLoad %6 %335 +%337 = OpIAdd %6 %336 %23 +%338 = OpAccessChain %334 %305 %23 +OpStore %338 %337 +%339 = OpAccessChain %334 %305 %23 +%340 = OpLoad %6 %339 +%341 = OpISub %6 %340 %23 +%342 = OpAccessChain %334 %305 %23 +OpStore %342 %341 OpReturn OpFunctionEnd -%339 = OpFunction %2 None %97 -%338 = OpLabel -OpBranch %340 -%340 = OpLabel -%341 = OpSNegate %6 %21 -%342 = OpSNegate %6 %21 -%343 = OpSNegate %6 %342 -%344 = OpSNegate %6 %21 -%345 = OpSNegate %6 %344 -%346 = OpSNegate %6 %21 -%347 = OpSNegate %6 %346 -%348 = OpSNegate %6 %21 -%349 = OpSNegate %6 %348 +%344 = OpFunction %2 None %104 +%343 = OpLabel +OpBranch %345 +%345 = OpLabel +%346 = OpSNegate %6 %23 +%347 = OpSNegate %6 %23 +%348 = OpSNegate %6 %347 +%349 = OpSNegate %6 %23 %350 = OpSNegate %6 %349 -%351 = OpSNegate %6 %21 +%351 = OpSNegate %6 %23 %352 = OpSNegate %6 %351 -%353 = OpSNegate %6 %352 +%353 = OpSNegate %6 %23 %354 = OpSNegate %6 %353 -%355 = OpSNegate %6 %21 -%356 = OpSNegate %6 %355 +%355 = OpSNegate %6 %354 +%356 = OpSNegate %6 %23 %357 = OpSNegate %6 %356 %358 = OpSNegate %6 %357 %359 = OpSNegate %6 %358 -%360 = OpSNegate %6 %21 +%360 = OpSNegate %6 %23 %361 = OpSNegate %6 %360 %362 = OpSNegate %6 %361 %363 = OpSNegate %6 %362 %364 = OpSNegate %6 %363 +%365 = OpSNegate %6 %23 +%366 = OpSNegate %6 %365 +%367 = OpSNegate %6 %366 +%368 = OpSNegate %6 %367 +%369 = OpSNegate %6 %368 OpReturn OpFunctionEnd -%366 = OpFunction %2 None %97 -%365 = OpLabel -OpBranch %368 -%368 = OpLabel -%369 = OpFunctionCall %3 %24 -%370 = OpFunctionCall %3 %52 -%371 = OpFunctionCall %10 %87 %367 -%372 = OpFunctionCall %2 %96 -%373 = OpFunctionCall %2 %114 -%374 = OpFunctionCall %2 %229 -%375 = OpFunctionCall %2 %256 -%376 = OpFunctionCall %2 %295 +%374 = OpFunction %2 None %104 +%370 = OpLabel +%373 = OpLoad %15 %371 +OpBranch %376 +%376 = OpLabel +%377 = OpFunctionCall %3 %26 +%378 = OpCompositeExtract %16 %373 0 +%379 = OpConvertUToF %4 %378 +%380 = OpCompositeExtract %16 %373 1 +%381 = OpBitcast %6 %380 +%382 = OpFunctionCall %3 %56 %379 %381 +%383 = OpFunctionCall %10 %94 %375 +%384 = OpFunctionCall %2 %103 +%385 = OpFunctionCall %2 %121 +%386 = OpFunctionCall %2 %234 +%387 = OpFunctionCall %2 %261 +%388 = OpFunctionCall %2 %300 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/const-exprs.wgsl b/naga/tests/out/wgsl/const-exprs.wgsl index 66cf2a745f..896b1e22fc 100644 --- a/naga/tests/out/wgsl/const-exprs.wgsl +++ b/naga/tests/out/wgsl/const-exprs.wgsl @@ -10,6 +10,8 @@ const DIV: vec4 = vec4(0.44444445, 0.0, 0.0, 0.0); const TEXTURE_KIND_REGULAR: i32 = 0; const TEXTURE_KIND_WARP: i32 = 1; const TEXTURE_KIND_SKY: i32 = 2; +const add_vec: vec2 = vec2(4.0, 5.0); +const compare_vec: vec2 = vec2(true, false); fn swizzle_of_compose() { var out: vec4 = vec4(4, 3, 2, 1); diff --git a/naga/tests/out/wgsl/operators.wgsl b/naga/tests/out/wgsl/operators.wgsl index 95621e32e2..a107f2b501 100644 --- a/naga/tests/out/wgsl/operators.wgsl +++ b/naga/tests/out/wgsl/operators.wgsl @@ -15,9 +15,9 @@ fn builtins() -> vec4 { return (((((vec4((vec4(s1_) + v_i32_zero)) + s2_) + m1_) + m2_) + vec4(b1_)) + b2_); } -fn splat() -> vec4 { - let a_2 = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0)); - let b = (vec4(5) % vec4(2)); +fn splat(m: f32, n: i32) -> vec4 { + let a_2 = (((vec2(2.0) + vec2(m)) - vec2(4.0)) / vec2(8.0)); + let b = (vec4(n) % vec4(2)); return (a_2.xyxy + vec4(b)); } @@ -241,10 +241,10 @@ fn negation_avoids_prefix_decrement() { } @compute @workgroup_size(1, 1, 1) -fn main() { - let _e0 = builtins(); - let _e1 = splat(); - let _e6 = bool_cast(vec3(1.0, 1.0, 1.0)); +fn main(@builtin(workgroup_id) id: vec3) { + let _e1 = builtins(); + let _e6 = splat(f32(id.x), i32(id.y)); + let _e11 = bool_cast(vec3(1.0, 1.0, 1.0)); logical(); arithmetic(); bit();