From 0b0accf26ffb858098431721d45ff697b6b0cadb Mon Sep 17 00:00:00 2001 From: Schell Carl Scivally Date: Tue, 10 Dec 2024 11:04:38 +1300 Subject: [PATCH 1/4] chore: [spv-in] clean up atomic upgrade tests --- naga/src/front/spv/mod.rs | 86 --------------------------------------- naga/tests/snapshots.rs | 8 +++- 2 files changed, 7 insertions(+), 87 deletions(-) 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/snapshots.rs b/naga/tests/snapshots.rs index 2460a69365..3c3149fa05 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -1071,7 +1071,13 @@ fn convert_spv_all() { false, Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ); - convert_spv("atomic_i_increment", false, Targets::IR); + let atomic_targets = Targets::IR | Targets::METAL | Targets::WGSL; + convert_spv("atomic_i_increment", false, atomic_targets); + convert_spv("atomic_load_and_store", false, atomic_targets); + convert_spv("atomic_exchange", false, atomic_targets); + convert_spv("atomic_compare_exchange", false, atomic_targets); + convert_spv("atomic_i_decrement", false, atomic_targets); + convert_spv("atomic_i_add_sub", false, atomic_targets); convert_spv( "fetch_depth", false, From 9809ca665b2b3a6f47af870ec4d945f331fcad8f Mon Sep 17 00:00:00 2001 From: Schell Carl Scivally Date: Tue, 10 Dec 2024 11:07:41 +1300 Subject: [PATCH 2/4] add output test files --- .../ir/atomic_compare_exchange.compact.ron | 495 +++++++++++++++ naga/tests/out/ir/atomic_compare_exchange.ron | 542 +++++++++++++++++ naga/tests/out/ir/atomic_exchange.compact.ron | 524 ++++++++++++++++ naga/tests/out/ir/atomic_exchange.ron | 564 ++++++++++++++++++ .../tests/out/ir/atomic_i_add_sub.compact.ron | 207 +++++++ naga/tests/out/ir/atomic_i_add_sub.ron | 262 ++++++++ .../out/ir/atomic_i_decrement.compact.ron | 273 +++++++++ naga/tests/out/ir/atomic_i_decrement.ron | 321 ++++++++++ .../out/ir/atomic_load_and_store.compact.ron | 471 +++++++++++++++ naga/tests/out/ir/atomic_load_and_store.ron | 504 ++++++++++++++++ .../tests/out/msl/atomic_compare_exchange.msl | 104 ++++ naga/tests/out/msl/atomic_exchange.msl | 89 +++ naga/tests/out/msl/atomic_i_add_sub.msl | 38 ++ naga/tests/out/msl/atomic_i_decrement.msl | 54 ++ naga/tests/out/msl/atomic_i_increment.msl | 52 ++ naga/tests/out/msl/atomic_load_and_store.msl | 81 +++ .../out/wgsl/atomic_compare_exchange.wgsl | 66 ++ naga/tests/out/wgsl/atomic_exchange.wgsl | 80 +++ naga/tests/out/wgsl/atomic_i_add_sub.wgsl | 26 + naga/tests/out/wgsl/atomic_i_decrement.wgsl | 37 ++ naga/tests/out/wgsl/atomic_i_increment.wgsl | 42 ++ .../tests/out/wgsl/atomic_load_and_store.wgsl | 72 +++ 22 files changed, 4904 insertions(+) create mode 100644 naga/tests/out/ir/atomic_compare_exchange.compact.ron create mode 100644 naga/tests/out/ir/atomic_compare_exchange.ron create mode 100644 naga/tests/out/ir/atomic_exchange.compact.ron create mode 100644 naga/tests/out/ir/atomic_exchange.ron create mode 100644 naga/tests/out/ir/atomic_i_add_sub.compact.ron create mode 100644 naga/tests/out/ir/atomic_i_add_sub.ron create mode 100644 naga/tests/out/ir/atomic_i_decrement.compact.ron create mode 100644 naga/tests/out/ir/atomic_i_decrement.ron create mode 100644 naga/tests/out/ir/atomic_load_and_store.compact.ron create mode 100644 naga/tests/out/ir/atomic_load_and_store.ron create mode 100644 naga/tests/out/msl/atomic_compare_exchange.msl create mode 100644 naga/tests/out/msl/atomic_exchange.msl create mode 100644 naga/tests/out/msl/atomic_i_add_sub.msl create mode 100644 naga/tests/out/msl/atomic_i_decrement.msl create mode 100644 naga/tests/out/msl/atomic_i_increment.msl create mode 100644 naga/tests/out/msl/atomic_load_and_store.msl create mode 100644 naga/tests/out/wgsl/atomic_compare_exchange.wgsl create mode 100644 naga/tests/out/wgsl/atomic_exchange.wgsl create mode 100644 naga/tests/out/wgsl/atomic_i_add_sub.wgsl create mode 100644 naga/tests/out/wgsl/atomic_i_decrement.wgsl create mode 100644 naga/tests/out/wgsl/atomic_i_increment.wgsl create mode 100644 naga/tests/out/wgsl/atomic_load_and_store.wgsl diff --git a/naga/tests/out/ir/atomic_compare_exchange.compact.ron b/naga/tests/out/ir/atomic_compare_exchange.compact.ron new file mode 100644 index 0000000000..aacc81e618 --- /dev/null +++ b/naga/tests/out/ir/atomic_compare_exchange.compact.ron @@ -0,0 +1,495 @@ +( + 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, + ), + ( + name: None, + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 0, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: Some("__atomic_compare_exchange_result"), + inner: Struct( + members: [ + ( + name: Some("old_value"), + ty: 0, + binding: None, + offset: 0, + ), + ( + name: Some("exchanged"), + ty: 1, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + 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: { + AtomicCompareExchangeWeakResult(( + kind: Uint, + width: 4, + )): 4, + }, + ), + constants: [ + ( + name: None, + ty: 0, + init: 0, + ), + ( + name: None, + ty: 1, + init: 1, + ), + ( + name: None, + ty: 0, + init: 2, + ), + ( + name: None, + ty: 2, + init: 3, + ), + ( + name: None, + ty: 0, + init: 4, + ), + ( + name: None, + ty: 1, + init: 5, + ), + ( + name: None, + ty: 1, + init: 6, + ), + ], + 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(Bool(false)), + Literal(U32(1)), + ZeroValue(2), + Literal(U32(3)), + ZeroValue(1), + Literal(Bool(true)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_33"), + ty: 2, + init: None, + ), + ( + name: Some("phi_34"), + ty: 2, + init: None, + ), + ( + name: Some("phi_49"), + ty: 2, + init: None, + ), + ( + name: Some("phi_63"), + ty: 1, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(5), + Constant(3), + Constant(1), + Constant(6), + Constant(4), + Constant(0), + Constant(2), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + Load( + pointer: 10, + ), + Compose( + ty: 2, + components: [ + 7, + 11, + ], + ), + LocalVariable(0), + Load( + pointer: 13, + ), + AccessIndex( + base: 14, + index: 0, + ), + AccessIndex( + base: 14, + index: 1, + ), + Binary( + op: Less, + left: 15, + right: 16, + ), + Binary( + op: Add, + left: 15, + right: 8, + ), + AccessIndex( + base: 14, + index: 1, + ), + Compose( + ty: 2, + components: [ + 18, + 19, + ], + ), + Compose( + ty: 2, + components: [ + 8, + 15, + ], + ), + AccessIndex( + base: 3, + index: 1, + ), + Compose( + ty: 2, + components: [ + 7, + 22, + ], + ), + LocalVariable(1), + Load( + pointer: 24, + ), + LocalVariable(2), + Load( + pointer: 26, + ), + AccessIndex( + base: 27, + index: 0, + ), + AccessIndex( + base: 27, + index: 1, + ), + As( + expr: 28, + kind: Sint, + convert: None, + ), + AtomicResult( + ty: 4, + comparison: true, + ), + AccessIndex( + base: 31, + index: 0, + ), + Binary( + op: Equal, + left: 32, + right: 6, + ), + Select( + condition: 33, + accept: 4, + reject: 5, + ), + LocalVariable(3), + Load( + pointer: 35, + ), + Unary( + op: LogicalNot, + expr: 36, + ), + LocalVariable(0), + LocalVariable(1), + LocalVariable(2), + LocalVariable(3), + ], + named_expressions: {}, + body: [ + Emit(( + start: 9, + end: 13, + )), + Store( + pointer: 38, + value: 12, + ), + Loop( + body: [ + Emit(( + start: 14, + end: 15, + )), + Emit(( + start: 15, + end: 18, + )), + If( + condition: 17, + accept: [ + Emit(( + start: 18, + end: 22, + )), + Store( + pointer: 39, + value: 20, + ), + Store( + pointer: 40, + value: 21, + ), + ], + reject: [ + Emit(( + start: 22, + end: 24, + )), + Store( + pointer: 39, + value: 14, + ), + Store( + pointer: 40, + value: 23, + ), + ], + ), + Emit(( + start: 25, + end: 26, + )), + Emit(( + start: 27, + end: 31, + )), + Switch( + selector: 30, + cases: [ + ( + value: I32(0), + body: [ + Store( + pointer: 41, + value: 4, + ), + Break, + ], + fall_through: false, + ), + ( + value: I32(1), + body: [ + Atomic( + pointer: 9, + fun: Exchange( + compare: Some(6), + ), + value: 29, + result: Some(31), + ), + Emit(( + start: 33, + end: 35, + )), + Store( + pointer: 41, + value: 34, + ), + Break, + ], + fall_through: false, + ), + ( + value: Default, + body: [ + Store( + pointer: 41, + value: 2, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 36, + end: 37, + )), + Continue, + ], + continuing: [ + Emit(( + start: 37, + end: 38, + )), + Store( + pointer: 38, + value: 25, + ), + ], + break_if: Some(37), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_compare_exchange", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_compare_exchange_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_compare_exchange.ron b/naga/tests/out/ir/atomic_compare_exchange.ron new file mode 100644 index 0000000000..4192ae8918 --- /dev/null +++ b/naga/tests/out/ir/atomic_compare_exchange.ron @@ -0,0 +1,542 @@ +( + 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, + ), + ( + name: None, + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Sint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 0, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 5, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + name: Some("__atomic_compare_exchange_result"), + inner: Struct( + members: [ + ( + name: Some("old_value"), + ty: 0, + binding: None, + offset: 0, + ), + ( + name: Some("exchanged"), + ty: 1, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + name: None, + inner: Atomic(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 8, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + predeclared_types: { + AtomicCompareExchangeWeakResult(( + kind: Uint, + width: 4, + )): 7, + }, + ), + 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: 3, + init: 4, + ), + ( + name: None, + ty: 0, + init: 5, + ), + ( + name: None, + ty: 1, + init: 6, + ), + ( + name: None, + ty: 1, + init: 7, + ), + ( + name: None, + ty: 0, + init: 8, + ), + ], + overrides: [], + global_variables: [ + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 9, + init: None, + ), + ( + name: None, + space: Storage( + access: ("LOAD"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 5, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(U32(2)), + Literal(Bool(false)), + Literal(U32(1)), + ZeroValue(3), + Literal(U32(3)), + ZeroValue(1), + Literal(Bool(true)), + Literal(U32(256)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_33"), + ty: 3, + init: None, + ), + ( + name: Some("phi_34"), + ty: 3, + init: None, + ), + ( + name: Some("phi_49"), + ty: 3, + init: None, + ), + ( + name: Some("phi_63"), + ty: 1, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(1), + Constant(8), + Constant(6), + Constant(4), + Constant(2), + Constant(7), + Constant(5), + Constant(0), + Constant(3), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + Load( + pointer: 12, + ), + Compose( + ty: 3, + components: [ + 9, + 13, + ], + ), + LocalVariable(0), + Load( + pointer: 15, + ), + AccessIndex( + base: 16, + index: 0, + ), + AccessIndex( + base: 16, + index: 1, + ), + Binary( + op: Less, + left: 17, + right: 18, + ), + Binary( + op: Add, + left: 17, + right: 10, + ), + AccessIndex( + base: 16, + index: 0, + ), + AccessIndex( + base: 16, + index: 1, + ), + Compose( + ty: 3, + components: [ + 20, + 22, + ], + ), + Compose( + ty: 3, + components: [ + 10, + 17, + ], + ), + AccessIndex( + base: 5, + index: 0, + ), + AccessIndex( + base: 5, + index: 1, + ), + Compose( + ty: 3, + components: [ + 9, + 26, + ], + ), + LocalVariable(1), + Load( + pointer: 28, + ), + LocalVariable(2), + Load( + pointer: 30, + ), + AccessIndex( + base: 31, + index: 0, + ), + AccessIndex( + base: 31, + index: 1, + ), + As( + expr: 32, + kind: Sint, + convert: None, + ), + AtomicResult( + ty: 7, + comparison: true, + ), + AccessIndex( + base: 35, + index: 0, + ), + Binary( + op: Equal, + left: 36, + right: 8, + ), + Select( + condition: 37, + accept: 6, + reject: 7, + ), + LocalVariable(3), + Load( + pointer: 39, + ), + Unary( + op: LogicalNot, + expr: 40, + ), + LocalVariable(0), + LocalVariable(1), + LocalVariable(2), + LocalVariable(3), + ], + named_expressions: {}, + body: [ + Emit(( + start: 11, + end: 15, + )), + Store( + pointer: 42, + value: 14, + ), + Loop( + body: [ + Emit(( + start: 16, + end: 17, + )), + Emit(( + start: 17, + end: 20, + )), + If( + condition: 19, + accept: [ + Emit(( + start: 20, + end: 25, + )), + Store( + pointer: 43, + value: 23, + ), + Store( + pointer: 44, + value: 24, + ), + ], + reject: [ + Emit(( + start: 25, + end: 28, + )), + Store( + pointer: 43, + value: 16, + ), + Store( + pointer: 44, + value: 27, + ), + ], + ), + Emit(( + start: 29, + end: 30, + )), + Emit(( + start: 31, + end: 35, + )), + Switch( + selector: 34, + cases: [ + ( + value: I32(0), + body: [ + Store( + pointer: 45, + value: 6, + ), + Break, + ], + fall_through: false, + ), + ( + value: I32(1), + body: [ + Atomic( + pointer: 11, + fun: Exchange( + compare: Some(8), + ), + value: 33, + result: Some(35), + ), + Emit(( + start: 37, + end: 39, + )), + Store( + pointer: 45, + value: 38, + ), + Break, + ], + fall_through: false, + ), + ( + value: Default, + body: [ + Store( + pointer: 45, + value: 4, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 40, + end: 41, + )), + Continue, + ], + continuing: [ + Emit(( + start: 41, + end: 42, + )), + Store( + pointer: 42, + value: 29, + ), + ], + break_if: Some(41), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_compare_exchange", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_compare_exchange_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_exchange.compact.ron b/naga/tests/out/ir/atomic_exchange.compact.ron new file mode 100644 index 0000000000..192bd6f913 --- /dev/null +++ b/naga/tests/out/ir/atomic_exchange.compact.ron @@ -0,0 +1,524 @@ +( + 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, + ), + ( + name: None, + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + 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: 4, + 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: 2, + 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: 5, + init: None, + ), + ( + name: None, + space: Storage( + access: ("LOAD"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 3, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(Bool(false)), + Literal(U32(1)), + ZeroValue(2), + Literal(Bool(true)), + ZeroValue(0), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_33"), + ty: 2, + init: None, + ), + ( + name: Some("phi_36"), + ty: 0, + init: None, + ), + ( + name: Some("phi_52"), + ty: 2, + init: None, + ), + ( + name: Some("phi_53"), + ty: 2, + init: None, + ), + ( + name: Some("phi_62"), + ty: 1, + init: None, + ), + ( + name: Some("phi_34"), + ty: 2, + init: None, + ), + ( + name: Some("phi_37"), + ty: 0, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(4), + Constant(2), + Constant(5), + Constant(0), + Constant(3), + Constant(1), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + Load( + pointer: 9, + ), + Compose( + ty: 2, + components: [ + 5, + 10, + ], + ), + LocalVariable(0), + Load( + pointer: 12, + ), + LocalVariable(1), + Load( + pointer: 14, + ), + AccessIndex( + base: 13, + index: 0, + ), + AccessIndex( + base: 13, + index: 1, + ), + Binary( + op: Less, + left: 16, + right: 17, + ), + Binary( + op: Add, + left: 16, + right: 3, + ), + AccessIndex( + base: 13, + index: 1, + ), + Compose( + ty: 2, + components: [ + 19, + 20, + ], + ), + Compose( + ty: 2, + components: [ + 3, + 16, + ], + ), + AccessIndex( + base: 6, + index: 1, + ), + Compose( + ty: 2, + components: [ + 5, + 23, + ], + ), + LocalVariable(2), + Load( + pointer: 25, + ), + LocalVariable(3), + Load( + pointer: 27, + ), + AccessIndex( + base: 28, + index: 0, + ), + As( + expr: 29, + kind: Sint, + convert: None, + ), + AtomicResult( + ty: 0, + comparison: false, + ), + Binary( + op: Add, + left: 15, + right: 31, + ), + LocalVariable(4), + Load( + pointer: 33, + ), + LocalVariable(5), + Load( + pointer: 35, + ), + LocalVariable(6), + Load( + pointer: 37, + ), + Unary( + op: LogicalNot, + expr: 34, + ), + LocalVariable(0), + LocalVariable(1), + LocalVariable(2), + LocalVariable(3), + LocalVariable(4), + LocalVariable(5), + LocalVariable(6), + ], + named_expressions: {}, + body: [ + Emit(( + start: 8, + end: 12, + )), + Store( + pointer: 40, + value: 11, + ), + Store( + pointer: 41, + value: 5, + ), + Loop( + body: [ + Emit(( + start: 13, + end: 14, + )), + Emit(( + start: 15, + end: 16, + )), + Emit(( + start: 16, + end: 19, + )), + If( + condition: 18, + accept: [ + Emit(( + start: 19, + end: 23, + )), + Store( + pointer: 42, + value: 21, + ), + Store( + pointer: 43, + value: 22, + ), + ], + reject: [ + Emit(( + start: 23, + end: 25, + )), + Store( + pointer: 42, + value: 13, + ), + Store( + pointer: 43, + value: 24, + ), + ], + ), + Emit(( + start: 26, + end: 27, + )), + Emit(( + start: 28, + end: 31, + )), + Switch( + selector: 30, + cases: [ + ( + value: I32(0), + body: [ + Store( + pointer: 44, + value: 7, + ), + Store( + pointer: 45, + value: 6, + ), + Store( + pointer: 46, + value: 4, + ), + Break, + ], + fall_through: false, + ), + ( + value: I32(1), + body: [ + Atomic( + pointer: 8, + fun: Exchange( + compare: None, + ), + value: 15, + result: Some(31), + ), + Emit(( + start: 32, + end: 33, + )), + Store( + pointer: 44, + value: 2, + ), + Store( + pointer: 45, + value: 26, + ), + Store( + pointer: 46, + value: 32, + ), + Break, + ], + fall_through: false, + ), + ( + value: Default, + body: [ + Store( + pointer: 44, + value: 7, + ), + Store( + pointer: 45, + value: 6, + ), + Store( + pointer: 46, + value: 4, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 34, + end: 35, + )), + Emit(( + start: 36, + end: 37, + )), + Emit(( + start: 38, + end: 39, + )), + Continue, + ], + continuing: [ + Emit(( + start: 39, + end: 40, + )), + Store( + pointer: 40, + value: 36, + ), + Store( + pointer: 41, + value: 38, + ), + ], + break_if: Some(39), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_exchange", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_exchange_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_exchange.ron b/naga/tests/out/ir/atomic_exchange.ron new file mode 100644 index 0000000000..8a96c1c1b4 --- /dev/null +++ b/naga/tests/out/ir/atomic_exchange.ron @@ -0,0 +1,564 @@ +( + 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, + ), + ( + name: None, + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Sint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 0, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 5, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + name: None, + inner: Atomic(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 7, + 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: 3, + init: 4, + ), + ( + name: None, + ty: 1, + init: 5, + ), + ( + name: None, + ty: 0, + init: 6, + ), + ], + overrides: [], + global_variables: [ + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 8, + init: None, + ), + ( + name: None, + space: Storage( + access: ("LOAD"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 5, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(U32(2)), + Literal(Bool(false)), + Literal(U32(1)), + ZeroValue(3), + Literal(Bool(true)), + ZeroValue(0), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_33"), + ty: 3, + init: None, + ), + ( + name: Some("phi_36"), + ty: 0, + init: None, + ), + ( + name: Some("phi_52"), + ty: 3, + init: None, + ), + ( + name: Some("phi_53"), + ty: 3, + init: None, + ), + ( + name: Some("phi_62"), + ty: 1, + init: None, + ), + ( + name: Some("phi_34"), + ty: 3, + init: None, + ), + ( + name: Some("phi_37"), + ty: 0, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(5), + Constant(3), + Constant(1), + Constant(6), + Constant(0), + Constant(4), + Constant(2), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + Load( + pointer: 10, + ), + Compose( + ty: 3, + components: [ + 6, + 11, + ], + ), + LocalVariable(0), + Load( + pointer: 13, + ), + LocalVariable(1), + Load( + pointer: 15, + ), + AccessIndex( + base: 14, + index: 0, + ), + AccessIndex( + base: 14, + index: 1, + ), + Binary( + op: Less, + left: 17, + right: 18, + ), + Binary( + op: Add, + left: 17, + right: 3, + ), + AccessIndex( + base: 14, + index: 0, + ), + AccessIndex( + base: 14, + index: 1, + ), + Compose( + ty: 3, + components: [ + 20, + 22, + ], + ), + Compose( + ty: 3, + components: [ + 3, + 17, + ], + ), + AccessIndex( + base: 7, + index: 0, + ), + AccessIndex( + base: 7, + index: 1, + ), + Compose( + ty: 3, + components: [ + 6, + 26, + ], + ), + LocalVariable(2), + Load( + pointer: 28, + ), + LocalVariable(3), + Load( + pointer: 30, + ), + AccessIndex( + base: 31, + index: 0, + ), + As( + expr: 32, + kind: Sint, + convert: None, + ), + AtomicResult( + ty: 0, + comparison: false, + ), + Binary( + op: Add, + left: 16, + right: 34, + ), + LocalVariable(4), + Load( + pointer: 36, + ), + LocalVariable(5), + Load( + pointer: 38, + ), + LocalVariable(6), + Load( + pointer: 40, + ), + Unary( + op: LogicalNot, + expr: 37, + ), + LocalVariable(0), + LocalVariable(1), + LocalVariable(2), + LocalVariable(3), + LocalVariable(4), + LocalVariable(5), + LocalVariable(6), + ], + named_expressions: {}, + body: [ + Emit(( + start: 9, + end: 13, + )), + Store( + pointer: 43, + value: 12, + ), + Store( + pointer: 44, + value: 6, + ), + Loop( + body: [ + Emit(( + start: 14, + end: 15, + )), + Emit(( + start: 16, + end: 17, + )), + Emit(( + start: 17, + end: 20, + )), + If( + condition: 19, + accept: [ + Emit(( + start: 20, + end: 25, + )), + Store( + pointer: 45, + value: 23, + ), + Store( + pointer: 46, + value: 24, + ), + ], + reject: [ + Emit(( + start: 25, + end: 28, + )), + Store( + pointer: 45, + value: 14, + ), + Store( + pointer: 46, + value: 27, + ), + ], + ), + Emit(( + start: 29, + end: 30, + )), + Emit(( + start: 31, + end: 34, + )), + Switch( + selector: 33, + cases: [ + ( + value: I32(0), + body: [ + Store( + pointer: 47, + value: 8, + ), + Store( + pointer: 48, + value: 7, + ), + Store( + pointer: 49, + value: 5, + ), + Break, + ], + fall_through: false, + ), + ( + value: I32(1), + body: [ + Atomic( + pointer: 9, + fun: Exchange( + compare: None, + ), + value: 16, + result: Some(34), + ), + Emit(( + start: 35, + end: 36, + )), + Store( + pointer: 47, + value: 2, + ), + Store( + pointer: 48, + value: 29, + ), + Store( + pointer: 49, + value: 35, + ), + Break, + ], + fall_through: false, + ), + ( + value: Default, + body: [ + Store( + pointer: 47, + value: 8, + ), + Store( + pointer: 48, + value: 7, + ), + Store( + pointer: 49, + value: 5, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 37, + end: 38, + )), + Emit(( + start: 39, + end: 40, + )), + Emit(( + start: 41, + end: 42, + )), + Continue, + ], + continuing: [ + Emit(( + start: 42, + end: 43, + )), + Store( + pointer: 43, + value: 39, + ), + Store( + pointer: 44, + value: 41, + ), + ], + break_if: Some(42), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_exchange", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_exchange_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_add_sub.compact.ron b/naga/tests/out/ir/atomic_i_add_sub.compact.ron new file mode 100644 index 0000000000..52702a1828 --- /dev/null +++ b/naga/tests/out/ir/atomic_i_add_sub.compact.ron @@ -0,0 +1,207 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 1, + 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, + ), + ], + 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 | STORE"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 2, + init: None, + ), + ], + global_expressions: [ + Literal(U32(2)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(0), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + ArrayLength(4), + AtomicResult( + ty: 0, + comparison: false, + ), + AtomicResult( + ty: 0, + comparison: false, + ), + Binary( + op: Less, + left: 6, + right: 5, + ), + AccessIndex( + base: 1, + index: 0, + ), + Access( + base: 9, + index: 6, + ), + ], + named_expressions: {}, + body: [ + Emit(( + start: 3, + end: 6, + )), + Atomic( + pointer: 3, + fun: Add, + value: 2, + result: Some(6), + ), + Atomic( + pointer: 3, + fun: Subtract, + value: 6, + result: Some(7), + ), + Emit(( + start: 8, + end: 9, + )), + If( + condition: 8, + accept: [ + Emit(( + start: 9, + end: 11, + )), + Store( + pointer: 10, + value: 7, + ), + ], + reject: [], + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_i_add_sub", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_i_add_sub_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_add_sub.ron b/naga/tests/out/ir/atomic_i_add_sub.ron new file mode 100644 index 0000000000..e8d41b6889 --- /dev/null +++ b/naga/tests/out/ir/atomic_i_add_sub.ron @@ -0,0 +1,262 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Scalar(( + kind: Bool, + width: 1, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 2, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 3, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + 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: 6, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + name: None, + inner: Atomic(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 8, + 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, + ), + ], + overrides: [], + global_variables: [ + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 9, + init: None, + ), + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 3, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(U32(2)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(1), + Constant(0), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + ArrayLength(5), + AtomicResult( + ty: 0, + comparison: false, + ), + AtomicResult( + ty: 0, + comparison: false, + ), + Binary( + op: Less, + left: 7, + right: 6, + ), + AccessIndex( + base: 1, + index: 0, + ), + Access( + base: 10, + index: 7, + ), + ], + named_expressions: {}, + body: [ + Emit(( + start: 4, + end: 7, + )), + Atomic( + pointer: 4, + fun: Add, + value: 2, + result: Some(7), + ), + Atomic( + pointer: 4, + fun: Subtract, + value: 7, + result: Some(8), + ), + Emit(( + start: 9, + end: 10, + )), + If( + condition: 9, + accept: [ + Emit(( + start: 10, + end: 12, + )), + Store( + pointer: 11, + value: 8, + ), + ], + reject: [], + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_i_add_sub", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_i_add_sub_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_decrement.compact.ron b/naga/tests/out/ir/atomic_i_decrement.compact.ron new file mode 100644 index 0000000000..5fec93a946 --- /dev/null +++ b/naga/tests/out/ir/atomic_i_decrement.compact.ron @@ -0,0 +1,273 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Scalar(( + kind: Bool, + width: 1, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 2, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Atomic(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 4, + 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: 1, + init: 2, + ), + ], + overrides: [], + global_variables: [ + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 5, + init: None, + ), + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 3, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(Bool(false)), + Literal(Bool(true)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_40"), + ty: 1, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(0), + Constant(1), + Constant(2), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + ArrayLength(6), + AtomicResult( + ty: 0, + comparison: false, + ), + Literal(U32(1)), + Binary( + op: Less, + left: 8, + right: 7, + ), + AccessIndex( + base: 1, + index: 0, + ), + Access( + base: 11, + index: 8, + ), + Binary( + op: Equal, + left: 8, + right: 2, + ), + Select( + condition: 13, + accept: 3, + reject: 4, + ), + LocalVariable(0), + Load( + pointer: 15, + ), + Unary( + op: LogicalNot, + expr: 16, + ), + LocalVariable(0), + ], + named_expressions: {}, + body: [ + Emit(( + start: 5, + end: 8, + )), + Loop( + body: [ + Atomic( + pointer: 5, + fun: Subtract, + value: 9, + result: Some(8), + ), + Emit(( + start: 10, + end: 11, + )), + If( + condition: 10, + accept: [ + Emit(( + start: 11, + end: 13, + )), + Store( + pointer: 12, + value: 8, + ), + Emit(( + start: 13, + end: 15, + )), + Store( + pointer: 18, + value: 14, + ), + ], + reject: [ + Store( + pointer: 18, + value: 3, + ), + ], + ), + Emit(( + start: 16, + end: 17, + )), + Continue, + ], + continuing: [ + Emit(( + start: 17, + end: 18, + )), + ], + break_if: Some(17), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_i_decrement", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_i_decrement_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_decrement.ron b/naga/tests/out/ir/atomic_i_decrement.ron new file mode 100644 index 0000000000..59be516b93 --- /dev/null +++ b/naga/tests/out/ir/atomic_i_decrement.ron @@ -0,0 +1,321 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Scalar(( + kind: Bool, + width: 1, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Dynamic, + stride: 4, + ), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 2, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 3, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + 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: 6, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + name: None, + inner: Atomic(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 8, + 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: 1, + init: 3, + ), + ], + overrides: [], + global_variables: [ + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 9, + init: None, + ), + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 3, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(U32(2)), + Literal(Bool(false)), + Literal(Bool(true)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_40"), + ty: 1, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(1), + Constant(0), + Constant(2), + Constant(3), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + ArrayLength(7), + AtomicResult( + ty: 0, + comparison: false, + ), + Literal(U32(1)), + Binary( + op: Less, + left: 9, + right: 8, + ), + AccessIndex( + base: 1, + index: 0, + ), + Access( + base: 12, + index: 9, + ), + Binary( + op: Equal, + left: 9, + right: 3, + ), + Select( + condition: 14, + accept: 4, + reject: 5, + ), + LocalVariable(0), + Load( + pointer: 16, + ), + Unary( + op: LogicalNot, + expr: 17, + ), + LocalVariable(0), + ], + named_expressions: {}, + body: [ + Emit(( + start: 6, + end: 9, + )), + Loop( + body: [ + Atomic( + pointer: 6, + fun: Subtract, + value: 10, + result: Some(9), + ), + Emit(( + start: 11, + end: 12, + )), + If( + condition: 11, + accept: [ + Emit(( + start: 12, + end: 14, + )), + Store( + pointer: 13, + value: 9, + ), + Emit(( + start: 14, + end: 16, + )), + Store( + pointer: 19, + value: 15, + ), + ], + reject: [ + Store( + pointer: 19, + value: 4, + ), + ], + ), + Emit(( + start: 17, + end: 18, + )), + Continue, + ], + continuing: [ + Emit(( + start: 18, + end: 19, + )), + ], + break_if: Some(18), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_i_decrement", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_i_decrement_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_load_and_store.compact.ron b/naga/tests/out/ir/atomic_load_and_store.compact.ron new file mode 100644 index 0000000000..6e5d94dde3 --- /dev/null +++ b/naga/tests/out/ir/atomic_load_and_store.compact.ron @@ -0,0 +1,471 @@ +( + 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, + ), + ( + name: None, + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + 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: 4, + 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: 2, + init: 4, + ), + ( + name: None, + ty: 1, + init: 5, + ), + ], + overrides: [], + global_variables: [ + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 5, + 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)), + ZeroValue(2), + Literal(Bool(true)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_32"), + ty: 2, + init: None, + ), + ( + name: Some("phi_49"), + ty: 2, + init: None, + ), + ( + name: Some("phi_50"), + ty: 2, + init: None, + ), + ( + name: Some("phi_59"), + ty: 1, + init: None, + ), + ( + name: Some("phi_33"), + ty: 2, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(5), + Constant(3), + Constant(1), + Constant(0), + Constant(4), + Constant(2), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + Load( + pointer: 9, + ), + Compose( + ty: 2, + components: [ + 5, + 10, + ], + ), + LocalVariable(0), + Load( + pointer: 12, + ), + AccessIndex( + base: 13, + index: 0, + ), + AccessIndex( + base: 13, + index: 1, + ), + Binary( + op: Less, + left: 14, + right: 15, + ), + Binary( + op: Add, + left: 14, + right: 3, + ), + AccessIndex( + base: 13, + index: 1, + ), + Compose( + ty: 2, + components: [ + 17, + 18, + ], + ), + Compose( + ty: 2, + components: [ + 3, + 14, + ], + ), + AccessIndex( + base: 6, + index: 1, + ), + Compose( + ty: 2, + components: [ + 5, + 21, + ], + ), + LocalVariable(1), + Load( + pointer: 23, + ), + LocalVariable(2), + Load( + pointer: 25, + ), + AccessIndex( + base: 26, + index: 0, + ), + As( + expr: 27, + kind: Sint, + convert: None, + ), + Load( + pointer: 8, + ), + Binary( + op: Add, + left: 29, + right: 4, + ), + LocalVariable(3), + Load( + pointer: 31, + ), + LocalVariable(4), + Load( + pointer: 33, + ), + Unary( + op: LogicalNot, + expr: 32, + ), + LocalVariable(0), + LocalVariable(1), + LocalVariable(2), + LocalVariable(3), + LocalVariable(4), + ], + named_expressions: {}, + body: [ + Emit(( + start: 8, + end: 12, + )), + Store( + pointer: 36, + value: 11, + ), + Loop( + body: [ + Emit(( + start: 13, + end: 14, + )), + Emit(( + start: 14, + end: 17, + )), + If( + condition: 16, + accept: [ + Emit(( + start: 17, + end: 21, + )), + Store( + pointer: 37, + value: 19, + ), + Store( + pointer: 38, + value: 20, + ), + ], + reject: [ + Emit(( + start: 21, + end: 23, + )), + Store( + pointer: 37, + value: 13, + ), + Store( + pointer: 38, + value: 22, + ), + ], + ), + Emit(( + start: 24, + end: 25, + )), + Emit(( + start: 26, + end: 29, + )), + Switch( + selector: 28, + cases: [ + ( + value: I32(0), + body: [ + Store( + pointer: 39, + value: 7, + ), + Store( + pointer: 40, + value: 6, + ), + Break, + ], + fall_through: false, + ), + ( + value: I32(1), + body: [ + Emit(( + start: 29, + end: 31, + )), + Store( + pointer: 8, + value: 30, + ), + Store( + pointer: 39, + value: 2, + ), + Store( + pointer: 40, + value: 24, + ), + Break, + ], + fall_through: false, + ), + ( + value: Default, + body: [ + Store( + pointer: 39, + value: 7, + ), + Store( + pointer: 40, + value: 6, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 32, + end: 33, + )), + Emit(( + start: 34, + end: 35, + )), + Continue, + ], + continuing: [ + Emit(( + start: 35, + end: 36, + )), + Store( + pointer: 36, + value: 34, + ), + ], + break_if: Some(35), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_load_and_store", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_load_and_store_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_load_and_store.ron b/naga/tests/out/ir/atomic_load_and_store.ron new file mode 100644 index 0000000000..005f23d883 --- /dev/null +++ b/naga/tests/out/ir/atomic_load_and_store.ron @@ -0,0 +1,504 @@ +( + 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, + ), + ( + name: None, + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Sint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 0, + binding: None, + offset: 0, + ), + ], + span: 4, + ), + ), + ( + name: None, + inner: Pointer( + base: 5, + space: Storage( + access: ("LOAD | STORE"), + ), + ), + ), + ( + name: None, + inner: Atomic(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Struct( + members: [ + ( + name: None, + ty: 7, + 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: 3, + init: 4, + ), + ( + name: None, + ty: 1, + init: 5, + ), + ], + overrides: [], + global_variables: [ + ( + name: None, + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 8, + init: None, + ), + ( + name: None, + space: Storage( + access: ("LOAD"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 5, + init: None, + ), + ], + global_expressions: [ + Literal(U32(0)), + Literal(U32(2)), + Literal(Bool(false)), + Literal(U32(1)), + ZeroValue(3), + Literal(Bool(true)), + ], + functions: [ + ( + name: None, + arguments: [], + result: None, + local_variables: [ + ( + name: Some("phi_32"), + ty: 3, + init: None, + ), + ( + name: Some("phi_49"), + ty: 3, + init: None, + ), + ( + name: Some("phi_50"), + ty: 3, + init: None, + ), + ( + name: Some("phi_59"), + ty: 1, + init: None, + ), + ( + name: Some("phi_33"), + ty: 3, + init: None, + ), + ], + expressions: [ + GlobalVariable(0), + GlobalVariable(1), + Constant(5), + Constant(3), + Constant(1), + Constant(0), + Constant(4), + Constant(2), + AccessIndex( + base: 0, + index: 0, + ), + AccessIndex( + base: 1, + index: 0, + ), + Load( + pointer: 9, + ), + Compose( + ty: 3, + components: [ + 5, + 10, + ], + ), + LocalVariable(0), + Load( + pointer: 12, + ), + AccessIndex( + base: 13, + index: 0, + ), + AccessIndex( + base: 13, + index: 1, + ), + Binary( + op: Less, + left: 14, + right: 15, + ), + Binary( + op: Add, + left: 14, + right: 3, + ), + AccessIndex( + base: 13, + index: 0, + ), + AccessIndex( + base: 13, + index: 1, + ), + Compose( + ty: 3, + components: [ + 17, + 19, + ], + ), + Compose( + ty: 3, + components: [ + 3, + 14, + ], + ), + AccessIndex( + base: 6, + index: 0, + ), + AccessIndex( + base: 6, + index: 1, + ), + Compose( + ty: 3, + components: [ + 5, + 23, + ], + ), + LocalVariable(1), + Load( + pointer: 25, + ), + LocalVariable(2), + Load( + pointer: 27, + ), + AccessIndex( + base: 28, + index: 0, + ), + As( + expr: 29, + kind: Sint, + convert: None, + ), + Load( + pointer: 8, + ), + Binary( + op: Add, + left: 31, + right: 4, + ), + LocalVariable(3), + Load( + pointer: 33, + ), + LocalVariable(4), + Load( + pointer: 35, + ), + Unary( + op: LogicalNot, + expr: 34, + ), + LocalVariable(0), + LocalVariable(1), + LocalVariable(2), + LocalVariable(3), + LocalVariable(4), + ], + named_expressions: {}, + body: [ + Emit(( + start: 8, + end: 12, + )), + Store( + pointer: 38, + value: 11, + ), + Loop( + body: [ + Emit(( + start: 13, + end: 14, + )), + Emit(( + start: 14, + end: 17, + )), + If( + condition: 16, + accept: [ + Emit(( + start: 17, + end: 22, + )), + Store( + pointer: 39, + value: 20, + ), + Store( + pointer: 40, + value: 21, + ), + ], + reject: [ + Emit(( + start: 22, + end: 25, + )), + Store( + pointer: 39, + value: 13, + ), + Store( + pointer: 40, + value: 24, + ), + ], + ), + Emit(( + start: 26, + end: 27, + )), + Emit(( + start: 28, + end: 31, + )), + Switch( + selector: 30, + cases: [ + ( + value: I32(0), + body: [ + Store( + pointer: 41, + value: 7, + ), + Store( + pointer: 42, + value: 6, + ), + Break, + ], + fall_through: false, + ), + ( + value: I32(1), + body: [ + Emit(( + start: 31, + end: 33, + )), + Store( + pointer: 8, + value: 32, + ), + Store( + pointer: 41, + value: 2, + ), + Store( + pointer: 42, + value: 26, + ), + Break, + ], + fall_through: false, + ), + ( + value: Default, + body: [ + Store( + pointer: 41, + value: 7, + ), + Store( + pointer: 42, + value: 6, + ), + Break, + ], + fall_through: false, + ), + ], + ), + Emit(( + start: 34, + end: 35, + )), + Emit(( + start: 36, + end: 37, + )), + Continue, + ], + continuing: [ + Emit(( + start: 37, + end: 38, + )), + Store( + pointer: 38, + value: 36, + ), + ], + break_if: Some(37), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "stage::test_atomic_load_and_store", + stage: Compute, + early_depth_test: None, + workgroup_size: (32, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("stage::test_atomic_load_and_store_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/msl/atomic_compare_exchange.msl b/naga/tests/out/msl/atomic_compare_exchange.msl new file mode 100644 index 0000000000..5b0fb97320 --- /dev/null +++ b/naga/tests/out/msl/atomic_compare_exchange.msl @@ -0,0 +1,104 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_2 { + uint member; + uint member_1; +}; +struct type_3 { + uint member; +}; +struct _atomic_compare_exchange_resultUint4_ { + uint old_value; + bool exchanged; +}; +struct type_5 { + metal::atomic_uint member; +}; + +template +_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit( + device A *atomic_ptr, + uint cmp, + uint v +) { + bool swapped = metal::atomic_compare_exchange_weak_explicit( + atomic_ptr, &cmp, v, + metal::memory_order_relaxed, metal::memory_order_relaxed + ); + return _atomic_compare_exchange_resultUint4_{cmp, swapped}; +} +template +_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit( + threadgroup A *atomic_ptr, + uint cmp, + uint v +) { + bool swapped = metal::atomic_compare_exchange_weak_explicit( + atomic_ptr, &cmp, v, + metal::memory_order_relaxed, metal::memory_order_relaxed + ); + return _atomic_compare_exchange_resultUint4_{cmp, swapped}; +} + +void function( + device type_5& global, + device type_3 const& global_1 +) { + type_2 phi_33_ = {}; + type_2 phi_34_ = {}; + type_2 phi_49_ = {}; + bool phi_63_ = {}; + uint _e11 = global_1.member; + phi_33_ = type_2 {0u, _e11}; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_33_ = phi_34_; + if (!(phi_63_)) { + break; + } + } + loop_init = false; + type_2 _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}; + } + type_2 _e25 = phi_34_; + type_2 _e27 = phi_49_; + switch(as_type(_e27.member)) { + case 0: { + phi_63_ = false; + break; + } + case 1: { + _atomic_compare_exchange_resultUint4_ _e31 = naga_atomic_compare_exchange_weak_explicit(&global.member, 3u, _e27.member_1); + phi_63_ = (_e31.old_value == 3u) ? false : true; + break; + } + default: { + phi_63_ = bool {}; + break; + } + } + bool _e36 = phi_63_; + continue; +#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } + LOOP_IS_BOUNDED + } + return; +} + +kernel void stagetest_atomic_compare_exchange( + device type_5& global [[user(fake0)]] +, device type_3 const& global_1 [[user(fake0)]] +) { + function(global, global_1); +} diff --git a/naga/tests/out/msl/atomic_exchange.msl b/naga/tests/out/msl/atomic_exchange.msl new file mode 100644 index 0000000000..7bc1e20c2b --- /dev/null +++ b/naga/tests/out/msl/atomic_exchange.msl @@ -0,0 +1,89 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_2 { + uint member; + uint member_1; +}; +struct type_3 { + uint member; +}; +struct type_5 { + metal::atomic_uint member; +}; + +void function( + device type_5& global, + device type_3 const& global_1 +) { + type_2 phi_33_ = {}; + uint phi_36_ = {}; + type_2 phi_52_ = {}; + type_2 phi_53_ = {}; + bool phi_62_ = {}; + type_2 phi_34_ = {}; + uint phi_37_ = {}; + uint _e10 = global_1.member; + phi_33_ = type_2 {0u, _e10}; + phi_36_ = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_33_ = phi_34_; + phi_36_ = phi_37_; + if (!(phi_62_)) { + break; + } + } + loop_init = false; + type_2 _e13 = phi_33_; + uint _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}; + } + type_2 _e26 = phi_52_; + type_2 _e28 = phi_53_; + switch(as_type(_e28.member)) { + case 0: { + phi_62_ = false; + phi_34_ = type_2 {}; + phi_37_ = uint {}; + break; + } + case 1: { + uint _e31 = metal::atomic_exchange_explicit(&global.member, _e15, metal::memory_order_relaxed); + phi_62_ = true; + phi_34_ = _e26; + phi_37_ = _e15 + _e31; + break; + } + default: { + phi_62_ = false; + phi_34_ = type_2 {}; + phi_37_ = uint {}; + break; + } + } + bool _e34 = phi_62_; + type_2 _e36 = phi_34_; + uint _e38 = phi_37_; + continue; +#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } + LOOP_IS_BOUNDED + } + return; +} + +kernel void stagetest_atomic_exchange( + device type_5& global [[user(fake0)]] +, device type_3 const& global_1 [[user(fake0)]] +) { + function(global, global_1); +} diff --git a/naga/tests/out/msl/atomic_i_add_sub.msl b/naga/tests/out/msl/atomic_i_add_sub.msl new file mode 100644 index 0000000000..31427767d3 --- /dev/null +++ b/naga/tests/out/msl/atomic_i_add_sub.msl @@ -0,0 +1,38 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct _mslBufferSizes { + uint size1; +}; + +typedef uint type_1[1]; +struct type_2 { + type_1 member; +}; +struct type_4 { + metal::atomic_uint member; +}; + +void function( + device type_4& global, + device type_2& global_1, + constant _mslBufferSizes& _buffer_sizes +) { + uint _e6 = metal::atomic_fetch_add_explicit(&global.member, 2u, metal::memory_order_relaxed); + uint _e7 = metal::atomic_fetch_sub_explicit(&global.member, _e6, metal::memory_order_relaxed); + if (_e6 < (1 + (_buffer_sizes.size1 - 0 - 4) / 4)) { + global_1.member[_e6] = _e7; + } + return; +} + +kernel void stagetest_atomic_i_add_sub( + device type_4& global [[user(fake0)]] +, device type_2& global_1 [[user(fake0)]] +, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] +) { + function(global, global_1, _buffer_sizes); +} diff --git a/naga/tests/out/msl/atomic_i_decrement.msl b/naga/tests/out/msl/atomic_i_decrement.msl new file mode 100644 index 0000000000..cc12aa98a7 --- /dev/null +++ b/naga/tests/out/msl/atomic_i_decrement.msl @@ -0,0 +1,54 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct _mslBufferSizes { + uint size1; +}; + +typedef uint type_2[1]; +struct type_3 { + type_2 member; +}; +struct type_5 { + metal::atomic_uint member; +}; + +void function( + device type_5& global, + device type_3& global_1, + constant _mslBufferSizes& _buffer_sizes +) { + bool phi_40_ = {}; + bool loop_init = true; + while(true) { + if (!loop_init) { + if (!(phi_40_)) { + break; + } + } + loop_init = false; + uint _e8 = metal::atomic_fetch_sub_explicit(&global.member, 1u, metal::memory_order_relaxed); + if (_e8 < (1 + (_buffer_sizes.size1 - 0 - 4) / 4)) { + global_1.member[_e8] = _e8; + phi_40_ = (_e8 == 0u) ? false : true; + } else { + phi_40_ = false; + } + bool _e16 = phi_40_; + continue; +#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } + LOOP_IS_BOUNDED + } + return; +} + +kernel void stagetest_atomic_i_decrement( + device type_5& global [[user(fake0)]] +, device type_3& global_1 [[user(fake0)]] +, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] +) { + function(global, global_1, _buffer_sizes); +} diff --git a/naga/tests/out/msl/atomic_i_increment.msl b/naga/tests/out/msl/atomic_i_increment.msl new file mode 100644 index 0000000000..507185da5f --- /dev/null +++ b/naga/tests/out/msl/atomic_i_increment.msl @@ -0,0 +1,52 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_2 { + uint member; +}; +struct type_4 { + metal::atomic_uint member; +}; + +void function( + device type_4& global, + device type_2 const& global_1 +) { + uint phi_23_ = {}; + uint phi_24_ = {}; + phi_23_ = 0u; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_23_ = phi_24_; + if (!(((phi_23_ >= global_1.member) ? false : true))) { + break; + } + } + loop_init = false; + uint _e10 = phi_23_; + uint _e11 = global_1.member; + bool _e12 = _e10 >= _e11; + if (_e12) { + phi_24_ = uint {}; + } else { + uint _e13 = metal::atomic_fetch_add_explicit(&global.member, 1u, metal::memory_order_relaxed); + phi_24_ = _e10 + 1u; + } + uint _e17 = phi_24_; + continue; +#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } + LOOP_IS_BOUNDED + } + return; +} + +kernel void stagetest_atomic_i_increment( + device type_4& global [[user(fake0)]] +, device type_2 const& global_1 [[user(fake0)]] +) { + function(global, global_1); +} diff --git a/naga/tests/out/msl/atomic_load_and_store.msl b/naga/tests/out/msl/atomic_load_and_store.msl new file mode 100644 index 0000000000..327980c5f6 --- /dev/null +++ b/naga/tests/out/msl/atomic_load_and_store.msl @@ -0,0 +1,81 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_2 { + uint member; + uint member_1; +}; +struct type_3 { + uint member; +}; +struct type_5 { + metal::atomic_uint member; +}; + +void function( + device type_5& global, + device type_3 const& global_1 +) { + type_2 phi_32_ = {}; + type_2 phi_49_ = {}; + type_2 phi_50_ = {}; + bool phi_59_ = {}; + type_2 phi_33_ = {}; + uint _e10 = global_1.member; + phi_32_ = type_2 {0u, _e10}; + bool loop_init = true; + while(true) { + if (!loop_init) { + phi_32_ = phi_33_; + if (!(phi_59_)) { + break; + } + } + loop_init = false; + type_2 _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}; + } + type_2 _e24 = phi_49_; + type_2 _e26 = phi_50_; + switch(as_type(_e26.member)) { + case 0: { + phi_59_ = false; + phi_33_ = type_2 {}; + break; + } + case 1: { + uint _e29 = metal::atomic_load_explicit(&global.member, metal::memory_order_relaxed); + metal::atomic_store_explicit(&global.member, _e29 + 2u, metal::memory_order_relaxed); + phi_59_ = true; + phi_33_ = _e24; + break; + } + default: { + phi_59_ = false; + phi_33_ = type_2 {}; + break; + } + } + bool _e32 = phi_59_; + type_2 _e34 = phi_33_; + continue; +#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } + LOOP_IS_BOUNDED + } + return; +} + +kernel void stagetest_atomic_load_and_store( + device type_5& global [[user(fake0)]] +, device type_3 const& global_1 [[user(fake0)]] +) { + function(global, global_1); +} 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(); +} From 1b88bae056b13e9864e5e7195ce50d4f26b6abd5 Mon Sep 17 00:00:00 2001 From: Schell Carl Scivally Date: Tue, 10 Dec 2024 11:16:48 +1300 Subject: [PATCH 3/4] update changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) 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). From 2ee7561a5c3b18c228e1102c44f8a6b36656d660 Mon Sep 17 00:00:00 2001 From: Schell Carl Scivally Date: Tue, 10 Dec 2024 15:36:06 +1300 Subject: [PATCH 4/4] remove extraneous snapshot out files --- .../ir/atomic_compare_exchange.compact.ron | 495 --------------- naga/tests/out/ir/atomic_compare_exchange.ron | 542 ----------------- naga/tests/out/ir/atomic_exchange.compact.ron | 524 ---------------- naga/tests/out/ir/atomic_exchange.ron | 564 ------------------ .../tests/out/ir/atomic_i_add_sub.compact.ron | 207 ------- naga/tests/out/ir/atomic_i_add_sub.ron | 262 -------- .../out/ir/atomic_i_decrement.compact.ron | 273 --------- naga/tests/out/ir/atomic_i_decrement.ron | 321 ---------- .../out/ir/atomic_i_increment.compact.ron | 287 --------- naga/tests/out/ir/atomic_i_increment.ron | 312 ---------- .../out/ir/atomic_load_and_store.compact.ron | 471 --------------- naga/tests/out/ir/atomic_load_and_store.ron | 504 ---------------- .../tests/out/msl/atomic_compare_exchange.msl | 104 ---- naga/tests/out/msl/atomic_exchange.msl | 89 --- naga/tests/out/msl/atomic_i_add_sub.msl | 38 -- naga/tests/out/msl/atomic_i_decrement.msl | 54 -- naga/tests/out/msl/atomic_i_increment.msl | 52 -- naga/tests/out/msl/atomic_load_and_store.msl | 81 --- naga/tests/snapshots.rs | 13 +- 19 files changed, 6 insertions(+), 5187 deletions(-) delete mode 100644 naga/tests/out/ir/atomic_compare_exchange.compact.ron delete mode 100644 naga/tests/out/ir/atomic_compare_exchange.ron delete mode 100644 naga/tests/out/ir/atomic_exchange.compact.ron delete mode 100644 naga/tests/out/ir/atomic_exchange.ron delete mode 100644 naga/tests/out/ir/atomic_i_add_sub.compact.ron delete mode 100644 naga/tests/out/ir/atomic_i_add_sub.ron delete mode 100644 naga/tests/out/ir/atomic_i_decrement.compact.ron delete mode 100644 naga/tests/out/ir/atomic_i_decrement.ron delete mode 100644 naga/tests/out/ir/atomic_i_increment.compact.ron delete mode 100644 naga/tests/out/ir/atomic_i_increment.ron delete mode 100644 naga/tests/out/ir/atomic_load_and_store.compact.ron delete mode 100644 naga/tests/out/ir/atomic_load_and_store.ron delete mode 100644 naga/tests/out/msl/atomic_compare_exchange.msl delete mode 100644 naga/tests/out/msl/atomic_exchange.msl delete mode 100644 naga/tests/out/msl/atomic_i_add_sub.msl delete mode 100644 naga/tests/out/msl/atomic_i_decrement.msl delete mode 100644 naga/tests/out/msl/atomic_i_increment.msl delete mode 100644 naga/tests/out/msl/atomic_load_and_store.msl diff --git a/naga/tests/out/ir/atomic_compare_exchange.compact.ron b/naga/tests/out/ir/atomic_compare_exchange.compact.ron deleted file mode 100644 index aacc81e618..0000000000 --- a/naga/tests/out/ir/atomic_compare_exchange.compact.ron +++ /dev/null @@ -1,495 +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, - ), - ( - name: None, - ty: 0, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 0, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: Some("__atomic_compare_exchange_result"), - inner: Struct( - members: [ - ( - name: Some("old_value"), - ty: 0, - binding: None, - offset: 0, - ), - ( - name: Some("exchanged"), - ty: 1, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - 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: { - AtomicCompareExchangeWeakResult(( - kind: Uint, - width: 4, - )): 4, - }, - ), - constants: [ - ( - name: None, - ty: 0, - init: 0, - ), - ( - name: None, - ty: 1, - init: 1, - ), - ( - name: None, - ty: 0, - init: 2, - ), - ( - name: None, - ty: 2, - init: 3, - ), - ( - name: None, - ty: 0, - init: 4, - ), - ( - name: None, - ty: 1, - init: 5, - ), - ( - name: None, - ty: 1, - init: 6, - ), - ], - 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(Bool(false)), - Literal(U32(1)), - ZeroValue(2), - Literal(U32(3)), - ZeroValue(1), - Literal(Bool(true)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_33"), - ty: 2, - init: None, - ), - ( - name: Some("phi_34"), - ty: 2, - init: None, - ), - ( - name: Some("phi_49"), - ty: 2, - init: None, - ), - ( - name: Some("phi_63"), - ty: 1, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(5), - Constant(3), - Constant(1), - Constant(6), - Constant(4), - Constant(0), - Constant(2), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - Load( - pointer: 10, - ), - Compose( - ty: 2, - components: [ - 7, - 11, - ], - ), - LocalVariable(0), - Load( - pointer: 13, - ), - AccessIndex( - base: 14, - index: 0, - ), - AccessIndex( - base: 14, - index: 1, - ), - Binary( - op: Less, - left: 15, - right: 16, - ), - Binary( - op: Add, - left: 15, - right: 8, - ), - AccessIndex( - base: 14, - index: 1, - ), - Compose( - ty: 2, - components: [ - 18, - 19, - ], - ), - Compose( - ty: 2, - components: [ - 8, - 15, - ], - ), - AccessIndex( - base: 3, - index: 1, - ), - Compose( - ty: 2, - components: [ - 7, - 22, - ], - ), - LocalVariable(1), - Load( - pointer: 24, - ), - LocalVariable(2), - Load( - pointer: 26, - ), - AccessIndex( - base: 27, - index: 0, - ), - AccessIndex( - base: 27, - index: 1, - ), - As( - expr: 28, - kind: Sint, - convert: None, - ), - AtomicResult( - ty: 4, - comparison: true, - ), - AccessIndex( - base: 31, - index: 0, - ), - Binary( - op: Equal, - left: 32, - right: 6, - ), - Select( - condition: 33, - accept: 4, - reject: 5, - ), - LocalVariable(3), - Load( - pointer: 35, - ), - Unary( - op: LogicalNot, - expr: 36, - ), - LocalVariable(0), - LocalVariable(1), - LocalVariable(2), - LocalVariable(3), - ], - named_expressions: {}, - body: [ - Emit(( - start: 9, - end: 13, - )), - Store( - pointer: 38, - value: 12, - ), - Loop( - body: [ - Emit(( - start: 14, - end: 15, - )), - Emit(( - start: 15, - end: 18, - )), - If( - condition: 17, - accept: [ - Emit(( - start: 18, - end: 22, - )), - Store( - pointer: 39, - value: 20, - ), - Store( - pointer: 40, - value: 21, - ), - ], - reject: [ - Emit(( - start: 22, - end: 24, - )), - Store( - pointer: 39, - value: 14, - ), - Store( - pointer: 40, - value: 23, - ), - ], - ), - Emit(( - start: 25, - end: 26, - )), - Emit(( - start: 27, - end: 31, - )), - Switch( - selector: 30, - cases: [ - ( - value: I32(0), - body: [ - Store( - pointer: 41, - value: 4, - ), - Break, - ], - fall_through: false, - ), - ( - value: I32(1), - body: [ - Atomic( - pointer: 9, - fun: Exchange( - compare: Some(6), - ), - value: 29, - result: Some(31), - ), - Emit(( - start: 33, - end: 35, - )), - Store( - pointer: 41, - value: 34, - ), - Break, - ], - fall_through: false, - ), - ( - value: Default, - body: [ - Store( - pointer: 41, - value: 2, - ), - Break, - ], - fall_through: false, - ), - ], - ), - Emit(( - start: 36, - end: 37, - )), - Continue, - ], - continuing: [ - Emit(( - start: 37, - end: 38, - )), - Store( - pointer: 38, - value: 25, - ), - ], - break_if: Some(37), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_compare_exchange", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_compare_exchange_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_compare_exchange.ron b/naga/tests/out/ir/atomic_compare_exchange.ron deleted file mode 100644 index 4192ae8918..0000000000 --- a/naga/tests/out/ir/atomic_compare_exchange.ron +++ /dev/null @@ -1,542 +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, - ), - ( - name: None, - ty: 0, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - name: None, - inner: Scalar(( - kind: Sint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 0, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Pointer( - base: 5, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - name: Some("__atomic_compare_exchange_result"), - inner: Struct( - members: [ - ( - name: Some("old_value"), - ty: 0, - binding: None, - offset: 0, - ), - ( - name: Some("exchanged"), - ty: 1, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 8, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ], - special_types: ( - ray_desc: None, - ray_intersection: None, - predeclared_types: { - AtomicCompareExchangeWeakResult(( - kind: Uint, - width: 4, - )): 7, - }, - ), - 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: 3, - init: 4, - ), - ( - name: None, - ty: 0, - init: 5, - ), - ( - name: None, - ty: 1, - init: 6, - ), - ( - name: None, - ty: 1, - init: 7, - ), - ( - name: None, - ty: 0, - init: 8, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 9, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 5, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(U32(2)), - Literal(Bool(false)), - Literal(U32(1)), - ZeroValue(3), - Literal(U32(3)), - ZeroValue(1), - Literal(Bool(true)), - Literal(U32(256)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_33"), - ty: 3, - init: None, - ), - ( - name: Some("phi_34"), - ty: 3, - init: None, - ), - ( - name: Some("phi_49"), - ty: 3, - init: None, - ), - ( - name: Some("phi_63"), - ty: 1, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(1), - Constant(8), - Constant(6), - Constant(4), - Constant(2), - Constant(7), - Constant(5), - Constant(0), - Constant(3), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - Load( - pointer: 12, - ), - Compose( - ty: 3, - components: [ - 9, - 13, - ], - ), - LocalVariable(0), - Load( - pointer: 15, - ), - AccessIndex( - base: 16, - index: 0, - ), - AccessIndex( - base: 16, - index: 1, - ), - Binary( - op: Less, - left: 17, - right: 18, - ), - Binary( - op: Add, - left: 17, - right: 10, - ), - AccessIndex( - base: 16, - index: 0, - ), - AccessIndex( - base: 16, - index: 1, - ), - Compose( - ty: 3, - components: [ - 20, - 22, - ], - ), - Compose( - ty: 3, - components: [ - 10, - 17, - ], - ), - AccessIndex( - base: 5, - index: 0, - ), - AccessIndex( - base: 5, - index: 1, - ), - Compose( - ty: 3, - components: [ - 9, - 26, - ], - ), - LocalVariable(1), - Load( - pointer: 28, - ), - LocalVariable(2), - Load( - pointer: 30, - ), - AccessIndex( - base: 31, - index: 0, - ), - AccessIndex( - base: 31, - index: 1, - ), - As( - expr: 32, - kind: Sint, - convert: None, - ), - AtomicResult( - ty: 7, - comparison: true, - ), - AccessIndex( - base: 35, - index: 0, - ), - Binary( - op: Equal, - left: 36, - right: 8, - ), - Select( - condition: 37, - accept: 6, - reject: 7, - ), - LocalVariable(3), - Load( - pointer: 39, - ), - Unary( - op: LogicalNot, - expr: 40, - ), - LocalVariable(0), - LocalVariable(1), - LocalVariable(2), - LocalVariable(3), - ], - named_expressions: {}, - body: [ - Emit(( - start: 11, - end: 15, - )), - Store( - pointer: 42, - value: 14, - ), - Loop( - body: [ - Emit(( - start: 16, - end: 17, - )), - Emit(( - start: 17, - end: 20, - )), - If( - condition: 19, - accept: [ - Emit(( - start: 20, - end: 25, - )), - Store( - pointer: 43, - value: 23, - ), - Store( - pointer: 44, - value: 24, - ), - ], - reject: [ - Emit(( - start: 25, - end: 28, - )), - Store( - pointer: 43, - value: 16, - ), - Store( - pointer: 44, - value: 27, - ), - ], - ), - Emit(( - start: 29, - end: 30, - )), - Emit(( - start: 31, - end: 35, - )), - Switch( - selector: 34, - cases: [ - ( - value: I32(0), - body: [ - Store( - pointer: 45, - value: 6, - ), - Break, - ], - fall_through: false, - ), - ( - value: I32(1), - body: [ - Atomic( - pointer: 11, - fun: Exchange( - compare: Some(8), - ), - value: 33, - result: Some(35), - ), - Emit(( - start: 37, - end: 39, - )), - Store( - pointer: 45, - value: 38, - ), - Break, - ], - fall_through: false, - ), - ( - value: Default, - body: [ - Store( - pointer: 45, - value: 4, - ), - Break, - ], - fall_through: false, - ), - ], - ), - Emit(( - start: 40, - end: 41, - )), - Continue, - ], - continuing: [ - Emit(( - start: 41, - end: 42, - )), - Store( - pointer: 42, - value: 29, - ), - ], - break_if: Some(41), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_compare_exchange", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_compare_exchange_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_exchange.compact.ron b/naga/tests/out/ir/atomic_exchange.compact.ron deleted file mode 100644 index 192bd6f913..0000000000 --- a/naga/tests/out/ir/atomic_exchange.compact.ron +++ /dev/null @@ -1,524 +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, - ), - ( - name: None, - ty: 0, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - 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: 4, - 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: 2, - 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: 5, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 3, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(Bool(false)), - Literal(U32(1)), - ZeroValue(2), - Literal(Bool(true)), - ZeroValue(0), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_33"), - ty: 2, - init: None, - ), - ( - name: Some("phi_36"), - ty: 0, - init: None, - ), - ( - name: Some("phi_52"), - ty: 2, - init: None, - ), - ( - name: Some("phi_53"), - ty: 2, - init: None, - ), - ( - name: Some("phi_62"), - ty: 1, - init: None, - ), - ( - name: Some("phi_34"), - ty: 2, - init: None, - ), - ( - name: Some("phi_37"), - ty: 0, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(4), - Constant(2), - Constant(5), - Constant(0), - Constant(3), - Constant(1), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - Load( - pointer: 9, - ), - Compose( - ty: 2, - components: [ - 5, - 10, - ], - ), - LocalVariable(0), - Load( - pointer: 12, - ), - LocalVariable(1), - Load( - pointer: 14, - ), - AccessIndex( - base: 13, - index: 0, - ), - AccessIndex( - base: 13, - index: 1, - ), - Binary( - op: Less, - left: 16, - right: 17, - ), - Binary( - op: Add, - left: 16, - right: 3, - ), - AccessIndex( - base: 13, - index: 1, - ), - Compose( - ty: 2, - components: [ - 19, - 20, - ], - ), - Compose( - ty: 2, - components: [ - 3, - 16, - ], - ), - AccessIndex( - base: 6, - index: 1, - ), - Compose( - ty: 2, - components: [ - 5, - 23, - ], - ), - LocalVariable(2), - Load( - pointer: 25, - ), - LocalVariable(3), - Load( - pointer: 27, - ), - AccessIndex( - base: 28, - index: 0, - ), - As( - expr: 29, - kind: Sint, - convert: None, - ), - AtomicResult( - ty: 0, - comparison: false, - ), - Binary( - op: Add, - left: 15, - right: 31, - ), - LocalVariable(4), - Load( - pointer: 33, - ), - LocalVariable(5), - Load( - pointer: 35, - ), - LocalVariable(6), - Load( - pointer: 37, - ), - Unary( - op: LogicalNot, - expr: 34, - ), - LocalVariable(0), - LocalVariable(1), - LocalVariable(2), - LocalVariable(3), - LocalVariable(4), - LocalVariable(5), - LocalVariable(6), - ], - named_expressions: {}, - body: [ - Emit(( - start: 8, - end: 12, - )), - Store( - pointer: 40, - value: 11, - ), - Store( - pointer: 41, - value: 5, - ), - Loop( - body: [ - Emit(( - start: 13, - end: 14, - )), - Emit(( - start: 15, - end: 16, - )), - Emit(( - start: 16, - end: 19, - )), - If( - condition: 18, - accept: [ - Emit(( - start: 19, - end: 23, - )), - Store( - pointer: 42, - value: 21, - ), - Store( - pointer: 43, - value: 22, - ), - ], - reject: [ - Emit(( - start: 23, - end: 25, - )), - Store( - pointer: 42, - value: 13, - ), - Store( - pointer: 43, - value: 24, - ), - ], - ), - Emit(( - start: 26, - end: 27, - )), - Emit(( - start: 28, - end: 31, - )), - Switch( - selector: 30, - cases: [ - ( - value: I32(0), - body: [ - Store( - pointer: 44, - value: 7, - ), - Store( - pointer: 45, - value: 6, - ), - Store( - pointer: 46, - value: 4, - ), - Break, - ], - fall_through: false, - ), - ( - value: I32(1), - body: [ - Atomic( - pointer: 8, - fun: Exchange( - compare: None, - ), - value: 15, - result: Some(31), - ), - Emit(( - start: 32, - end: 33, - )), - Store( - pointer: 44, - value: 2, - ), - Store( - pointer: 45, - value: 26, - ), - Store( - pointer: 46, - value: 32, - ), - Break, - ], - fall_through: false, - ), - ( - value: Default, - body: [ - Store( - pointer: 44, - value: 7, - ), - Store( - pointer: 45, - value: 6, - ), - Store( - pointer: 46, - value: 4, - ), - Break, - ], - fall_through: false, - ), - ], - ), - Emit(( - start: 34, - end: 35, - )), - Emit(( - start: 36, - end: 37, - )), - Emit(( - start: 38, - end: 39, - )), - Continue, - ], - continuing: [ - Emit(( - start: 39, - end: 40, - )), - Store( - pointer: 40, - value: 36, - ), - Store( - pointer: 41, - value: 38, - ), - ], - break_if: Some(39), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_exchange", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_exchange_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_exchange.ron b/naga/tests/out/ir/atomic_exchange.ron deleted file mode 100644 index 8a96c1c1b4..0000000000 --- a/naga/tests/out/ir/atomic_exchange.ron +++ /dev/null @@ -1,564 +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, - ), - ( - name: None, - ty: 0, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - name: None, - inner: Scalar(( - kind: Sint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 0, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Pointer( - base: 5, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 7, - 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: 3, - init: 4, - ), - ( - name: None, - ty: 1, - init: 5, - ), - ( - name: None, - ty: 0, - init: 6, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 8, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 5, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(U32(2)), - Literal(Bool(false)), - Literal(U32(1)), - ZeroValue(3), - Literal(Bool(true)), - ZeroValue(0), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_33"), - ty: 3, - init: None, - ), - ( - name: Some("phi_36"), - ty: 0, - init: None, - ), - ( - name: Some("phi_52"), - ty: 3, - init: None, - ), - ( - name: Some("phi_53"), - ty: 3, - init: None, - ), - ( - name: Some("phi_62"), - ty: 1, - init: None, - ), - ( - name: Some("phi_34"), - ty: 3, - init: None, - ), - ( - name: Some("phi_37"), - ty: 0, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(5), - Constant(3), - Constant(1), - Constant(6), - Constant(0), - Constant(4), - Constant(2), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - Load( - pointer: 10, - ), - Compose( - ty: 3, - components: [ - 6, - 11, - ], - ), - LocalVariable(0), - Load( - pointer: 13, - ), - LocalVariable(1), - Load( - pointer: 15, - ), - AccessIndex( - base: 14, - index: 0, - ), - AccessIndex( - base: 14, - index: 1, - ), - Binary( - op: Less, - left: 17, - right: 18, - ), - Binary( - op: Add, - left: 17, - right: 3, - ), - AccessIndex( - base: 14, - index: 0, - ), - AccessIndex( - base: 14, - index: 1, - ), - Compose( - ty: 3, - components: [ - 20, - 22, - ], - ), - Compose( - ty: 3, - components: [ - 3, - 17, - ], - ), - AccessIndex( - base: 7, - index: 0, - ), - AccessIndex( - base: 7, - index: 1, - ), - Compose( - ty: 3, - components: [ - 6, - 26, - ], - ), - LocalVariable(2), - Load( - pointer: 28, - ), - LocalVariable(3), - Load( - pointer: 30, - ), - AccessIndex( - base: 31, - index: 0, - ), - As( - expr: 32, - kind: Sint, - convert: None, - ), - AtomicResult( - ty: 0, - comparison: false, - ), - Binary( - op: Add, - left: 16, - right: 34, - ), - LocalVariable(4), - Load( - pointer: 36, - ), - LocalVariable(5), - Load( - pointer: 38, - ), - LocalVariable(6), - Load( - pointer: 40, - ), - Unary( - op: LogicalNot, - expr: 37, - ), - LocalVariable(0), - LocalVariable(1), - LocalVariable(2), - LocalVariable(3), - LocalVariable(4), - LocalVariable(5), - LocalVariable(6), - ], - named_expressions: {}, - body: [ - Emit(( - start: 9, - end: 13, - )), - Store( - pointer: 43, - value: 12, - ), - Store( - pointer: 44, - value: 6, - ), - Loop( - body: [ - Emit(( - start: 14, - end: 15, - )), - Emit(( - start: 16, - end: 17, - )), - Emit(( - start: 17, - end: 20, - )), - If( - condition: 19, - accept: [ - Emit(( - start: 20, - end: 25, - )), - Store( - pointer: 45, - value: 23, - ), - Store( - pointer: 46, - value: 24, - ), - ], - reject: [ - Emit(( - start: 25, - end: 28, - )), - Store( - pointer: 45, - value: 14, - ), - Store( - pointer: 46, - value: 27, - ), - ], - ), - Emit(( - start: 29, - end: 30, - )), - Emit(( - start: 31, - end: 34, - )), - Switch( - selector: 33, - cases: [ - ( - value: I32(0), - body: [ - Store( - pointer: 47, - value: 8, - ), - Store( - pointer: 48, - value: 7, - ), - Store( - pointer: 49, - value: 5, - ), - Break, - ], - fall_through: false, - ), - ( - value: I32(1), - body: [ - Atomic( - pointer: 9, - fun: Exchange( - compare: None, - ), - value: 16, - result: Some(34), - ), - Emit(( - start: 35, - end: 36, - )), - Store( - pointer: 47, - value: 2, - ), - Store( - pointer: 48, - value: 29, - ), - Store( - pointer: 49, - value: 35, - ), - Break, - ], - fall_through: false, - ), - ( - value: Default, - body: [ - Store( - pointer: 47, - value: 8, - ), - Store( - pointer: 48, - value: 7, - ), - Store( - pointer: 49, - value: 5, - ), - Break, - ], - fall_through: false, - ), - ], - ), - Emit(( - start: 37, - end: 38, - )), - Emit(( - start: 39, - end: 40, - )), - Emit(( - start: 41, - end: 42, - )), - Continue, - ], - continuing: [ - Emit(( - start: 42, - end: 43, - )), - Store( - pointer: 43, - value: 39, - ), - Store( - pointer: 44, - value: 41, - ), - ], - break_if: Some(42), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_exchange", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_exchange_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_add_sub.compact.ron b/naga/tests/out/ir/atomic_i_add_sub.compact.ron deleted file mode 100644 index 52702a1828..0000000000 --- a/naga/tests/out/ir/atomic_i_add_sub.compact.ron +++ /dev/null @@ -1,207 +0,0 @@ -( - types: [ - ( - name: None, - inner: Scalar(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Array( - base: 0, - size: Dynamic, - stride: 4, - ), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 1, - 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, - ), - ], - 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 | STORE"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 2, - init: None, - ), - ], - global_expressions: [ - Literal(U32(2)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(0), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - ArrayLength(4), - AtomicResult( - ty: 0, - comparison: false, - ), - AtomicResult( - ty: 0, - comparison: false, - ), - Binary( - op: Less, - left: 6, - right: 5, - ), - AccessIndex( - base: 1, - index: 0, - ), - Access( - base: 9, - index: 6, - ), - ], - named_expressions: {}, - body: [ - Emit(( - start: 3, - end: 6, - )), - Atomic( - pointer: 3, - fun: Add, - value: 2, - result: Some(6), - ), - Atomic( - pointer: 3, - fun: Subtract, - value: 6, - result: Some(7), - ), - Emit(( - start: 8, - end: 9, - )), - If( - condition: 8, - accept: [ - Emit(( - start: 9, - end: 11, - )), - Store( - pointer: 10, - value: 7, - ), - ], - reject: [], - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_i_add_sub", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_i_add_sub_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_add_sub.ron b/naga/tests/out/ir/atomic_i_add_sub.ron deleted file mode 100644 index e8d41b6889..0000000000 --- a/naga/tests/out/ir/atomic_i_add_sub.ron +++ /dev/null @@ -1,262 +0,0 @@ -( - types: [ - ( - name: None, - inner: Scalar(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Scalar(( - kind: Bool, - width: 1, - )), - ), - ( - name: None, - inner: Array( - base: 0, - size: Dynamic, - stride: 4, - ), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 2, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Pointer( - base: 3, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - 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: 6, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 8, - 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, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 9, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 3, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(U32(2)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(1), - Constant(0), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - ArrayLength(5), - AtomicResult( - ty: 0, - comparison: false, - ), - AtomicResult( - ty: 0, - comparison: false, - ), - Binary( - op: Less, - left: 7, - right: 6, - ), - AccessIndex( - base: 1, - index: 0, - ), - Access( - base: 10, - index: 7, - ), - ], - named_expressions: {}, - body: [ - Emit(( - start: 4, - end: 7, - )), - Atomic( - pointer: 4, - fun: Add, - value: 2, - result: Some(7), - ), - Atomic( - pointer: 4, - fun: Subtract, - value: 7, - result: Some(8), - ), - Emit(( - start: 9, - end: 10, - )), - If( - condition: 9, - accept: [ - Emit(( - start: 10, - end: 12, - )), - Store( - pointer: 11, - value: 8, - ), - ], - reject: [], - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_i_add_sub", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_i_add_sub_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_decrement.compact.ron b/naga/tests/out/ir/atomic_i_decrement.compact.ron deleted file mode 100644 index 5fec93a946..0000000000 --- a/naga/tests/out/ir/atomic_i_decrement.compact.ron +++ /dev/null @@ -1,273 +0,0 @@ -( - types: [ - ( - name: None, - inner: Scalar(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Scalar(( - kind: Bool, - width: 1, - )), - ), - ( - name: None, - inner: Array( - base: 0, - size: Dynamic, - stride: 4, - ), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 2, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 4, - 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: 1, - init: 2, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 5, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 3, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(Bool(false)), - Literal(Bool(true)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_40"), - ty: 1, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(0), - Constant(1), - Constant(2), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - ArrayLength(6), - AtomicResult( - ty: 0, - comparison: false, - ), - Literal(U32(1)), - Binary( - op: Less, - left: 8, - right: 7, - ), - AccessIndex( - base: 1, - index: 0, - ), - Access( - base: 11, - index: 8, - ), - Binary( - op: Equal, - left: 8, - right: 2, - ), - Select( - condition: 13, - accept: 3, - reject: 4, - ), - LocalVariable(0), - Load( - pointer: 15, - ), - Unary( - op: LogicalNot, - expr: 16, - ), - LocalVariable(0), - ], - named_expressions: {}, - body: [ - Emit(( - start: 5, - end: 8, - )), - Loop( - body: [ - Atomic( - pointer: 5, - fun: Subtract, - value: 9, - result: Some(8), - ), - Emit(( - start: 10, - end: 11, - )), - If( - condition: 10, - accept: [ - Emit(( - start: 11, - end: 13, - )), - Store( - pointer: 12, - value: 8, - ), - Emit(( - start: 13, - end: 15, - )), - Store( - pointer: 18, - value: 14, - ), - ], - reject: [ - Store( - pointer: 18, - value: 3, - ), - ], - ), - Emit(( - start: 16, - end: 17, - )), - Continue, - ], - continuing: [ - Emit(( - start: 17, - end: 18, - )), - ], - break_if: Some(17), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_i_decrement", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_i_decrement_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_decrement.ron b/naga/tests/out/ir/atomic_i_decrement.ron deleted file mode 100644 index 59be516b93..0000000000 --- a/naga/tests/out/ir/atomic_i_decrement.ron +++ /dev/null @@ -1,321 +0,0 @@ -( - types: [ - ( - name: None, - inner: Scalar(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Scalar(( - kind: Bool, - width: 1, - )), - ), - ( - name: None, - inner: Array( - base: 0, - size: Dynamic, - stride: 4, - ), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 2, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Pointer( - base: 3, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - 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: 6, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 8, - 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: 1, - init: 3, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 9, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 3, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(U32(2)), - Literal(Bool(false)), - Literal(Bool(true)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_40"), - ty: 1, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(1), - Constant(0), - Constant(2), - Constant(3), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - ArrayLength(7), - AtomicResult( - ty: 0, - comparison: false, - ), - Literal(U32(1)), - Binary( - op: Less, - left: 9, - right: 8, - ), - AccessIndex( - base: 1, - index: 0, - ), - Access( - base: 12, - index: 9, - ), - Binary( - op: Equal, - left: 9, - right: 3, - ), - Select( - condition: 14, - accept: 4, - reject: 5, - ), - LocalVariable(0), - Load( - pointer: 16, - ), - Unary( - op: LogicalNot, - expr: 17, - ), - LocalVariable(0), - ], - named_expressions: {}, - body: [ - Emit(( - start: 6, - end: 9, - )), - Loop( - body: [ - Atomic( - pointer: 6, - fun: Subtract, - value: 10, - result: Some(9), - ), - Emit(( - start: 11, - end: 12, - )), - If( - condition: 11, - accept: [ - Emit(( - start: 12, - end: 14, - )), - Store( - pointer: 13, - value: 9, - ), - Emit(( - start: 14, - end: 16, - )), - Store( - pointer: 19, - value: 15, - ), - ], - reject: [ - Store( - pointer: 19, - value: 4, - ), - ], - ), - Emit(( - start: 17, - end: 18, - )), - Continue, - ], - continuing: [ - Emit(( - start: 18, - end: 19, - )), - ], - break_if: Some(18), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_i_decrement", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_i_decrement_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.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/ir/atomic_load_and_store.compact.ron b/naga/tests/out/ir/atomic_load_and_store.compact.ron deleted file mode 100644 index 6e5d94dde3..0000000000 --- a/naga/tests/out/ir/atomic_load_and_store.compact.ron +++ /dev/null @@ -1,471 +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, - ), - ( - name: None, - ty: 0, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - 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: 4, - 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: 2, - init: 4, - ), - ( - name: None, - ty: 1, - init: 5, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 5, - 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)), - ZeroValue(2), - Literal(Bool(true)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_32"), - ty: 2, - init: None, - ), - ( - name: Some("phi_49"), - ty: 2, - init: None, - ), - ( - name: Some("phi_50"), - ty: 2, - init: None, - ), - ( - name: Some("phi_59"), - ty: 1, - init: None, - ), - ( - name: Some("phi_33"), - ty: 2, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(5), - Constant(3), - Constant(1), - Constant(0), - Constant(4), - Constant(2), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - Load( - pointer: 9, - ), - Compose( - ty: 2, - components: [ - 5, - 10, - ], - ), - LocalVariable(0), - Load( - pointer: 12, - ), - AccessIndex( - base: 13, - index: 0, - ), - AccessIndex( - base: 13, - index: 1, - ), - Binary( - op: Less, - left: 14, - right: 15, - ), - Binary( - op: Add, - left: 14, - right: 3, - ), - AccessIndex( - base: 13, - index: 1, - ), - Compose( - ty: 2, - components: [ - 17, - 18, - ], - ), - Compose( - ty: 2, - components: [ - 3, - 14, - ], - ), - AccessIndex( - base: 6, - index: 1, - ), - Compose( - ty: 2, - components: [ - 5, - 21, - ], - ), - LocalVariable(1), - Load( - pointer: 23, - ), - LocalVariable(2), - Load( - pointer: 25, - ), - AccessIndex( - base: 26, - index: 0, - ), - As( - expr: 27, - kind: Sint, - convert: None, - ), - Load( - pointer: 8, - ), - Binary( - op: Add, - left: 29, - right: 4, - ), - LocalVariable(3), - Load( - pointer: 31, - ), - LocalVariable(4), - Load( - pointer: 33, - ), - Unary( - op: LogicalNot, - expr: 32, - ), - LocalVariable(0), - LocalVariable(1), - LocalVariable(2), - LocalVariable(3), - LocalVariable(4), - ], - named_expressions: {}, - body: [ - Emit(( - start: 8, - end: 12, - )), - Store( - pointer: 36, - value: 11, - ), - Loop( - body: [ - Emit(( - start: 13, - end: 14, - )), - Emit(( - start: 14, - end: 17, - )), - If( - condition: 16, - accept: [ - Emit(( - start: 17, - end: 21, - )), - Store( - pointer: 37, - value: 19, - ), - Store( - pointer: 38, - value: 20, - ), - ], - reject: [ - Emit(( - start: 21, - end: 23, - )), - Store( - pointer: 37, - value: 13, - ), - Store( - pointer: 38, - value: 22, - ), - ], - ), - Emit(( - start: 24, - end: 25, - )), - Emit(( - start: 26, - end: 29, - )), - Switch( - selector: 28, - cases: [ - ( - value: I32(0), - body: [ - Store( - pointer: 39, - value: 7, - ), - Store( - pointer: 40, - value: 6, - ), - Break, - ], - fall_through: false, - ), - ( - value: I32(1), - body: [ - Emit(( - start: 29, - end: 31, - )), - Store( - pointer: 8, - value: 30, - ), - Store( - pointer: 39, - value: 2, - ), - Store( - pointer: 40, - value: 24, - ), - Break, - ], - fall_through: false, - ), - ( - value: Default, - body: [ - Store( - pointer: 39, - value: 7, - ), - Store( - pointer: 40, - value: 6, - ), - Break, - ], - fall_through: false, - ), - ], - ), - Emit(( - start: 32, - end: 33, - )), - Emit(( - start: 34, - end: 35, - )), - Continue, - ], - continuing: [ - Emit(( - start: 35, - end: 36, - )), - Store( - pointer: 36, - value: 34, - ), - ], - break_if: Some(35), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_load_and_store", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_load_and_store_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_load_and_store.ron b/naga/tests/out/ir/atomic_load_and_store.ron deleted file mode 100644 index 005f23d883..0000000000 --- a/naga/tests/out/ir/atomic_load_and_store.ron +++ /dev/null @@ -1,504 +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, - ), - ( - name: None, - ty: 0, - binding: None, - offset: 4, - ), - ], - span: 8, - ), - ), - ( - name: None, - inner: Scalar(( - kind: Sint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 0, - binding: None, - offset: 0, - ), - ], - span: 4, - ), - ), - ( - name: None, - inner: Pointer( - base: 5, - space: Storage( - access: ("LOAD | STORE"), - ), - ), - ), - ( - name: None, - inner: Atomic(( - kind: Uint, - width: 4, - )), - ), - ( - name: None, - inner: Struct( - members: [ - ( - name: None, - ty: 7, - 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: 3, - init: 4, - ), - ( - name: None, - ty: 1, - init: 5, - ), - ], - overrides: [], - global_variables: [ - ( - name: None, - space: Storage( - access: ("LOAD | STORE"), - ), - binding: Some(( - group: 0, - binding: 0, - )), - ty: 8, - init: None, - ), - ( - name: None, - space: Storage( - access: ("LOAD"), - ), - binding: Some(( - group: 0, - binding: 1, - )), - ty: 5, - init: None, - ), - ], - global_expressions: [ - Literal(U32(0)), - Literal(U32(2)), - Literal(Bool(false)), - Literal(U32(1)), - ZeroValue(3), - Literal(Bool(true)), - ], - functions: [ - ( - name: None, - arguments: [], - result: None, - local_variables: [ - ( - name: Some("phi_32"), - ty: 3, - init: None, - ), - ( - name: Some("phi_49"), - ty: 3, - init: None, - ), - ( - name: Some("phi_50"), - ty: 3, - init: None, - ), - ( - name: Some("phi_59"), - ty: 1, - init: None, - ), - ( - name: Some("phi_33"), - ty: 3, - init: None, - ), - ], - expressions: [ - GlobalVariable(0), - GlobalVariable(1), - Constant(5), - Constant(3), - Constant(1), - Constant(0), - Constant(4), - Constant(2), - AccessIndex( - base: 0, - index: 0, - ), - AccessIndex( - base: 1, - index: 0, - ), - Load( - pointer: 9, - ), - Compose( - ty: 3, - components: [ - 5, - 10, - ], - ), - LocalVariable(0), - Load( - pointer: 12, - ), - AccessIndex( - base: 13, - index: 0, - ), - AccessIndex( - base: 13, - index: 1, - ), - Binary( - op: Less, - left: 14, - right: 15, - ), - Binary( - op: Add, - left: 14, - right: 3, - ), - AccessIndex( - base: 13, - index: 0, - ), - AccessIndex( - base: 13, - index: 1, - ), - Compose( - ty: 3, - components: [ - 17, - 19, - ], - ), - Compose( - ty: 3, - components: [ - 3, - 14, - ], - ), - AccessIndex( - base: 6, - index: 0, - ), - AccessIndex( - base: 6, - index: 1, - ), - Compose( - ty: 3, - components: [ - 5, - 23, - ], - ), - LocalVariable(1), - Load( - pointer: 25, - ), - LocalVariable(2), - Load( - pointer: 27, - ), - AccessIndex( - base: 28, - index: 0, - ), - As( - expr: 29, - kind: Sint, - convert: None, - ), - Load( - pointer: 8, - ), - Binary( - op: Add, - left: 31, - right: 4, - ), - LocalVariable(3), - Load( - pointer: 33, - ), - LocalVariable(4), - Load( - pointer: 35, - ), - Unary( - op: LogicalNot, - expr: 34, - ), - LocalVariable(0), - LocalVariable(1), - LocalVariable(2), - LocalVariable(3), - LocalVariable(4), - ], - named_expressions: {}, - body: [ - Emit(( - start: 8, - end: 12, - )), - Store( - pointer: 38, - value: 11, - ), - Loop( - body: [ - Emit(( - start: 13, - end: 14, - )), - Emit(( - start: 14, - end: 17, - )), - If( - condition: 16, - accept: [ - Emit(( - start: 17, - end: 22, - )), - Store( - pointer: 39, - value: 20, - ), - Store( - pointer: 40, - value: 21, - ), - ], - reject: [ - Emit(( - start: 22, - end: 25, - )), - Store( - pointer: 39, - value: 13, - ), - Store( - pointer: 40, - value: 24, - ), - ], - ), - Emit(( - start: 26, - end: 27, - )), - Emit(( - start: 28, - end: 31, - )), - Switch( - selector: 30, - cases: [ - ( - value: I32(0), - body: [ - Store( - pointer: 41, - value: 7, - ), - Store( - pointer: 42, - value: 6, - ), - Break, - ], - fall_through: false, - ), - ( - value: I32(1), - body: [ - Emit(( - start: 31, - end: 33, - )), - Store( - pointer: 8, - value: 32, - ), - Store( - pointer: 41, - value: 2, - ), - Store( - pointer: 42, - value: 26, - ), - Break, - ], - fall_through: false, - ), - ( - value: Default, - body: [ - Store( - pointer: 41, - value: 7, - ), - Store( - pointer: 42, - value: 6, - ), - Break, - ], - fall_through: false, - ), - ], - ), - Emit(( - start: 34, - end: 35, - )), - Emit(( - start: 36, - end: 37, - )), - Continue, - ], - continuing: [ - Emit(( - start: 37, - end: 38, - )), - Store( - pointer: 38, - value: 36, - ), - ], - break_if: Some(37), - ), - Return( - value: None, - ), - ], - diagnostic_filter_leaf: None, - ), - ], - entry_points: [ - ( - name: "stage::test_atomic_load_and_store", - stage: Compute, - early_depth_test: None, - workgroup_size: (32, 1, 1), - workgroup_size_overrides: None, - function: ( - name: Some("stage::test_atomic_load_and_store_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/msl/atomic_compare_exchange.msl b/naga/tests/out/msl/atomic_compare_exchange.msl deleted file mode 100644 index 5b0fb97320..0000000000 --- a/naga/tests/out/msl/atomic_compare_exchange.msl +++ /dev/null @@ -1,104 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; - -struct type_2 { - uint member; - uint member_1; -}; -struct type_3 { - uint member; -}; -struct _atomic_compare_exchange_resultUint4_ { - uint old_value; - bool exchanged; -}; -struct type_5 { - metal::atomic_uint member; -}; - -template -_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit( - device A *atomic_ptr, - uint cmp, - uint v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_resultUint4_{cmp, swapped}; -} -template -_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit( - threadgroup A *atomic_ptr, - uint cmp, - uint v -) { - bool swapped = metal::atomic_compare_exchange_weak_explicit( - atomic_ptr, &cmp, v, - metal::memory_order_relaxed, metal::memory_order_relaxed - ); - return _atomic_compare_exchange_resultUint4_{cmp, swapped}; -} - -void function( - device type_5& global, - device type_3 const& global_1 -) { - type_2 phi_33_ = {}; - type_2 phi_34_ = {}; - type_2 phi_49_ = {}; - bool phi_63_ = {}; - uint _e11 = global_1.member; - phi_33_ = type_2 {0u, _e11}; - bool loop_init = true; - while(true) { - if (!loop_init) { - phi_33_ = phi_34_; - if (!(phi_63_)) { - break; - } - } - loop_init = false; - type_2 _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}; - } - type_2 _e25 = phi_34_; - type_2 _e27 = phi_49_; - switch(as_type(_e27.member)) { - case 0: { - phi_63_ = false; - break; - } - case 1: { - _atomic_compare_exchange_resultUint4_ _e31 = naga_atomic_compare_exchange_weak_explicit(&global.member, 3u, _e27.member_1); - phi_63_ = (_e31.old_value == 3u) ? false : true; - break; - } - default: { - phi_63_ = bool {}; - break; - } - } - bool _e36 = phi_63_; - continue; -#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } - LOOP_IS_BOUNDED - } - return; -} - -kernel void stagetest_atomic_compare_exchange( - device type_5& global [[user(fake0)]] -, device type_3 const& global_1 [[user(fake0)]] -) { - function(global, global_1); -} diff --git a/naga/tests/out/msl/atomic_exchange.msl b/naga/tests/out/msl/atomic_exchange.msl deleted file mode 100644 index 7bc1e20c2b..0000000000 --- a/naga/tests/out/msl/atomic_exchange.msl +++ /dev/null @@ -1,89 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; - -struct type_2 { - uint member; - uint member_1; -}; -struct type_3 { - uint member; -}; -struct type_5 { - metal::atomic_uint member; -}; - -void function( - device type_5& global, - device type_3 const& global_1 -) { - type_2 phi_33_ = {}; - uint phi_36_ = {}; - type_2 phi_52_ = {}; - type_2 phi_53_ = {}; - bool phi_62_ = {}; - type_2 phi_34_ = {}; - uint phi_37_ = {}; - uint _e10 = global_1.member; - phi_33_ = type_2 {0u, _e10}; - phi_36_ = 0u; - bool loop_init = true; - while(true) { - if (!loop_init) { - phi_33_ = phi_34_; - phi_36_ = phi_37_; - if (!(phi_62_)) { - break; - } - } - loop_init = false; - type_2 _e13 = phi_33_; - uint _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}; - } - type_2 _e26 = phi_52_; - type_2 _e28 = phi_53_; - switch(as_type(_e28.member)) { - case 0: { - phi_62_ = false; - phi_34_ = type_2 {}; - phi_37_ = uint {}; - break; - } - case 1: { - uint _e31 = metal::atomic_exchange_explicit(&global.member, _e15, metal::memory_order_relaxed); - phi_62_ = true; - phi_34_ = _e26; - phi_37_ = _e15 + _e31; - break; - } - default: { - phi_62_ = false; - phi_34_ = type_2 {}; - phi_37_ = uint {}; - break; - } - } - bool _e34 = phi_62_; - type_2 _e36 = phi_34_; - uint _e38 = phi_37_; - continue; -#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } - LOOP_IS_BOUNDED - } - return; -} - -kernel void stagetest_atomic_exchange( - device type_5& global [[user(fake0)]] -, device type_3 const& global_1 [[user(fake0)]] -) { - function(global, global_1); -} diff --git a/naga/tests/out/msl/atomic_i_add_sub.msl b/naga/tests/out/msl/atomic_i_add_sub.msl deleted file mode 100644 index 31427767d3..0000000000 --- a/naga/tests/out/msl/atomic_i_add_sub.msl +++ /dev/null @@ -1,38 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; - -struct _mslBufferSizes { - uint size1; -}; - -typedef uint type_1[1]; -struct type_2 { - type_1 member; -}; -struct type_4 { - metal::atomic_uint member; -}; - -void function( - device type_4& global, - device type_2& global_1, - constant _mslBufferSizes& _buffer_sizes -) { - uint _e6 = metal::atomic_fetch_add_explicit(&global.member, 2u, metal::memory_order_relaxed); - uint _e7 = metal::atomic_fetch_sub_explicit(&global.member, _e6, metal::memory_order_relaxed); - if (_e6 < (1 + (_buffer_sizes.size1 - 0 - 4) / 4)) { - global_1.member[_e6] = _e7; - } - return; -} - -kernel void stagetest_atomic_i_add_sub( - device type_4& global [[user(fake0)]] -, device type_2& global_1 [[user(fake0)]] -, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] -) { - function(global, global_1, _buffer_sizes); -} diff --git a/naga/tests/out/msl/atomic_i_decrement.msl b/naga/tests/out/msl/atomic_i_decrement.msl deleted file mode 100644 index cc12aa98a7..0000000000 --- a/naga/tests/out/msl/atomic_i_decrement.msl +++ /dev/null @@ -1,54 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; - -struct _mslBufferSizes { - uint size1; -}; - -typedef uint type_2[1]; -struct type_3 { - type_2 member; -}; -struct type_5 { - metal::atomic_uint member; -}; - -void function( - device type_5& global, - device type_3& global_1, - constant _mslBufferSizes& _buffer_sizes -) { - bool phi_40_ = {}; - bool loop_init = true; - while(true) { - if (!loop_init) { - if (!(phi_40_)) { - break; - } - } - loop_init = false; - uint _e8 = metal::atomic_fetch_sub_explicit(&global.member, 1u, metal::memory_order_relaxed); - if (_e8 < (1 + (_buffer_sizes.size1 - 0 - 4) / 4)) { - global_1.member[_e8] = _e8; - phi_40_ = (_e8 == 0u) ? false : true; - } else { - phi_40_ = false; - } - bool _e16 = phi_40_; - continue; -#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } - LOOP_IS_BOUNDED - } - return; -} - -kernel void stagetest_atomic_i_decrement( - device type_5& global [[user(fake0)]] -, device type_3& global_1 [[user(fake0)]] -, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] -) { - function(global, global_1, _buffer_sizes); -} diff --git a/naga/tests/out/msl/atomic_i_increment.msl b/naga/tests/out/msl/atomic_i_increment.msl deleted file mode 100644 index 507185da5f..0000000000 --- a/naga/tests/out/msl/atomic_i_increment.msl +++ /dev/null @@ -1,52 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; - -struct type_2 { - uint member; -}; -struct type_4 { - metal::atomic_uint member; -}; - -void function( - device type_4& global, - device type_2 const& global_1 -) { - uint phi_23_ = {}; - uint phi_24_ = {}; - phi_23_ = 0u; - bool loop_init = true; - while(true) { - if (!loop_init) { - phi_23_ = phi_24_; - if (!(((phi_23_ >= global_1.member) ? false : true))) { - break; - } - } - loop_init = false; - uint _e10 = phi_23_; - uint _e11 = global_1.member; - bool _e12 = _e10 >= _e11; - if (_e12) { - phi_24_ = uint {}; - } else { - uint _e13 = metal::atomic_fetch_add_explicit(&global.member, 1u, metal::memory_order_relaxed); - phi_24_ = _e10 + 1u; - } - uint _e17 = phi_24_; - continue; -#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } - LOOP_IS_BOUNDED - } - return; -} - -kernel void stagetest_atomic_i_increment( - device type_4& global [[user(fake0)]] -, device type_2 const& global_1 [[user(fake0)]] -) { - function(global, global_1); -} diff --git a/naga/tests/out/msl/atomic_load_and_store.msl b/naga/tests/out/msl/atomic_load_and_store.msl deleted file mode 100644 index 327980c5f6..0000000000 --- a/naga/tests/out/msl/atomic_load_and_store.msl +++ /dev/null @@ -1,81 +0,0 @@ -// language: metal1.0 -#include -#include - -using metal::uint; - -struct type_2 { - uint member; - uint member_1; -}; -struct type_3 { - uint member; -}; -struct type_5 { - metal::atomic_uint member; -}; - -void function( - device type_5& global, - device type_3 const& global_1 -) { - type_2 phi_32_ = {}; - type_2 phi_49_ = {}; - type_2 phi_50_ = {}; - bool phi_59_ = {}; - type_2 phi_33_ = {}; - uint _e10 = global_1.member; - phi_32_ = type_2 {0u, _e10}; - bool loop_init = true; - while(true) { - if (!loop_init) { - phi_32_ = phi_33_; - if (!(phi_59_)) { - break; - } - } - loop_init = false; - type_2 _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}; - } - type_2 _e24 = phi_49_; - type_2 _e26 = phi_50_; - switch(as_type(_e26.member)) { - case 0: { - phi_59_ = false; - phi_33_ = type_2 {}; - break; - } - case 1: { - uint _e29 = metal::atomic_load_explicit(&global.member, metal::memory_order_relaxed); - metal::atomic_store_explicit(&global.member, _e29 + 2u, metal::memory_order_relaxed); - phi_59_ = true; - phi_33_ = _e24; - break; - } - default: { - phi_59_ = false; - phi_33_ = type_2 {}; - break; - } - } - bool _e32 = phi_59_; - type_2 _e34 = phi_33_; - continue; -#define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; } - LOOP_IS_BOUNDED - } - return; -} - -kernel void stagetest_atomic_load_and_store( - device type_5& global [[user(fake0)]] -, device type_3 const& global_1 [[user(fake0)]] -) { - function(global, global_1); -} diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 3c3149fa05..72ce323585 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -1071,13 +1071,12 @@ fn convert_spv_all() { false, Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ); - let atomic_targets = Targets::IR | Targets::METAL | Targets::WGSL; - convert_spv("atomic_i_increment", false, atomic_targets); - convert_spv("atomic_load_and_store", false, atomic_targets); - convert_spv("atomic_exchange", false, atomic_targets); - convert_spv("atomic_compare_exchange", false, atomic_targets); - convert_spv("atomic_i_decrement", false, atomic_targets); - convert_spv("atomic_i_add_sub", false, atomic_targets); + 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,