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,