From 2f255edc60e9669c8c737464c59af10d59a31126 Mon Sep 17 00:00:00 2001 From: Samson <16504129+sagudev@users.noreply.github.com> Date: Mon, 24 Feb 2025 16:24:37 +0100 Subject: [PATCH] [naga] Use const ctx instead of global ctx for type resolution (#6935) Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> --- naga/src/front/wgsl/lower/construction.rs | 8 +- naga/src/front/wgsl/lower/mod.rs | 93 +++--- naga/tests/in/const-exprs.wgsl | 6 + .../out/glsl/const-exprs.main.Compute.glsl | 6 + naga/tests/out/hlsl/const-exprs.hlsl | 8 + naga/tests/out/msl/const-exprs.msl | 10 + naga/tests/out/spv/const-exprs.spvasm | 296 +++++++++--------- naga/tests/out/wgsl/const-exprs.wgsl | 7 + 8 files changed, 244 insertions(+), 190 deletions(-) diff --git a/naga/src/front/wgsl/lower/construction.rs b/naga/src/front/wgsl/lower/construction.rs index ee973fb9f1..21317545bb 100644 --- a/naga/src/front/wgsl/lower/construction.rs +++ b/naga/src/front/wgsl/lower/construction.rs @@ -579,7 +579,7 @@ impl<'source> Lowerer<'source, '_> { } ast::ConstructorType::PartialVector { size } => Constructor::PartialVector { size }, ast::ConstructorType::Vector { size, ty, ty_span } => { - let ty = self.resolve_ast_type(ty, &mut ctx.as_global())?; + let ty = self.resolve_ast_type(ty, &mut ctx.as_const())?; let scalar = match ctx.module.types[ty].inner { crate::TypeInner::Scalar(sc) => sc, _ => return Err(Error::UnknownScalarType(ty_span)), @@ -596,7 +596,7 @@ impl<'source> Lowerer<'source, '_> { ty, ty_span, } => { - let ty = self.resolve_ast_type(ty, &mut ctx.as_global())?; + let ty = self.resolve_ast_type(ty, &mut ctx.as_const())?; let scalar = match ctx.module.types[ty].inner { crate::TypeInner::Scalar(sc) => sc, _ => return Err(Error::UnknownScalarType(ty_span)), @@ -613,8 +613,8 @@ impl<'source> Lowerer<'source, '_> { } ast::ConstructorType::PartialArray => Constructor::PartialArray, ast::ConstructorType::Array { base, size } => { - let base = self.resolve_ast_type(base, &mut ctx.as_global())?; - let size = self.array_size(size, &mut ctx.as_global())?; + let base = self.resolve_ast_type(base, &mut ctx.as_const())?; + let size = self.array_size(size, &mut ctx.as_const())?; ctx.layouter.update(ctx.module.to_ctx()).unwrap(); let stride = ctx.layouter[base].to_stride(); diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 7cd3ef90f8..70900e0db3 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -244,6 +244,7 @@ impl<'a, 'temp> StatementContext<'a, 'temp, '_> { } } + #[allow(dead_code)] fn as_global(&mut self) -> GlobalContext<'a, '_, '_> { GlobalContext { ast_expressions: self.ast_expressions, @@ -468,29 +469,28 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> { .map_err(|e| Error::ConstantEvaluatorError(e.into(), span)) } - fn const_access(&self, handle: Handle) -> Option { + fn const_eval_expr_to_u32( + &self, + handle: Handle, + ) -> Result { match self.expr_type { ExpressionContextType::Runtime(ref ctx) => { if !ctx.local_expression_kind_tracker.is_const(handle) { - return None; + return Err(crate::proc::U32EvalError::NonConst); } self.module .to_ctx() .eval_expr_to_u32_from(handle, &ctx.function.expressions) - .ok() } ExpressionContextType::Constant(Some(ref ctx)) => { assert!(ctx.local_expression_kind_tracker.is_const(handle)); self.module .to_ctx() .eval_expr_to_u32_from(handle, &ctx.function.expressions) - .ok() } - ExpressionContextType::Constant(None) => { - self.module.to_ctx().eval_expr_to_u32(handle).ok() - } - ExpressionContextType::Override => None, + ExpressionContextType::Constant(None) => self.module.to_ctx().eval_expr_to_u32(handle), + ExpressionContextType::Override => Err(crate::proc::U32EvalError::NonConst), } } @@ -1069,7 +1069,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { } ast::GlobalDeclKind::Var(ref v) => { let explicit_ty = - v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx)) + v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const())) .transpose()?; let (ty, initializer) = @@ -1102,7 +1102,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let mut ectx = ctx.as_const(); let explicit_ty = - c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx.as_global())) + c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx)) .transpose()?; let (ty, init) = @@ -1123,7 +1123,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { } ast::GlobalDeclKind::Override(ref o) => { let explicit_ty = - o.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx)) + o.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const())) .transpose()?; let mut ectx = ctx.as_override(); @@ -1165,7 +1165,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let ty = self.resolve_named_ast_type( alias.ty, Some(alias.name.name.to_string()), - &mut ctx, + &mut ctx.as_const(), )?; ctx.globals .insert(alias.name.name, LoweredGlobalDecl::Type(ty)); @@ -1263,7 +1263,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { .iter() .enumerate() .map(|(i, arg)| -> Result<_, Error<'_>> { - let ty = self.resolve_ast_type(arg.ty, ctx)?; + let ty = self.resolve_ast_type(arg.ty, &mut ctx.as_const())?; let expr = expressions .append(crate::Expression::FunctionArgument(i as u32), arg.name.span); local_table.insert(arg.handle, Declared::Runtime(Typed::Plain(expr))); @@ -1282,7 +1282,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { .result .as_ref() .map(|res| -> Result<_, Error<'_>> { - let ty = self.resolve_ast_type(res.ty, ctx)?; + let ty = self.resolve_ast_type(res.ty, &mut ctx.as_const())?; Ok(crate::FunctionResult { ty, binding: self.binding(&res.binding, ty, ctx)?, @@ -1440,9 +1440,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { // optimization. ctx.local_expression_kind_tracker.force_non_const(value); - let explicit_ty = - l.ty.map(|ty| self.resolve_ast_type(ty, &mut ctx.as_global())) - .transpose()?; + let explicit_ty = l + .ty + .map(|ty| self.resolve_ast_type(ty, &mut ctx.as_const(block, &mut emitter))) + .transpose()?; if let Some(ty) = explicit_ty { let mut ctx = ctx.as_expression(block, &mut emitter); @@ -1469,12 +1470,15 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { return Ok(()); } ast::LocalDecl::Var(ref v) => { - let explicit_ty = - v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_global())) - .transpose()?; - let mut emitter = Emitter::default(); emitter.start(&ctx.function.expressions); + + let explicit_ty = + v.ty.map(|ast| { + self.resolve_ast_type(ast, &mut ctx.as_const(block, &mut emitter)) + }) + .transpose()?; + let mut ectx = ctx.as_expression(block, &mut emitter); let (ty, initializer) = self.type_and_init(v.name, v.init, explicit_ty, &mut ectx)?; @@ -1533,11 +1537,15 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let ectx = &mut ctx.as_const(block, &mut emitter); let explicit_ty = - c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx.as_global())) + c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx.as_const())) .transpose()?; - let (_ty, init) = - self.type_and_init(c.name, Some(c.init), explicit_ty, ectx)?; + let (_ty, init) = self.type_and_init( + c.name, + Some(c.init), + explicit_ty, + &mut ectx.as_const(), + )?; let init = init.expect("Local const must have init"); block.extend(emitter.finish(&ctx.function.expressions)); @@ -1992,7 +2000,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { } } - lowered_base.map(|base| match ctx.const_access(index) { + lowered_base.map(|base| match ctx.const_eval_expr_to_u32(index).ok() { Some(index) => crate::Expression::AccessIndex { base, index }, None => crate::Expression::Access { base, index }, }) @@ -2069,7 +2077,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { } ast::Expression::Bitcast { expr, to, ty_span } => { let expr = self.expression(expr, ctx)?; - let to_resolved = self.resolve_ast_type(to, &mut ctx.as_global())?; + let to_resolved = self.resolve_ast_type(to, &mut ctx.as_const())?; let element_scalar = match ctx.module.types[to_resolved].inner { crate::TypeInner::Scalar(scalar) => scalar, @@ -3051,7 +3059,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let mut members = Vec::with_capacity(s.members.len()); for member in s.members.iter() { - let ty = self.resolve_ast_type(member.ty, ctx)?; + let ty = self.resolve_ast_type(member.ty, &mut ctx.as_const())?; ctx.layouter.update(ctx.module.to_ctx()).unwrap(); @@ -3138,7 +3146,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { fn array_size( &mut self, size: ast::ArraySize<'source>, - ctx: &mut GlobalContext<'source, '_, '_>, + ctx: &mut ExpressionContext<'source, '_, '_>, ) -> Result> { Ok(match size { ast::ArraySize::Constant(expr) => { @@ -3146,17 +3154,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let const_expr = self.expression(expr, &mut ctx.as_const()); match const_expr { Ok(value) => { - let len = - ctx.module.to_ctx().eval_expr_to_u32(value).map_err( - |err| match err { - crate::proc::U32EvalError::NonConst => { - Error::ExpectedConstExprConcreteIntegerScalar(span) - } - crate::proc::U32EvalError::Negative => { - Error::ExpectedPositiveArrayLength(span) - } - }, - )?; + let len = ctx.const_eval_expr_to_u32(value).map_err(|err| match err { + crate::proc::U32EvalError::NonConst => { + Error::ExpectedConstExprConcreteIntegerScalar(span) + } + crate::proc::U32EvalError::Negative => { + Error::ExpectedPositiveArrayLength(span) + } + })?; let size = NonZeroU32::new(len).ok_or(Error::ExpectedPositiveArrayLength(span))?; crate::ArraySize::Constant(size) @@ -3167,7 +3172,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { crate::proc::ConstantEvaluatorError::OverrideExpr => { crate::ArraySize::Pending(self.array_size_override( expr, - &mut ctx.as_override(), + &mut ctx.as_global().as_override(), span, )?) } @@ -3219,7 +3224,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { &mut self, handle: Handle>, name: Option, - ctx: &mut GlobalContext<'source, '_, '_>, + ctx: &mut ExpressionContext<'source, '_, '_>, ) -> Result, Error<'source>> { let inner = match ctx.types[handle] { ast::Type::Scalar(scalar) => scalar.to_inner_scalar(), @@ -3257,7 +3262,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { crate::TypeInner::Pointer { base, space } } ast::Type::Array { base, size } => { - let base = self.resolve_ast_type(base, ctx)?; + let base = self.resolve_ast_type(base, &mut ctx.as_const())?; let size = self.array_size(size, ctx)?; ctx.layouter.update(ctx.module.to_ctx()).unwrap(); @@ -3297,14 +3302,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { } }; - Ok(ctx.ensure_type_exists(name, inner)) + Ok(ctx.as_global().ensure_type_exists(name, inner)) } /// Return a Naga `Handle` representing the front-end type `handle`. fn resolve_ast_type( &mut self, handle: Handle>, - ctx: &mut GlobalContext<'source, '_, '_>, + ctx: &mut ExpressionContext<'source, '_, '_>, ) -> Result, Error<'source>> { self.resolve_named_ast_type(handle, None, ctx) } diff --git a/naga/tests/in/const-exprs.wgsl b/naga/tests/in/const-exprs.wgsl index 72e5726355..c8ab84cd3f 100644 --- a/naga/tests/in/const-exprs.wgsl +++ b/naga/tests/in/const-exprs.wgsl @@ -12,6 +12,7 @@ fn main() { splat_of_constant(); compose_of_constant(); compose_of_splat(); + test_local_const(); } // Swizzle the value of nested Compose expressions. @@ -109,3 +110,8 @@ fn relational() { var vec_all_false = all(vec4(vec3(vec2(), TRUE), false)); var vec_all_true = all(vec4(true)); } + +fn test_local_const() { + const local_const = 2; + var arr: array; +} diff --git a/naga/tests/out/glsl/const-exprs.main.Compute.glsl b/naga/tests/out/glsl/const-exprs.main.Compute.glsl index f2607ce0eb..e4d50b297a 100644 --- a/naga/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/naga/tests/out/glsl/const-exprs.main.Compute.glsl @@ -71,6 +71,11 @@ void compose_of_splat() { return; } +void test_local_const() { + float arr[2] = float[2](0.0, 0.0); + return; +} + uint map_texture_kind(int texture_kind) { switch(texture_kind) { case 0: { @@ -115,6 +120,7 @@ void main() { splat_of_constant(); compose_of_constant(); compose_of_splat(); + test_local_const(); return; } diff --git a/naga/tests/out/hlsl/const-exprs.hlsl b/naga/tests/out/hlsl/const-exprs.hlsl index e0c4da5c77..40222d54e4 100644 --- a/naga/tests/out/hlsl/const-exprs.hlsl +++ b/naga/tests/out/hlsl/const-exprs.hlsl @@ -77,6 +77,13 @@ void compose_of_splat() return; } +void test_local_const() +{ + float arr[2] = (float[2])0; + + return; +} + uint map_texture_kind(int texture_kind) { switch(texture_kind) { @@ -128,5 +135,6 @@ void main() splat_of_constant(); compose_of_constant(); compose_of_splat(); + test_local_const(); return; } diff --git a/naga/tests/out/msl/const-exprs.msl b/naga/tests/out/msl/const-exprs.msl index 71dd868f4e..fd22ef72ac 100644 --- a/naga/tests/out/msl/const-exprs.msl +++ b/naga/tests/out/msl/const-exprs.msl @@ -4,6 +4,9 @@ using metal::uint; +struct type_6 { + float inner[2]; +}; constant uint TWO = 2u; constant int THREE = 3; constant bool TRUE = true; @@ -76,6 +79,12 @@ void compose_of_splat( return; } +void test_local_const( +) { + type_6 arr = {}; + return; +} + uint map_texture_kind( int texture_kind ) { @@ -125,5 +134,6 @@ kernel void main_( splat_of_constant(); compose_of_constant(); compose_of_splat(); + test_local_const(); return; } diff --git a/naga/tests/out/spv/const-exprs.spvasm b/naga/tests/out/spv/const-exprs.spvasm index 83777667d9..ddd837b2f0 100644 --- a/naga/tests/out/spv/const-exprs.spvasm +++ b/naga/tests/out/spv/const-exprs.spvasm @@ -1,12 +1,13 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 132 +; Bound: 140 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %123 "main" -OpExecutionMode %123 LocalSize 2 3 1 +OpEntryPoint GLCompute %130 "main" +OpExecutionMode %130 LocalSize 2 3 1 +OpDecorate %9 ArrayStride 4 %2 = OpTypeVoid %3 = OpTypeInt 32 0 %4 = OpTypeInt 32 1 @@ -14,168 +15,179 @@ OpExecutionMode %123 LocalSize 2 3 1 %6 = OpTypeVector %4 4 %7 = OpTypeFloat 32 %8 = OpTypeVector %7 4 -%9 = OpTypeVector %7 2 -%10 = OpTypeVector %5 2 -%11 = OpTypeVector %4 3 -%12 = OpConstant %3 2 -%13 = OpConstant %4 3 -%14 = OpConstantTrue %5 -%15 = OpConstantFalse %5 -%16 = OpConstant %4 4 -%17 = OpConstant %4 8 -%18 = OpConstant %7 3.141 -%19 = OpConstant %7 6.282 -%20 = OpConstant %7 0.44444445 -%21 = OpConstant %7 0.0 -%22 = OpConstantComposite %8 %20 %21 %21 %21 -%23 = OpConstant %4 0 -%24 = OpConstant %4 1 -%25 = OpConstant %4 2 -%26 = OpConstant %7 4.0 -%27 = OpConstant %7 5.0 -%28 = OpConstantComposite %9 %26 %27 -%29 = OpConstantComposite %10 %14 %15 -%32 = OpTypeFunction %2 -%33 = OpConstantComposite %6 %16 %13 %25 %24 -%35 = OpTypePointer Function %6 -%40 = OpTypePointer Function %4 -%44 = OpConstant %4 6 -%49 = OpConstant %4 30 -%50 = OpConstant %4 70 -%53 = OpConstantNull %4 -%55 = OpConstantNull %4 -%58 = OpConstantNull %6 -%69 = OpConstant %4 -4 -%70 = OpConstantComposite %6 %69 %69 %69 %69 -%79 = OpConstant %7 1.0 -%80 = OpConstant %7 2.0 -%81 = OpConstantComposite %8 %80 %79 %79 %79 -%83 = OpTypePointer Function %8 -%88 = OpTypeFunction %3 %4 -%89 = OpConstant %3 10 -%90 = OpConstant %3 20 -%91 = OpConstant %3 30 -%92 = OpConstant %3 0 -%99 = OpConstantNull %3 -%102 = OpConstantComposite %11 %24 %24 %24 -%103 = OpConstantComposite %11 %23 %24 %25 -%104 = OpConstantComposite %11 %24 %23 %25 -%106 = OpTypePointer Function %11 -%113 = OpTypePointer Function %5 -%31 = OpFunction %2 None %32 -%30 = OpLabel -%34 = OpVariable %35 Function %33 -OpBranch %36 -%36 = OpLabel -OpReturn -OpFunctionEnd -%38 = OpFunction %2 None %32 +%10 = OpConstant %3 2 +%9 = OpTypeArray %7 %10 +%11 = OpTypeVector %7 2 +%12 = OpTypeVector %5 2 +%13 = OpTypeVector %4 3 +%14 = OpConstant %4 3 +%15 = OpConstantTrue %5 +%16 = OpConstantFalse %5 +%17 = OpConstant %4 4 +%18 = OpConstant %4 8 +%19 = OpConstant %7 3.141 +%20 = OpConstant %7 6.282 +%21 = OpConstant %7 0.44444445 +%22 = OpConstant %7 0.0 +%23 = OpConstantComposite %8 %21 %22 %22 %22 +%24 = OpConstant %4 0 +%25 = OpConstant %4 1 +%26 = OpConstant %4 2 +%27 = OpConstant %7 4.0 +%28 = OpConstant %7 5.0 +%29 = OpConstantComposite %11 %27 %28 +%30 = OpConstantComposite %12 %15 %16 +%33 = OpTypeFunction %2 +%34 = OpConstantComposite %6 %17 %14 %26 %25 +%36 = OpTypePointer Function %6 +%41 = OpTypePointer Function %4 +%45 = OpConstant %4 6 +%50 = OpConstant %4 30 +%51 = OpConstant %4 70 +%54 = OpConstantNull %4 +%56 = OpConstantNull %4 +%59 = OpConstantNull %6 +%70 = OpConstant %4 -4 +%71 = OpConstantComposite %6 %70 %70 %70 %70 +%80 = OpConstant %7 1.0 +%81 = OpConstant %7 2.0 +%82 = OpConstantComposite %8 %81 %80 %80 %80 +%84 = OpTypePointer Function %8 +%89 = OpTypePointer Function %9 +%90 = OpConstantNull %9 +%95 = OpTypeFunction %3 %4 +%96 = OpConstant %3 10 +%97 = OpConstant %3 20 +%98 = OpConstant %3 30 +%99 = OpConstant %3 0 +%106 = OpConstantNull %3 +%109 = OpConstantComposite %13 %25 %25 %25 +%110 = OpConstantComposite %13 %24 %25 %26 +%111 = OpConstantComposite %13 %25 %24 %26 +%113 = OpTypePointer Function %13 +%120 = OpTypePointer Function %5 +%32 = OpFunction %2 None %33 +%31 = OpLabel +%35 = OpVariable %36 Function %34 +OpBranch %37 %37 = OpLabel -%39 = OpVariable %40 Function %25 -OpBranch %41 -%41 = OpLabel OpReturn OpFunctionEnd -%43 = OpFunction %2 None %32 +%39 = OpFunction %2 None %33 +%38 = OpLabel +%40 = OpVariable %41 Function %26 +OpBranch %42 %42 = OpLabel -%45 = OpVariable %40 Function %44 -OpBranch %46 -%46 = OpLabel OpReturn OpFunctionEnd -%48 = OpFunction %2 None %32 +%44 = OpFunction %2 None %33 +%43 = OpLabel +%46 = OpVariable %41 Function %45 +OpBranch %47 %47 = OpLabel -%57 = OpVariable %35 Function %58 -%52 = OpVariable %40 Function %53 -%56 = OpVariable %40 Function %50 -%51 = OpVariable %40 Function %49 -%54 = OpVariable %40 Function %55 -OpBranch %59 -%59 = OpLabel -%60 = OpLoad %4 %51 -OpStore %52 %60 -%61 = OpLoad %4 %52 -OpStore %54 %61 -%62 = OpLoad %4 %51 -%63 = OpLoad %4 %52 -%64 = OpLoad %4 %54 -%65 = OpLoad %4 %56 -%66 = OpCompositeConstruct %6 %62 %63 %64 %65 -OpStore %57 %66 OpReturn OpFunctionEnd -%68 = OpFunction %2 None %32 -%67 = OpLabel -%71 = OpVariable %35 Function %70 -OpBranch %72 -%72 = OpLabel +%49 = OpFunction %2 None %33 +%48 = OpLabel +%58 = OpVariable %36 Function %59 +%53 = OpVariable %41 Function %54 +%57 = OpVariable %41 Function %51 +%52 = OpVariable %41 Function %50 +%55 = OpVariable %41 Function %56 +OpBranch %60 +%60 = OpLabel +%61 = OpLoad %4 %52 +OpStore %53 %61 +%62 = OpLoad %4 %53 +OpStore %55 %62 +%63 = OpLoad %4 %52 +%64 = OpLoad %4 %53 +%65 = OpLoad %4 %55 +%66 = OpLoad %4 %57 +%67 = OpCompositeConstruct %6 %63 %64 %65 %66 +OpStore %58 %67 OpReturn OpFunctionEnd -%74 = OpFunction %2 None %32 +%69 = OpFunction %2 None %33 +%68 = OpLabel +%72 = OpVariable %36 Function %71 +OpBranch %73 %73 = OpLabel -%75 = OpVariable %35 Function %70 -OpBranch %76 -%76 = OpLabel OpReturn OpFunctionEnd -%78 = OpFunction %2 None %32 +%75 = OpFunction %2 None %33 +%74 = OpLabel +%76 = OpVariable %36 Function %71 +OpBranch %77 %77 = OpLabel -%82 = OpVariable %83 Function %81 -OpBranch %84 -%84 = OpLabel OpReturn OpFunctionEnd -%87 = OpFunction %3 None %88 -%86 = OpFunctionParameter %4 +%79 = OpFunction %2 None %33 +%78 = OpLabel +%83 = OpVariable %84 Function %82 +OpBranch %85 %85 = OpLabel -OpBranch %93 -%93 = OpLabel -OpSelectionMerge %94 None -OpSwitch %86 %98 0 %95 1 %96 2 %97 -%95 = OpLabel -OpReturnValue %89 -%96 = OpLabel -OpReturnValue %90 -%97 = OpLabel -OpReturnValue %91 -%98 = OpLabel -OpReturnValue %92 -%94 = OpLabel -OpReturnValue %99 +OpReturn +OpFunctionEnd +%87 = OpFunction %2 None %33 +%86 = OpLabel +%88 = OpVariable %89 Function %90 +OpBranch %91 +%91 = OpLabel +OpReturn OpFunctionEnd -%101 = OpFunction %2 None %32 +%94 = OpFunction %3 None %95 +%93 = OpFunctionParameter %4 +%92 = OpLabel +OpBranch %100 %100 = OpLabel -%105 = OpVariable %106 Function %102 -%107 = OpVariable %106 Function %103 -%108 = OpVariable %106 Function %104 -OpBranch %109 -%109 = OpLabel +OpSelectionMerge %101 None +OpSwitch %93 %105 0 %102 1 %103 2 %104 +%102 = OpLabel +OpReturnValue %96 +%103 = OpLabel +OpReturnValue %97 +%104 = OpLabel +OpReturnValue %98 +%105 = OpLabel +OpReturnValue %99 +%101 = OpLabel +OpReturnValue %106 +OpFunctionEnd +%108 = OpFunction %2 None %33 +%107 = OpLabel +%112 = OpVariable %113 Function %109 +%114 = OpVariable %113 Function %110 +%115 = OpVariable %113 Function %111 +OpBranch %116 +%116 = OpLabel OpReturn OpFunctionEnd -%111 = OpFunction %2 None %32 -%110 = OpLabel -%119 = OpVariable %113 Function %15 -%116 = OpVariable %113 Function %14 -%112 = OpVariable %113 Function %15 -%120 = OpVariable %113 Function %14 -%117 = OpVariable %113 Function %15 -%114 = OpVariable %113 Function %14 -%118 = OpVariable %113 Function %14 -%115 = OpVariable %113 Function %15 -OpBranch %121 -%121 = OpLabel +%118 = OpFunction %2 None %33 +%117 = OpLabel +%126 = OpVariable %120 Function %16 +%123 = OpVariable %120 Function %15 +%119 = OpVariable %120 Function %16 +%127 = OpVariable %120 Function %15 +%124 = OpVariable %120 Function %16 +%121 = OpVariable %120 Function %15 +%125 = OpVariable %120 Function %15 +%122 = OpVariable %120 Function %16 +OpBranch %128 +%128 = OpLabel OpReturn OpFunctionEnd -%123 = OpFunction %2 None %32 -%122 = OpLabel -OpBranch %124 -%124 = OpLabel -%125 = OpFunctionCall %2 %31 -%126 = OpFunctionCall %2 %38 -%127 = OpFunctionCall %2 %43 -%128 = OpFunctionCall %2 %48 -%129 = OpFunctionCall %2 %68 -%130 = OpFunctionCall %2 %74 -%131 = OpFunctionCall %2 %78 +%130 = OpFunction %2 None %33 +%129 = OpLabel +OpBranch %131 +%131 = OpLabel +%132 = OpFunctionCall %2 %32 +%133 = OpFunctionCall %2 %39 +%134 = OpFunctionCall %2 %44 +%135 = OpFunctionCall %2 %49 +%136 = OpFunctionCall %2 %69 +%137 = OpFunctionCall %2 %75 +%138 = OpFunctionCall %2 %79 +%139 = OpFunctionCall %2 %87 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 9fcbb588c9..1e94167c24 100644 --- a/naga/tests/out/wgsl/const-exprs.wgsl +++ b/naga/tests/out/wgsl/const-exprs.wgsl @@ -70,6 +70,12 @@ fn compose_of_splat() { return; } +fn test_local_const() { + var arr: array; + + return; +} + fn map_texture_kind(texture_kind: i32) -> u32 { switch texture_kind { case 0: { @@ -117,5 +123,6 @@ fn main() { splat_of_constant(); compose_of_constant(); compose_of_splat(); + test_local_const(); return; }