From 3912120977b819019ade77efb642346b56bd34ae Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Thu, 12 Jun 2025 16:43:54 -0400 Subject: [PATCH 1/3] Add force-loop-bounding support to naga-cli --- naga-cli/src/bin/naga.rs | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/naga-cli/src/bin/naga.rs b/naga-cli/src/bin/naga.rs index 6f95e429f68..97cb05eb66c 100644 --- a/naga-cli/src/bin/naga.rs +++ b/naga-cli/src/bin/naga.rs @@ -75,6 +75,10 @@ struct Args { #[argh(switch)] keep_coordinate_space: bool, + /// force loop bounding if the backend supports it. + #[argh(switch)] + force_loop_bounding: bool, + /// in dot output, include only the control flow graph #[argh(switch)] dot_cfg_only: bool, @@ -427,6 +431,10 @@ fn run() -> anyhow::Result<()> { block_ctx_dump_prefix: args.block_ctx_dir.clone().map(std::path::PathBuf::from), }; + params.spv_out.force_loop_bounding = args.force_loop_bounding; + params.msl.force_loop_bounding = args.force_loop_bounding; + params.hlsl.force_loop_bounding = args.force_loop_bounding; + params.entry_point.clone_from(&args.entry_point); if let Some(ref version) = args.profile { params.glsl.version = version.0; @@ -706,7 +714,7 @@ fn write_output( let file = fs::File::create(output_path)?; bincode::serialize_into(file, module)?; } - "metal" => { + "metal" | "msl" => { use naga::back::msl; let mut options = params.msl.clone(); From 165abcbbcdd61ab93aaedef8e1f8d372a958ee3b Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Thu, 12 Jun 2025 16:43:54 -0400 Subject: [PATCH 2/3] Test outputs --- naga/tests/in/spv/test.slang | 19 ++ naga/tests/in/spv/test.spv | Bin 0 -> 1012 bytes naga/tests/in/spv/test.spvasm | 77 +++++ naga/tests/in/spv/test.toml | 10 + naga/tests/in/wgsl/test.toml | 10 + naga/tests/in/wgsl/test.wgsl | 11 + .../tests/out/glsl/spv-test.main.Compute.glsl | 48 +++ .../out/glsl/wgsl-test.main.Compute.glsl | 33 ++ naga/tests/out/hlsl/spv-test.hlsl | 41 +++ naga/tests/out/hlsl/spv-test.ron | 12 + naga/tests/out/hlsl/wgsl-test.hlsl | 27 ++ naga/tests/out/hlsl/wgsl-test.ron | 12 + naga/tests/out/ir/spv-test.compact.ron | 273 ++++++++++++++++ naga/tests/out/ir/spv-test.ron | 298 ++++++++++++++++++ naga/tests/out/ir/wgsl-test.compact.ron | 157 +++++++++ naga/tests/out/ir/wgsl-test.ron | 157 +++++++++ naga/tests/out/msl/spv-test.metal | 57 ++++ naga/tests/out/msl/spv-test.msl | 57 ++++ naga/tests/out/msl/wgsl-test.metal | 37 +++ naga/tests/out/msl/wgsl-test.msl | 37 +++ naga/tests/out/spv/wgsl-test.spvasm | 87 +++++ naga/tests/out/wgsl/spv-test.wgsl | 43 +++ naga/tests/out/wgsl/wgsl-test.wgsl | 25 ++ 23 files changed, 1528 insertions(+) create mode 100644 naga/tests/in/spv/test.slang create mode 100644 naga/tests/in/spv/test.spv create mode 100644 naga/tests/in/spv/test.spvasm create mode 100644 naga/tests/in/spv/test.toml create mode 100644 naga/tests/in/wgsl/test.toml create mode 100644 naga/tests/in/wgsl/test.wgsl create mode 100644 naga/tests/out/glsl/spv-test.main.Compute.glsl create mode 100644 naga/tests/out/glsl/wgsl-test.main.Compute.glsl create mode 100644 naga/tests/out/hlsl/spv-test.hlsl create mode 100644 naga/tests/out/hlsl/spv-test.ron create mode 100644 naga/tests/out/hlsl/wgsl-test.hlsl create mode 100644 naga/tests/out/hlsl/wgsl-test.ron create mode 100644 naga/tests/out/ir/spv-test.compact.ron create mode 100644 naga/tests/out/ir/spv-test.ron create mode 100644 naga/tests/out/ir/wgsl-test.compact.ron create mode 100644 naga/tests/out/ir/wgsl-test.ron create mode 100644 naga/tests/out/msl/spv-test.metal create mode 100644 naga/tests/out/msl/spv-test.msl create mode 100644 naga/tests/out/msl/wgsl-test.metal create mode 100644 naga/tests/out/msl/wgsl-test.msl create mode 100644 naga/tests/out/spv/wgsl-test.spvasm create mode 100644 naga/tests/out/wgsl/spv-test.wgsl create mode 100644 naga/tests/out/wgsl/wgsl-test.wgsl diff --git a/naga/tests/in/spv/test.slang b/naga/tests/in/spv/test.slang new file mode 100644 index 00000000000..b027a173c85 --- /dev/null +++ b/naga/tests/in/spv/test.slang @@ -0,0 +1,19 @@ +[vk::binding(0, 0)] +RWStructuredBuffer data; + +[shader("compute")] +void main() +{ + for (uint32_t i = 0; i < 4; i++) + { + if (data[i] == 1) + { + continue; + } + if (data[i] == 2) + { + break; + } + data[i] = i * 2; // Example operation: double each element + } +} diff --git a/naga/tests/in/spv/test.spv b/naga/tests/in/spv/test.spv new file mode 100644 index 0000000000000000000000000000000000000000..5c35682a467ea7f61aacc14a7b09d4c06e9b1e97 GIT binary patch literal 1012 zcmYk4OHUM05QU2whT);SK~TXSP%&|DNDL&nfrJf#1eTK)21CLlnaAD*zlN26PDKrg z-`CR@<|b#V&aJ9bRkx>FS+7TJC?Y<@2EU$}sDwA;F%d@x$Ibn{!)B*@)6T!OnqPb8 z=dJc5J-y639cw0|swwKGVkBy%f0bWcdtU+5+^ahSv+DO>Gte9^bk6c_PI|%BNIiCm55q@*b{8e5viBa@{Rt!1;U4piH465`*GfJ|)IAzh zs;R#~O&yy7&&QIF_G!#hGd6HPj_)kixnl;L**f_%;K|_6v2yiS=xLq)D!#MJ&6~uU z=UwVQg`ae-nFj9mxy;wM<`Y|_Zv!p%x+l|b@&`MMm1}>}U9CCbe&)=B1@H^q#nQ(V zwJ)&NYJW4`k^#;?e0whgibUCyuy0^qHj$88d#@wM*GT*?XZFGwgS?Q z>TOZqf4hF{uI%@Q*edY7?3ZCZ_A}N18Bkx}e(oXnKI!@&v1iitKjoe8uI^i`M_ps@ zfOqQX5ANapE%Q0=d*J-)Y@=fj?*`{^p6_M8-@| data: array; + +@compute @workgroup_size(1) +fn main() +{ + for (var i = 0u; i < 4; i++) + { + data[i] = i * 2; // Example operation: double each element + } +} diff --git a/naga/tests/out/glsl/spv-test.main.Compute.glsl b/naga/tests/out/glsl/spv-test.main.Compute.glsl new file mode 100644 index 00000000000..3ac8a7c9146 --- /dev/null +++ b/naga/tests/out/glsl/spv-test.main.Compute.glsl @@ -0,0 +1,48 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(std430) buffer RWStructuredBuffer_block_0Compute { + uint member[]; +} _group_0_binding_0_cs; + + +void main_1() { + uint phi_19_ = 0u; + phi_19_ = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_19_ = (phi_19_ + 1u); + } + loop_init = false; + uint _e7 = phi_19_; + bool should_continue = false; + do { + if ((_e7 < 4u)) { + } else { + break; + } + uint _e11 = _group_0_binding_0_cs.member[_e7]; + if ((_e11 == 1u)) { + break; + } + uint _e13 = _group_0_binding_0_cs.member[_e7]; + if ((_e13 == 2u)) { + break; + } + _group_0_binding_0_cs.member[_e7] = (_e7 * 2u); + break; + } while(false); + continue; + } + return; +} + +void main() { + main_1(); +} + diff --git a/naga/tests/out/glsl/wgsl-test.main.Compute.glsl b/naga/tests/out/glsl/wgsl-test.main.Compute.glsl new file mode 100644 index 00000000000..40c04f1cd98 --- /dev/null +++ b/naga/tests/out/glsl/wgsl-test.main.Compute.glsl @@ -0,0 +1,33 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(std430) buffer type_1_block_0Compute { uint _group_0_binding_0_cs[]; }; + + +void main() { + uint i = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + uint _e12 = i; + i = (_e12 + 1u); + } + loop_init = false; + uint _e2 = i; + if ((_e2 < 4u)) { + } else { + break; + } + { + uint _e6 = i; + uint _e8 = i; + _group_0_binding_0_cs[_e6] = (_e8 * 2u); + } + } + return; +} + diff --git a/naga/tests/out/hlsl/spv-test.hlsl b/naga/tests/out/hlsl/spv-test.hlsl new file mode 100644 index 00000000000..189e695daa9 --- /dev/null +++ b/naga/tests/out/hlsl/spv-test.hlsl @@ -0,0 +1,41 @@ +RWByteAddressBuffer data : register(u0); + +void main_1() +{ + uint phi_19_ = (uint)0; + + phi_19_ = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_19_ = (phi_19_ + 1u); + } + loop_init = false; + uint _e7 = phi_19_; + bool should_continue = false; + do { + if ((_e7 < 4u)) { + } else { + break; + } + uint _e11 = asuint(data.Load(_e7*4+0)); + if ((_e11 == 1u)) { + break; + } + uint _e13 = asuint(data.Load(_e7*4+0)); + if ((_e13 == 2u)) { + break; + } + data.Store(_e7*4+0, asuint((_e7 * 2u))); + break; + } while(false); + continue; + } + return; +} + +[numthreads(1, 1, 1)] +void main() +{ + main_1(); +} diff --git a/naga/tests/out/hlsl/spv-test.ron b/naga/tests/out/hlsl/spv-test.ron new file mode 100644 index 00000000000..a07b03300b1 --- /dev/null +++ b/naga/tests/out/hlsl/spv-test.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/hlsl/wgsl-test.hlsl b/naga/tests/out/hlsl/wgsl-test.hlsl new file mode 100644 index 00000000000..f3815df26be --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-test.hlsl @@ -0,0 +1,27 @@ +RWByteAddressBuffer data : register(u0); + +[numthreads(1, 1, 1)] +void main() +{ + uint i = 0u; + + bool loop_init = true; + while(true) { + if (!loop_init) { + uint _e12 = i; + i = (_e12 + 1u); + } + loop_init = false; + uint _e2 = i; + if ((_e2 < 4u)) { + } else { + break; + } + { + uint _e6 = i; + uint _e8 = i; + data.Store(_e6*4, asuint((_e8 * 2u))); + } + } + return; +} diff --git a/naga/tests/out/hlsl/wgsl-test.ron b/naga/tests/out/hlsl/wgsl-test.ron new file mode 100644 index 00000000000..a07b03300b1 --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-test.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/ir/spv-test.compact.ron b/naga/tests/out/ir/spv-test.compact.ron new file mode 100644 index 00000000000..b072cf16def --- /dev/null +++ b/naga/tests/out/ir/spv-test.compact.ron @@ -0,0 +1,273 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Scalar(( + kind: Sint, + width: 4, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ( + name: Some("RWStructuredBuffer"), + inner: Struct( + members: [ + ( + name: None, + ty: 2, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [ + ( + name: None, + ty: 0, + init: 0, + ), + ( + name: None, + ty: 0, + init: 1, + ), + ( + name: None, + ty: 1, + init: 2, + ), + ( + name: None, + ty: 0, + init: 3, + ), + ( + name: None, + ty: 0, + init: 4, + ), + ], + overrides: [], + global_variables: [ + ( + name: Some("data"), + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 3, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(U32(4)), + Literal(I32(0)), + Literal(U32(1)), + Literal(U32(2)), + ], + functions: [ + ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_19"), + ty: 0, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + Constant(0), + Constant(4), + Constant(1), + Constant(3), + Constant(2), + LocalVariable(0), + Load( + pointer: 6, + ), + Binary( + op: Less, + left: 7, + right: 3, + ), + AccessIndex( + base: 0, + index: 0, + ), + Access( + base: 9, + index: 7, + ), + Load( + pointer: 10, + ), + Binary( + op: Equal, + left: 11, + right: 4, + ), + Load( + pointer: 10, + ), + Binary( + op: Equal, + left: 13, + right: 2, + ), + Binary( + op: Multiply, + left: 7, + right: 2, + ), + Binary( + op: Add, + left: 7, + right: 4, + ), + LocalVariable(0), + ], + named_expressions: {}, + body: [ + Store( + pointer: 17, + value: 1, + ), + Loop( + body: [ + Emit(( + start: 7, + end: 8, + )), + Switch( + selector: 5, + cases: [ + ( + value: Default, + body: [ + Emit(( + start: 8, + end: 9, + )), + If( + condition: 8, + accept: [], + reject: [ + Break, + ], + ), + Emit(( + start: 9, + end: 13, + )), + If( + condition: 12, + accept: [ + Break, + ], + reject: [], + ), + Emit(( + start: 13, + end: 15, + )), + If( + condition: 14, + accept: [ + Break, + ], + reject: [], + ), + Emit(( + start: 15, + end: 16, + )), + Store( + pointer: 10, + value: 15, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 16, + end: 17, + )), + Continue, + ], + continuing: [ + Store( + pointer: 17, + value: 16, + ), + ], + break_if: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main_wrap"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [], + result: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/spv-test.ron b/naga/tests/out/ir/spv-test.ron new file mode 100644 index 00000000000..c35ee52bb93 --- /dev/null +++ b/naga/tests/out/ir/spv-test.ron @@ -0,0 +1,298 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Scalar(( + kind: Bool, + width: 1, + )), + ), + ( + name: None, + inner: Scalar(( + kind: Sint, + width: 4, + )), + ), + ( + name: None, + inner: Pointer( + base: 0, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ( + name: Some("RWStructuredBuffer"), + inner: Struct( + members: [ + ( + name: None, + ty: 4, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 5, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [ + ( + name: None, + ty: 0, + init: 0, + ), + ( + name: None, + ty: 0, + init: 1, + ), + ( + name: None, + ty: 2, + init: 2, + ), + ( + name: None, + ty: 0, + init: 3, + ), + ( + name: None, + ty: 0, + init: 4, + ), + ], + overrides: [], + global_variables: [ + ( + name: Some("data"), + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 5, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(U32(4)), + Literal(I32(0)), + Literal(U32(1)), + Literal(U32(2)), + ], + functions: [ + ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_19"), + ty: 0, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + Constant(0), + Constant(4), + Constant(1), + Constant(3), + Constant(2), + LocalVariable(0), + Load( + pointer: 6, + ), + Binary( + op: Less, + left: 7, + right: 3, + ), + AccessIndex( + base: 0, + index: 0, + ), + Access( + base: 9, + index: 7, + ), + Load( + pointer: 10, + ), + Binary( + op: Equal, + left: 11, + right: 4, + ), + Load( + pointer: 10, + ), + Binary( + op: Equal, + left: 13, + right: 2, + ), + Binary( + op: Multiply, + left: 7, + right: 2, + ), + Binary( + op: Add, + left: 7, + right: 4, + ), + LocalVariable(0), + ], + named_expressions: {}, + body: [ + Store( + pointer: 17, + value: 1, + ), + Loop( + body: [ + Emit(( + start: 7, + end: 8, + )), + Switch( + selector: 5, + cases: [ + ( + value: Default, + body: [ + Emit(( + start: 8, + end: 9, + )), + If( + condition: 8, + accept: [], + reject: [ + Break, + ], + ), + Emit(( + start: 9, + end: 13, + )), + If( + condition: 12, + accept: [ + Break, + ], + reject: [], + ), + Emit(( + start: 13, + end: 15, + )), + If( + condition: 14, + accept: [ + Break, + ], + reject: [], + ), + Emit(( + start: 15, + end: 16, + )), + Store( + pointer: 10, + value: 15, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 16, + end: 17, + )), + Continue, + ], + continuing: [ + Store( + pointer: 17, + value: 16, + ), + ], + break_if: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main_wrap"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [], + result: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/wgsl-test.compact.ron b/naga/tests/out/ir/wgsl-test.compact.ron new file mode 100644 index 00000000000..681d476056d --- /dev/null +++ b/naga/tests/out/ir/wgsl-test.compact.ron @@ -0,0 +1,157 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("data"), + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 1, + init: None, + ), + ], + global_expressions: [], + functions: [], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("i"), + ty: 0, + init: Some(0), + ), + ], + expressions: [ + Literal(U32(0)), + LocalVariable(0), + Load( + pointer: 1, + ), + Literal(U32(4)), + Binary( + op: Less, + left: 2, + right: 3, + ), + GlobalVariable(0), + Load( + pointer: 1, + ), + Access( + base: 5, + index: 6, + ), + Load( + pointer: 1, + ), + Literal(U32(2)), + Binary( + op: Multiply, + left: 8, + right: 9, + ), + Literal(U32(1)), + Load( + pointer: 1, + ), + Binary( + op: Add, + left: 12, + right: 11, + ), + ], + named_expressions: {}, + body: [ + Loop( + body: [ + Emit(( + start: 2, + end: 3, + )), + Emit(( + start: 4, + end: 5, + )), + If( + condition: 4, + accept: [], + reject: [ + Break, + ], + ), + Block([ + Emit(( + start: 6, + end: 9, + )), + Emit(( + start: 10, + end: 11, + )), + Store( + pointer: 7, + value: 10, + ), + ]), + ], + continuing: [ + Emit(( + start: 12, + end: 14, + )), + Store( + pointer: 1, + value: 13, + ), + ], + break_if: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/wgsl-test.ron b/naga/tests/out/ir/wgsl-test.ron new file mode 100644 index 00000000000..681d476056d --- /dev/null +++ b/naga/tests/out/ir/wgsl-test.ron @@ -0,0 +1,157 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("data"), + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 1, + init: None, + ), + ], + global_expressions: [], + functions: [], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("i"), + ty: 0, + init: Some(0), + ), + ], + expressions: [ + Literal(U32(0)), + LocalVariable(0), + Load( + pointer: 1, + ), + Literal(U32(4)), + Binary( + op: Less, + left: 2, + right: 3, + ), + GlobalVariable(0), + Load( + pointer: 1, + ), + Access( + base: 5, + index: 6, + ), + Load( + pointer: 1, + ), + Literal(U32(2)), + Binary( + op: Multiply, + left: 8, + right: 9, + ), + Literal(U32(1)), + Load( + pointer: 1, + ), + Binary( + op: Add, + left: 12, + right: 11, + ), + ], + named_expressions: {}, + body: [ + Loop( + body: [ + Emit(( + start: 2, + end: 3, + )), + Emit(( + start: 4, + end: 5, + )), + If( + condition: 4, + accept: [], + reject: [ + Break, + ], + ), + Block([ + Emit(( + start: 6, + end: 9, + )), + Emit(( + start: 10, + end: 11, + )), + Store( + pointer: 7, + value: 10, + ), + ]), + ], + continuing: [ + Emit(( + start: 12, + end: 14, + )), + Store( + pointer: 1, + value: 13, + ), + ], + break_if: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/msl/spv-test.metal b/naga/tests/out/msl/spv-test.metal new file mode 100644 index 00000000000..802db4e5f65 --- /dev/null +++ b/naga/tests/out/msl/spv-test.metal @@ -0,0 +1,57 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct _mslBufferSizes { + uint size0; +}; + +typedef uint type_4[1]; +struct RWStructuredBuffer { + type_4 member; +}; + +void main_1( + device RWStructuredBuffer& data, + constant _mslBufferSizes& _buffer_sizes +) { + uint phi_19_ = {}; + phi_19_ = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_19_ = phi_19_ + 1u; + } + loop_init = false; + uint _e7 = phi_19_; + switch(0) { + default: { + if (_e7 < 4u) { + } else { + break; + } + uint _e11 = data.member[_e7]; + if (_e11 == 1u) { + break; + } + uint _e13 = data.member[_e7]; + if (_e13 == 2u) { + break; + } + data.member[_e7] = _e7 * 2u; + break; + } + } + continue; + } + return; +} + +kernel void main_( + device RWStructuredBuffer& data [[user(fake0)]] +, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] +) { + main_1(data, _buffer_sizes); +} diff --git a/naga/tests/out/msl/spv-test.msl b/naga/tests/out/msl/spv-test.msl new file mode 100644 index 00000000000..6a38aacadaa --- /dev/null +++ b/naga/tests/out/msl/spv-test.msl @@ -0,0 +1,57 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct _mslBufferSizes { + uint size0; +}; + +typedef uint type_2[1]; +struct RWStructuredBuffer { + type_2 member; +}; + +void main_1( + device RWStructuredBuffer& data, + constant _mslBufferSizes& _buffer_sizes +) { + uint phi_19_ = {}; + phi_19_ = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_19_ = phi_19_ + 1u; + } + loop_init = false; + uint _e7 = phi_19_; + switch(0) { + default: { + if (_e7 < 4u) { + } else { + break; + } + uint _e11 = data.member[_e7]; + if (_e11 == 1u) { + break; + } + uint _e13 = data.member[_e7]; + if (_e13 == 2u) { + break; + } + data.member[_e7] = _e7 * 2u; + break; + } + } + continue; + } + return; +} + +kernel void main_( + device RWStructuredBuffer& data [[user(fake0)]] +, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] +) { + main_1(data, _buffer_sizes); +} diff --git a/naga/tests/out/msl/wgsl-test.metal b/naga/tests/out/msl/wgsl-test.metal new file mode 100644 index 00000000000..1fd8fb498e3 --- /dev/null +++ b/naga/tests/out/msl/wgsl-test.metal @@ -0,0 +1,37 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct _mslBufferSizes { + uint size0; +}; + +typedef uint type_1[1]; + +kernel void main_( + device type_1& data [[user(fake0)]] +, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] +) { + uint i = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + uint _e12 = i; + i = _e12 + 1u; + } + loop_init = false; + uint _e2 = i; + if (_e2 < 4u) { + } else { + break; + } + { + uint _e6 = i; + uint _e8 = i; + data[_e6] = _e8 * 2u; + } + } + return; +} diff --git a/naga/tests/out/msl/wgsl-test.msl b/naga/tests/out/msl/wgsl-test.msl new file mode 100644 index 00000000000..1fd8fb498e3 --- /dev/null +++ b/naga/tests/out/msl/wgsl-test.msl @@ -0,0 +1,37 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct _mslBufferSizes { + uint size0; +}; + +typedef uint type_1[1]; + +kernel void main_( + device type_1& data [[user(fake0)]] +, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] +) { + uint i = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + uint _e12 = i; + i = _e12 + 1u; + } + loop_init = false; + uint _e2 = i; + if (_e2 < 4u) { + } else { + break; + } + { + uint _e6 = i; + uint _e8 = i; + data[_e6] = _e8 * 2u; + } + } + return; +} diff --git a/naga/tests/out/spv/wgsl-test.spvasm b/naga/tests/out/spv/wgsl-test.spvasm new file mode 100644 index 00000000000..156df3d78a6 --- /dev/null +++ b/naga/tests/out/spv/wgsl-test.spvasm @@ -0,0 +1,87 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 55 +OpCapability Shader +OpExtension "SPV_KHR_storage_buffer_storage_class" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %9 "main" +OpExecutionMode %9 LocalSize 1 1 1 +OpDecorate %4 ArrayStride 4 +OpDecorate %5 DescriptorSet 0 +OpDecorate %5 Binding 0 +OpDecorate %6 Block +OpMemberDecorate %6 0 Offset 0 +%2 = OpTypeVoid +%3 = OpTypeInt 32 0 +%4 = OpTypeRuntimeArray %3 +%6 = OpTypeStruct %4 +%7 = OpTypePointer StorageBuffer %6 +%5 = OpVariable %7 StorageBuffer +%10 = OpTypeFunction %2 +%11 = OpTypePointer StorageBuffer %4 +%12 = OpConstant %3 0 +%14 = OpConstant %3 4 +%15 = OpConstant %3 2 +%16 = OpConstant %3 1 +%18 = OpTypePointer Function %3 +%24 = OpTypeVector %3 2 +%25 = OpTypePointer Function %24 +%26 = OpTypeBool +%27 = OpTypeVector %26 2 +%28 = OpConstantComposite %24 %12 %12 +%29 = OpConstant %3 4294967295 +%30 = OpConstantComposite %24 %29 %29 +%49 = OpTypePointer StorageBuffer %3 +%9 = OpFunction %2 None %10 +%8 = OpLabel +%17 = OpVariable %18 Function %12 +%31 = OpVariable %25 Function %30 +%13 = OpAccessChain %11 %5 %12 +OpBranch %19 +%19 = OpLabel +OpBranch %20 +%20 = OpLabel +OpLoopMerge %21 %23 None +OpBranch %32 +%32 = OpLabel +%33 = OpLoad %24 %31 +%34 = OpIEqual %27 %28 %33 +%35 = OpAll %26 %34 +OpSelectionMerge %36 None +OpBranchConditional %35 %21 %36 +%36 = OpLabel +%37 = OpCompositeExtract %3 %33 1 +%38 = OpIEqual %26 %37 %12 +%39 = OpSelect %3 %38 %16 %12 +%40 = OpCompositeConstruct %24 %39 %16 +%41 = OpISub %24 %33 %40 +OpStore %31 %41 +OpBranch %22 +%22 = OpLabel +%42 = OpLoad %3 %17 +%43 = OpULessThan %26 %42 %14 +OpSelectionMerge %44 None +OpBranchConditional %43 %44 %45 +%45 = OpLabel +OpBranch %21 +%44 = OpLabel +OpBranch %46 +%46 = OpLabel +%48 = OpLoad %3 %17 +%50 = OpLoad %3 %17 +%51 = OpIMul %3 %50 %15 +%52 = OpAccessChain %49 %13 %48 +OpStore %52 %51 +OpBranch %47 +%47 = OpLabel +OpBranch %23 +%23 = OpLabel +%53 = OpLoad %3 %17 +%54 = OpIAdd %3 %53 %16 +OpStore %17 %54 +OpBranch %20 +%21 = OpLabel +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/spv-test.wgsl b/naga/tests/out/wgsl/spv-test.wgsl new file mode 100644 index 00000000000..23df6e5f33f --- /dev/null +++ b/naga/tests/out/wgsl/spv-test.wgsl @@ -0,0 +1,43 @@ +struct RWStructuredBuffer { + member: array, +} + +@group(0) @binding(0) +var data: RWStructuredBuffer; + +fn main_1() { + var phi_19_: u32; + + phi_19_ = 0u; + loop { + let _e7 = phi_19_; + switch 0i { + default: { + if (_e7 < 4u) { + } else { + break; + } + let _e11 = data.member[_e7]; + if (_e11 == 1u) { + break; + } + let _e13 = data.member[_e7]; + if (_e13 == 2u) { + break; + } + data.member[_e7] = (_e7 * 2u); + break; + } + } + continue; + continuing { + phi_19_ = (_e7 + 1u); + } + } + return; +} + +@compute @workgroup_size(1, 1, 1) +fn main() { + main_1(); +} diff --git a/naga/tests/out/wgsl/wgsl-test.wgsl b/naga/tests/out/wgsl/wgsl-test.wgsl new file mode 100644 index 00000000000..4944dae5d57 --- /dev/null +++ b/naga/tests/out/wgsl/wgsl-test.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) +var data: array; + +@compute @workgroup_size(1, 1, 1) +fn main() { + var i: u32 = 0u; + + loop { + let _e2 = i; + if (_e2 < 4u) { + } else { + break; + } + { + let _e6 = i; + let _e8 = i; + data[_e6] = (_e8 * 2u); + } + continuing { + let _e12 = i; + i = (_e12 + 1u); + } + } + return; +} From dc4863a52f337dbd464f41509f8586bb6c9b423d Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Thu, 12 Jun 2025 18:54:12 -0400 Subject: [PATCH 3/3] Start writing load baking routine --- naga/src/back/load_baking.rs | 166 ++++++++++++++++++ naga/src/back/mod.rs | 2 + naga/tests/in/spv/load-elimination.spvasm | 40 +++++ naga/tests/in/wgsl/load-elimination.toml | 1 + naga/tests/in/wgsl/load-elimination.wgsl | 15 ++ .../wgsl-load-elimination.main.Compute.glsl | 25 +++ .../tests/out/hlsl/wgsl-load-elimination.hlsl | 20 +++ naga/tests/out/hlsl/wgsl-load-elimination.ron | 12 ++ .../out/ir/wgsl-load-elimination.compact.ron | 117 ++++++++++++ naga/tests/out/ir/wgsl-load-elimination.ron | 117 ++++++++++++ naga/tests/out/msl/wgsl-load-elimination.msl | 25 +++ .../out/spv/wgsl-load-elimination.spvasm | 37 ++++ naga/tests/out/wgsl/spv-load-elimination.wgsl | 22 +++ .../tests/out/wgsl/wgsl-load-elimination.wgsl | 18 ++ 14 files changed, 617 insertions(+) create mode 100644 naga/src/back/load_baking.rs create mode 100644 naga/tests/in/spv/load-elimination.spvasm create mode 100644 naga/tests/in/wgsl/load-elimination.toml create mode 100644 naga/tests/in/wgsl/load-elimination.wgsl create mode 100644 naga/tests/out/glsl/wgsl-load-elimination.main.Compute.glsl create mode 100644 naga/tests/out/hlsl/wgsl-load-elimination.hlsl create mode 100644 naga/tests/out/hlsl/wgsl-load-elimination.ron create mode 100644 naga/tests/out/ir/wgsl-load-elimination.compact.ron create mode 100644 naga/tests/out/ir/wgsl-load-elimination.ron create mode 100644 naga/tests/out/msl/wgsl-load-elimination.msl create mode 100644 naga/tests/out/spv/wgsl-load-elimination.spvasm create mode 100644 naga/tests/out/wgsl/spv-load-elimination.wgsl create mode 100644 naga/tests/out/wgsl/wgsl-load-elimination.wgsl diff --git a/naga/src/back/load_baking.rs b/naga/src/back/load_baking.rs new file mode 100644 index 00000000000..92bf81d6bcd --- /dev/null +++ b/naga/src/back/load_baking.rs @@ -0,0 +1,166 @@ +use core::{cell::Cell, mem}; +use std::vec::Vec; + +use crate::{ir, valid, FastHashMap, FastHashSet, Handle}; + +enum VariableState { + Uninitialized, + Initialized, + Written, + ReRead, +} + +impl VariableState { + fn write(&mut self) { + match self { + VariableState::Uninitialized => *self = VariableState::Initialized, + VariableState::Initialized => *self = VariableState::Written, + VariableState::ReRead => {} + VariableState::Written => {} + } + } + + fn read(&mut self) { + match self { + VariableState::Uninitialized => unreachable!(), + VariableState::Initialized => {} + VariableState::Written => *self = VariableState::ReRead, + VariableState::ReRead => {} + } + } +} + +struct LoadBaker<'a> { + pub module: &'a ir::Module, + pub function: &'a ir::Function, + pub function_info: &'a valid::FunctionInfo, + + requires_bake: FastHashSet>, + + states: FastHashMap, (usize, VariableState)>, + + loop_levels: Vec, +} + +impl<'a> LoadBaker<'a> { + pub fn new( + module: &'a ir::Module, + function: &'a ir::Function, + function_info: &'a valid::FunctionInfo, + ) -> Self { + let requires_bake = FastHashSet::default(); + let states = FastHashMap::default(); + + Self { + module, + function, + function_info, + requires_bake, + states, + loop_levels: Vec::new(), + } + } + + pub fn evaluate(&mut self) { + let current_depth = ScopedDepth::new(); + self.evaluate_block(¤t_depth, &self.function.body); + } + + fn evaluate_block(&mut self, depth: &ScopedDepth, block: &'a ir::Block) { + let _guard = depth.enter(); + for statement in block { + match *statement { + ir::Statement::Store { pointer, value } => { + self.evaluate_expression(depth, value); + self.register_write(depth, pointer); + } + ir::Statement::Emit(ref range) => { + for expression in range.clone() { + self.evaluate_expression(depth, expression); + } + } + ir::Statement::Block(ref block) => { + self.evaluate_block(&depth, block); + } + ir::Statement::Loop { + ref body, + ref continuing, + break_if, + } => { + self.loop_levels.push(depth.get()); + + self.evaluate_block(&depth, body); + self.evaluate_block(&depth, continuing); + if let Some(condition) = break_if { + self.evaluate_expression(depth, condition); + } + + self.loop_levels.pop(); + } + _ => {} + } + } + } + + fn evaluate_expression(&mut self, depth: &ScopedDepth, expression: Handle) { + let expression = &self.function.expressions[expression]; + + match *expression { + ir::Expression::Load { pointer } => { + let exp = &self.function.expressions[pointer]; + + let is_local_variable = matches!(exp, ir::Expression::LocalVariable(_)); + + if is_local_variable { + self.register_load(depth, pointer); + } + } + _ => {} + } + } + + fn register_load(&mut self, depth: &ScopedDepth, pointer: Handle) { + // Register the load as happening at the depth of the variable. + if let Some((_, state)) = self.states.get_mut(&pointer) { + state.read(); + return; + } + } + + fn register_write(&mut self, depth: &ScopedDepth, pointer: Handle) {} +} + +struct ScopedDepth { + current: Cell, +} + +impl ScopedDepth { + pub fn new() -> Self { + Self { + current: Cell::new(0), + } + } + + pub fn enter(&self) -> DepthGuard<'_> { + self.current.set(self.current.get() + 1); + DepthGuard { + manager: self, + depth: self.current.get(), + } + } + + pub fn get(&self) -> usize { + self.current.get() + } +} + +struct DepthGuard<'a> { + manager: &'a ScopedDepth, + depth: usize, +} + +impl Drop for DepthGuard<'_> { + fn drop(&mut self) { + self.manager.current.set(self.manager.current.get() - 1); + } +} diff --git a/naga/src/back/mod.rs b/naga/src/back/mod.rs index d7b14475bff..0e35ab74ea4 100644 --- a/naga/src/back/mod.rs +++ b/naga/src/back/mod.rs @@ -30,6 +30,8 @@ pub mod pipeline_constants; #[cfg(any(hlsl_out, glsl_out))] mod continue_forward; +mod load_baking; + /// Names of vector components. pub const COMPONENTS: &[char] = &['x', 'y', 'z', 'w']; /// Indent for backends. diff --git a/naga/tests/in/spv/load-elimination.spvasm b/naga/tests/in/spv/load-elimination.spvasm new file mode 100644 index 00000000000..50e69392c11 --- /dev/null +++ b/naga/tests/in/spv/load-elimination.spvasm @@ -0,0 +1,40 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 21 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %18 "main" +OpExecutionMode %18 LocalSize 1 1 1 +OpName %5 "sink" +OpName %11 "a" +OpName %16 "b" +%2 = OpTypeVoid +%3 = OpTypeInt 32 0 +%4 = OpConstant %3 0 +%6 = OpTypePointer Private %3 +%5 = OpVariable %6 Private %4 +%9 = OpTypeFunction %2 +%10 = OpConstant %3 2 +%12 = OpTypePointer Function %3 +%13 = OpConstantNull %3 +%8 = OpFunction %2 None %9 +%7 = OpLabel +%11 = OpVariable %12 Function %13 +OpBranch %14 +%14 = OpLabel +%15 = OpLoad %3 %5 +OpStore %11 %15 +%16 = OpLoad %3 %11 +OpStore %11 %10 +OpStore %5 %16 +OpReturn +OpFunctionEnd +%18 = OpFunction %2 None %9 +%17 = OpLabel +OpBranch %19 +%19 = OpLabel +%20 = OpFunctionCall %2 %8 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/in/wgsl/load-elimination.toml b/naga/tests/in/wgsl/load-elimination.toml new file mode 100644 index 00000000000..024bc76ea71 --- /dev/null +++ b/naga/tests/in/wgsl/load-elimination.toml @@ -0,0 +1 @@ +targets = "SPIRV | METAL | GLSL | WGSL | HLSL | IR" diff --git a/naga/tests/in/wgsl/load-elimination.wgsl b/naga/tests/in/wgsl/load-elimination.wgsl new file mode 100644 index 00000000000..a095afbaab2 --- /dev/null +++ b/naga/tests/in/wgsl/load-elimination.wgsl @@ -0,0 +1,15 @@ +var sink: u32 = 0; + +fn simple() { + var a = sink; + let b = a; + + a = 2u; + + sink = b; +} + +@compute @workgroup_size(1) +fn main() { + simple(); +} \ No newline at end of file diff --git a/naga/tests/out/glsl/wgsl-load-elimination.main.Compute.glsl b/naga/tests/out/glsl/wgsl-load-elimination.main.Compute.glsl new file mode 100644 index 00000000000..902efb7e6a9 --- /dev/null +++ b/naga/tests/out/glsl/wgsl-load-elimination.main.Compute.glsl @@ -0,0 +1,25 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +uint sink = 0u; + + +void simple() { + uint a = 0u; + uint _e1 = sink; + a = _e1; + uint b = a; + a = 2u; + sink = b; + return; +} + +void main() { + simple(); + return; +} + diff --git a/naga/tests/out/hlsl/wgsl-load-elimination.hlsl b/naga/tests/out/hlsl/wgsl-load-elimination.hlsl new file mode 100644 index 00000000000..a6734d09200 --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-load-elimination.hlsl @@ -0,0 +1,20 @@ +static uint sink = 0u; + +void simple() +{ + uint a = (uint)0; + + uint _e1 = sink; + a = _e1; + uint b = a; + a = 2u; + sink = b; + return; +} + +[numthreads(1, 1, 1)] +void main() +{ + simple(); + return; +} diff --git a/naga/tests/out/hlsl/wgsl-load-elimination.ron b/naga/tests/out/hlsl/wgsl-load-elimination.ron new file mode 100644 index 00000000000..a07b03300b1 --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-load-elimination.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/ir/wgsl-load-elimination.compact.ron b/naga/tests/out/ir/wgsl-load-elimination.compact.ron new file mode 100644 index 00000000000..a4b5f5eac9b --- /dev/null +++ b/naga/tests/out/ir/wgsl-load-elimination.compact.ron @@ -0,0 +1,117 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("sink"), + space: Private, + binding: None, + ty: 0, + init: Some(0), + ), + ], + global_expressions: [ + Literal(U32(0)), + ], + functions: [ + ( + name: Some("simple"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("a"), + ty: 0, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + Load( + pointer: 0, + ), + LocalVariable(0), + Load( + pointer: 2, + ), + Literal(U32(2)), + GlobalVariable(0), + ], + named_expressions: { + 3: "b", + }, + body: [ + Emit(( + start: 1, + end: 2, + )), + Store( + pointer: 2, + value: 1, + ), + Emit(( + start: 3, + end: 4, + )), + Store( + pointer: 2, + value: 4, + ), + Store( + pointer: 5, + value: 3, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [], + result: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/wgsl-load-elimination.ron b/naga/tests/out/ir/wgsl-load-elimination.ron new file mode 100644 index 00000000000..a4b5f5eac9b --- /dev/null +++ b/naga/tests/out/ir/wgsl-load-elimination.ron @@ -0,0 +1,117 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("sink"), + space: Private, + binding: None, + ty: 0, + init: Some(0), + ), + ], + global_expressions: [ + Literal(U32(0)), + ], + functions: [ + ( + name: Some("simple"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("a"), + ty: 0, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + Load( + pointer: 0, + ), + LocalVariable(0), + Load( + pointer: 2, + ), + Literal(U32(2)), + GlobalVariable(0), + ], + named_expressions: { + 3: "b", + }, + body: [ + Emit(( + start: 1, + end: 2, + )), + Store( + pointer: 2, + value: 1, + ), + Emit(( + start: 3, + end: 4, + )), + Store( + pointer: 2, + value: 4, + ), + Store( + pointer: 5, + value: 3, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [], + result: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/msl/wgsl-load-elimination.msl b/naga/tests/out/msl/wgsl-load-elimination.msl new file mode 100644 index 00000000000..12c10e12e58 --- /dev/null +++ b/naga/tests/out/msl/wgsl-load-elimination.msl @@ -0,0 +1,25 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + + +void simple( + thread uint& sink +) { + uint a = {}; + uint _e1 = sink; + a = _e1; + uint b = a; + a = 2u; + sink = b; + return; +} + +kernel void main_( +) { + uint sink = 0u; + simple(sink); + return; +} diff --git a/naga/tests/out/spv/wgsl-load-elimination.spvasm b/naga/tests/out/spv/wgsl-load-elimination.spvasm new file mode 100644 index 00000000000..029c1559068 --- /dev/null +++ b/naga/tests/out/spv/wgsl-load-elimination.spvasm @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 21 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %18 "main" +OpExecutionMode %18 LocalSize 1 1 1 +%2 = OpTypeVoid +%3 = OpTypeInt 32 0 +%4 = OpConstant %3 0 +%6 = OpTypePointer Private %3 +%5 = OpVariable %6 Private %4 +%9 = OpTypeFunction %2 +%10 = OpConstant %3 2 +%12 = OpTypePointer Function %3 +%13 = OpConstantNull %3 +%8 = OpFunction %2 None %9 +%7 = OpLabel +%11 = OpVariable %12 Function %13 +OpBranch %14 +%14 = OpLabel +%15 = OpLoad %3 %5 +OpStore %11 %15 +%16 = OpLoad %3 %11 +OpStore %11 %10 +OpStore %5 %16 +OpReturn +OpFunctionEnd +%18 = OpFunction %2 None %9 +%17 = OpLabel +OpBranch %19 +%19 = OpLabel +%20 = OpFunctionCall %2 %8 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/spv-load-elimination.wgsl b/naga/tests/out/wgsl/spv-load-elimination.wgsl new file mode 100644 index 00000000000..bb85fbdc056 --- /dev/null +++ b/naga/tests/out/wgsl/spv-load-elimination.wgsl @@ -0,0 +1,22 @@ +var sink: u32 = 0u; + +fn function() { + var a: u32 = u32(); + + let _e4 = sink; + a = _e4; + let _e5 = a; + a = 2u; + sink = _e5; + return; +} + +fn function_1() { + function(); + return; +} + +@compute @workgroup_size(1, 1, 1) +fn main() { + function_1(); +} diff --git a/naga/tests/out/wgsl/wgsl-load-elimination.wgsl b/naga/tests/out/wgsl/wgsl-load-elimination.wgsl new file mode 100644 index 00000000000..5c9d19941b3 --- /dev/null +++ b/naga/tests/out/wgsl/wgsl-load-elimination.wgsl @@ -0,0 +1,18 @@ +var sink: u32 = 0u; + +fn simple() { + var a: u32; + + let _e1 = sink; + a = _e1; + let b = a; + a = 2u; + sink = b; + return; +} + +@compute @workgroup_size(1, 1, 1) +fn main() { + simple(); + return; +}