diff --git a/CHANGELOG.md b/CHANGELOG.md index b763d71ef0..6717a9d330 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -107,6 +107,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148] #### Naga +- Clean up tests for atomic operations support in SPIR-V frontend. By @schell in [#6692](https://github.com/gfx-rs/wgpu/pull/6692) - Fix an issue where `naga` CLI would incorrectly skip the first positional argument when `--stdin-file-path` was specified. By @ErichDonGubler in [#6480](https://github.com/gfx-rs/wgpu/pull/6480). - Fix textureNumLevels in the GLSL backend. By @magcius in [#6483](https://github.com/gfx-rs/wgpu/pull/6483). - Support 64-bit hex literals and unary operations in constants [#6616](https://github.com/gfx-rs/wgpu/pull/6616). diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index 63406d3220..839f37ab5f 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -6057,89 +6057,3 @@ mod test { let _ = super::parse_u8_slice(&bin, &Default::default()).unwrap(); } } - -#[cfg(all(test, feature = "wgsl-in", wgsl_out))] -mod test_atomic { - fn atomic_test(bytes: &[u8]) { - let _ = env_logger::builder().is_test(true).try_init(); - let m = crate::front::spv::parse_u8_slice(bytes, &Default::default()).unwrap(); - - let mut wgsl = String::new(); - - for (vflags, name) in [ - (crate::valid::ValidationFlags::empty(), "empty"), - (crate::valid::ValidationFlags::all(), "all"), - ] { - log::info!("validating with flags - {name}"); - let mut validator = crate::valid::Validator::new(vflags, Default::default()); - match validator.validate(&m) { - Err(e) => { - log::error!("SPIR-V validation {}", e.emit_to_string("")); - log::info!("types: {:#?}", m.types); - panic!("validation error"); - } - Ok(i) => { - wgsl = crate::back::wgsl::write_string( - &m, - &i, - crate::back::wgsl::WriterFlags::empty(), - ) - .unwrap(); - log::info!("wgsl-out:\n{wgsl}"); - } - }; - } - - let m = match crate::front::wgsl::parse_str(&wgsl) { - Ok(m) => m, - Err(e) => { - log::error!("round trip WGSL validation {}", e.emit_to_string(&wgsl)); - panic!("invalid module"); - } - }; - let mut validator = - crate::valid::Validator::new(crate::valid::ValidationFlags::all(), Default::default()); - if let Err(e) = validator.validate(&m) { - log::error!("{}", e.emit_to_string(&wgsl)); - panic!("invalid generated wgsl"); - } - } - - #[test] - fn atomic_i_inc() { - atomic_test(include_bytes!( - "../../../tests/in/spv/atomic_i_increment.spv" - )); - } - - #[test] - fn atomic_load_and_store() { - atomic_test(include_bytes!( - "../../../tests/in/spv/atomic_load_and_store.spv" - )); - } - - #[test] - fn atomic_exchange() { - atomic_test(include_bytes!("../../../tests/in/spv/atomic_exchange.spv")); - } - - #[test] - fn atomic_compare_exchange() { - atomic_test(include_bytes!( - "../../../tests/in/spv/atomic_compare_exchange.spv" - )); - } - - #[test] - fn atomic_i_decrement() { - atomic_test(include_bytes!( - "../../../tests/in/spv/atomic_i_decrement.spv" - )); - } - - #[test] - fn atomic_i_add_and_sub() { - atomic_test(include_bytes!("../../../tests/in/spv/atomic_i_add_sub.spv")); - } -} diff --git a/naga/tests/out/ir/atomic_i_increment.compact.ron b/naga/tests/out/ir/atomic_i_increment.compact.ron deleted file mode 100644 index 5bb6820258..0000000000 --- a/naga/tests/out/ir/atomic_i_increment.compact.ron +++ /dev/null @@ -1,287 +0,0 @@ -( - types: [ - ( - name: None, - inner: Scalar(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Scalar(( - kind: Bool, - width: 1, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 0, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 3, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ], - special_types: ( - ray_desc: None, - ray_intersection: None, - predeclared_types: {}, - ), - constants: [ - ( - name: None, - ty: 0, - init: 0, - ), - ( - name: None, - ty: 1, - init: 1, - ), - ( - name: None, - ty: 0, - init: 2, - ), - ( - name: None, - ty: 1, - init: 3, - ), - ( - name: None, - ty: 0, - init: 4, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 4, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 2, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(Bool(false)), - Literal(U32(1)), - Literal(Bool(true)), - ZeroValue(0), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_23"), - ty: 0, - init: None, - ), - ( - name: Some("phi_24"), - ty: 0, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(3), - Constant(1), - Constant(4), - Constant(2), - Constant(0), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - LocalVariable(0), - Load( - pointer: 9, - ), - Load( - pointer: 8, - ), - Binary( - op: GreaterEqual, - left: 10, - right: 11, - ), - AtomicResult( - ty: 0, - comparison: false, - ), - Literal(U32(1)), - Binary( - op: Add, - left: 10, - right: 5, - ), - LocalVariable(1), - Load( - pointer: 16, - ), - Select( - condition: 12, - accept: 3, - reject: 2, - ), - Unary( - op: LogicalNot, - expr: 18, - ), - LocalVariable(0), - LocalVariable(1), - ], - named_expressions: {}, - body: [ - Emit(( - start: 7, - end: 9, - )), - Store( - pointer: 20, - value: 6, - ), - Loop( - body: [ - Emit(( - start: 10, - end: 11, - )), - Emit(( - start: 11, - end: 13, - )), - If( - condition: 12, - accept: [ - Store( - pointer: 21, - value: 4, - ), - ], - reject: [ - Atomic( - pointer: 7, - fun: Add, - value: 14, - result: Some(13), - ), - Emit(( - start: 15, - end: 16, - )), - Store( - pointer: 21, - value: 15, - ), - ], - ), - Emit(( - start: 17, - end: 19, - )), - Continue, - ], - continuing: [ - Emit(( - start: 19, - end: 20, - )), - Store( - pointer: 20, - value: 17, - ), - ], - break_if: Some(19), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_i_increment", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_i_increment_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, -) \ No newline at end of file diff --git a/naga/tests/out/ir/atomic_i_increment.ron b/naga/tests/out/ir/atomic_i_increment.ron deleted file mode 100644 index ae14821330..0000000000 --- a/naga/tests/out/ir/atomic_i_increment.ron +++ /dev/null @@ -1,312 +0,0 @@ -( - types: [ - ( - name: None, - inner: Scalar(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Scalar(( - kind: Bool, - width: 1, - )), - ), - ( - name: None, - inner: Pointer( - base: 0, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 0, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Pointer( - base: 3, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 5, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ], - special_types: ( - ray_desc: None, - ray_intersection: 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: 1, - init: 4, - ), - ( - name: None, - ty: 0, - init: 5, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 6, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 3, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(U32(2)), - Literal(Bool(false)), - Literal(U32(1)), - Literal(Bool(true)), - ZeroValue(0), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_23"), - ty: 0, - init: None, - ), - ( - name: Some("phi_24"), - ty: 0, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(4), - Constant(2), - Constant(5), - Constant(3), - Constant(1), - Constant(0), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - LocalVariable(0), - Load( - pointer: 10, - ), - Load( - pointer: 9, - ), - Binary( - op: GreaterEqual, - left: 11, - right: 12, - ), - AtomicResult( - ty: 0, - comparison: false, - ), - Literal(U32(1)), - Binary( - op: Add, - left: 11, - right: 5, - ), - LocalVariable(1), - Load( - pointer: 17, - ), - Select( - condition: 13, - accept: 3, - reject: 2, - ), - Unary( - op: LogicalNot, - expr: 19, - ), - LocalVariable(0), - LocalVariable(1), - ], - named_expressions: {}, - body: [ - Emit(( - start: 8, - end: 10, - )), - Store( - pointer: 21, - value: 7, - ), - Loop( - body: [ - Emit(( - start: 11, - end: 12, - )), - Emit(( - start: 12, - end: 14, - )), - If( - condition: 13, - accept: [ - Store( - pointer: 22, - value: 4, - ), - ], - reject: [ - Atomic( - pointer: 8, - fun: Add, - value: 15, - result: Some(14), - ), - Emit(( - start: 16, - end: 17, - )), - Store( - pointer: 22, - value: 16, - ), - ], - ), - Emit(( - start: 18, - end: 20, - )), - Continue, - ], - continuing: [ - Emit(( - start: 20, - end: 21, - )), - Store( - pointer: 21, - value: 18, - ), - ], - break_if: Some(20), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_i_increment", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_i_increment_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, -) \ No newline at end of file diff --git a/naga/tests/out/wgsl/atomic_compare_exchange.wgsl b/naga/tests/out/wgsl/atomic_compare_exchange.wgsl new file mode 100644 index 0000000000..00d3b6008c --- /dev/null +++ b/naga/tests/out/wgsl/atomic_compare_exchange.wgsl @@ -0,0 +1,66 @@ +struct type_2 { + member: u32, + member_1: u32, +} + +struct type_3 { + member: u32, +} + +struct type_5 { + member: atomic, +} + +@group(0) @binding(0) +var global: type_5; +@group(0) @binding(1) +var global_1: type_3; + +fn function() { + var phi_33_: type_2; + var phi_34_: type_2; + var phi_49_: type_2; + var phi_63_: bool; + + let _e11 = global_1.member; + phi_33_ = type_2(0u, _e11); + loop { + let _e14 = phi_33_; + if (_e14.member < _e14.member_1) { + phi_34_ = type_2((_e14.member + 1u), _e14.member_1); + phi_49_ = type_2(1u, _e14.member); + } else { + phi_34_ = _e14; + phi_49_ = type_2(0u, type_2().member_1); + } + let _e25 = phi_34_; + let _e27 = phi_49_; + switch bitcast(_e27.member) { + case 0: { + phi_63_ = false; + break; + } + case 1: { + let _e31 = atomicCompareExchangeWeak((&global.member), 3u, _e27.member_1); + phi_63_ = select(true, false, (_e31.old_value == 3u)); + break; + } + default: { + phi_63_ = bool(); + break; + } + } + let _e36 = phi_63_; + continue; + continuing { + phi_33_ = _e25; + break if !(_e36); + } + } + return; +} + +@compute @workgroup_size(32, 1, 1) +fn stagetest_atomic_compare_exchange() { + function(); +} diff --git a/naga/tests/out/wgsl/atomic_exchange.wgsl b/naga/tests/out/wgsl/atomic_exchange.wgsl new file mode 100644 index 0000000000..bbdb369ae8 --- /dev/null +++ b/naga/tests/out/wgsl/atomic_exchange.wgsl @@ -0,0 +1,80 @@ +struct type_2 { + member: u32, + member_1: u32, +} + +struct type_3 { + member: u32, +} + +struct type_5 { + member: atomic, +} + +@group(0) @binding(0) +var global: type_5; +@group(0) @binding(1) +var global_1: type_3; + +fn function() { + var phi_33_: type_2; + var phi_36_: u32; + var phi_52_: type_2; + var phi_53_: type_2; + var phi_62_: bool; + var phi_34_: type_2; + var phi_37_: u32; + + let _e10 = global_1.member; + phi_33_ = type_2(0u, _e10); + phi_36_ = 0u; + loop { + let _e13 = phi_33_; + let _e15 = phi_36_; + if (_e13.member < _e13.member_1) { + phi_52_ = type_2((_e13.member + 1u), _e13.member_1); + phi_53_ = type_2(1u, _e13.member); + } else { + phi_52_ = _e13; + phi_53_ = type_2(0u, type_2().member_1); + } + let _e26 = phi_52_; + let _e28 = phi_53_; + switch bitcast(_e28.member) { + case 0: { + phi_62_ = false; + phi_34_ = type_2(); + phi_37_ = u32(); + break; + } + case 1: { + let _e31 = atomicExchange((&global.member), _e15); + phi_62_ = true; + phi_34_ = _e26; + phi_37_ = (_e15 + _e31); + break; + } + default: { + phi_62_ = false; + phi_34_ = type_2(); + phi_37_ = u32(); + break; + } + } + let _e34 = phi_62_; + let _e36 = phi_34_; + let _e38 = phi_37_; + continue; + continuing { + phi_33_ = _e36; + phi_36_ = _e38; + break if !(_e34); + } + } + return; +} + +@compute @workgroup_size(32, 1, 1) +fn stagetest_atomic_exchange() { + function(); +} diff --git a/naga/tests/out/wgsl/atomic_i_add_sub.wgsl b/naga/tests/out/wgsl/atomic_i_add_sub.wgsl new file mode 100644 index 0000000000..d7798bd411 --- /dev/null +++ b/naga/tests/out/wgsl/atomic_i_add_sub.wgsl @@ -0,0 +1,26 @@ +struct type_2 { + member: array, +} + +struct type_4 { + member: atomic, +} + +@group(0) @binding(0) +var global: type_4; +@group(0) @binding(1) +var global_1: type_2; + +fn function() { + let _e6 = atomicAdd((&global.member), 2u); + let _e7 = atomicSub((&global.member), _e6); + if (_e6 < arrayLength((&global_1.member))) { + global_1.member[_e6] = _e7; + } + return; +} + +@compute @workgroup_size(32, 1, 1) +fn stagetest_atomic_i_add_sub() { + function(); +} diff --git a/naga/tests/out/wgsl/atomic_i_decrement.wgsl b/naga/tests/out/wgsl/atomic_i_decrement.wgsl new file mode 100644 index 0000000000..fdd067cd67 --- /dev/null +++ b/naga/tests/out/wgsl/atomic_i_decrement.wgsl @@ -0,0 +1,37 @@ +struct type_3 { + member: array, +} + +struct type_5 { + member: atomic, +} + +@group(0) @binding(0) +var global: type_5; +@group(0) @binding(1) +var global_1: type_3; + +fn function() { + var phi_40_: bool; + + loop { + let _e8 = atomicSub((&global.member), 1u); + if (_e8 < arrayLength((&global_1.member))) { + global_1.member[_e8] = _e8; + phi_40_ = select(true, false, (_e8 == 0u)); + } else { + phi_40_ = false; + } + let _e16 = phi_40_; + continue; + continuing { + break if !(_e16); + } + } + return; +} + +@compute @workgroup_size(32, 1, 1) +fn stagetest_atomic_i_decrement() { + function(); +} diff --git a/naga/tests/out/wgsl/atomic_i_increment.wgsl b/naga/tests/out/wgsl/atomic_i_increment.wgsl new file mode 100644 index 0000000000..da88e803e0 --- /dev/null +++ b/naga/tests/out/wgsl/atomic_i_increment.wgsl @@ -0,0 +1,42 @@ +struct type_2 { + member: u32, +} + +struct type_4 { + member: atomic, +} + +@group(0) @binding(0) +var global: type_4; +@group(0) @binding(1) +var global_1: type_2; + +fn function() { + var phi_23_: u32; + var phi_24_: u32; + + phi_23_ = 0u; + loop { + let _e10 = phi_23_; + let _e11 = global_1.member; + let _e12 = (_e10 >= _e11); + if _e12 { + phi_24_ = u32(); + } else { + let _e13 = atomicAdd((&global.member), 1u); + phi_24_ = (_e10 + 1u); + } + let _e17 = phi_24_; + continue; + continuing { + phi_23_ = _e17; + break if !(select(true, false, _e12)); + } + } + return; +} + +@compute @workgroup_size(32, 1, 1) +fn stagetest_atomic_i_increment() { + function(); +} diff --git a/naga/tests/out/wgsl/atomic_load_and_store.wgsl b/naga/tests/out/wgsl/atomic_load_and_store.wgsl new file mode 100644 index 0000000000..5102d8ebb1 --- /dev/null +++ b/naga/tests/out/wgsl/atomic_load_and_store.wgsl @@ -0,0 +1,72 @@ +struct type_2 { + member: u32, + member_1: u32, +} + +struct type_3 { + member: u32, +} + +struct type_5 { + member: atomic, +} + +@group(0) @binding(0) +var global: type_5; +@group(0) @binding(1) +var global_1: type_3; + +fn function() { + var phi_32_: type_2; + var phi_49_: type_2; + var phi_50_: type_2; + var phi_59_: bool; + var phi_33_: type_2; + + let _e10 = global_1.member; + phi_32_ = type_2(0u, _e10); + loop { + let _e13 = phi_32_; + if (_e13.member < _e13.member_1) { + phi_49_ = type_2((_e13.member + 1u), _e13.member_1); + phi_50_ = type_2(1u, _e13.member); + } else { + phi_49_ = _e13; + phi_50_ = type_2(0u, type_2().member_1); + } + let _e24 = phi_49_; + let _e26 = phi_50_; + switch bitcast(_e26.member) { + case 0: { + phi_59_ = false; + phi_33_ = type_2(); + break; + } + case 1: { + let _e29 = atomicLoad((&global.member)); + atomicStore((&global.member), (_e29 + 2u)); + phi_59_ = true; + phi_33_ = _e24; + break; + } + default: { + phi_59_ = false; + phi_33_ = type_2(); + break; + } + } + let _e32 = phi_59_; + let _e34 = phi_33_; + continue; + continuing { + phi_32_ = _e34; + break if !(_e32); + } + } + return; +} + +@compute @workgroup_size(32, 1, 1) +fn stagetest_atomic_load_and_store() { + function(); +} diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 2460a69365..72ce323585 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -1071,7 +1071,12 @@ fn convert_spv_all() { false, Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ); - convert_spv("atomic_i_increment", false, Targets::IR); + convert_spv("atomic_i_increment", false, Targets::WGSL); + convert_spv("atomic_load_and_store", false, Targets::WGSL); + convert_spv("atomic_exchange", false, Targets::WGSL); + convert_spv("atomic_compare_exchange", false, Targets::WGSL); + convert_spv("atomic_i_decrement", false, Targets::WGSL); + convert_spv("atomic_i_add_sub", false, Targets::WGSL); convert_spv( "fetch_depth", false,