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 wgsl-in] Do not eagerly concretize global const declarations of abstract types

Browse Source 
Instead allow the const to be converted and 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.

A consequence of this is that abstract types may now find their way to
the validation stage, which we don't want. We therefore additionally
now ensure that the compact pass removes global constants of abstract
types. This will have no *functional* effect on shaders generated by
the backends, as the expressions belonging to the abstract consts in
the IR will not actually be used, as any usage in the input shader
will have been const-evaluated away. Certain unused const declarations
will now be removed, however, as can be seen by the effect on the
snapshot outputs.
This commit is contained in:
Jamie Nicol
2025-02-06 10:09:01 +00:00
committed by Jamie Nicol
parent da90d7aaa6
commit dd2d53814d
19 changed files with 178 additions and 243 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
CHANGELOG.md
View File

@@ -182,6 +182,7 @@ By @brodycj in [#6924](https://github.com/gfx-rs/wgpu/pull/6924).
- Error if structs have two fields with the same name. By @SparkyPotato in [#7088](https://github.com/gfx-rs/wgpu/pull/7088).
- Forward '--keep-coordinate-space' flag to GLSL backend in naga-cli. By @cloone8 in [#7206](https://github.com/gfx-rs/wgpu/pull/7206).
- Allow template lists to have a trailing comma. By @KentSlaney in [#7142](https://github.com/gfx-rs/wgpu/pull/7142).
- Allow WGSL const declarations to have abstract types. By @jamienicol in [#7055](https://github.com/gfx-rs/wgpu/pull/7055).
#### General

21
naga/src/compact/mod.rs
View File

@@ -110,14 +110,27 @@ pub fn compact(module: &mut crate::Module) {
log::trace!("tracing special types");
module_tracer.trace_special_types(&module.special_types);
// We treat all named constants as used by definition.
// We treat all named constants as used by definition, unless they have an
// abstract type as we do not want those reaching the validator.
log::trace!("tracing named constants");
for (handle, constant) in module.constants.iter() {
if constant.name.is_some() {
log::trace!("tracing constant {:?}", constant.name.as_ref().unwrap());
module_tracer.constants_used.insert(handle);
module_tracer.types_used.insert(constant.ty);
module_tracer.global_expressions_used.insert(constant.init);
// If the type is an array (of an array, etc) then we must check whether the
// type of the innermost array's base type is abstract.
let mut ty = constant.ty;
while let crate::TypeInner::Array { base, .. } = module.types[ty].inner {
ty = base;
}
if !module.types[ty]
.inner
.scalar()
.is_some_and(|s| s.is_abstract())
{
module_tracer.constants_used.insert(handle);
module_tracer.types_used.insert(constant.ty);
module_tracer.global_expressions_used.insert(constant.init);
}
}
}

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

@@ -1072,8 +1072,13 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const()))
.transpose()?;
let (ty, initializer) =
self.type_and_init(v.name, v.init, explicit_ty, &mut ctx.as_override())?;
let (ty, initializer) = self.type_and_init(
v.name,
v.init,
explicit_ty,
false,
&mut ctx.as_override(),
)?;
let binding = if let Some(ref binding) = v.binding {
Some(crate::ResourceBinding {
@@ -1106,7 +1111,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
.transpose()?;
let (ty, init) =
self.type_and_init(c.name, Some(c.init), explicit_ty, &mut ectx)?;
self.type_and_init(c.name, Some(c.init), explicit_ty, true, &mut ectx)?;
let init = init.expect("Global const must have init");
let handle = ctx.module.constants.append(
@@ -1128,7 +1133,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let mut ectx = ctx.as_override();
let (ty, init) = self.type_and_init(o.name, o.init, explicit_ty, &mut ectx)?;
let (ty, init) =
self.type_and_init(o.name, o.init, explicit_ty, false, &mut ectx)?;
let id =
o.id.map(|id| self.const_u32(id, &mut ctx.as_const()))
@@ -1201,6 +1207,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
name: ast::Ident<'source>,
init: Option<Handle<ast::Expression<'source>>>,
explicit_ty: Option<Handle<crate::Type>>,
allow_abstract: bool,
ectx: &mut ExpressionContext<'source, '_, '_>,
) -> Result<(Handle<crate::Type>, Option<Handle<crate::Expression>>), Error<'source>> {
let ty;
@@ -1234,9 +1241,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
initializer = Some(init);
}
(Some(init), None) => {
let concretized = self.expression(init, ectx)?;
ty = ectx.register_type(concretized)?;
initializer = Some(concretized);
let mut init = self.expression_for_abstract(init, ectx)?;
if !allow_abstract {
init = ectx.concretize(init)?;
}
ty = ectx.register_type(init)?;
initializer = Some(init);
}
(None, Some(explicit_ty)) => {
ty = explicit_ty;
@@ -1481,7 +1491,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let mut ectx = ctx.as_expression(block, &mut emitter);
let (ty, initializer) =
self.type_and_init(v.name, v.init, explicit_ty, &mut ectx)?;
self.type_and_init(v.name, v.init, explicit_ty, false, &mut ectx)?;
let (const_initializer, initializer) = {
match initializer {
@@ -1544,6 +1554,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
c.name,
Some(c.init),
explicit_ty,
false,
&mut ectx.as_const(),
)?;
let init = init.expect("Local const must have init");

2
naga/tests/out/glsl/constructors.main.Compute.glsl
View File

@@ -9,7 +9,6 @@ struct Foo {
vec4 a;
int b;
};
const vec3 const2_ = vec3(0.0, 1.0, 2.0);
const mat2x2 const3_ = mat2x2(vec2(0.0, 1.0), vec2(2.0, 3.0));
const mat2x2 const4_[1] = mat2x2[1](mat2x2(vec2(0.0, 1.0), vec2(2.0, 3.0)));
const bool cz0_ = false;
@@ -20,7 +19,6 @@ const uvec2 cz4_ = uvec2(0u);
const mat2x2 cz5_ = mat2x2(0.0);
const Foo cz6_[3] = Foo[3](Foo(vec4(0.0), 0), Foo(vec4(0.0), 0), Foo(vec4(0.0), 0));
const Foo cz7_ = Foo(vec4(0.0), 0);
const int cp3_[4] = int[4](0, 1, 2, 3);
void main() {

14
naga/tests/out/hlsl/constructors.hlsl
View File

@@ -12,12 +12,6 @@ ret_Constructarray1_float2x2_ Constructarray1_float2x2_(float2x2 arg0) {
return ret;
}
typedef int ret_Constructarray4_int_[4];
ret_Constructarray4_int_ Constructarray4_int_(int arg0, int arg1, int arg2, int arg3) {
int ret[4] = { arg0, arg1, arg2, arg3 };
return ret;
}
bool ZeroValuebool() {
return (bool)0;
}
@@ -51,7 +45,6 @@ Foo ZeroValueFoo() {
return (Foo)0;
}
static const float3 const2_ = float3(0.0, 1.0, 2.0);
static const float2x2 const3_ = float2x2(float2(0.0, 1.0), float2(2.0, 3.0));
static const float2x2 const4_[1] = Constructarray1_float2x2_(float2x2(float2(0.0, 1.0), float2(2.0, 3.0)));
static const bool cz0_ = ZeroValuebool();
@@ -62,7 +55,6 @@ static const uint2 cz4_ = ZeroValueuint2();
static const float2x2 cz5_ = ZeroValuefloat2x2();
static const Foo cz6_[3] = ZeroValuearray3_Foo_();
static const Foo cz7_ = ZeroValueFoo();
static const int cp3_[4] = Constructarray4_int_(int(0), int(1), int(2), int(3));
Foo ConstructFoo(float4 arg0, int arg1) {
Foo ret = (Foo)0;
@@ -71,6 +63,12 @@ Foo ConstructFoo(float4 arg0, int arg1) {
return ret;
}
typedef int ret_Constructarray4_int_[4];
ret_Constructarray4_int_ Constructarray4_int_(int arg0, int arg1, int arg2, int arg3) {
int ret[4] = { arg0, arg1, arg2, arg3 };
return ret;
}
float2x3 ZeroValuefloat2x3() {
return (float2x3)0;
}

28
naga/tests/out/ir/const_assert.compact.ron
View File

@@ -1,36 +1,14 @@
(
types: [
(
name: None,
inner: Scalar((
kind: Sint,
width: 4,
)),
),
],
types: [],
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [
(
name: Some("x"),
ty: 0,
init: 0,
),
(
name: Some("y"),
ty: 0,
init: 1,
),
],
constants: [],
overrides: [],
global_variables: [],
global_expressions: [
Literal(I32(1)),
Literal(I32(2)),
],
global_expressions: [],
functions: [
(
name: Some("foo"),

28
naga/tests/out/ir/const_assert.ron
View File

@@ -1,36 +1,14 @@
(
types: [
(
name: None,
inner: Scalar((
kind: Sint,
width: 4,
)),
),
],
types: [],
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [
(
name: Some("x"),
ty: 0,
init: 0,
),
(
name: Some("y"),
ty: 0,
init: 1,
),
],
constants: [],
overrides: [],
global_variables: [],
global_expressions: [
Literal(I32(1)),
Literal(I32(2)),
],
global_expressions: [],
functions: [
(
name: Some("foo"),

59
naga/tests/out/ir/local-const.compact.ron
View File

@@ -39,52 +39,27 @@
),
constants: [
(
name: Some("ga"),
name: Some("gb"),
ty: 0,
init: 0,
),
(
name: Some("gb"),
ty: 0,
init: 1,
),
(
name: Some("gc"),
ty: 1,
init: 2,
init: 1,
),
(
name: Some("gd"),
ty: 2,
init: 3,
),
(
name: Some("ge"),
ty: 3,
init: 4,
),
(
name: Some("gf"),
ty: 2,
init: 5,
init: 2,
),
],
overrides: [],
global_variables: [],
global_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)),
],
functions: [
(
@@ -106,12 +81,22 @@
],
),
Literal(F32(2.0)),
Literal(I32(4)),
Constant(0),
Constant(1),
Constant(2),
Constant(3),
Constant(4),
Constant(5),
Literal(I32(4)),
Literal(I32(4)),
Literal(I32(4)),
Compose(
ty: 3,
components: [
10,
11,
12,
],
),
Literal(F32(2.0)),
],
named_expressions: {
0: "a",
@@ -124,14 +109,22 @@
7: "bg",
8: "cg",
9: "dg",
10: "eg",
11: "fg",
13: "eg",
14: "fg",
},
body: [
Emit((
start: 4,
end: 5,
)),
Emit((
start: 0,
end: 0,
)),
Emit((
start: 13,
end: 14,
)),
Return(
value: None,
),

59
naga/tests/out/ir/local-const.ron
View File

@@ -39,52 +39,27 @@
),
constants: [
(
name: Some("ga"),
name: Some("gb"),
ty: 0,
init: 0,
),
(
name: Some("gb"),
ty: 0,
init: 1,
),
(
name: Some("gc"),
ty: 1,
init: 2,
init: 1,
),
(
name: Some("gd"),
ty: 2,
init: 3,
),
(
name: Some("ge"),
ty: 3,
init: 4,
),
(
name: Some("gf"),
ty: 2,
init: 5,
init: 2,
),
],
overrides: [],
global_variables: [],
global_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)),
],
functions: [
(
@@ -106,12 +81,22 @@
],
),
Literal(F32(2.0)),
Literal(I32(4)),
Constant(0),
Constant(1),
Constant(2),
Constant(3),
Constant(4),
Constant(5),
Literal(I32(4)),
Literal(I32(4)),
Literal(I32(4)),
Compose(
ty: 3,
components: [
10,
11,
12,
],
),
Literal(F32(2.0)),
],
named_expressions: {
0: "a",
@@ -124,14 +109,22 @@
7: "bg",
8: "cg",
9: "dg",
10: "eg",
11: "fg",
13: "eg",
14: "fg",
},
body: [
Emit((
start: 4,
end: 5,
)),
Emit((
start: 0,
end: 0,
)),
Emit((
start: 13,
end: 14,
)),
Return(
value: None,
),

8
naga/tests/out/msl/abstract-types-const.msl
View File

@@ -23,20 +23,12 @@ constant metal::float2x2 xmfpafaiaiai = metal::float2x2(metal::float2(1.0, 2.0),
constant metal::float2x2 xmfpaiafaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 xmfpaiaiafai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 xmfpaiaiaiaf = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 imfpaiaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 imfpafaiaiai = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::float2x2 imfpafafafaf = metal::float2x2(metal::float2(1.0, 2.0), metal::float2(3.0, 4.0));
constant metal::int2 ivispai = metal::int2(1);
constant metal::float2 ivfspaf = metal::float2(1.0);
constant metal::int2 ivis_ai = metal::int2(1);
constant metal::uint2 ivus_ai = metal::uint2(1u);
constant metal::float2 ivfs_ai = metal::float2(1.0);
constant metal::float2 ivfs_af = metal::float2(1.0);
constant type_7 iafafaf = type_7 {1.0, 2.0};
constant type_7 iafaiai = type_7 {1.0, 2.0};
constant type_7 iafpafaf = type_7 {1.0, 2.0};
constant type_7 iafpaiaf = type_7 {1.0, 2.0};
constant type_7 iafpafai = type_7 {1.0, 2.0};
constant type_7 xafpafaf = type_7 {1.0, 2.0};
constant S s_f_i_u = S {1.0, 1, 1u};
constant S s_f_iai = S {1.0, 1, 1u};

1
naga/tests/out/msl/abstract-types-operators.msl
View File

@@ -37,7 +37,6 @@ constant int shr_iai_u = 0;
constant uint shr_uaiai = 0u;
constant uint shr_uai_u = 0u;
constant int wgpu_4492_ = -2147483648;
constant int wgpu_4492_2_ = -2147483648;
void runtime_values(
) {

10
naga/tests/out/msl/constructors.msl
View File

@@ -8,27 +8,25 @@ struct Foo {
metal::float4 a;
int b;
};
struct type_6 {
struct type_5 {
metal::float2x2 inner[1];
};
struct type_10 {
struct type_9 {
Foo inner[3];
};
struct type_11 {
int inner[4];
};
constant metal::float3 const2_ = metal::float3(0.0, 1.0, 2.0);
constant metal::float2x2 const3_ = metal::float2x2(metal::float2(0.0, 1.0), metal::float2(2.0, 3.0));
constant type_6 const4_ = type_6 {metal::float2x2(metal::float2(0.0, 1.0), metal::float2(2.0, 3.0))};
constant type_5 const4_ = type_5 {metal::float2x2(metal::float2(0.0, 1.0), metal::float2(2.0, 3.0))};
constant bool cz0_ = bool {};
constant int cz1_ = int {};
constant uint cz2_ = uint {};
constant float cz3_ = float {};
constant metal::uint2 cz4_ = metal::uint2 {};
constant metal::float2x2 cz5_ = metal::float2x2 {};
constant type_10 cz6_ = type_10 {};
constant type_9 cz6_ = type_9 {};
constant Foo cz7_ = Foo {};
constant type_11 cp3_ = type_11 {0, 1, 2, 3};
kernel void main_(
) {

18
naga/tests/out/spv/abstract-types-const.spvasm
View File

@@ -16,11 +16,11 @@ OpMemberDecorate %12 2 Offset 8
%5 = OpTypeFloat 32
%6 = OpTypeVector %5 2
%7 = OpTypeMatrix %6 2
%9 = OpTypeInt 32 1
%8 = OpTypeVector %9 2
%8 = OpTypeInt 32 1
%9 = OpTypeVector %8 2
%11 = OpConstant %3 2
%10 = OpTypeArray %5 %11
%12 = OpTypeStruct %5 %9 %3
%12 = OpTypeStruct %5 %8 %3
%13 = OpTypeVector %5 3
%14 = OpConstant %3 42
%15 = OpConstant %3 43
@@ -35,12 +35,12 @@ OpMemberDecorate %12 2 Offset 8
%24 = OpConstant %5 4.0
%25 = OpConstantComposite %6 %23 %24
%26 = OpConstantComposite %7 %22 %25
%27 = OpConstant %9 1
%28 = OpConstantComposite %8 %27 %27
%29 = OpConstantComposite %6 %20 %20
%30 = OpConstant %3 1
%31 = OpConstantComposite %4 %30 %30
%27 = OpConstant %8 1
%28 = OpConstantComposite %9 %27 %27
%29 = OpConstant %3 1
%30 = OpConstantComposite %4 %29 %29
%31 = OpConstantComposite %6 %20 %20
%32 = OpConstantComposite %10 %20 %21
%33 = OpConstantComposite %12 %20 %27 %30
%33 = OpConstantComposite %12 %20 %27 %29
%34 = OpConstantComposite %13 %20 %21 %23
%35 = OpConstantComposite %6 %21 %23

127
naga/tests/out/spv/constructors.spvasm
View File

@@ -1,84 +1,83 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 68
; Bound: 67
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %44 "main"
OpExecutionMode %44 LocalSize 1 1 1
OpEntryPoint GLCompute %38 "main"
OpExecutionMode %38 LocalSize 1 1 1
OpMemberDecorate %6 0 Offset 0
OpMemberDecorate %6 1 Offset 16
OpDecorate %10 ArrayStride 16
OpDecorate %15 ArrayStride 32
OpDecorate %9 ArrayStride 16
OpDecorate %14 ArrayStride 32
OpDecorate %17 ArrayStride 4
%2 = OpTypeVoid
%3 = OpTypeFloat 32
%4 = OpTypeVector %3 4
%5 = OpTypeInt 32 1
%6 = OpTypeStruct %4 %5
%7 = OpTypeVector %3 3
%9 = OpTypeVector %3 2
%8 = OpTypeMatrix %9 2
%12 = OpTypeInt 32 0
%11 = OpConstant %12 1
%10 = OpTypeArray %8 %11
%13 = OpTypeBool
%14 = OpTypeVector %12 2
%16 = OpConstant %12 3
%15 = OpTypeArray %6 %16
%18 = OpConstant %12 4
%8 = OpTypeVector %3 2
%7 = OpTypeMatrix %8 2
%11 = OpTypeInt 32 0
%10 = OpConstant %11 1
%9 = OpTypeArray %7 %10
%12 = OpTypeBool
%13 = OpTypeVector %11 2
%15 = OpConstant %11 3
%14 = OpTypeArray %6 %15
%16 = OpTypeMatrix %4 4
%18 = OpConstant %11 4
%17 = OpTypeArray %5 %18
%19 = OpTypeMatrix %4 4
%20 = OpTypeMatrix %7 2
%20 = OpTypeVector %3 3
%19 = OpTypeMatrix %20 2
%21 = OpConstant %3 0.0
%22 = OpConstant %3 1.0
%23 = OpConstant %3 2.0
%24 = OpConstantComposite %7 %21 %22 %23
%25 = OpConstant %3 3.0
%26 = OpConstantComposite %9 %21 %22
%27 = OpConstantComposite %9 %23 %25
%28 = OpConstantComposite %8 %26 %27
%29 = OpConstantComposite %10 %28
%30 = OpConstantNull %13
%31 = OpConstantNull %5
%32 = OpConstantNull %12
%33 = OpConstantNull %3
%34 = OpConstantNull %14
%35 = OpConstantNull %8
%36 = OpConstantNull %15
%37 = OpConstantNull %6
%38 = OpConstant %5 0
%39 = OpConstant %5 1
%40 = OpConstant %5 2
%41 = OpConstant %5 3
%42 = OpConstantComposite %17 %38 %39 %40 %41
%45 = OpTypeFunction %2
%46 = OpConstantComposite %4 %22 %22 %22 %22
%47 = OpConstantComposite %6 %46 %39
%48 = OpConstantComposite %9 %22 %21
%49 = OpConstantComposite %8 %48 %26
%50 = OpConstantComposite %4 %22 %21 %21 %21
%51 = OpConstantComposite %4 %21 %22 %21 %21
%52 = OpConstantComposite %4 %21 %21 %22 %21
%53 = OpConstantComposite %4 %21 %21 %21 %22
%54 = OpConstantComposite %19 %50 %51 %52 %53
%55 = OpConstant %12 0
%56 = OpConstantComposite %14 %55 %55
%57 = OpConstantComposite %9 %21 %21
%58 = OpConstantComposite %8 %57 %57
%59 = OpConstantFalse %13
%60 = OpConstantComposite %14 %55 %55
%61 = OpConstantComposite %7 %21 %21 %21
%62 = OpConstantComposite %20 %61 %61
%63 = OpConstantNull %20
%65 = OpTypePointer Function %6
%66 = OpConstantNull %6
%44 = OpFunction %2 None %45
%43 = OpLabel
%64 = OpVariable %65 Function %66
OpBranch %67
%67 = OpLabel
OpStore %64 %47
%24 = OpConstant %3 3.0
%25 = OpConstantComposite %8 %21 %22
%26 = OpConstantComposite %8 %23 %24
%27 = OpConstantComposite %7 %25 %26
%28 = OpConstantComposite %9 %27
%29 = OpConstantNull %12
%30 = OpConstantNull %5
%31 = OpConstantNull %11
%32 = OpConstantNull %3
%33 = OpConstantNull %13
%34 = OpConstantNull %7
%35 = OpConstantNull %14
%36 = OpConstantNull %6
%39 = OpTypeFunction %2
%40 = OpConstantComposite %4 %22 %22 %22 %22
%41 = OpConstant %5 1
%42 = OpConstantComposite %6 %40 %41
%43 = OpConstantComposite %8 %22 %21
%44 = OpConstantComposite %7 %43 %25
%45 = OpConstantComposite %4 %22 %21 %21 %21
%46 = OpConstantComposite %4 %21 %22 %21 %21
%47 = OpConstantComposite %4 %21 %21 %22 %21
%48 = OpConstantComposite %4 %21 %21 %21 %22
%49 = OpConstantComposite %16 %45 %46 %47 %48
%50 = OpConstant %11 0
%51 = OpConstantComposite %13 %50 %50
%52 = OpConstantComposite %8 %21 %21
%53 = OpConstantComposite %7 %52 %52
%54 = OpConstant %5 0
%55 = OpConstant %5 2
%56 = OpConstant %5 3
%57 = OpConstantComposite %17 %54 %41 %55 %56
%58 = OpConstantFalse %12
%59 = OpConstantComposite %13 %50 %50
%60 = OpConstantComposite %20 %21 %21 %21
%61 = OpConstantComposite %19 %60 %60
%62 = OpConstantNull %19
%64 = OpTypePointer Function %6
%65 = OpConstantNull %6
%38 = OpFunction %2 None %39
%37 = OpLabel
%63 = OpVariable %64 Function %65
OpBranch %66
%66 = OpLabel
OpStore %63 %42
OpReturn
OpFunctionEnd

8
naga/tests/out/wgsl/abstract-types-const.wgsl
View File

@@ -15,20 +15,12 @@ const xmfpafaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4
const xmfpaiafaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4f));
const xmfpaiaiafai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4f));
const xmfpaiaiaiaf: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4f));
const imfpaiaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4f));
const imfpafaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4f));
const imfpafafafaf: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4f));
const ivispai: vec2<i32> = vec2(1i);
const ivfspaf: vec2<f32> = vec2(1f);
const ivis_ai: vec2<i32> = vec2(1i);
const ivus_ai: vec2<u32> = vec2(1u);
const ivfs_ai: vec2<f32> = vec2(1f);
const ivfs_af: vec2<f32> = vec2(1f);
const iafafaf: array<f32, 2> = array<f32, 2>(1f, 2f);
const iafaiai: array<f32, 2> = array<f32, 2>(1f, 2f);
const iafpafaf: array<f32, 2> = array<f32, 2>(1f, 2f);
const iafpaiaf: array<f32, 2> = array<f32, 2>(1f, 2f);
const iafpafai: array<f32, 2> = array<f32, 2>(1f, 2f);
const xafpafaf: array<f32, 2> = array<f32, 2>(1f, 2f);
const s_f_i_u: S = S(1f, 1i, 1u);
const s_f_iai: S = S(1f, 1i, 1u);

1
naga/tests/out/wgsl/abstract-types-operators.wgsl
View File

@@ -28,7 +28,6 @@ const shr_iai_u: i32 = 0i;
const shr_uaiai: u32 = 0u;
const shr_uai_u: u32 = 0u;
const wgpu_4492_: i32 = i32(-2147483648);
const wgpu_4492_2_: i32 = i32(-2147483648);
var<workgroup> a: array<u32, 64>;

3
naga/tests/out/wgsl/const_assert.wgsl
View File

@@ -1,6 +1,3 @@
const x: i32 = 1i;
const y: i32 = 2i;
fn foo() {
return;
}

2
naga/tests/out/wgsl/constructors.wgsl
View File

@@ -3,7 +3,6 @@ struct Foo {
b: i32,
}
const const2_: vec3<f32> = vec3<f32>(0f, 1f, 2f);
const const3_: mat2x2<f32> = mat2x2<f32>(vec2<f32>(0f, 1f), vec2<f32>(2f, 3f));
const const4_: array<mat2x2<f32>, 1> = array<mat2x2<f32>, 1>(mat2x2<f32>(vec2<f32>(0f, 1f), vec2<f32>(2f, 3f)));
const cz0_: bool = bool();
@@ -14,7 +13,6 @@ const cz4_: vec2<u32> = vec2<u32>();
const cz5_: mat2x2<f32> = mat2x2<f32>();
const cz6_: array<Foo, 3> = array<Foo, 3>();
const cz7_: Foo = Foo();
const cp3_: array<i32, 4> = array<i32, 4>(0i, 1i, 2i, 3i);
@compute @workgroup_size(1, 1, 1)
fn main() {

4
naga/tests/out/wgsl/local-const.wgsl
View File

@@ -1,12 +1,10 @@
const ga: i32 = 4i;
const gb: i32 = 4i;
const gc: u32 = 4u;
const gd: f32 = 4f;
const ge: vec3<i32> = vec3<i32>(4i, 4i, 4i);
const gf: f32 = 2f;
fn const_in_fn() {
const e = vec3<i32>(4i, 4i, 4i);
const eg = vec3<i32>(4i, 4i, 4i);
return;
}