diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index f9e52914f7..a3810c5dab 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1856,14 +1856,37 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } }; - let var_handle = self.fill_access_chain(module, pointer, func_ctx)?; - // working around the borrow checker in `self.write_expr` - let chain = mem::take(&mut self.temp_access_chain); - let var_name = &self.names[&NameKey::GlobalVariable(var_handle)]; + // Validation ensures that `pointer` has a `Pointer` type. + let pointer_space = func_ctx.info[pointer] + .ty + .inner_with(&module.types) + .pointer_space() + .unwrap(); let fun_str = fun.to_hlsl_suffix(); - write!(self.out, " {res_name}; {var_name}.Interlocked{fun_str}(")?; - self.write_storage_address(module, &chain, func_ctx)?; + write!(self.out, " {res_name}; ")?; + match pointer_space { + crate::AddressSpace::WorkGroup => { + write!(self.out, "Interlocked{fun_str}(")?; + self.write_expr(module, pointer, func_ctx)?; + } + crate::AddressSpace::Storage { .. } => { + let var_handle = self.fill_access_chain(module, pointer, func_ctx)?; + // The call to `self.write_storage_address` wants + // mutable access to all of `self`, so temporarily take + // ownership of our reusable access chain buffer. + let chain = mem::take(&mut self.temp_access_chain); + let var_name = &self.names[&NameKey::GlobalVariable(var_handle)]; + write!(self.out, "{var_name}.Interlocked{fun_str}(")?; + self.write_storage_address(module, &chain, func_ctx)?; + self.temp_access_chain = chain; + } + ref other => { + return Err(Error::Custom(format!( + "invalid address space {other:?} for atomic statement" + ))) + } + } write!(self.out, ", ")?; // handle the special cases match *fun { @@ -1878,7 +1901,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } self.write_expr(module, value, func_ctx)?; writeln!(self.out, ", {res_name});")?; - self.temp_access_chain = chain; self.named_expressions.insert(result, res_name); } Statement::Switch { diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index a9c27ee78c..ed4988532d 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -151,23 +151,6 @@ fn foo_frag() -> @location(0) vec4 { return vec4(0.0); } -@compute @workgroup_size(1) -fn atomics() { - var tmp: i32; - let value = atomicLoad(&bar.atom); - tmp = atomicAdd(&bar.atom, 5); - tmp = atomicSub(&bar.atom, 5); - tmp = atomicAnd(&bar.atom, 5); - tmp = atomicOr(&bar.atom, 5); - tmp = atomicXor(&bar.atom, 5); - tmp = atomicMin(&bar.atom, 5); - tmp = atomicMax(&bar.atom, 5); - tmp = atomicExchange(&bar.atom, 5); - // https://github.com/gpuweb/gpuweb/issues/2021 - // tmp = atomicCompareExchangeWeak(&bar.atom, 5, 5); - atomicStore(&bar.atom, value); -} - var val: u32; fn assign_through_ptr_fn(p: ptr) { @@ -184,4 +167,4 @@ fn assign_through_ptr() { assign_through_ptr_fn(&val); assign_array_through_ptr_fn(&arr); -} \ No newline at end of file +} diff --git a/tests/in/atomicOps.wgsl b/tests/in/atomicOps.wgsl new file mode 100644 index 0000000000..111a9abbc8 --- /dev/null +++ b/tests/in/atomicOps.wgsl @@ -0,0 +1,141 @@ +// This test covers the cross product of: +// +// * All atomic operations. +// * On all applicable scopes (storage read-write, workgroup). +// * For all shapes of modeling atomic data. + +struct Struct { + atomic_scalar: atomic, + atomic_arr: array, 2>, +} + +@group(0) @binding(0) +var storage_atomic_scalar: atomic; +@group(0) @binding(1) +var storage_atomic_arr: array, 2>; +@group(0) @binding(2) +var storage_struct: Struct; + +var workgroup_atomic_scalar: atomic; +var workgroup_atomic_arr: array, 2>; +var workgroup_struct: Struct; + +@compute +@workgroup_size(2) +fn cs_main(@builtin(local_invocation_id) id: vec3) { + atomicStore(&storage_atomic_scalar, 1u); + atomicStore(&storage_atomic_arr[1], 1i); + atomicStore(&storage_struct.atomic_scalar, 1u); + atomicStore(&storage_struct.atomic_arr[1], 1i); + atomicStore(&workgroup_atomic_scalar, 1u); + atomicStore(&workgroup_atomic_arr[1], 1i); + atomicStore(&workgroup_struct.atomic_scalar, 1u); + atomicStore(&workgroup_struct.atomic_arr[1], 1i); + + workgroupBarrier(); + + atomicLoad(&storage_atomic_scalar); + atomicLoad(&storage_atomic_arr[1]); + atomicLoad(&storage_struct.atomic_scalar); + atomicLoad(&storage_struct.atomic_arr[1]); + atomicLoad(&workgroup_atomic_scalar); + atomicLoad(&workgroup_atomic_arr[1]); + atomicLoad(&workgroup_struct.atomic_scalar); + atomicLoad(&workgroup_struct.atomic_arr[1]); + + workgroupBarrier(); + + atomicAdd(&storage_atomic_scalar, 1u); + atomicAdd(&storage_atomic_arr[1], 1i); + atomicAdd(&storage_struct.atomic_scalar, 1u); + atomicAdd(&storage_struct.atomic_arr[1], 1i); + atomicAdd(&workgroup_atomic_scalar, 1u); + atomicAdd(&workgroup_atomic_arr[1], 1i); + atomicAdd(&workgroup_struct.atomic_scalar, 1u); + atomicAdd(&workgroup_struct.atomic_arr[1], 1i); + + workgroupBarrier(); + + atomicSub(&storage_atomic_scalar, 1u); + atomicSub(&storage_atomic_arr[1], 1i); + atomicSub(&storage_struct.atomic_scalar, 1u); + atomicSub(&storage_struct.atomic_arr[1], 1i); + atomicSub(&workgroup_atomic_scalar, 1u); + atomicSub(&workgroup_atomic_arr[1], 1i); + atomicSub(&workgroup_struct.atomic_scalar, 1u); + atomicSub(&workgroup_struct.atomic_arr[1], 1i); + + workgroupBarrier(); + + atomicMax(&storage_atomic_scalar, 1u); + atomicMax(&storage_atomic_arr[1], 1i); + atomicMax(&storage_struct.atomic_scalar, 1u); + atomicMax(&storage_struct.atomic_arr[1], 1i); + atomicMax(&workgroup_atomic_scalar, 1u); + atomicMax(&workgroup_atomic_arr[1], 1i); + atomicMax(&workgroup_struct.atomic_scalar, 1u); + atomicMax(&workgroup_struct.atomic_arr[1], 1i); + + workgroupBarrier(); + + atomicMin(&storage_atomic_scalar, 1u); + atomicMin(&storage_atomic_arr[1], 1i); + atomicMin(&storage_struct.atomic_scalar, 1u); + atomicMin(&storage_struct.atomic_arr[1], 1i); + atomicMin(&workgroup_atomic_scalar, 1u); + atomicMin(&workgroup_atomic_arr[1], 1i); + atomicMin(&workgroup_struct.atomic_scalar, 1u); + atomicMin(&workgroup_struct.atomic_arr[1], 1i); + + workgroupBarrier(); + + atomicAnd(&storage_atomic_scalar, 1u); + atomicAnd(&storage_atomic_arr[1], 1i); + atomicAnd(&storage_struct.atomic_scalar, 1u); + atomicAnd(&storage_struct.atomic_arr[1], 1i); + atomicAnd(&workgroup_atomic_scalar, 1u); + atomicAnd(&workgroup_atomic_arr[1], 1i); + atomicAnd(&workgroup_struct.atomic_scalar, 1u); + atomicAnd(&workgroup_struct.atomic_arr[1], 1i); + + workgroupBarrier(); + + atomicOr(&storage_atomic_scalar, 1u); + atomicOr(&storage_atomic_arr[1], 1i); + atomicOr(&storage_struct.atomic_scalar, 1u); + atomicOr(&storage_struct.atomic_arr[1], 1i); + atomicOr(&workgroup_atomic_scalar, 1u); + atomicOr(&workgroup_atomic_arr[1], 1i); + atomicOr(&workgroup_struct.atomic_scalar, 1u); + atomicOr(&workgroup_struct.atomic_arr[1], 1i); + + workgroupBarrier(); + + atomicXor(&storage_atomic_scalar, 1u); + atomicXor(&storage_atomic_arr[1], 1i); + atomicXor(&storage_struct.atomic_scalar, 1u); + atomicXor(&storage_struct.atomic_arr[1], 1i); + atomicXor(&workgroup_atomic_scalar, 1u); + atomicXor(&workgroup_atomic_arr[1], 1i); + atomicXor(&workgroup_struct.atomic_scalar, 1u); + atomicXor(&workgroup_struct.atomic_arr[1], 1i); + + atomicExchange(&storage_atomic_scalar, 1u); + atomicExchange(&storage_atomic_arr[1], 1i); + atomicExchange(&storage_struct.atomic_scalar, 1u); + atomicExchange(&storage_struct.atomic_arr[1], 1i); + atomicExchange(&workgroup_atomic_scalar, 1u); + atomicExchange(&workgroup_atomic_arr[1], 1i); + atomicExchange(&workgroup_struct.atomic_scalar, 1u); + atomicExchange(&workgroup_struct.atomic_arr[1], 1i); + + // // TODO: https://github.com/gpuweb/gpuweb/issues/2021 + // atomicCompareExchangeWeak(&storage_atomic_scalar, 1u); + // atomicCompareExchangeWeak(&storage_atomic_arr[1], 1i); + // atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1u); + // atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1i); + // atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1u); + // atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1i); + // atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1u); + // atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1i); +} diff --git a/tests/out/analysis/access.info.ron b/tests/out/analysis/access.info.ron index 01e08a63ba..8c5097f65a 100644 --- a/tests/out/analysis/access.info.ron +++ b/tests/out/analysis/access.info.ron @@ -5123,633 +5123,6 @@ ], sampling: [], ), - ( - flags: ( - bits: 63, - ), - available_stages: ( - bits: 7, - ), - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - may_kill: false, - sampling_set: [], - global_uses: [ - ( - bits: 0, - ), - ( - bits: 3, - ), - ( - bits: 0, - ), - ( - bits: 0, - ), - ( - bits: 0, - ), - ( - bits: 0, - ), - ], - expressions: [ - ( - uniformity: ( - non_uniform_result: Some(1), - requirements: ( - bits: 0, - ), - ), - ref_count: 8, - assignable_global: None, - ty: Value(Pointer( - base: 3, - space: Function, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(2), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(2), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(2), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(5), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(5), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(8), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(9), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(9), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(12), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(13), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(13), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(16), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(17), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(17), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(20), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(21), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(21), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(24), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(25), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(25), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(28), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(29), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(29), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(32), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(33), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(33), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: None, - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Value(Scalar( - kind: Sint, - width: 4, - )), - ), - ( - uniformity: ( - non_uniform_result: Some(36), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: None, - ty: Handle(3), - ), - ( - uniformity: ( - non_uniform_result: Some(37), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 14, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ( - uniformity: ( - non_uniform_result: Some(37), - requirements: ( - bits: 0, - ), - ), - ref_count: 1, - assignable_global: Some(2), - ty: Value(Pointer( - base: 9, - space: Storage( - access: ( - bits: 3, - ), - ), - )), - ), - ], - sampling: [], - ), ( flags: ( bits: 63, diff --git a/tests/out/glsl/atomicOps.cs_main.Compute.glsl b/tests/out/glsl/atomicOps.cs_main.Compute.glsl new file mode 100644 index 0000000000..30086fa8ef --- /dev/null +++ b/tests/out/glsl/atomicOps.cs_main.Compute.glsl @@ -0,0 +1,132 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in; + +struct Struct { + uint atomic_scalar; + int atomic_arr[2]; +}; +layout(std430) buffer type_block_0Compute { uint _group_0_binding_0_cs; }; + +layout(std430) buffer type_2_block_1Compute { int _group_0_binding_1_cs[2]; }; + +layout(std430) buffer Struct_block_2Compute { Struct _group_0_binding_2_cs; }; + +shared uint workgroup_atomic_scalar; + +shared int workgroup_atomic_arr[2]; + +shared Struct workgroup_struct; + + +void main() { + if (gl_LocalInvocationID == uvec3(0u)) { + workgroup_atomic_scalar = 0u; + workgroup_atomic_arr = int[2](0, 0); + workgroup_struct = Struct(0u, int[2](0, 0)); + } + memoryBarrierShared(); + barrier(); + uvec3 id = gl_LocalInvocationID; + _group_0_binding_0_cs = 1u; + _group_0_binding_1_cs[1] = 1; + _group_0_binding_2_cs.atomic_scalar = 1u; + _group_0_binding_2_cs.atomic_arr[1] = 1; + workgroup_atomic_scalar = 1u; + workgroup_atomic_arr[1] = 1; + workgroup_struct.atomic_scalar = 1u; + workgroup_struct.atomic_arr[1] = 1; + memoryBarrierShared(); + barrier(); + uint unnamed = _group_0_binding_0_cs; + int unnamed_1 = _group_0_binding_1_cs[1]; + uint unnamed_2 = _group_0_binding_2_cs.atomic_scalar; + int unnamed_3 = _group_0_binding_2_cs.atomic_arr[1]; + uint unnamed_4 = workgroup_atomic_scalar; + int unnamed_5 = workgroup_atomic_arr[1]; + uint unnamed_6 = workgroup_struct.atomic_scalar; + int unnamed_7 = workgroup_struct.atomic_arr[1]; + memoryBarrierShared(); + barrier(); + uint _e59 = atomicAdd(_group_0_binding_0_cs, 1u); + int _e64 = atomicAdd(_group_0_binding_1_cs[1], 1); + uint _e68 = atomicAdd(_group_0_binding_2_cs.atomic_scalar, 1u); + int _e74 = atomicAdd(_group_0_binding_2_cs.atomic_arr[1], 1); + uint _e77 = atomicAdd(workgroup_atomic_scalar, 1u); + int _e82 = atomicAdd(workgroup_atomic_arr[1], 1); + uint _e86 = atomicAdd(workgroup_struct.atomic_scalar, 1u); + int _e92 = atomicAdd(workgroup_struct.atomic_arr[1], 1); + memoryBarrierShared(); + barrier(); + uint _e95 = atomicAdd(_group_0_binding_0_cs, -1u); + int _e100 = atomicAdd(_group_0_binding_1_cs[1], -1); + uint _e104 = atomicAdd(_group_0_binding_2_cs.atomic_scalar, -1u); + int _e110 = atomicAdd(_group_0_binding_2_cs.atomic_arr[1], -1); + uint _e113 = atomicAdd(workgroup_atomic_scalar, -1u); + int _e118 = atomicAdd(workgroup_atomic_arr[1], -1); + uint _e122 = atomicAdd(workgroup_struct.atomic_scalar, -1u); + int _e128 = atomicAdd(workgroup_struct.atomic_arr[1], -1); + memoryBarrierShared(); + barrier(); + uint _e131 = atomicMax(_group_0_binding_0_cs, 1u); + int _e136 = atomicMax(_group_0_binding_1_cs[1], 1); + uint _e140 = atomicMax(_group_0_binding_2_cs.atomic_scalar, 1u); + int _e146 = atomicMax(_group_0_binding_2_cs.atomic_arr[1], 1); + uint _e149 = atomicMax(workgroup_atomic_scalar, 1u); + int _e154 = atomicMax(workgroup_atomic_arr[1], 1); + uint _e158 = atomicMax(workgroup_struct.atomic_scalar, 1u); + int _e164 = atomicMax(workgroup_struct.atomic_arr[1], 1); + memoryBarrierShared(); + barrier(); + uint _e167 = atomicMin(_group_0_binding_0_cs, 1u); + int _e172 = atomicMin(_group_0_binding_1_cs[1], 1); + uint _e176 = atomicMin(_group_0_binding_2_cs.atomic_scalar, 1u); + int _e182 = atomicMin(_group_0_binding_2_cs.atomic_arr[1], 1); + uint _e185 = atomicMin(workgroup_atomic_scalar, 1u); + int _e190 = atomicMin(workgroup_atomic_arr[1], 1); + uint _e194 = atomicMin(workgroup_struct.atomic_scalar, 1u); + int _e200 = atomicMin(workgroup_struct.atomic_arr[1], 1); + memoryBarrierShared(); + barrier(); + uint _e203 = atomicAnd(_group_0_binding_0_cs, 1u); + int _e208 = atomicAnd(_group_0_binding_1_cs[1], 1); + uint _e212 = atomicAnd(_group_0_binding_2_cs.atomic_scalar, 1u); + int _e218 = atomicAnd(_group_0_binding_2_cs.atomic_arr[1], 1); + uint _e221 = atomicAnd(workgroup_atomic_scalar, 1u); + int _e226 = atomicAnd(workgroup_atomic_arr[1], 1); + uint _e230 = atomicAnd(workgroup_struct.atomic_scalar, 1u); + int _e236 = atomicAnd(workgroup_struct.atomic_arr[1], 1); + memoryBarrierShared(); + barrier(); + uint _e239 = atomicOr(_group_0_binding_0_cs, 1u); + int _e244 = atomicOr(_group_0_binding_1_cs[1], 1); + uint _e248 = atomicOr(_group_0_binding_2_cs.atomic_scalar, 1u); + int _e254 = atomicOr(_group_0_binding_2_cs.atomic_arr[1], 1); + uint _e257 = atomicOr(workgroup_atomic_scalar, 1u); + int _e262 = atomicOr(workgroup_atomic_arr[1], 1); + uint _e266 = atomicOr(workgroup_struct.atomic_scalar, 1u); + int _e272 = atomicOr(workgroup_struct.atomic_arr[1], 1); + memoryBarrierShared(); + barrier(); + uint _e275 = atomicXor(_group_0_binding_0_cs, 1u); + int _e280 = atomicXor(_group_0_binding_1_cs[1], 1); + uint _e284 = atomicXor(_group_0_binding_2_cs.atomic_scalar, 1u); + int _e290 = atomicXor(_group_0_binding_2_cs.atomic_arr[1], 1); + uint _e293 = atomicXor(workgroup_atomic_scalar, 1u); + int _e298 = atomicXor(workgroup_atomic_arr[1], 1); + uint _e302 = atomicXor(workgroup_struct.atomic_scalar, 1u); + int _e308 = atomicXor(workgroup_struct.atomic_arr[1], 1); + uint _e311 = atomicExchange(_group_0_binding_0_cs, 1u); + int _e316 = atomicExchange(_group_0_binding_1_cs[1], 1); + uint _e320 = atomicExchange(_group_0_binding_2_cs.atomic_scalar, 1u); + int _e326 = atomicExchange(_group_0_binding_2_cs.atomic_arr[1], 1); + uint _e329 = atomicExchange(workgroup_atomic_scalar, 1u); + int _e334 = atomicExchange(workgroup_atomic_arr[1], 1); + uint _e338 = atomicExchange(workgroup_struct.atomic_scalar, 1u); + int _e344 = atomicExchange(workgroup_struct.atomic_arr[1], 1); + return; +} + diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index f20cb726db..27786717d4 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -311,32 +311,6 @@ float4 foo_frag() : SV_Target0 return (0.0).xxxx; } -[numthreads(1, 1, 1)] -void atomics() -{ - int tmp = (int)0; - - int value_1 = asint(bar.Load(96)); - int _e7; bar.InterlockedAdd(96, 5, _e7); - tmp = _e7; - int _e11; bar.InterlockedAdd(96, -5, _e11); - tmp = _e11; - int _e15; bar.InterlockedAnd(96, 5, _e15); - tmp = _e15; - int _e19; bar.InterlockedOr(96, 5, _e19); - tmp = _e19; - int _e23; bar.InterlockedXor(96, 5, _e23); - tmp = _e23; - int _e27; bar.InterlockedMin(96, 5, _e27); - tmp = _e27; - int _e31; bar.InterlockedMax(96, 5, _e31); - tmp = _e31; - int _e35; bar.InterlockedExchange(96, 5, _e35); - tmp = _e35; - bar.Store(96, asuint(value_1)); - return; -} - [numthreads(1, 1, 1)] void assign_through_ptr(uint3 __local_invocation_id : SV_GroupThreadID) { diff --git a/tests/out/hlsl/access.hlsl.config b/tests/out/hlsl/access.hlsl.config index b68dda72d3..de7719bdff 100644 --- a/tests/out/hlsl/access.hlsl.config +++ b/tests/out/hlsl/access.hlsl.config @@ -1,3 +1,3 @@ vertex=(foo_vert:vs_5_1 ) fragment=(foo_frag:ps_5_1 ) -compute=(atomics:cs_5_1 assign_through_ptr:cs_5_1 ) +compute=(assign_through_ptr:cs_5_1 ) diff --git a/tests/out/hlsl/atomicOps.hlsl b/tests/out/hlsl/atomicOps.hlsl new file mode 100644 index 0000000000..f6322cd92f --- /dev/null +++ b/tests/out/hlsl/atomicOps.hlsl @@ -0,0 +1,112 @@ + +struct Struct { + uint atomic_scalar; + int atomic_arr[2]; +}; + +RWByteAddressBuffer storage_atomic_scalar : register(u0); +RWByteAddressBuffer storage_atomic_arr : register(u1); +RWByteAddressBuffer storage_struct : register(u2); +groupshared uint workgroup_atomic_scalar; +groupshared int workgroup_atomic_arr[2]; +groupshared Struct workgroup_struct; + +[numthreads(2, 1, 1)] +void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_GroupThreadID) +{ + if (all(__local_invocation_id == uint3(0u, 0u, 0u))) { + workgroup_atomic_scalar = (uint)0; + workgroup_atomic_arr = (int[2])0; + workgroup_struct = (Struct)0; + } + GroupMemoryBarrierWithGroupSync(); + storage_atomic_scalar.Store(0, asuint(1u)); + storage_atomic_arr.Store(4, asuint(1)); + storage_struct.Store(0, asuint(1u)); + storage_struct.Store(4+4, asuint(1)); + workgroup_atomic_scalar = 1u; + workgroup_atomic_arr[1] = 1; + workgroup_struct.atomic_scalar = 1u; + workgroup_struct.atomic_arr[1] = 1; + GroupMemoryBarrierWithGroupSync(); + uint unnamed = asuint(storage_atomic_scalar.Load(0)); + int unnamed_1 = asint(storage_atomic_arr.Load(4)); + uint unnamed_2 = asuint(storage_struct.Load(0)); + int unnamed_3 = asint(storage_struct.Load(4+4)); + uint unnamed_4 = workgroup_atomic_scalar; + int unnamed_5 = workgroup_atomic_arr[1]; + uint unnamed_6 = workgroup_struct.atomic_scalar; + int unnamed_7 = workgroup_struct.atomic_arr[1]; + GroupMemoryBarrierWithGroupSync(); + uint _e59; storage_atomic_scalar.InterlockedAdd(0, 1u, _e59); + int _e64; storage_atomic_arr.InterlockedAdd(4, 1, _e64); + uint _e68; storage_struct.InterlockedAdd(0, 1u, _e68); + int _e74; storage_struct.InterlockedAdd(4+4, 1, _e74); + uint _e77; InterlockedAdd(workgroup_atomic_scalar, 1u, _e77); + int _e82; InterlockedAdd(workgroup_atomic_arr[1], 1, _e82); + uint _e86; InterlockedAdd(workgroup_struct.atomic_scalar, 1u, _e86); + int _e92; InterlockedAdd(workgroup_struct.atomic_arr[1], 1, _e92); + GroupMemoryBarrierWithGroupSync(); + uint _e95; storage_atomic_scalar.InterlockedAdd(0, -1u, _e95); + int _e100; storage_atomic_arr.InterlockedAdd(4, -1, _e100); + uint _e104; storage_struct.InterlockedAdd(0, -1u, _e104); + int _e110; storage_struct.InterlockedAdd(4+4, -1, _e110); + uint _e113; InterlockedAdd(workgroup_atomic_scalar, -1u, _e113); + int _e118; InterlockedAdd(workgroup_atomic_arr[1], -1, _e118); + uint _e122; InterlockedAdd(workgroup_struct.atomic_scalar, -1u, _e122); + int _e128; InterlockedAdd(workgroup_struct.atomic_arr[1], -1, _e128); + GroupMemoryBarrierWithGroupSync(); + uint _e131; storage_atomic_scalar.InterlockedMax(0, 1u, _e131); + int _e136; storage_atomic_arr.InterlockedMax(4, 1, _e136); + uint _e140; storage_struct.InterlockedMax(0, 1u, _e140); + int _e146; storage_struct.InterlockedMax(4+4, 1, _e146); + uint _e149; InterlockedMax(workgroup_atomic_scalar, 1u, _e149); + int _e154; InterlockedMax(workgroup_atomic_arr[1], 1, _e154); + uint _e158; InterlockedMax(workgroup_struct.atomic_scalar, 1u, _e158); + int _e164; InterlockedMax(workgroup_struct.atomic_arr[1], 1, _e164); + GroupMemoryBarrierWithGroupSync(); + uint _e167; storage_atomic_scalar.InterlockedMin(0, 1u, _e167); + int _e172; storage_atomic_arr.InterlockedMin(4, 1, _e172); + uint _e176; storage_struct.InterlockedMin(0, 1u, _e176); + int _e182; storage_struct.InterlockedMin(4+4, 1, _e182); + uint _e185; InterlockedMin(workgroup_atomic_scalar, 1u, _e185); + int _e190; InterlockedMin(workgroup_atomic_arr[1], 1, _e190); + uint _e194; InterlockedMin(workgroup_struct.atomic_scalar, 1u, _e194); + int _e200; InterlockedMin(workgroup_struct.atomic_arr[1], 1, _e200); + GroupMemoryBarrierWithGroupSync(); + uint _e203; storage_atomic_scalar.InterlockedAnd(0, 1u, _e203); + int _e208; storage_atomic_arr.InterlockedAnd(4, 1, _e208); + uint _e212; storage_struct.InterlockedAnd(0, 1u, _e212); + int _e218; storage_struct.InterlockedAnd(4+4, 1, _e218); + uint _e221; InterlockedAnd(workgroup_atomic_scalar, 1u, _e221); + int _e226; InterlockedAnd(workgroup_atomic_arr[1], 1, _e226); + uint _e230; InterlockedAnd(workgroup_struct.atomic_scalar, 1u, _e230); + int _e236; InterlockedAnd(workgroup_struct.atomic_arr[1], 1, _e236); + GroupMemoryBarrierWithGroupSync(); + uint _e239; storage_atomic_scalar.InterlockedOr(0, 1u, _e239); + int _e244; storage_atomic_arr.InterlockedOr(4, 1, _e244); + uint _e248; storage_struct.InterlockedOr(0, 1u, _e248); + int _e254; storage_struct.InterlockedOr(4+4, 1, _e254); + uint _e257; InterlockedOr(workgroup_atomic_scalar, 1u, _e257); + int _e262; InterlockedOr(workgroup_atomic_arr[1], 1, _e262); + uint _e266; InterlockedOr(workgroup_struct.atomic_scalar, 1u, _e266); + int _e272; InterlockedOr(workgroup_struct.atomic_arr[1], 1, _e272); + GroupMemoryBarrierWithGroupSync(); + uint _e275; storage_atomic_scalar.InterlockedXor(0, 1u, _e275); + int _e280; storage_atomic_arr.InterlockedXor(4, 1, _e280); + uint _e284; storage_struct.InterlockedXor(0, 1u, _e284); + int _e290; storage_struct.InterlockedXor(4+4, 1, _e290); + uint _e293; InterlockedXor(workgroup_atomic_scalar, 1u, _e293); + int _e298; InterlockedXor(workgroup_atomic_arr[1], 1, _e298); + uint _e302; InterlockedXor(workgroup_struct.atomic_scalar, 1u, _e302); + int _e308; InterlockedXor(workgroup_struct.atomic_arr[1], 1, _e308); + uint _e311; storage_atomic_scalar.InterlockedExchange(0, 1u, _e311); + int _e316; storage_atomic_arr.InterlockedExchange(4, 1, _e316); + uint _e320; storage_struct.InterlockedExchange(0, 1u, _e320); + int _e326; storage_struct.InterlockedExchange(4+4, 1, _e326); + uint _e329; InterlockedExchange(workgroup_atomic_scalar, 1u, _e329); + int _e334; InterlockedExchange(workgroup_atomic_arr[1], 1, _e334); + uint _e338; InterlockedExchange(workgroup_struct.atomic_scalar, 1u, _e338); + int _e344; InterlockedExchange(workgroup_struct.atomic_arr[1], 1, _e344); + return; +} diff --git a/tests/out/hlsl/atomicOps.hlsl.config b/tests/out/hlsl/atomicOps.hlsl.config new file mode 100644 index 0000000000..522d99ed1a --- /dev/null +++ b/tests/out/hlsl/atomicOps.hlsl.config @@ -0,0 +1,3 @@ +vertex=() +fragment=() +compute=(cs_main:cs_5_1 ) diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index 41772b9332..0c1f9df1cb 100644 --- a/tests/out/ir/access.ron +++ b/tests/out/ir/access.ron @@ -2535,254 +2535,6 @@ ], ), ), - ( - name: "atomics", - stage: Compute, - early_depth_test: None, - workgroup_size: (1, 1, 1), - function: ( - name: Some("atomics"), - arguments: [], - result: None, - local_variables: [ - ( - name: Some("tmp"), - ty: 3, - init: None, - ), - ], - expressions: [ - LocalVariable(1), - GlobalVariable(2), - AccessIndex( - base: 2, - index: 2, - ), - Load( - pointer: 3, - ), - GlobalVariable(2), - AccessIndex( - base: 5, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 9, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 13, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 17, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 21, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 25, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 29, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 33, - index: 2, - ), - Constant(26), - AtomicResult( - ty: 3, - comparison: false, - ), - GlobalVariable(2), - AccessIndex( - base: 37, - index: 2, - ), - ], - named_expressions: { - 4: "value", - }, - body: [ - Emit(( - start: 2, - end: 4, - )), - Emit(( - start: 5, - end: 6, - )), - Atomic( - pointer: 6, - fun: Add, - value: 7, - result: 8, - ), - Store( - pointer: 1, - value: 8, - ), - Emit(( - start: 9, - end: 10, - )), - Atomic( - pointer: 10, - fun: Subtract, - value: 11, - result: 12, - ), - Store( - pointer: 1, - value: 12, - ), - Emit(( - start: 13, - end: 14, - )), - Atomic( - pointer: 14, - fun: And, - value: 15, - result: 16, - ), - Store( - pointer: 1, - value: 16, - ), - Emit(( - start: 17, - end: 18, - )), - Atomic( - pointer: 18, - fun: InclusiveOr, - value: 19, - result: 20, - ), - Store( - pointer: 1, - value: 20, - ), - Emit(( - start: 21, - end: 22, - )), - Atomic( - pointer: 22, - fun: ExclusiveOr, - value: 23, - result: 24, - ), - Store( - pointer: 1, - value: 24, - ), - Emit(( - start: 25, - end: 26, - )), - Atomic( - pointer: 26, - fun: Min, - value: 27, - result: 28, - ), - Store( - pointer: 1, - value: 28, - ), - Emit(( - start: 29, - end: 30, - )), - Atomic( - pointer: 30, - fun: Max, - value: 31, - result: 32, - ), - Store( - pointer: 1, - value: 32, - ), - Emit(( - start: 33, - end: 34, - )), - Atomic( - pointer: 34, - fun: Exchange( - compare: None, - ), - value: 35, - result: 36, - ), - Store( - pointer: 1, - value: 36, - ), - Emit(( - start: 37, - end: 38, - )), - Store( - pointer: 38, - value: 4, - ), - Return( - value: None, - ), - ], - ), - ), ( name: "assign_through_ptr", stage: Compute, diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index 8ad40973f4..a4967117ff 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -222,33 +222,6 @@ fragment foo_fragOutput foo_frag( } -kernel void atomics( - device Bar& bar [[buffer(0)]] -, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] -) { - int tmp = {}; - int value_1 = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed); - int _e7 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e7; - int _e11 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e11; - int _e15 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e15; - int _e19 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e19; - int _e23 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e23; - int _e27 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e27; - int _e31 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e31; - int _e35 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e35; - metal::atomic_store_explicit(&bar.atom, value_1, metal::memory_order_relaxed); - return; -} - - kernel void assign_through_ptr( metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]] , threadgroup uint& val diff --git a/tests/out/msl/atomicOps.msl b/tests/out/msl/atomicOps.msl new file mode 100644 index 0000000000..feed8fb743 --- /dev/null +++ b/tests/out/msl/atomicOps.msl @@ -0,0 +1,126 @@ +// language: metal2.0 +#include +#include + +using metal::uint; + +struct type_2 { + metal::atomic_int inner[2]; +}; +struct Struct { + metal::atomic_uint atomic_scalar; + type_2 atomic_arr; +}; + +struct cs_mainInput { +}; +kernel void cs_main( + metal::uint3 id [[thread_position_in_threadgroup]] +, device metal::atomic_uint& storage_atomic_scalar [[user(fake0)]] +, device type_2& storage_atomic_arr [[user(fake0)]] +, device Struct& storage_struct [[user(fake0)]] +, threadgroup metal::atomic_uint& workgroup_atomic_scalar +, threadgroup type_2& workgroup_atomic_arr +, threadgroup Struct& workgroup_struct +) { + if (metal::all(id == metal::uint3(0u))) { + metal::atomic_store_explicit(&workgroup_atomic_scalar, 0, metal::memory_order_relaxed); + for (int __i0 = 0; __i0 < 2; __i0++) { + metal::atomic_store_explicit(&workgroup_atomic_arr.inner[__i0], 0, metal::memory_order_relaxed); + } + metal::atomic_store_explicit(&workgroup_struct.atomic_scalar, 0, metal::memory_order_relaxed); + for (int __i0 = 0; __i0 < 2; __i0++) { + metal::atomic_store_explicit(&workgroup_struct.atomic_arr.inner[__i0], 0, metal::memory_order_relaxed); + } + } + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + metal::atomic_store_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + metal::atomic_store_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::atomic_store_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + metal::atomic_store_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::atomic_store_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + metal::atomic_store_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::atomic_store_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + metal::atomic_store_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint unnamed = metal::atomic_load_explicit(&storage_atomic_scalar, metal::memory_order_relaxed); + int unnamed_1 = metal::atomic_load_explicit(&storage_atomic_arr.inner[1], metal::memory_order_relaxed); + uint unnamed_2 = metal::atomic_load_explicit(&storage_struct.atomic_scalar, metal::memory_order_relaxed); + int unnamed_3 = metal::atomic_load_explicit(&storage_struct.atomic_arr.inner[1], metal::memory_order_relaxed); + uint unnamed_4 = metal::atomic_load_explicit(&workgroup_atomic_scalar, metal::memory_order_relaxed); + int unnamed_5 = metal::atomic_load_explicit(&workgroup_atomic_arr.inner[1], metal::memory_order_relaxed); + uint unnamed_6 = metal::atomic_load_explicit(&workgroup_struct.atomic_scalar, metal::memory_order_relaxed); + int unnamed_7 = metal::atomic_load_explicit(&workgroup_struct.atomic_arr.inner[1], metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint _e59 = metal::atomic_fetch_add_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e64 = metal::atomic_fetch_add_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e68 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e74 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e77 = metal::atomic_fetch_add_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e82 = metal::atomic_fetch_add_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e86 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e92 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint _e95 = metal::atomic_fetch_sub_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e100 = metal::atomic_fetch_sub_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e104 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e110 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e113 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e118 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e122 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e128 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint _e131 = metal::atomic_fetch_max_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e136 = metal::atomic_fetch_max_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e140 = metal::atomic_fetch_max_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e146 = metal::atomic_fetch_max_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e149 = metal::atomic_fetch_max_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e154 = metal::atomic_fetch_max_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e158 = metal::atomic_fetch_max_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e164 = metal::atomic_fetch_max_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint _e167 = metal::atomic_fetch_min_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e172 = metal::atomic_fetch_min_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e176 = metal::atomic_fetch_min_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e182 = metal::atomic_fetch_min_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e185 = metal::atomic_fetch_min_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e190 = metal::atomic_fetch_min_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e194 = metal::atomic_fetch_min_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e200 = metal::atomic_fetch_min_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint _e203 = metal::atomic_fetch_and_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e208 = metal::atomic_fetch_and_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e212 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e218 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e221 = metal::atomic_fetch_and_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e226 = metal::atomic_fetch_and_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e230 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e236 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint _e239 = metal::atomic_fetch_or_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e244 = metal::atomic_fetch_or_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e248 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e254 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e257 = metal::atomic_fetch_or_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e262 = metal::atomic_fetch_or_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e266 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e272 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + uint _e275 = metal::atomic_fetch_xor_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e280 = metal::atomic_fetch_xor_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e284 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e290 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e293 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e298 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e302 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e308 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e311 = metal::atomic_exchange_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e316 = metal::atomic_exchange_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e320 = metal::atomic_exchange_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e326 = metal::atomic_exchange_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e329 = metal::atomic_exchange_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e334 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e338 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e344 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + return; +} diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index a882322922..fa828b4d88 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,18 +1,16 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 354 +; Bound: 328 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %247 "foo_vert" %242 %245 OpEntryPoint Fragment %289 "foo_frag" %288 -OpEntryPoint GLCompute %308 "atomics" -OpEntryPoint GLCompute %334 "assign_through_ptr" %337 +OpEntryPoint GLCompute %308 "assign_through_ptr" %311 OpExecutionMode %289 OriginUpperLeft OpExecutionMode %308 LocalSize 1 1 1 -OpExecutionMode %334 LocalSize 1 1 1 OpSource GLSL 450 OpMemberName %36 0 "a" OpMemberName %36 1 "b" @@ -56,10 +54,8 @@ OpName %238 "c2" OpName %242 "vi" OpName %247 "foo_vert" OpName %289 "foo_frag" -OpName %305 "tmp" -OpName %308 "atomics" -OpName %331 "arr" -OpName %334 "assign_through_ptr" +OpName %305 "arr" +OpName %308 "assign_through_ptr" OpMemberDecorate %36 0 Offset 0 OpMemberDecorate %36 1 Offset 16 OpMemberDecorate %36 2 Offset 28 @@ -107,7 +103,7 @@ OpMemberDecorate %83 0 Offset 0 OpDecorate %242 BuiltIn VertexIndex OpDecorate %245 BuiltIn Position OpDecorate %288 Location 0 -OpDecorate %337 BuiltIn LocalInvocationId +OpDecorate %311 BuiltIn LocalInvocationId %2 = OpTypeVoid %4 = OpTypeInt 32 0 %3 = OpConstant %4 0 @@ -233,17 +229,14 @@ OpDecorate %337 BuiltIn LocalInvocationId %270 = OpTypePointer StorageBuffer %37 %271 = OpConstant %4 5 %288 = OpVariable %246 Output -%306 = OpConstantNull %6 -%310 = OpTypePointer StorageBuffer %6 -%313 = OpConstant %4 64 -%332 = OpConstantNull %62 -%336 = OpConstantNull %4 -%338 = OpTypePointer Input %35 -%337 = OpVariable %338 Input -%340 = OpConstantNull %35 -%342 = OpTypeBool -%341 = OpTypeVector %342 3 -%347 = OpConstant %4 264 +%306 = OpConstantNull %62 +%310 = OpConstantNull %4 +%312 = OpTypePointer Input %35 +%311 = OpVariable %312 Input +%314 = OpConstantNull %35 +%316 = OpTypeBool +%315 = OpTypeVector %316 3 +%321 = OpConstant %4 264 %93 = OpFunction %2 None %94 %92 = OpLabel %86 = OpVariable %87 Function %88 @@ -484,61 +477,26 @@ OpReturn OpFunctionEnd %308 = OpFunction %2 None %94 %307 = OpLabel -%305 = OpVariable %87 Function %306 +%305 = OpVariable %63 Function %306 OpBranch %309 %309 = OpLabel -%311 = OpAccessChain %310 %74 %30 -%312 = OpAtomicLoad %6 %311 %9 %313 -%315 = OpAccessChain %310 %74 %30 -%314 = OpAtomicIAdd %6 %315 %9 %313 %26 -OpStore %305 %314 -%317 = OpAccessChain %310 %74 %30 -%316 = OpAtomicISub %6 %317 %9 %313 %26 -OpStore %305 %316 -%319 = OpAccessChain %310 %74 %30 -%318 = OpAtomicAnd %6 %319 %9 %313 %26 -OpStore %305 %318 -%321 = OpAccessChain %310 %74 %30 -%320 = OpAtomicOr %6 %321 %9 %313 %26 -OpStore %305 %320 -%323 = OpAccessChain %310 %74 %30 -%322 = OpAtomicXor %6 %323 %9 %313 %26 -OpStore %305 %322 -%325 = OpAccessChain %310 %74 %30 -%324 = OpAtomicSMin %6 %325 %9 %313 %26 -OpStore %305 %324 -%327 = OpAccessChain %310 %74 %30 -%326 = OpAtomicSMax %6 %327 %9 %313 %26 -OpStore %305 %326 -%329 = OpAccessChain %310 %74 %30 -%328 = OpAtomicExchange %6 %329 %9 %313 %26 -OpStore %305 %328 -%330 = OpAccessChain %310 %74 %30 -OpAtomicStore %330 %9 %313 %312 -OpReturn -OpFunctionEnd -%334 = OpFunction %2 None %94 -%333 = OpLabel -%331 = OpVariable %63 Function %332 -OpBranch %335 -%335 = OpLabel -%339 = OpLoad %35 %337 -%343 = OpIEqual %341 %339 %340 -%344 = OpAll %342 %343 -OpSelectionMerge %345 None -OpBranchConditional %344 %346 %345 -%346 = OpLabel -OpStore %85 %336 -OpBranch %345 -%345 = OpLabel -OpControlBarrier %30 %30 %347 -OpBranch %348 -%348 = OpLabel -%349 = OpCompositeConstruct %57 %14 %14 %14 %14 -%350 = OpCompositeConstruct %57 %25 %25 %25 %25 -%351 = OpCompositeConstruct %62 %349 %350 -OpStore %331 %351 -%352 = OpFunctionCall %2 %225 %85 -%353 = OpFunctionCall %2 %230 %331 +%313 = OpLoad %35 %311 +%317 = OpIEqual %315 %313 %314 +%318 = OpAll %316 %317 +OpSelectionMerge %319 None +OpBranchConditional %318 %320 %319 +%320 = OpLabel +OpStore %85 %310 +OpBranch %319 +%319 = OpLabel +OpControlBarrier %30 %30 %321 +OpBranch %322 +%322 = OpLabel +%323 = OpCompositeConstruct %57 %14 %14 %14 %14 +%324 = OpCompositeConstruct %57 %25 %25 %25 %25 +%325 = OpCompositeConstruct %62 %323 %324 +OpStore %305 %325 +%326 = OpFunctionCall %2 %225 %85 +%327 = OpFunctionCall %2 %230 %305 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/atomicOps.spvasm b/tests/out/spv/atomicOps.spvasm new file mode 100644 index 0000000000..8e9bad660f --- /dev/null +++ b/tests/out/spv/atomicOps.spvasm @@ -0,0 +1,240 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 189 +OpCapability Shader +OpExtension "SPV_KHR_storage_buffer_storage_class" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %30 "cs_main" %27 +OpExecutionMode %30 LocalSize 2 1 1 +OpDecorate %8 ArrayStride 4 +OpMemberDecorate %9 0 Offset 0 +OpMemberDecorate %9 1 Offset 4 +OpDecorate %11 DescriptorSet 0 +OpDecorate %11 Binding 0 +OpDecorate %12 Block +OpMemberDecorate %12 0 Offset 0 +OpDecorate %14 DescriptorSet 0 +OpDecorate %14 Binding 1 +OpDecorate %15 Block +OpMemberDecorate %15 0 Offset 0 +OpDecorate %17 DescriptorSet 0 +OpDecorate %17 Binding 2 +OpDecorate %18 Block +OpMemberDecorate %18 0 Offset 0 +OpDecorate %27 BuiltIn LocalInvocationId +%2 = OpTypeVoid +%4 = OpTypeInt 32 1 +%3 = OpConstant %4 2 +%6 = OpTypeInt 32 0 +%5 = OpConstant %6 1 +%7 = OpConstant %4 1 +%8 = OpTypeArray %4 %3 +%9 = OpTypeStruct %6 %8 +%10 = OpTypeVector %6 3 +%12 = OpTypeStruct %6 +%13 = OpTypePointer StorageBuffer %12 +%11 = OpVariable %13 StorageBuffer +%15 = OpTypeStruct %8 +%16 = OpTypePointer StorageBuffer %15 +%14 = OpVariable %16 StorageBuffer +%18 = OpTypeStruct %9 +%19 = OpTypePointer StorageBuffer %18 +%17 = OpVariable %19 StorageBuffer +%21 = OpTypePointer Workgroup %6 +%20 = OpVariable %21 Workgroup +%23 = OpTypePointer Workgroup %8 +%22 = OpVariable %23 Workgroup +%25 = OpTypePointer Workgroup %9 +%24 = OpVariable %25 Workgroup +%28 = OpTypePointer Input %10 +%27 = OpVariable %28 Input +%31 = OpTypeFunction %2 +%32 = OpTypePointer StorageBuffer %6 +%33 = OpConstant %6 0 +%35 = OpTypePointer StorageBuffer %8 +%37 = OpTypePointer StorageBuffer %9 +%40 = OpConstantNull %6 +%41 = OpConstantNull %8 +%42 = OpConstantNull %9 +%43 = OpConstantNull %10 +%45 = OpTypeBool +%44 = OpTypeVector %45 3 +%50 = OpConstant %6 2 +%51 = OpConstant %6 264 +%53 = OpConstant %6 64 +%54 = OpTypePointer StorageBuffer %4 +%58 = OpConstant %6 256 +%59 = OpTypePointer Workgroup %4 +%30 = OpFunction %2 None %31 +%26 = OpLabel +%29 = OpLoad %10 %27 +%34 = OpAccessChain %32 %11 %33 +%36 = OpAccessChain %35 %14 %33 +%38 = OpAccessChain %37 %17 %33 +OpBranch %39 +%39 = OpLabel +%46 = OpIEqual %44 %29 %43 +%47 = OpAll %45 %46 +OpSelectionMerge %48 None +OpBranchConditional %47 %49 %48 +%49 = OpLabel +OpStore %20 %40 +OpStore %22 %41 +OpStore %24 %42 +OpBranch %48 +%48 = OpLabel +OpControlBarrier %50 %50 %51 +OpBranch %52 +%52 = OpLabel +OpAtomicStore %34 %7 %53 %5 +%55 = OpAccessChain %54 %36 %5 +OpAtomicStore %55 %7 %53 %7 +%56 = OpAccessChain %32 %38 %33 +OpAtomicStore %56 %7 %53 %5 +%57 = OpAccessChain %54 %38 %5 %5 +OpAtomicStore %57 %7 %53 %7 +OpAtomicStore %20 %3 %58 %5 +%60 = OpAccessChain %59 %22 %5 +OpAtomicStore %60 %3 %58 %7 +%61 = OpAccessChain %21 %24 %33 +OpAtomicStore %61 %3 %58 %5 +%62 = OpAccessChain %59 %24 %5 %5 +OpAtomicStore %62 %3 %58 %7 +OpControlBarrier %50 %50 %51 +%63 = OpAtomicLoad %6 %34 %7 %53 +%64 = OpAccessChain %54 %36 %5 +%65 = OpAtomicLoad %4 %64 %7 %53 +%66 = OpAccessChain %32 %38 %33 +%67 = OpAtomicLoad %6 %66 %7 %53 +%68 = OpAccessChain %54 %38 %5 %5 +%69 = OpAtomicLoad %4 %68 %7 %53 +%70 = OpAtomicLoad %6 %20 %3 %58 +%71 = OpAccessChain %59 %22 %5 +%72 = OpAtomicLoad %4 %71 %3 %58 +%73 = OpAccessChain %21 %24 %33 +%74 = OpAtomicLoad %6 %73 %3 %58 +%75 = OpAccessChain %59 %24 %5 %5 +%76 = OpAtomicLoad %4 %75 %3 %58 +OpControlBarrier %50 %50 %51 +%77 = OpAtomicIAdd %6 %34 %7 %53 %5 +%79 = OpAccessChain %54 %36 %5 +%78 = OpAtomicIAdd %4 %79 %7 %53 %7 +%81 = OpAccessChain %32 %38 %33 +%80 = OpAtomicIAdd %6 %81 %7 %53 %5 +%83 = OpAccessChain %54 %38 %5 %5 +%82 = OpAtomicIAdd %4 %83 %7 %53 %7 +%84 = OpAtomicIAdd %6 %20 %3 %58 %5 +%86 = OpAccessChain %59 %22 %5 +%85 = OpAtomicIAdd %4 %86 %3 %58 %7 +%88 = OpAccessChain %21 %24 %33 +%87 = OpAtomicIAdd %6 %88 %3 %58 %5 +%90 = OpAccessChain %59 %24 %5 %5 +%89 = OpAtomicIAdd %4 %90 %3 %58 %7 +OpControlBarrier %50 %50 %51 +%91 = OpAtomicISub %6 %34 %7 %53 %5 +%93 = OpAccessChain %54 %36 %5 +%92 = OpAtomicISub %4 %93 %7 %53 %7 +%95 = OpAccessChain %32 %38 %33 +%94 = OpAtomicISub %6 %95 %7 %53 %5 +%97 = OpAccessChain %54 %38 %5 %5 +%96 = OpAtomicISub %4 %97 %7 %53 %7 +%98 = OpAtomicISub %6 %20 %3 %58 %5 +%100 = OpAccessChain %59 %22 %5 +%99 = OpAtomicISub %4 %100 %3 %58 %7 +%102 = OpAccessChain %21 %24 %33 +%101 = OpAtomicISub %6 %102 %3 %58 %5 +%104 = OpAccessChain %59 %24 %5 %5 +%103 = OpAtomicISub %4 %104 %3 %58 %7 +OpControlBarrier %50 %50 %51 +%105 = OpAtomicUMax %6 %34 %7 %53 %5 +%107 = OpAccessChain %54 %36 %5 +%106 = OpAtomicSMax %4 %107 %7 %53 %7 +%109 = OpAccessChain %32 %38 %33 +%108 = OpAtomicUMax %6 %109 %7 %53 %5 +%111 = OpAccessChain %54 %38 %5 %5 +%110 = OpAtomicSMax %4 %111 %7 %53 %7 +%112 = OpAtomicUMax %6 %20 %3 %58 %5 +%114 = OpAccessChain %59 %22 %5 +%113 = OpAtomicSMax %4 %114 %3 %58 %7 +%116 = OpAccessChain %21 %24 %33 +%115 = OpAtomicUMax %6 %116 %3 %58 %5 +%118 = OpAccessChain %59 %24 %5 %5 +%117 = OpAtomicSMax %4 %118 %3 %58 %7 +OpControlBarrier %50 %50 %51 +%119 = OpAtomicUMin %6 %34 %7 %53 %5 +%121 = OpAccessChain %54 %36 %5 +%120 = OpAtomicSMin %4 %121 %7 %53 %7 +%123 = OpAccessChain %32 %38 %33 +%122 = OpAtomicUMin %6 %123 %7 %53 %5 +%125 = OpAccessChain %54 %38 %5 %5 +%124 = OpAtomicSMin %4 %125 %7 %53 %7 +%126 = OpAtomicUMin %6 %20 %3 %58 %5 +%128 = OpAccessChain %59 %22 %5 +%127 = OpAtomicSMin %4 %128 %3 %58 %7 +%130 = OpAccessChain %21 %24 %33 +%129 = OpAtomicUMin %6 %130 %3 %58 %5 +%132 = OpAccessChain %59 %24 %5 %5 +%131 = OpAtomicSMin %4 %132 %3 %58 %7 +OpControlBarrier %50 %50 %51 +%133 = OpAtomicAnd %6 %34 %7 %53 %5 +%135 = OpAccessChain %54 %36 %5 +%134 = OpAtomicAnd %4 %135 %7 %53 %7 +%137 = OpAccessChain %32 %38 %33 +%136 = OpAtomicAnd %6 %137 %7 %53 %5 +%139 = OpAccessChain %54 %38 %5 %5 +%138 = OpAtomicAnd %4 %139 %7 %53 %7 +%140 = OpAtomicAnd %6 %20 %3 %58 %5 +%142 = OpAccessChain %59 %22 %5 +%141 = OpAtomicAnd %4 %142 %3 %58 %7 +%144 = OpAccessChain %21 %24 %33 +%143 = OpAtomicAnd %6 %144 %3 %58 %5 +%146 = OpAccessChain %59 %24 %5 %5 +%145 = OpAtomicAnd %4 %146 %3 %58 %7 +OpControlBarrier %50 %50 %51 +%147 = OpAtomicOr %6 %34 %7 %53 %5 +%149 = OpAccessChain %54 %36 %5 +%148 = OpAtomicOr %4 %149 %7 %53 %7 +%151 = OpAccessChain %32 %38 %33 +%150 = OpAtomicOr %6 %151 %7 %53 %5 +%153 = OpAccessChain %54 %38 %5 %5 +%152 = OpAtomicOr %4 %153 %7 %53 %7 +%154 = OpAtomicOr %6 %20 %3 %58 %5 +%156 = OpAccessChain %59 %22 %5 +%155 = OpAtomicOr %4 %156 %3 %58 %7 +%158 = OpAccessChain %21 %24 %33 +%157 = OpAtomicOr %6 %158 %3 %58 %5 +%160 = OpAccessChain %59 %24 %5 %5 +%159 = OpAtomicOr %4 %160 %3 %58 %7 +OpControlBarrier %50 %50 %51 +%161 = OpAtomicXor %6 %34 %7 %53 %5 +%163 = OpAccessChain %54 %36 %5 +%162 = OpAtomicXor %4 %163 %7 %53 %7 +%165 = OpAccessChain %32 %38 %33 +%164 = OpAtomicXor %6 %165 %7 %53 %5 +%167 = OpAccessChain %54 %38 %5 %5 +%166 = OpAtomicXor %4 %167 %7 %53 %7 +%168 = OpAtomicXor %6 %20 %3 %58 %5 +%170 = OpAccessChain %59 %22 %5 +%169 = OpAtomicXor %4 %170 %3 %58 %7 +%172 = OpAccessChain %21 %24 %33 +%171 = OpAtomicXor %6 %172 %3 %58 %5 +%174 = OpAccessChain %59 %24 %5 %5 +%173 = OpAtomicXor %4 %174 %3 %58 %7 +%175 = OpAtomicExchange %6 %34 %7 %53 %5 +%177 = OpAccessChain %54 %36 %5 +%176 = OpAtomicExchange %4 %177 %7 %53 %7 +%179 = OpAccessChain %32 %38 %33 +%178 = OpAtomicExchange %6 %179 %7 %53 %5 +%181 = OpAccessChain %54 %38 %5 %5 +%180 = OpAtomicExchange %4 %181 %7 %53 %7 +%182 = OpAtomicExchange %6 %20 %3 %58 %5 +%184 = OpAccessChain %59 %22 %5 +%183 = OpAtomicExchange %4 %184 %3 %58 %7 +%186 = OpAccessChain %21 %24 %33 +%185 = OpAtomicExchange %6 %186 %3 %58 %5 +%188 = OpAccessChain %59 %24 %5 %5 +%187 = OpAtomicExchange %4 %188 %3 %58 %7 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 7133f53d69..0bbd1a7759 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -165,31 +165,6 @@ fn foo_frag() -> @location(0) vec4 { return vec4(0.0); } -@compute @workgroup_size(1, 1, 1) -fn atomics() { - var tmp: i32; - - let value_1 = atomicLoad((&bar.atom)); - let _e7 = atomicAdd((&bar.atom), 5); - tmp = _e7; - let _e11 = atomicSub((&bar.atom), 5); - tmp = _e11; - let _e15 = atomicAnd((&bar.atom), 5); - tmp = _e15; - let _e19 = atomicOr((&bar.atom), 5); - tmp = _e19; - let _e23 = atomicXor((&bar.atom), 5); - tmp = _e23; - let _e27 = atomicMin((&bar.atom), 5); - tmp = _e27; - let _e31 = atomicMax((&bar.atom), 5); - tmp = _e31; - let _e35 = atomicExchange((&bar.atom), 5); - tmp = _e35; - atomicStore((&bar.atom), value_1); - return; -} - @compute @workgroup_size(1, 1, 1) fn assign_through_ptr() { var arr: array,2>; diff --git a/tests/out/wgsl/atomicOps.wgsl b/tests/out/wgsl/atomicOps.wgsl new file mode 100644 index 0000000000..80d1776548 --- /dev/null +++ b/tests/out/wgsl/atomicOps.wgsl @@ -0,0 +1,107 @@ +struct Struct { + atomic_scalar: atomic, + atomic_arr: array,2>, +} + +@group(0) @binding(0) +var storage_atomic_scalar: atomic; +@group(0) @binding(1) +var storage_atomic_arr: array,2>; +@group(0) @binding(2) +var storage_struct: Struct; +var workgroup_atomic_scalar: atomic; +var workgroup_atomic_arr: array,2>; +var workgroup_struct: Struct; + +@compute @workgroup_size(2, 1, 1) +fn cs_main(@builtin(local_invocation_id) id: vec3) { + atomicStore((&storage_atomic_scalar), 1u); + atomicStore((&storage_atomic_arr[1]), 1); + atomicStore((&storage_struct.atomic_scalar), 1u); + atomicStore((&storage_struct.atomic_arr[1]), 1); + atomicStore((&workgroup_atomic_scalar), 1u); + atomicStore((&workgroup_atomic_arr[1]), 1); + atomicStore((&workgroup_struct.atomic_scalar), 1u); + atomicStore((&workgroup_struct.atomic_arr[1]), 1); + workgroupBarrier(); + _ = atomicLoad((&storage_atomic_scalar)); + _ = atomicLoad((&storage_atomic_arr[1])); + _ = atomicLoad((&storage_struct.atomic_scalar)); + _ = atomicLoad((&storage_struct.atomic_arr[1])); + _ = atomicLoad((&workgroup_atomic_scalar)); + _ = atomicLoad((&workgroup_atomic_arr[1])); + _ = atomicLoad((&workgroup_struct.atomic_scalar)); + _ = atomicLoad((&workgroup_struct.atomic_arr[1])); + workgroupBarrier(); + let _e59 = atomicAdd((&storage_atomic_scalar), 1u); + let _e64 = atomicAdd((&storage_atomic_arr[1]), 1); + let _e68 = atomicAdd((&storage_struct.atomic_scalar), 1u); + let _e74 = atomicAdd((&storage_struct.atomic_arr[1]), 1); + let _e77 = atomicAdd((&workgroup_atomic_scalar), 1u); + let _e82 = atomicAdd((&workgroup_atomic_arr[1]), 1); + let _e86 = atomicAdd((&workgroup_struct.atomic_scalar), 1u); + let _e92 = atomicAdd((&workgroup_struct.atomic_arr[1]), 1); + workgroupBarrier(); + let _e95 = atomicSub((&storage_atomic_scalar), 1u); + let _e100 = atomicSub((&storage_atomic_arr[1]), 1); + let _e104 = atomicSub((&storage_struct.atomic_scalar), 1u); + let _e110 = atomicSub((&storage_struct.atomic_arr[1]), 1); + let _e113 = atomicSub((&workgroup_atomic_scalar), 1u); + let _e118 = atomicSub((&workgroup_atomic_arr[1]), 1); + let _e122 = atomicSub((&workgroup_struct.atomic_scalar), 1u); + let _e128 = atomicSub((&workgroup_struct.atomic_arr[1]), 1); + workgroupBarrier(); + let _e131 = atomicMax((&storage_atomic_scalar), 1u); + let _e136 = atomicMax((&storage_atomic_arr[1]), 1); + let _e140 = atomicMax((&storage_struct.atomic_scalar), 1u); + let _e146 = atomicMax((&storage_struct.atomic_arr[1]), 1); + let _e149 = atomicMax((&workgroup_atomic_scalar), 1u); + let _e154 = atomicMax((&workgroup_atomic_arr[1]), 1); + let _e158 = atomicMax((&workgroup_struct.atomic_scalar), 1u); + let _e164 = atomicMax((&workgroup_struct.atomic_arr[1]), 1); + workgroupBarrier(); + let _e167 = atomicMin((&storage_atomic_scalar), 1u); + let _e172 = atomicMin((&storage_atomic_arr[1]), 1); + let _e176 = atomicMin((&storage_struct.atomic_scalar), 1u); + let _e182 = atomicMin((&storage_struct.atomic_arr[1]), 1); + let _e185 = atomicMin((&workgroup_atomic_scalar), 1u); + let _e190 = atomicMin((&workgroup_atomic_arr[1]), 1); + let _e194 = atomicMin((&workgroup_struct.atomic_scalar), 1u); + let _e200 = atomicMin((&workgroup_struct.atomic_arr[1]), 1); + workgroupBarrier(); + let _e203 = atomicAnd((&storage_atomic_scalar), 1u); + let _e208 = atomicAnd((&storage_atomic_arr[1]), 1); + let _e212 = atomicAnd((&storage_struct.atomic_scalar), 1u); + let _e218 = atomicAnd((&storage_struct.atomic_arr[1]), 1); + let _e221 = atomicAnd((&workgroup_atomic_scalar), 1u); + let _e226 = atomicAnd((&workgroup_atomic_arr[1]), 1); + let _e230 = atomicAnd((&workgroup_struct.atomic_scalar), 1u); + let _e236 = atomicAnd((&workgroup_struct.atomic_arr[1]), 1); + workgroupBarrier(); + let _e239 = atomicOr((&storage_atomic_scalar), 1u); + let _e244 = atomicOr((&storage_atomic_arr[1]), 1); + let _e248 = atomicOr((&storage_struct.atomic_scalar), 1u); + let _e254 = atomicOr((&storage_struct.atomic_arr[1]), 1); + let _e257 = atomicOr((&workgroup_atomic_scalar), 1u); + let _e262 = atomicOr((&workgroup_atomic_arr[1]), 1); + let _e266 = atomicOr((&workgroup_struct.atomic_scalar), 1u); + let _e272 = atomicOr((&workgroup_struct.atomic_arr[1]), 1); + workgroupBarrier(); + let _e275 = atomicXor((&storage_atomic_scalar), 1u); + let _e280 = atomicXor((&storage_atomic_arr[1]), 1); + let _e284 = atomicXor((&storage_struct.atomic_scalar), 1u); + let _e290 = atomicXor((&storage_struct.atomic_arr[1]), 1); + let _e293 = atomicXor((&workgroup_atomic_scalar), 1u); + let _e298 = atomicXor((&workgroup_atomic_arr[1]), 1); + let _e302 = atomicXor((&workgroup_struct.atomic_scalar), 1u); + let _e308 = atomicXor((&workgroup_struct.atomic_arr[1]), 1); + let _e311 = atomicExchange((&storage_atomic_scalar), 1u); + let _e316 = atomicExchange((&storage_atomic_arr[1]), 1); + let _e320 = atomicExchange((&storage_struct.atomic_scalar), 1u); + let _e326 = atomicExchange((&storage_struct.atomic_arr[1]), 1); + let _e329 = atomicExchange((&workgroup_atomic_scalar), 1u); + let _e334 = atomicExchange((&workgroup_atomic_arr[1]), 1); + let _e338 = atomicExchange((&workgroup_struct.atomic_scalar), 1u); + let _e344 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1); + return; +} diff --git a/tests/snapshots.rs b/tests/snapshots.rs index d968a0dfc1..877b2597b4 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -505,6 +505,10 @@ fn convert_wgsl() { | Targets::IR | Targets::ANALYSIS, ), + ( + "atomicOps", + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, + ), ("atomicCompareExchange", Targets::SPIRV | Targets::WGSL), ( "padding",