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(); 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/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 00000000000..5c35682a467 Binary files /dev/null and b/naga/tests/in/spv/test.spv differ diff --git a/naga/tests/in/spv/test.spvasm b/naga/tests/in/spv/test.spvasm new file mode 100644 index 00000000000..03e47c2c65f --- /dev/null +++ b/naga/tests/in/spv/test.spvasm @@ -0,0 +1,77 @@ +; SPIR-V +; Version: 1.5 +; Generator: Khronos Slang Compiler; 0 +; Bound: 39 +; Schema: 0 + OpCapability Shader + OpExtension "SPV_KHR_storage_buffer_storage_class" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %data + OpExecutionMode %main LocalSize 1 1 1 + OpSource Slang 1 + OpName %RWStructuredBuffer "RWStructuredBuffer" + OpName %data "data" + OpName %main "main" + OpDecorate %_runtimearr_uint ArrayStride 4 + OpDecorate %RWStructuredBuffer Block + OpMemberDecorate %RWStructuredBuffer 0 Offset 0 + OpDecorate %data Binding 0 + OpDecorate %data DescriptorSet 0 + %void = OpTypeVoid + %6 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %bool = OpTypeBool + %uint_4 = OpConstant %uint 4 + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint +%_runtimearr_uint = OpTypeRuntimeArray %uint +%RWStructuredBuffer = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer_RWStructuredBuffer = OpTypePointer StorageBuffer %RWStructuredBuffer + %uint_1 = OpConstant %uint 1 + %uint_2 = OpConstant %uint 2 + %data = OpVariable %_ptr_StorageBuffer_RWStructuredBuffer StorageBuffer + %main = OpFunction %void None %6 + %17 = OpLabel + OpBranch %18 + %18 = OpLabel + %19 = OpPhi %uint %uint_0 %17 %20 %21 + OpLoopMerge %22 %21 None + OpBranch %23 + %23 = OpLabel + OpSelectionMerge %24 None + OpSwitch %int_0 %25 + %25 = OpLabel + %26 = OpULessThan %bool %19 %uint_4 + OpSelectionMerge %27 None + OpBranchConditional %26 %27 %28 + %28 = OpLabel + OpBranch %22 + %27 = OpLabel + %29 = OpAccessChain %_ptr_StorageBuffer_uint %data %int_0 %19 + %30 = OpLoad %uint %29 + %31 = OpIEqual %bool %30 %uint_1 + OpSelectionMerge %32 None + OpBranchConditional %31 %33 %32 + %33 = OpLabel + OpBranch %24 + %32 = OpLabel + %34 = OpLoad %uint %29 + %35 = OpIEqual %bool %34 %uint_2 + OpSelectionMerge %36 None + OpBranchConditional %35 %37 %36 + %37 = OpLabel + OpBranch %22 + %36 = OpLabel + %38 = OpIMul %uint %19 %uint_2 + OpStore %29 %38 + OpBranch %24 + %24 = OpLabel + %20 = OpIAdd %uint %19 %uint_1 + OpBranch %21 + %21 = OpLabel + OpBranch %18 + %22 = OpLabel + OpReturn + OpFunctionEnd diff --git a/naga/tests/in/spv/test.toml b/naga/tests/in/spv/test.toml new file mode 100644 index 00000000000..192e8e5c3d9 --- /dev/null +++ b/naga/tests/in/spv/test.toml @@ -0,0 +1,10 @@ +targets = "HLSL | METAL | GLSL | WGSL | IR" + +[spv] +force_loop_bounding = false + +[hlsl] +force_loop_bounding = false + +[msl] +force_loop_bounding = false 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/in/wgsl/test.toml b/naga/tests/in/wgsl/test.toml new file mode 100644 index 00000000000..16e776f872a --- /dev/null +++ b/naga/tests/in/wgsl/test.toml @@ -0,0 +1,10 @@ +targets = "SPIRV | METAL | GLSL | WGSL| HLSL | IR" + +[spv] +force_loop_bounding = false + +[hlsl] +force_loop_bounding = false + +[msl] +force_loop_bounding = false diff --git a/naga/tests/in/wgsl/test.wgsl b/naga/tests/in/wgsl/test.wgsl new file mode 100644 index 00000000000..125593c96ef --- /dev/null +++ b/naga/tests/in/wgsl/test.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) +var 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-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/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-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/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-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/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-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/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-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/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-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/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-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; +} 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; +}