github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

[naga] Use const ctx instead of global ctx for type resolution (#6935)

Browse Source 
Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>
This commit is contained in:
Samson
2025-02-24 16:24:37 +01:00
committed by GitHub
parent e95f6d632c
commit 2f255edc60
8 changed files with 244 additions and 190 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
naga/src/front/wgsl/lower/construction.rs
View File

@@ -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();

93
naga/src/front/wgsl/lower/mod.rs
View File

@@ -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<crate::Expression>) -> Option<u32> {
fn const_eval_expr_to_u32(
&self,
handle: Handle<crate::Expression>,
) -> Result<u32, crate::proc::U32EvalError> {
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<crate::ArraySize, Error<'source>> {
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<ast::Type<'source>>,
name: Option<String>,
ctx: &mut GlobalContext<'source, '_, '_>,
ctx: &mut ExpressionContext<'source, '_, '_>,
) -> Result<Handle<crate::Type>, 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<Type>` representing the front-end type `handle`.
fn resolve_ast_type(
&mut self,
handle: Handle<ast::Type<'source>>,
ctx: &mut GlobalContext<'source, '_, '_>,
ctx: &mut ExpressionContext<'source, '_, '_>,
) -> Result<Handle<crate::Type>, Error<'source>> {
self.resolve_named_ast_type(handle, None, ctx)
}

6
naga/tests/in/const-exprs.wgsl
View File

@@ -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<bool>(), TRUE), false));
var vec_all_true = all(vec4(true));
}
fn test_local_const() {
const local_const = 2;
var arr: array<f32, local_const>;
}

6
naga/tests/out/glsl/const-exprs.main.Compute.glsl
View File

@@ -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;
}

8
naga/tests/out/hlsl/const-exprs.hlsl
View File

@@ -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;
}

10
naga/tests/out/msl/const-exprs.msl
View File

@@ -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;
}

296
naga/tests/out/spv/const-exprs.spvasm
View File

@@ -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
OpReturn
OpFunctionEnd
%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 %54 %61
%62 = OpLoad %4 %51
OpStore %53 %61
%62 = OpLoad %4 %53
OpStore %55 %62
%63 = OpLoad %4 %52
%64 = OpLoad %4 %54
%65 = OpLoad %4 %56
%66 = OpCompositeConstruct %6 %62 %63 %64 %65
OpStore %57 %66
%64 = OpLoad %4 %53
%65 = OpLoad %4 %55
%66 = OpLoad %4 %57
%67 = OpCompositeConstruct %6 %63 %64 %65 %66
OpStore %58 %67
OpReturn
OpFunctionEnd
%68 = OpFunction %2 None %32
%67 = OpLabel
%71 = OpVariable %35 Function %70
OpBranch %72
%72 = OpLabel
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
%101 = OpFunction %2 None %32
%87 = OpFunction %2 None %33
%86 = OpLabel
%88 = OpVariable %89 Function %90
OpBranch %91
%91 = OpLabel
OpReturn
OpFunctionEnd
%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

7
naga/tests/out/wgsl/const-exprs.wgsl
View File

@@ -70,6 +70,12 @@ fn compose_of_splat() {
return;
}
fn test_local_const() {
var arr: array<f32, 2>;
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;
}