From a546e60d6dd0005812cca7ebdb968642f5a9055a Mon Sep 17 00:00:00 2001 From: Jamie Nicol Date: Wed, 12 Feb 2025 14:46:13 +0000 Subject: [PATCH] [naga wgsl-in] Ensure correct array stride when converting types during const evaluation (#7112) When converting the underlying scalar type of an array during const evaluation, we currently use the resulting base type's size as the array stride for the resulting type. For certain types, this may not match the required alignment and will therefore result in a validation error. For example, `array, N>` should have a stride of 16. But if declared with an abstract initializer, eg `array(vec3(0.0))` we will incorrectly determine the stride to be 12. To solve this, we use the proc::Layouter struct to determine the required array stride during const evaluation. To avoid repeating layouting work, we reuse the Lowerer's layouter, passing it through the various *Contexts through to the ConstantEvaluator. --- naga/src/back/pipeline_constants.rs | 8 +- naga/src/front/glsl/context.rs | 10 +- naga/src/front/wgsl/lower/construction.rs | 8 +- naga/src/front/wgsl/lower/mod.rs | 35 +- naga/src/proc/constant_evaluator.rs | 24 +- naga/tests/in/abstract-types-var.wgsl | 16 + naga/tests/out/msl/abstract-types-var.msl | 12 + naga/tests/out/spv/abstract-types-var.spvasm | 698 ++++++++++--------- naga/tests/out/wgsl/abstract-types-var.wgsl | 12 + 9 files changed, 466 insertions(+), 357 deletions(-) diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index ef498b6554..7261c1d858 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -75,6 +75,7 @@ pub fn process_overrides<'a>( let mut adjusted_constant_initializers = HashSet::with_capacity(module.constants.len()); let mut global_expression_kind_tracker = crate::proc::ExpressionKindTracker::new(); + let mut layouter = crate::proc::Layouter::default(); // An iterator through the original overrides table, consumed in // approximate tandem with the global expressions. @@ -147,6 +148,7 @@ pub fn process_overrides<'a>( let mut evaluator = ConstantEvaluator::for_wgsl_module( &mut module, &mut global_expression_kind_tracker, + &mut layouter, false, ); adjust_expr(&adjusted_global_expressions, &mut expr); @@ -186,13 +188,13 @@ pub fn process_overrides<'a>( let mut functions = mem::take(&mut module.functions); for (_, function) in functions.iter_mut() { - process_function(&mut module, &override_map, function)?; + process_function(&mut module, &override_map, &mut layouter, function)?; } module.functions = functions; let mut entry_points = mem::take(&mut module.entry_points); for ep in entry_points.iter_mut() { - process_function(&mut module, &override_map, &mut ep.function)?; + process_function(&mut module, &override_map, &mut layouter, &mut ep.function)?; process_workgroup_size_override(&mut module, &adjusted_global_expressions, ep)?; } module.entry_points = entry_points; @@ -365,6 +367,7 @@ fn process_override( fn process_function( module: &mut Module, override_map: &HandleVec>, + layouter: &mut crate::proc::Layouter, function: &mut Function, ) -> Result<(), ConstantEvaluatorError> { // A map from original local expression handles to @@ -389,6 +392,7 @@ fn process_function( module, &mut function.expressions, &mut local_expression_kind_tracker, + layouter, &mut emitter, &mut block, false, diff --git a/naga/src/front/glsl/context.rs b/naga/src/front/glsl/context.rs index b4cb1c874e..04e89e8397 100644 --- a/naga/src/front/glsl/context.rs +++ b/naga/src/front/glsl/context.rs @@ -8,9 +8,9 @@ use super::{ Frontend, Result, }; use crate::{ - front::Typifier, proc::Emitter, AddressSpace, Arena, BinaryOperator, Block, Expression, - FastHashMap, FunctionArgument, Handle, Literal, LocalVariable, RelationalFunction, Scalar, - Span, Statement, Type, TypeInner, VectorSize, + front::Typifier, proc::Emitter, proc::Layouter, AddressSpace, Arena, BinaryOperator, Block, + Expression, FastHashMap, FunctionArgument, Handle, Literal, LocalVariable, RelationalFunction, + Scalar, Span, Statement, Type, TypeInner, VectorSize, }; use std::ops::Index; @@ -72,6 +72,7 @@ pub struct Context<'a> { pub const_typifier: Typifier, pub typifier: Typifier, + layouter: Layouter, emitter: Emitter, stmt_ctx: Option, pub body: Block, @@ -103,6 +104,7 @@ impl<'a> Context<'a> { const_typifier: Typifier::new(), typifier: Typifier::new(), + layouter: Layouter::default(), emitter: Emitter::default(), stmt_ctx: Some(StmtContext::new()), body: Block::new(), @@ -260,12 +262,14 @@ impl<'a> Context<'a> { crate::proc::ConstantEvaluator::for_glsl_module( self.module, self.global_expression_kind_tracker, + &mut self.layouter, ) } else { crate::proc::ConstantEvaluator::for_glsl_function( self.module, &mut self.expressions, &mut self.local_expression_kind_tracker, + &mut self.layouter, &mut self.emitter, &mut self.body, ) diff --git a/naga/src/front/wgsl/lower/construction.rs b/naga/src/front/wgsl/lower/construction.rs index e52d4776ab..ee973fb9f1 100644 --- a/naga/src/front/wgsl/lower/construction.rs +++ b/naga/src/front/wgsl/lower/construction.rs @@ -490,8 +490,8 @@ impl<'source> Lowerer<'source, '_> { NonZeroU32::new(u32::try_from(components.len()).unwrap()).unwrap(), ), stride: { - self.layouter.update(ctx.module.to_ctx()).unwrap(); - self.layouter[base].to_stride() + ctx.layouter.update(ctx.module.to_ctx()).unwrap(); + ctx.layouter[base].to_stride() }, }; let ty = ctx.ensure_type_exists(inner); @@ -616,8 +616,8 @@ impl<'source> Lowerer<'source, '_> { let base = self.resolve_ast_type(base, &mut ctx.as_global())?; let size = self.array_size(size, &mut ctx.as_global())?; - self.layouter.update(ctx.module.to_ctx()).unwrap(); - let stride = self.layouter[base].to_stride(); + ctx.layouter.update(ctx.module.to_ctx()).unwrap(); + let stride = ctx.layouter[base].to_stride(); let ty = ctx.ensure_type_exists(crate::TypeInner::Array { base, size, stride }); Constructor::Type(ty) diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index e405d3e56f..a1bb65533f 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -87,6 +87,8 @@ pub struct GlobalContext<'source, 'temp, 'out> { const_typifier: &'temp mut Typifier, + layouter: &'temp mut Layouter, + global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker, } @@ -98,6 +100,7 @@ impl<'source> GlobalContext<'source, '_, '_> { types: self.types, module: self.module, const_typifier: self.const_typifier, + layouter: self.layouter, expr_type: ExpressionContextType::Constant(None), global_expression_kind_tracker: self.global_expression_kind_tracker, } @@ -110,6 +113,7 @@ impl<'source> GlobalContext<'source, '_, '_> { types: self.types, module: self.module, const_typifier: self.const_typifier, + layouter: self.layouter, expr_type: ExpressionContextType::Override, global_expression_kind_tracker: self.global_expression_kind_tracker, } @@ -165,6 +169,7 @@ pub struct StatementContext<'source, 'temp, 'out> { const_typifier: &'temp mut Typifier, typifier: &'temp mut Typifier, + layouter: &'temp mut Layouter, function: &'out mut crate::Function, /// Stores the names of expressions that are assigned in `let` statement /// Also stores the spans of the names, for use in errors. @@ -198,6 +203,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> { types: self.types, ast_expressions: self.ast_expressions, const_typifier: self.const_typifier, + layouter: self.layouter, global_expression_kind_tracker: self.global_expression_kind_tracker, module: self.module, expr_type: ExpressionContextType::Constant(Some(LocalExpressionContext { @@ -224,6 +230,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> { types: self.types, ast_expressions: self.ast_expressions, const_typifier: self.const_typifier, + layouter: self.layouter, global_expression_kind_tracker: self.global_expression_kind_tracker, module: self.module, expr_type: ExpressionContextType::Runtime(LocalExpressionContext { @@ -244,6 +251,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> { types: self.types, module: self.module, const_typifier: self.const_typifier, + layouter: self.layouter, global_expression_kind_tracker: self.global_expression_kind_tracker, } } @@ -364,6 +372,7 @@ pub struct ExpressionContext<'source, 'temp, 'out> { /// /// [`module::global_expressions`]: crate::Module::global_expressions const_typifier: &'temp mut Typifier, + layouter: &'temp mut Layouter, global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker, /// Whether we are lowering a constant expression or a general @@ -379,6 +388,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { types: self.types, ast_expressions: self.ast_expressions, const_typifier: self.const_typifier, + layouter: self.layouter, module: self.module, expr_type: ExpressionContextType::Constant(match self.expr_type { ExpressionContextType::Runtime(ref mut local_expression_context) @@ -406,6 +416,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { types: self.types, module: self.module, const_typifier: self.const_typifier, + layouter: self.layouter, global_expression_kind_tracker: self.global_expression_kind_tracker, } } @@ -416,6 +427,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { self.module, &mut rctx.function.expressions, rctx.local_expression_kind_tracker, + self.layouter, rctx.emitter, rctx.block, false, @@ -425,6 +437,7 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { self.module, &mut rctx.function.expressions, rctx.local_expression_kind_tracker, + self.layouter, rctx.emitter, rctx.block, true, @@ -433,11 +446,13 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { ExpressionContextType::Constant(None) => ConstantEvaluator::for_wgsl_module( self.module, self.global_expression_kind_tracker, + self.layouter, false, ), ExpressionContextType::Override => ConstantEvaluator::for_wgsl_module( self.module, self.global_expression_kind_tracker, + self.layouter, true, ), } @@ -1016,15 +1031,11 @@ impl SubgroupGather { pub struct Lowerer<'source, 'temp> { index: &'temp Index<'source>, - layouter: Layouter, } impl<'source, 'temp> Lowerer<'source, 'temp> { - pub fn new(index: &'temp Index<'source>) -> Self { - Self { - index, - layouter: Layouter::default(), - } + pub const fn new(index: &'temp Index<'source>) -> Self { + Self { index } } pub fn lower( @@ -1043,6 +1054,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { types: &tu.types, module: &mut module, const_typifier: &mut Typifier::new(), + layouter: &mut Layouter::default(), global_expression_kind_tracker: &mut crate::proc::ExpressionKindTracker::new(), }; @@ -1299,6 +1311,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { ast_expressions: ctx.ast_expressions, const_typifier: ctx.const_typifier, typifier: &mut typifier, + layouter: ctx.layouter, function: &mut function, named_expressions: &mut named_expressions, types: ctx.types, @@ -3051,10 +3064,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { for member in s.members.iter() { let ty = self.resolve_ast_type(member.ty, ctx)?; - self.layouter.update(ctx.module.to_ctx()).unwrap(); + ctx.layouter.update(ctx.module.to_ctx()).unwrap(); - let member_min_size = self.layouter[ty].size; - let member_min_alignment = self.layouter[ty].alignment; + let member_min_size = ctx.layouter[ty].size; + let member_min_alignment = ctx.layouter[ty].alignment; let member_size = if let Some(size_expr) = member.size { let (size, span) = self.const_u32(size_expr, &mut ctx.as_const())?; @@ -3258,8 +3271,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let base = self.resolve_ast_type(base, ctx)?; let size = self.array_size(size, ctx)?; - self.layouter.update(ctx.module.to_ctx()).unwrap(); - let stride = self.layouter[base].to_stride(); + ctx.layouter.update(ctx.module.to_ctx()).unwrap(); + let stride = ctx.layouter[base].to_stride(); crate::TypeInner::Array { base, size, stride } } diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index c8911077b7..bf97bedc34 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -312,6 +312,8 @@ pub struct ConstantEvaluator<'a> { /// Tracks the constness of expressions residing in [`Self::expressions`] expression_kind_tracker: &'a mut ExpressionKindTracker, + + layouter: &'a mut crate::proc::Layouter, } #[derive(Debug)] @@ -594,6 +596,7 @@ impl<'a> ConstantEvaluator<'a> { pub fn for_wgsl_module( module: &'a mut crate::Module, global_expression_kind_tracker: &'a mut ExpressionKindTracker, + layouter: &'a mut crate::proc::Layouter, in_override_ctx: bool, ) -> Self { Self::for_module( @@ -604,6 +607,7 @@ impl<'a> ConstantEvaluator<'a> { }), module, global_expression_kind_tracker, + layouter, ) } @@ -614,11 +618,13 @@ impl<'a> ConstantEvaluator<'a> { pub fn for_glsl_module( module: &'a mut crate::Module, global_expression_kind_tracker: &'a mut ExpressionKindTracker, + layouter: &'a mut crate::proc::Layouter, ) -> Self { Self::for_module( Behavior::Glsl(GlslRestrictions::Const), module, global_expression_kind_tracker, + layouter, ) } @@ -626,6 +632,7 @@ impl<'a> ConstantEvaluator<'a> { behavior: Behavior<'a>, module: &'a mut crate::Module, global_expression_kind_tracker: &'a mut ExpressionKindTracker, + layouter: &'a mut crate::proc::Layouter, ) -> Self { Self { behavior, @@ -634,6 +641,7 @@ impl<'a> ConstantEvaluator<'a> { overrides: &module.overrides, expressions: &mut module.global_expressions, expression_kind_tracker: global_expression_kind_tracker, + layouter, } } @@ -645,6 +653,7 @@ impl<'a> ConstantEvaluator<'a> { module: &'a mut crate::Module, expressions: &'a mut Arena, local_expression_kind_tracker: &'a mut ExpressionKindTracker, + layouter: &'a mut crate::proc::Layouter, emitter: &'a mut super::Emitter, block: &'a mut crate::Block, is_const: bool, @@ -665,6 +674,7 @@ impl<'a> ConstantEvaluator<'a> { overrides: &module.overrides, expressions, expression_kind_tracker: local_expression_kind_tracker, + layouter, } } @@ -676,6 +686,7 @@ impl<'a> ConstantEvaluator<'a> { module: &'a mut crate::Module, expressions: &'a mut Arena, local_expression_kind_tracker: &'a mut ExpressionKindTracker, + layouter: &'a mut crate::proc::Layouter, emitter: &'a mut super::Emitter, block: &'a mut crate::Block, ) -> Self { @@ -690,6 +701,7 @@ impl<'a> ConstantEvaluator<'a> { overrides: &module.overrides, expressions, expression_kind_tracker: local_expression_kind_tracker, + layouter, } } @@ -1718,7 +1730,11 @@ impl<'a> ConstantEvaluator<'a> { self.types.insert(Type { name: None, inner }, span) } }; - let new_base_stride = self.types[new_base].inner.size(self.to_ctx()); + let mut layouter = std::mem::take(self.layouter); + layouter.update(self.to_ctx()).unwrap(); + *self.layouter = layouter; + + let new_base_stride = self.layouter[new_base].to_stride(); let new_array_ty = self.types.insert( Type { name: None, @@ -2567,6 +2583,7 @@ mod tests { overrides: &overrides, expressions: &mut global_expressions, expression_kind_tracker, + layouter: &mut crate::proc::Layouter::default(), }; let res1 = solver @@ -2653,6 +2670,7 @@ mod tests { overrides: &overrides, expressions: &mut global_expressions, expression_kind_tracker, + layouter: &mut crate::proc::Layouter::default(), }; let res = solver @@ -2771,6 +2789,7 @@ mod tests { overrides: &overrides, expressions: &mut global_expressions, expression_kind_tracker, + layouter: &mut crate::proc::Layouter::default(), }; let root1 = Expression::AccessIndex { base, index: 1 }; @@ -2864,6 +2883,7 @@ mod tests { overrides: &overrides, expressions: &mut global_expressions, expression_kind_tracker, + layouter: &mut crate::proc::Layouter::default(), }; let solved_compose = solver @@ -2946,6 +2966,7 @@ mod tests { overrides: &overrides, expressions: &mut global_expressions, expression_kind_tracker, + layouter: &mut crate::proc::Layouter::default(), }; let solved_compose = solver @@ -3034,6 +3055,7 @@ mod tests { overrides: &overrides, expressions: &mut global_expressions, expression_kind_tracker, + layouter: &mut crate::proc::Layouter::default(), }; let solved_add = solver diff --git a/naga/tests/in/abstract-types-var.wgsl b/naga/tests/in/abstract-types-var.wgsl index 4f0bfe44bc..e12a14b372 100644 --- a/naga/tests/in/abstract-types-var.wgsl +++ b/naga/tests/in/abstract-types-var.wgsl @@ -43,6 +43,10 @@ var xafpaiaf: array = array(1, 2.0); var xafpafai: array = array(1.0, 2); var xafpafaf: array = array(1.0, 2.0); +var xavipai: array, 1> = array(vec3(1)); +var xavfpai: array, 1> = array(vec3(1)); +var xavfpaf: array, 1> = array(vec3(1.0)); + var ivispai = vec2(1); var ivfspaf = vec2(1.0); var ivis_ai = vec2(1); @@ -58,6 +62,10 @@ var iafpafaf = array(1.0, 2.0); var iafpaiaf = array(1, 2.0); var iafpafai = array(1.0, 2); +var iavipai = array(vec3(1)); +var iavfpai = array(vec3(1)); +var iavfpaf = array(vec3(1.0)); + fn all_constant_arguments() { var xvipaiai: vec2 = vec2(42, 43); var xvupaiai: vec2 = vec2(44, 45); @@ -100,6 +108,10 @@ fn all_constant_arguments() { var xafpafai: array = array(1.0, 2); var xafpafaf: array = array(1.0, 2.0); + var xavipai: array, 1> = array(vec3(1)); + var xavfpai: array, 1> = array(vec3(1)); + var xavfpaf: array, 1> = array(vec3(1.0)); + var iaipaiai = array(1, 2); var iafpaiaf = array(1, 2.0); var iafpafai = array(1.0, 2); @@ -147,6 +159,10 @@ fn all_constant_arguments() { xafpafai = array(1.0, 2); xafpafaf = array(1.0, 2.0); + xavipai = array(vec3(1)); + xavfpai = array(vec3(1)); + xavfpaf = array(vec3(1.0)); + iaipaiai = array(1, 2); iafpaiaf = array(1, 2.0); iafpafai = array(1.0, 2); diff --git a/naga/tests/out/msl/abstract-types-var.msl b/naga/tests/out/msl/abstract-types-var.msl index 1c9d6fcb25..da88a396e8 100644 --- a/naga/tests/out/msl/abstract-types-var.msl +++ b/naga/tests/out/msl/abstract-types-var.msl @@ -10,6 +10,12 @@ struct type_7 { struct type_8 { int inner[2]; }; +struct type_10 { + metal::int3 inner[1]; +}; +struct type_12 { + metal::float3 inner[1]; +}; void all_constant_arguments( ) { @@ -46,6 +52,9 @@ void all_constant_arguments( type_7 xafpaiaf = type_7 {1.0, 2.0}; type_7 xafpafai = type_7 {1.0, 2.0}; type_7 xafpafaf = type_7 {1.0, 2.0}; + type_10 xavipai = type_10 {metal::int3(1)}; + type_12 xavfpai = type_12 {metal::float3(1.0)}; + type_12 xavfpaf = type_12 {metal::float3(1.0)}; type_8 iaipaiai = type_8 {1, 2}; type_7 iafpaiaf = type_7 {1.0, 2.0}; type_7 iafpafai = type_7 {1.0, 2.0}; @@ -83,6 +92,9 @@ void all_constant_arguments( xafpaiaf = type_7 {1.0, 2.0}; xafpafai = type_7 {1.0, 2.0}; xafpafaf = type_7 {1.0, 2.0}; + xavipai = type_10 {metal::int3(1)}; + xavfpai = type_12 {metal::float3(1.0)}; + xavfpaf = type_12 {metal::float3(1.0)}; iaipaiai = type_8 {1, 2}; iafpaiaf = type_7 {1.0, 2.0}; iafpafai = type_7 {1.0, 2.0}; diff --git a/naga/tests/out/spv/abstract-types-var.spvasm b/naga/tests/out/spv/abstract-types-var.spvasm index 47249bf278..3c3f42b74b 100644 --- a/naga/tests/out/spv/abstract-types-var.spvasm +++ b/naga/tests/out/spv/abstract-types-var.spvasm @@ -1,13 +1,15 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 269 +; Bound: 290 OpCapability Shader OpCapability Linkage %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpDecorate %10 ArrayStride 4 OpDecorate %12 ArrayStride 4 +OpDecorate %14 ArrayStride 16 +OpDecorate %17 ArrayStride 16 %2 = OpTypeVoid %3 = OpTypeInt 32 1 %4 = OpTypeVector %3 2 @@ -19,342 +21,366 @@ OpDecorate %12 ArrayStride 4 %11 = OpConstant %5 2 %10 = OpTypeArray %7 %11 %12 = OpTypeArray %3 %11 -%13 = OpConstant %3 42 -%14 = OpConstant %3 43 -%15 = OpConstantComposite %4 %13 %14 -%16 = OpConstant %5 44 -%17 = OpConstant %5 45 -%18 = OpConstantComposite %6 %16 %17 -%19 = OpConstant %7 46.0 -%20 = OpConstant %7 47.0 -%21 = OpConstantComposite %8 %19 %20 -%22 = OpConstant %5 42 -%23 = OpConstant %5 43 -%24 = OpConstantComposite %6 %22 %23 -%25 = OpConstant %7 1.0 -%26 = OpConstant %7 2.0 -%27 = OpConstantComposite %8 %25 %26 -%28 = OpConstant %7 3.0 -%29 = OpConstant %7 4.0 -%30 = OpConstantComposite %8 %28 %29 -%31 = OpConstantComposite %9 %27 %30 -%32 = OpConstant %3 1 -%33 = OpConstantComposite %4 %32 %32 -%34 = OpConstantComposite %8 %25 %25 -%35 = OpConstant %5 1 -%36 = OpConstantComposite %6 %35 %35 -%37 = OpConstantComposite %10 %25 %26 -%38 = OpConstant %3 2 -%39 = OpConstantComposite %12 %32 %38 -%41 = OpTypePointer Private %4 -%40 = OpVariable %41 Private %15 -%43 = OpTypePointer Private %6 -%42 = OpVariable %43 Private %18 -%45 = OpTypePointer Private %8 -%44 = OpVariable %45 Private %21 -%46 = OpVariable %43 Private %24 -%47 = OpVariable %43 Private %24 -%48 = OpVariable %43 Private %24 -%49 = OpVariable %43 Private %24 -%51 = OpTypePointer Private %9 -%50 = OpVariable %51 Private %31 -%52 = OpVariable %51 Private %31 -%53 = OpVariable %51 Private %31 -%54 = OpVariable %51 Private %31 -%55 = OpVariable %51 Private %31 -%56 = OpVariable %41 Private %33 -%57 = OpVariable %45 Private %34 -%58 = OpVariable %41 Private %33 -%59 = OpVariable %43 Private %36 -%60 = OpVariable %45 Private %34 -%61 = OpVariable %45 Private %34 -%63 = OpTypePointer Private %10 -%62 = OpVariable %63 Private %37 -%64 = OpVariable %63 Private %37 -%66 = OpTypePointer Private %12 -%65 = OpVariable %66 Private %39 -%67 = OpVariable %63 Private %37 -%68 = OpVariable %63 Private %37 -%69 = OpVariable %63 Private %37 -%70 = OpVariable %41 Private %33 -%71 = OpVariable %45 Private %34 -%72 = OpVariable %41 Private %33 -%73 = OpVariable %43 Private %36 -%74 = OpVariable %45 Private %34 -%75 = OpVariable %45 Private %34 -%76 = OpVariable %63 Private %37 -%77 = OpVariable %63 Private %37 -%78 = OpVariable %66 Private %39 -%79 = OpVariable %63 Private %37 -%80 = OpVariable %63 Private %37 -%81 = OpVariable %63 Private %37 -%84 = OpTypeFunction %2 -%86 = OpTypePointer Function %4 -%88 = OpTypePointer Function %6 -%90 = OpTypePointer Function %8 -%96 = OpTypePointer Function %9 -%112 = OpTypePointer Function %10 -%117 = OpTypePointer Function %12 -%132 = OpTypePointer Function %5 -%133 = OpConstantNull %5 -%135 = OpTypePointer Function %3 -%136 = OpConstantNull %3 -%138 = OpTypePointer Function %7 -%139 = OpConstantNull %7 -%141 = OpConstantNull %6 -%143 = OpConstantNull %6 -%145 = OpConstantNull %6 -%147 = OpConstantNull %6 -%149 = OpConstantNull %9 -%151 = OpConstantNull %9 -%153 = OpConstantNull %9 -%155 = OpConstantNull %9 -%157 = OpConstantNull %10 -%159 = OpConstantNull %10 -%161 = OpConstantNull %10 -%163 = OpConstantNull %10 -%165 = OpConstantNull %12 -%167 = OpConstantNull %12 -%169 = OpConstantNull %10 -%171 = OpConstantNull %10 -%173 = OpConstantNull %10 -%175 = OpConstantNull %10 -%177 = OpConstantNull %12 -%179 = OpConstantNull %12 -%83 = OpFunction %2 None %84 -%82 = OpLabel -%127 = OpVariable %112 Function %37 -%124 = OpVariable %117 Function %39 -%121 = OpVariable %112 Function %37 -%118 = OpVariable %117 Function %39 -%114 = OpVariable %112 Function %37 -%110 = OpVariable %90 Function %34 -%107 = OpVariable %86 Function %33 -%104 = OpVariable %96 Function %31 -%101 = OpVariable %96 Function %31 -%98 = OpVariable %96 Function %31 -%94 = OpVariable %88 Function %24 -%91 = OpVariable %88 Function %24 -%85 = OpVariable %86 Function %15 -%125 = OpVariable %112 Function %37 -%122 = OpVariable %112 Function %37 -%119 = OpVariable %117 Function %39 -%115 = OpVariable %112 Function %37 -%111 = OpVariable %112 Function %37 -%108 = OpVariable %88 Function %36 -%105 = OpVariable %86 Function %33 -%102 = OpVariable %96 Function %31 -%99 = OpVariable %96 Function %31 -%95 = OpVariable %96 Function %31 -%92 = OpVariable %88 Function %24 -%87 = OpVariable %88 Function %18 -%126 = OpVariable %112 Function %37 -%123 = OpVariable %112 Function %37 -%120 = OpVariable %112 Function %37 -%116 = OpVariable %117 Function %39 -%113 = OpVariable %112 Function %37 -%109 = OpVariable %90 Function %34 -%106 = OpVariable %90 Function %34 -%103 = OpVariable %96 Function %31 -%100 = OpVariable %96 Function %31 -%97 = OpVariable %96 Function %31 -%93 = OpVariable %88 Function %24 -%89 = OpVariable %90 Function %21 -OpBranch %128 -%128 = OpLabel -OpStore %85 %15 -OpStore %87 %18 -OpStore %89 %21 -OpStore %91 %24 -OpStore %92 %24 -OpStore %93 %24 -OpStore %94 %24 -OpStore %95 %31 -OpStore %97 %31 -OpStore %98 %31 -OpStore %99 %31 -OpStore %100 %31 -OpStore %101 %31 -OpStore %102 %31 -OpStore %103 %31 -OpStore %104 %31 -OpStore %105 %33 -OpStore %106 %34 -OpStore %107 %33 -OpStore %108 %36 -OpStore %109 %34 -OpStore %110 %34 -OpStore %111 %37 -OpStore %113 %37 -OpStore %114 %37 -OpStore %115 %37 -OpStore %116 %39 -OpStore %118 %39 -OpStore %119 %39 -OpStore %120 %37 -OpStore %121 %37 -OpStore %122 %37 -OpStore %123 %37 -OpStore %124 %39 -OpStore %125 %37 -OpStore %126 %37 -OpStore %127 %37 +%13 = OpTypeVector %3 3 +%15 = OpConstant %5 1 +%14 = OpTypeArray %13 %15 +%16 = OpTypeVector %7 3 +%17 = OpTypeArray %16 %15 +%18 = OpConstant %3 42 +%19 = OpConstant %3 43 +%20 = OpConstantComposite %4 %18 %19 +%21 = OpConstant %5 44 +%22 = OpConstant %5 45 +%23 = OpConstantComposite %6 %21 %22 +%24 = OpConstant %7 46.0 +%25 = OpConstant %7 47.0 +%26 = OpConstantComposite %8 %24 %25 +%27 = OpConstant %5 42 +%28 = OpConstant %5 43 +%29 = OpConstantComposite %6 %27 %28 +%30 = OpConstant %7 1.0 +%31 = OpConstant %7 2.0 +%32 = OpConstantComposite %8 %30 %31 +%33 = OpConstant %7 3.0 +%34 = OpConstant %7 4.0 +%35 = OpConstantComposite %8 %33 %34 +%36 = OpConstantComposite %9 %32 %35 +%37 = OpConstant %3 1 +%38 = OpConstantComposite %4 %37 %37 +%39 = OpConstantComposite %8 %30 %30 +%40 = OpConstantComposite %6 %15 %15 +%41 = OpConstantComposite %10 %30 %31 +%42 = OpConstant %3 2 +%43 = OpConstantComposite %12 %37 %42 +%44 = OpConstantComposite %13 %37 %37 %37 +%45 = OpConstantComposite %14 %44 +%46 = OpConstantComposite %16 %30 %30 %30 +%47 = OpConstantComposite %17 %46 +%49 = OpTypePointer Private %4 +%48 = OpVariable %49 Private %20 +%51 = OpTypePointer Private %6 +%50 = OpVariable %51 Private %23 +%53 = OpTypePointer Private %8 +%52 = OpVariable %53 Private %26 +%54 = OpVariable %51 Private %29 +%55 = OpVariable %51 Private %29 +%56 = OpVariable %51 Private %29 +%57 = OpVariable %51 Private %29 +%59 = OpTypePointer Private %9 +%58 = OpVariable %59 Private %36 +%60 = OpVariable %59 Private %36 +%61 = OpVariable %59 Private %36 +%62 = OpVariable %59 Private %36 +%63 = OpVariable %59 Private %36 +%64 = OpVariable %49 Private %38 +%65 = OpVariable %53 Private %39 +%66 = OpVariable %49 Private %38 +%67 = OpVariable %51 Private %40 +%68 = OpVariable %53 Private %39 +%69 = OpVariable %53 Private %39 +%71 = OpTypePointer Private %10 +%70 = OpVariable %71 Private %41 +%72 = OpVariable %71 Private %41 +%74 = OpTypePointer Private %12 +%73 = OpVariable %74 Private %43 +%75 = OpVariable %71 Private %41 +%76 = OpVariable %71 Private %41 +%77 = OpVariable %71 Private %41 +%79 = OpTypePointer Private %14 +%78 = OpVariable %79 Private %45 +%81 = OpTypePointer Private %17 +%80 = OpVariable %81 Private %47 +%82 = OpVariable %81 Private %47 +%83 = OpVariable %49 Private %38 +%84 = OpVariable %53 Private %39 +%85 = OpVariable %49 Private %38 +%86 = OpVariable %51 Private %40 +%87 = OpVariable %53 Private %39 +%88 = OpVariable %53 Private %39 +%89 = OpVariable %71 Private %41 +%90 = OpVariable %71 Private %41 +%91 = OpVariable %74 Private %43 +%92 = OpVariable %71 Private %41 +%93 = OpVariable %71 Private %41 +%94 = OpVariable %71 Private %41 +%95 = OpVariable %79 Private %45 +%96 = OpVariable %79 Private %45 +%97 = OpVariable %81 Private %47 +%100 = OpTypeFunction %2 +%102 = OpTypePointer Function %4 +%104 = OpTypePointer Function %6 +%106 = OpTypePointer Function %8 +%112 = OpTypePointer Function %9 +%128 = OpTypePointer Function %10 +%133 = OpTypePointer Function %12 +%141 = OpTypePointer Function %14 +%143 = OpTypePointer Function %17 +%153 = OpTypePointer Function %5 +%154 = OpConstantNull %5 +%156 = OpTypePointer Function %3 +%157 = OpConstantNull %3 +%159 = OpTypePointer Function %7 +%160 = OpConstantNull %7 +%162 = OpConstantNull %6 +%164 = OpConstantNull %6 +%166 = OpConstantNull %6 +%168 = OpConstantNull %6 +%170 = OpConstantNull %9 +%172 = OpConstantNull %9 +%174 = OpConstantNull %9 +%176 = OpConstantNull %9 +%178 = OpConstantNull %10 +%180 = OpConstantNull %10 +%182 = OpConstantNull %10 +%184 = OpConstantNull %10 +%186 = OpConstantNull %12 +%188 = OpConstantNull %12 +%190 = OpConstantNull %10 +%192 = OpConstantNull %10 +%194 = OpConstantNull %10 +%196 = OpConstantNull %10 +%198 = OpConstantNull %12 +%200 = OpConstantNull %12 +%99 = OpFunction %2 None %100 +%98 = OpLabel +%148 = OpVariable %128 Function %41 +%145 = OpVariable %133 Function %43 +%140 = OpVariable %141 Function %45 +%137 = OpVariable %128 Function %41 +%134 = OpVariable %133 Function %43 +%130 = OpVariable %128 Function %41 +%126 = OpVariable %106 Function %39 +%123 = OpVariable %102 Function %38 +%120 = OpVariable %112 Function %36 +%117 = OpVariable %112 Function %36 +%114 = OpVariable %112 Function %36 +%110 = OpVariable %104 Function %29 +%107 = OpVariable %104 Function %29 +%101 = OpVariable %102 Function %20 +%146 = OpVariable %128 Function %41 +%142 = OpVariable %143 Function %47 +%138 = OpVariable %128 Function %41 +%135 = OpVariable %133 Function %43 +%131 = OpVariable %128 Function %41 +%127 = OpVariable %128 Function %41 +%124 = OpVariable %104 Function %40 +%121 = OpVariable %102 Function %38 +%118 = OpVariable %112 Function %36 +%115 = OpVariable %112 Function %36 +%111 = OpVariable %112 Function %36 +%108 = OpVariable %104 Function %29 +%103 = OpVariable %104 Function %23 +%147 = OpVariable %128 Function %41 +%144 = OpVariable %143 Function %47 +%139 = OpVariable %128 Function %41 +%136 = OpVariable %128 Function %41 +%132 = OpVariable %133 Function %43 +%129 = OpVariable %128 Function %41 +%125 = OpVariable %106 Function %39 +%122 = OpVariable %106 Function %39 +%119 = OpVariable %112 Function %36 +%116 = OpVariable %112 Function %36 +%113 = OpVariable %112 Function %36 +%109 = OpVariable %104 Function %29 +%105 = OpVariable %106 Function %26 +OpBranch %149 +%149 = OpLabel +OpStore %101 %20 +OpStore %103 %23 +OpStore %105 %26 +OpStore %107 %29 +OpStore %108 %29 +OpStore %109 %29 +OpStore %110 %29 +OpStore %111 %36 +OpStore %113 %36 +OpStore %114 %36 +OpStore %115 %36 +OpStore %116 %36 +OpStore %117 %36 +OpStore %118 %36 +OpStore %119 %36 +OpStore %120 %36 +OpStore %121 %38 +OpStore %122 %39 +OpStore %123 %38 +OpStore %124 %40 +OpStore %125 %39 +OpStore %126 %39 +OpStore %127 %41 +OpStore %129 %41 +OpStore %130 %41 +OpStore %131 %41 +OpStore %132 %43 +OpStore %134 %43 +OpStore %135 %43 +OpStore %136 %41 +OpStore %137 %41 +OpStore %138 %41 +OpStore %139 %41 +OpStore %140 %45 +OpStore %142 %47 +OpStore %144 %47 +OpStore %145 %43 +OpStore %146 %41 +OpStore %147 %41 +OpStore %148 %41 OpReturn OpFunctionEnd -%130 = OpFunction %2 None %84 -%129 = OpLabel -%178 = OpVariable %117 Function %179 -%172 = OpVariable %112 Function %173 -%166 = OpVariable %117 Function %167 -%160 = OpVariable %112 Function %161 -%154 = OpVariable %96 Function %155 -%148 = OpVariable %96 Function %149 -%142 = OpVariable %88 Function %143 -%134 = OpVariable %135 Function %136 -%176 = OpVariable %117 Function %177 -%170 = OpVariable %112 Function %171 -%164 = OpVariable %117 Function %165 -%158 = OpVariable %112 Function %159 -%152 = OpVariable %96 Function %153 -%146 = OpVariable %88 Function %147 -%140 = OpVariable %88 Function %141 -%131 = OpVariable %132 Function %133 -%174 = OpVariable %112 Function %175 -%168 = OpVariable %112 Function %169 -%162 = OpVariable %112 Function %163 -%156 = OpVariable %112 Function %157 -%150 = OpVariable %96 Function %151 -%144 = OpVariable %88 Function %145 -%137 = OpVariable %138 Function %139 -OpBranch %180 -%180 = OpLabel -%181 = OpLoad %5 %131 -%182 = OpCompositeConstruct %6 %181 %23 -OpStore %140 %182 -%183 = OpLoad %5 %131 -%184 = OpCompositeConstruct %6 %22 %183 -OpStore %142 %184 -%185 = OpLoad %5 %131 -%186 = OpCompositeConstruct %6 %185 %23 -OpStore %144 %186 -%187 = OpLoad %5 %131 -%188 = OpCompositeConstruct %6 %22 %187 -OpStore %146 %188 -%189 = OpLoad %7 %137 -%190 = OpCompositeConstruct %8 %189 %26 -%191 = OpCompositeConstruct %9 %190 %30 -OpStore %148 %191 -%192 = OpLoad %7 %137 -%193 = OpCompositeConstruct %8 %25 %192 -%194 = OpCompositeConstruct %9 %193 %30 -OpStore %150 %194 -%195 = OpLoad %7 %137 -%196 = OpCompositeConstruct %8 %195 %29 -%197 = OpCompositeConstruct %9 %27 %196 -OpStore %152 %197 -%198 = OpLoad %7 %137 -%199 = OpCompositeConstruct %8 %28 %198 -%200 = OpCompositeConstruct %9 %27 %199 -OpStore %154 %200 -%201 = OpLoad %7 %137 -%202 = OpCompositeConstruct %10 %201 %26 -OpStore %156 %202 -%203 = OpLoad %7 %137 -%204 = OpCompositeConstruct %10 %25 %203 -OpStore %158 %204 -%205 = OpLoad %7 %137 -%206 = OpCompositeConstruct %10 %205 %26 -OpStore %160 %206 -%207 = OpLoad %7 %137 -%208 = OpCompositeConstruct %10 %25 %207 -OpStore %162 %208 -%209 = OpLoad %3 %134 -%210 = OpCompositeConstruct %12 %209 %38 -OpStore %164 %210 -%211 = OpLoad %3 %134 -%212 = OpCompositeConstruct %12 %32 %211 -OpStore %166 %212 -%213 = OpLoad %7 %137 -%214 = OpCompositeConstruct %10 %213 %26 -OpStore %168 %214 -%215 = OpLoad %7 %137 -%216 = OpCompositeConstruct %10 %25 %215 -OpStore %170 %216 -%217 = OpLoad %7 %137 -%218 = OpCompositeConstruct %10 %217 %26 -OpStore %172 %218 -%219 = OpLoad %7 %137 -%220 = OpCompositeConstruct %10 %25 %219 -OpStore %174 %220 -%221 = OpLoad %3 %134 -%222 = OpCompositeConstruct %12 %221 %38 -OpStore %176 %222 -%223 = OpLoad %3 %134 -%224 = OpCompositeConstruct %12 %32 %223 -OpStore %178 %224 -%225 = OpLoad %5 %131 -%226 = OpCompositeConstruct %6 %225 %23 -OpStore %140 %226 -%227 = OpLoad %5 %131 -%228 = OpCompositeConstruct %6 %22 %227 -OpStore %142 %228 -%229 = OpLoad %5 %131 -%230 = OpCompositeConstruct %6 %229 %23 -OpStore %144 %230 -%231 = OpLoad %5 %131 -%232 = OpCompositeConstruct %6 %22 %231 -OpStore %146 %232 -%233 = OpLoad %7 %137 -%234 = OpCompositeConstruct %8 %233 %26 -%235 = OpCompositeConstruct %9 %234 %30 -OpStore %148 %235 -%236 = OpLoad %7 %137 -%237 = OpCompositeConstruct %8 %25 %236 -%238 = OpCompositeConstruct %9 %237 %30 -OpStore %150 %238 -%239 = OpLoad %7 %137 -%240 = OpCompositeConstruct %8 %239 %29 -%241 = OpCompositeConstruct %9 %27 %240 -OpStore %152 %241 -%242 = OpLoad %7 %137 -%243 = OpCompositeConstruct %8 %28 %242 -%244 = OpCompositeConstruct %9 %27 %243 -OpStore %154 %244 -%245 = OpLoad %7 %137 -%246 = OpCompositeConstruct %10 %245 %26 -OpStore %156 %246 -%247 = OpLoad %7 %137 -%248 = OpCompositeConstruct %10 %25 %247 -OpStore %158 %248 -%249 = OpLoad %7 %137 -%250 = OpCompositeConstruct %10 %249 %26 -OpStore %160 %250 -%251 = OpLoad %7 %137 -%252 = OpCompositeConstruct %10 %25 %251 -OpStore %162 %252 -%253 = OpLoad %3 %134 -%254 = OpCompositeConstruct %12 %253 %38 -OpStore %164 %254 -%255 = OpLoad %3 %134 -%256 = OpCompositeConstruct %12 %32 %255 -OpStore %166 %256 -%257 = OpLoad %7 %137 -%258 = OpCompositeConstruct %10 %257 %26 -OpStore %168 %258 -%259 = OpLoad %7 %137 -%260 = OpCompositeConstruct %10 %25 %259 -OpStore %170 %260 -%261 = OpLoad %7 %137 -%262 = OpCompositeConstruct %10 %261 %26 -OpStore %172 %262 -%263 = OpLoad %7 %137 -%264 = OpCompositeConstruct %10 %25 %263 -OpStore %174 %264 -%265 = OpLoad %3 %134 -%266 = OpCompositeConstruct %12 %265 %38 -OpStore %176 %266 -%267 = OpLoad %3 %134 -%268 = OpCompositeConstruct %12 %32 %267 -OpStore %178 %268 +%151 = OpFunction %2 None %100 +%150 = OpLabel +%199 = OpVariable %133 Function %200 +%193 = OpVariable %128 Function %194 +%187 = OpVariable %133 Function %188 +%181 = OpVariable %128 Function %182 +%175 = OpVariable %112 Function %176 +%169 = OpVariable %112 Function %170 +%163 = OpVariable %104 Function %164 +%155 = OpVariable %156 Function %157 +%197 = OpVariable %133 Function %198 +%191 = OpVariable %128 Function %192 +%185 = OpVariable %133 Function %186 +%179 = OpVariable %128 Function %180 +%173 = OpVariable %112 Function %174 +%167 = OpVariable %104 Function %168 +%161 = OpVariable %104 Function %162 +%152 = OpVariable %153 Function %154 +%195 = OpVariable %128 Function %196 +%189 = OpVariable %128 Function %190 +%183 = OpVariable %128 Function %184 +%177 = OpVariable %128 Function %178 +%171 = OpVariable %112 Function %172 +%165 = OpVariable %104 Function %166 +%158 = OpVariable %159 Function %160 +OpBranch %201 +%201 = OpLabel +%202 = OpLoad %5 %152 +%203 = OpCompositeConstruct %6 %202 %28 +OpStore %161 %203 +%204 = OpLoad %5 %152 +%205 = OpCompositeConstruct %6 %27 %204 +OpStore %163 %205 +%206 = OpLoad %5 %152 +%207 = OpCompositeConstruct %6 %206 %28 +OpStore %165 %207 +%208 = OpLoad %5 %152 +%209 = OpCompositeConstruct %6 %27 %208 +OpStore %167 %209 +%210 = OpLoad %7 %158 +%211 = OpCompositeConstruct %8 %210 %31 +%212 = OpCompositeConstruct %9 %211 %35 +OpStore %169 %212 +%213 = OpLoad %7 %158 +%214 = OpCompositeConstruct %8 %30 %213 +%215 = OpCompositeConstruct %9 %214 %35 +OpStore %171 %215 +%216 = OpLoad %7 %158 +%217 = OpCompositeConstruct %8 %216 %34 +%218 = OpCompositeConstruct %9 %32 %217 +OpStore %173 %218 +%219 = OpLoad %7 %158 +%220 = OpCompositeConstruct %8 %33 %219 +%221 = OpCompositeConstruct %9 %32 %220 +OpStore %175 %221 +%222 = OpLoad %7 %158 +%223 = OpCompositeConstruct %10 %222 %31 +OpStore %177 %223 +%224 = OpLoad %7 %158 +%225 = OpCompositeConstruct %10 %30 %224 +OpStore %179 %225 +%226 = OpLoad %7 %158 +%227 = OpCompositeConstruct %10 %226 %31 +OpStore %181 %227 +%228 = OpLoad %7 %158 +%229 = OpCompositeConstruct %10 %30 %228 +OpStore %183 %229 +%230 = OpLoad %3 %155 +%231 = OpCompositeConstruct %12 %230 %42 +OpStore %185 %231 +%232 = OpLoad %3 %155 +%233 = OpCompositeConstruct %12 %37 %232 +OpStore %187 %233 +%234 = OpLoad %7 %158 +%235 = OpCompositeConstruct %10 %234 %31 +OpStore %189 %235 +%236 = OpLoad %7 %158 +%237 = OpCompositeConstruct %10 %30 %236 +OpStore %191 %237 +%238 = OpLoad %7 %158 +%239 = OpCompositeConstruct %10 %238 %31 +OpStore %193 %239 +%240 = OpLoad %7 %158 +%241 = OpCompositeConstruct %10 %30 %240 +OpStore %195 %241 +%242 = OpLoad %3 %155 +%243 = OpCompositeConstruct %12 %242 %42 +OpStore %197 %243 +%244 = OpLoad %3 %155 +%245 = OpCompositeConstruct %12 %37 %244 +OpStore %199 %245 +%246 = OpLoad %5 %152 +%247 = OpCompositeConstruct %6 %246 %28 +OpStore %161 %247 +%248 = OpLoad %5 %152 +%249 = OpCompositeConstruct %6 %27 %248 +OpStore %163 %249 +%250 = OpLoad %5 %152 +%251 = OpCompositeConstruct %6 %250 %28 +OpStore %165 %251 +%252 = OpLoad %5 %152 +%253 = OpCompositeConstruct %6 %27 %252 +OpStore %167 %253 +%254 = OpLoad %7 %158 +%255 = OpCompositeConstruct %8 %254 %31 +%256 = OpCompositeConstruct %9 %255 %35 +OpStore %169 %256 +%257 = OpLoad %7 %158 +%258 = OpCompositeConstruct %8 %30 %257 +%259 = OpCompositeConstruct %9 %258 %35 +OpStore %171 %259 +%260 = OpLoad %7 %158 +%261 = OpCompositeConstruct %8 %260 %34 +%262 = OpCompositeConstruct %9 %32 %261 +OpStore %173 %262 +%263 = OpLoad %7 %158 +%264 = OpCompositeConstruct %8 %33 %263 +%265 = OpCompositeConstruct %9 %32 %264 +OpStore %175 %265 +%266 = OpLoad %7 %158 +%267 = OpCompositeConstruct %10 %266 %31 +OpStore %177 %267 +%268 = OpLoad %7 %158 +%269 = OpCompositeConstruct %10 %30 %268 +OpStore %179 %269 +%270 = OpLoad %7 %158 +%271 = OpCompositeConstruct %10 %270 %31 +OpStore %181 %271 +%272 = OpLoad %7 %158 +%273 = OpCompositeConstruct %10 %30 %272 +OpStore %183 %273 +%274 = OpLoad %3 %155 +%275 = OpCompositeConstruct %12 %274 %42 +OpStore %185 %275 +%276 = OpLoad %3 %155 +%277 = OpCompositeConstruct %12 %37 %276 +OpStore %187 %277 +%278 = OpLoad %7 %158 +%279 = OpCompositeConstruct %10 %278 %31 +OpStore %189 %279 +%280 = OpLoad %7 %158 +%281 = OpCompositeConstruct %10 %30 %280 +OpStore %191 %281 +%282 = OpLoad %7 %158 +%283 = OpCompositeConstruct %10 %282 %31 +OpStore %193 %283 +%284 = OpLoad %7 %158 +%285 = OpCompositeConstruct %10 %30 %284 +OpStore %195 %285 +%286 = OpLoad %3 %155 +%287 = OpCompositeConstruct %12 %286 %42 +OpStore %197 %287 +%288 = OpLoad %3 %155 +%289 = OpCompositeConstruct %12 %37 %288 +OpStore %199 %289 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/abstract-types-var.wgsl b/naga/tests/out/wgsl/abstract-types-var.wgsl index cd35501732..58d5f414f2 100644 --- a/naga/tests/out/wgsl/abstract-types-var.wgsl +++ b/naga/tests/out/wgsl/abstract-types-var.wgsl @@ -22,6 +22,9 @@ var xafpaiai_1: array = array(1i, 2i); var xafpaiaf_1: array = array(1f, 2f); var xafpafai_1: array = array(1f, 2f); var xafpafaf_1: array = array(1f, 2f); +var xavipai_1: array, 1> = array, 1>(vec3(1i)); +var xavfpai_1: array, 1> = array, 1>(vec3(1f)); +var xavfpaf_1: array, 1> = array, 1>(vec3(1f)); var ivispai: vec2 = vec2(1i); var ivfspaf: vec2 = vec2(1f); var ivis_ai: vec2 = vec2(1i); @@ -34,6 +37,9 @@ var iaipaiai_1: array = array(1i, 2i); var iafpafaf_1: array = array(1f, 2f); var iafpaiaf_1: array = array(1f, 2f); var iafpafai_1: array = array(1f, 2f); +var iavipai: array, 1> = array, 1>(vec3(1i)); +var iavfpai: array, 1> = array, 1>(vec3(1i)); +var iavfpaf: array, 1> = array, 1>(vec3(1f)); fn all_constant_arguments() { var xvipaiai: vec2 = vec2(42i, 43i); @@ -69,6 +75,9 @@ fn all_constant_arguments() { var xafpaiaf: array = array(1f, 2f); var xafpafai: array = array(1f, 2f); var xafpafaf: array = array(1f, 2f); + var xavipai: array, 1> = array, 1>(vec3(1i)); + var xavfpai: array, 1> = array, 1>(vec3(1f)); + var xavfpaf: array, 1> = array, 1>(vec3(1f)); var iaipaiai: array = array(1i, 2i); var iafpaiaf: array = array(1f, 2f); var iafpafai: array = array(1f, 2f); @@ -107,6 +116,9 @@ fn all_constant_arguments() { xafpaiaf = array(1f, 2f); xafpafai = array(1f, 2f); xafpafaf = array(1f, 2f); + xavipai = array, 1>(vec3(1i)); + xavfpai = array, 1>(vec3(1f)); + xavfpaf = array, 1>(vec3(1f)); iaipaiai = array(1i, 2i); iafpaiaf = array(1f, 2f); iafpafai = array(1f, 2f);