Skip to content

Commit

Permalink
[naga wgsl-in] Do not eagerly concretize local const declarations of …
Browse files Browse the repository at this point in the history 
…abstract types

Instead allow the const to be converted each time it is const
evaluated as part of another expression. This allows an abstract const
to be used as a different type depending on the context.

As a result, abstract types may now find their way in to the IR, which
we don't want. This occurs because the compact pass treats named
expressions as used, mostly so that our snapshot tests are more
useful, and therefore does not remove them. To prevent this, we avoid
adding abstract-typed local consts to the named expressions list. This
will have no functional effect on any shaders produced by the
backends, but some unused local const declarations will no longer be
present.
  • Loading branch information
@jamienicol
jamienicol committed Feb 25, 2025
1 parent e0f0185 commit da8d10e
Show file tree
Hide file tree
Showing 13 changed files with 64 additions and 131 deletions.
14 changes: 9 additions & 5 deletions naga/src/front/wgsl/lower/mod.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1574,21 +1574,25 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx.as_const()))
.transpose()?;

let (_ty, init) = self.type_and_init(
let (ty, init) = self.type_and_init(
c.name,
Some(c.init),
explicit_ty,
AbstractRule::Concretize,
AbstractRule::Allow,
&mut ectx.as_const(),
)?;
let init = init.expect("Local const must have init");

block.extend(emitter.finish(&ctx.function.expressions));
ctx.local_table
.insert(c.handle, Declared::Const(Typed::Plain(init)));
ctx.named_expressions
.insert(init, (c.name.name.to_string(), c.name.span));

// Only add constants of non-abstract types to the named expressions
// to prevent abstract types ending up in the IR.
let is_abstract = ctx.module.types[ty].inner.is_abstract(&ctx.module.types);
if !is_abstract {
ctx.named_expressions
.insert(init, (c.name.name.to_string(), c.name.span));
}
return Ok(());
}
},
Expand Down
5 changes: 5 additions & 0 deletions naga/tests/in/wgsl/abstract-types-return.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -29,3 +29,8 @@ const one = 1;
fn return_const_f32_const_ai() -> f32 {
return one;
}

fn return_vec2f32_const_ai() -> vec2<f32> {
const vec_one = vec2(1);
return vec_one;
}
4 changes: 4 additions & 0 deletions naga/tests/in/wgsl/const_assert.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,4 +13,8 @@ fn foo() {
const z = x + y - 2;
const_assert z > 0; // valid in functions.
const_assert(z > 0);

const_assert z == 1i;
const_assert z > 0u;
const_assert z < 2.0f;
}
4 changes: 4 additions & 0 deletions naga/tests/out/glsl/abstract-types-return.main.Compute.glsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,10 @@ float return_const_f32_const_ai() {
return 1.0;
}

vec2 return_vec2f32_const_ai() {
return vec2(1.0);
}

void main() {
return;
}
Expand Down
5 changes: 5 additions & 0 deletions naga/tests/out/hlsl/abstract-types-return.hlsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,11 @@ float return_const_f32_const_ai()
return 1.0;
}

float2 return_vec2f32_const_ai()
{
return (1.0).xx;
}

[numthreads(1, 1, 1)]
void main()
{
Expand Down
8 changes: 2 additions & 6 deletions naga/tests/out/ir/const_assert.compact.ron
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15,12 +15,8 @@
arguments: [],
result: None,
local_variables: [],
expressions: [
Literal(I32(1)),
],
named_expressions: {
0: "z",
},
expressions: [],
named_expressions: {},
body: [
Return(
value: None,
Expand Down
8 changes: 2 additions & 6 deletions naga/tests/out/ir/const_assert.ron
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15,12 +15,8 @@
arguments: [],
result: None,
local_variables: [],
expressions: [
Literal(I32(1)),
],
named_expressions: {
0: "z",
},
expressions: [],
named_expressions: {},
body: [
Return(
value: None,
Expand Down
59 changes: 6 additions & 53 deletions naga/tests/out/ir/local-const.compact.ron
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,16 +21,6 @@
width: 4,
)),
),
(
name: None,
inner: Vector(
size: Tri,
scalar: (
kind: Sint,
width: 4,
),
),
),
],
special_types: (
ray_desc: None,
Expand Down Expand Up @@ -68,63 +58,26 @@
result: None,
local_variables: [],
expressions: [
Literal(I32(4)),
Literal(I32(4)),
Literal(U32(4)),
Literal(F32(4.0)),
Compose(
ty: 3,
components: [
0,
0,
0,
],
),
Literal(F32(2.0)),
Literal(I32(4)),
Constant(0),
Constant(1),
Constant(2),
Literal(I32(4)),
Literal(I32(4)),
Literal(I32(4)),
Compose(
ty: 3,
components: [
10,
11,
12,
],
),
Literal(F32(2.0)),
],
named_expressions: {
0: "a",
1: "b",
2: "c",
3: "d",
4: "e",
5: "f",
6: "ag",
7: "bg",
8: "cg",
9: "dg",
13: "eg",
14: "fg",
0: "b",
1: "c",
2: "d",
3: "bg",
4: "cg",
5: "dg",
},
body: [
Emit((
start: 4,
end: 5,
)),
Emit((
start: 0,
end: 0,
)),
Emit((
start: 13,
end: 14,
)),
Return(
value: None,
),
Expand Down
59 changes: 6 additions & 53 deletions naga/tests/out/ir/local-const.ron
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,16 +21,6 @@
width: 4,
)),
),
(
name: None,
inner: Vector(
size: Tri,
scalar: (
kind: Sint,
width: 4,
),
),
),
],
special_types: (
ray_desc: None,
Expand Down Expand Up @@ -68,63 +58,26 @@
result: None,
local_variables: [],
expressions: [
Literal(I32(4)),
Literal(I32(4)),
Literal(U32(4)),
Literal(F32(4.0)),
Compose(
ty: 3,
components: [
0,
0,
0,
],
),
Literal(F32(2.0)),
Literal(I32(4)),
Constant(0),
Constant(1),
Constant(2),
Literal(I32(4)),
Literal(I32(4)),
Literal(I32(4)),
Compose(
ty: 3,
components: [
10,
11,
12,
],
),
Literal(F32(2.0)),
],
named_expressions: {
0: "a",
1: "b",
2: "c",
3: "d",
4: "e",
5: "f",
6: "ag",
7: "bg",
8: "cg",
9: "dg",
13: "eg",
14: "fg",
0: "b",
1: "c",
2: "d",
3: "bg",
4: "cg",
5: "dg",
},
body: [
Emit((
start: 4,
end: 5,
)),
Emit((
start: 0,
end: 0,
)),
Emit((
start: 13,
end: 14,
)),
Return(
value: None,
),
Expand Down
5 changes: 5 additions & 0 deletions naga/tests/out/msl/abstract-types-return.msl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,11 @@ float return_const_f32_const_ai(
return 1.0;
}

metal::float2 return_vec2f32_const_ai(
) {
return metal::float2(1.0);
}

kernel void main_(
) {
return;
Expand Down
18 changes: 12 additions & 6 deletions naga/tests/out/spv/abstract-types-return.spvasm
View file
Original file line number Diff line number Diff line change
@@ -1,12 +1,12 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 44
; Bound: 47
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %41 "main"
OpExecutionMode %41 LocalSize 1 1 1
OpEntryPoint GLCompute %44 "main"
OpExecutionMode %44 LocalSize 1 1 1
OpDecorate %7 ArrayStride 4
%2 = OpTypeVoid
%3 = OpTypeInt 32 1
Expand All @@ -25,7 +25,7 @@ OpDecorate %7 ArrayStride 4
%30 = OpConstantComposite %6 %22 %22
%34 = OpTypeFunction %7
%35 = OpConstantComposite %7 %22 %22 %22 %22
%42 = OpTypeFunction %2
%45 = OpTypeFunction %2
%10 = OpFunction %3 None %11
%9 = OpLabel
OpBranch %13
Expand Down Expand Up @@ -68,9 +68,15 @@ OpBranch %39
%39 = OpLabel
OpReturnValue %22
OpFunctionEnd
%41 = OpFunction %2 None %42
%41 = OpFunction %6 None %29
%40 = OpLabel
OpBranch %43
OpBranch %42
%42 = OpLabel
OpReturnValue %30
OpFunctionEnd
%44 = OpFunction %2 None %45
%43 = OpLabel
OpBranch %46
%46 = OpLabel
OpReturn
OpFunctionEnd
4 changes: 4 additions & 0 deletions naga/tests/out/wgsl/abstract-types-return.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -26,6 +26,10 @@ fn return_const_f32_const_ai() -> f32 {
return 1f;
}

fn return_vec2f32_const_ai() -> vec2<f32> {
return vec2(1f);
}

@compute @workgroup_size(1, 1, 1)
fn main() {
return;
Expand Down
2 changes: 0 additions & 2 deletions naga/tests/out/wgsl/local-const.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,8 +3,6 @@ const gc: u32 = 4u;
const gd: f32 = 4f;

fn const_in_fn() {
const e = vec3<i32>(4i, 4i, 4i);
const eg = vec3<i32>(4i, 4i, 4i);
return;
}

0 comments on commit da8d10e

Please sign in to comment.