github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

fix(hlsl-out): use Interlocked<op> intrinsic for atomic integers (#2294)

Browse Source 
We currently assume that we are using raw `RWByteAddressBuffer` methods for all atomic operations (`<pointer>.Interlocked<op>(<raw_byte_offset>, …)`), which is only true when we use `var<storage, read_write>` globals. For `var<workgroup>` globals, we need `Interlocked<op>(<pointer>, …)`, using the original expression as the first argument.

Fix this by branching on the `pointer`'s address space in `Atomic` statements, and implementing the workgroup address space case with intrinsics.

Remove atomic ops from `access`, add new `atomicOps` test.

Fixes #2284
This commit is contained in:
Erich Gubler
2023-04-05 22:37:22 -04:00
committed by GitHub
parent 1158709747
commit 99a7773e65
17 changed files with 928 additions and 1053 deletions
Download Patch File Download Diff File Expand all files Collapse all files

36
src/back/hlsl/writer.rs
View File

@@ -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 {

19
tests/in/access.wgsl
View File

@@ -151,23 +151,6 @@ fn foo_frag() -> @location(0) vec4<f32> {
return vec4<f32>(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<workgroup> val: u32;
fn assign_through_ptr_fn(p: ptr<workgroup, u32>) {
@@ -184,4 +167,4 @@ fn assign_through_ptr() {
assign_through_ptr_fn(&val);
assign_array_through_ptr_fn(&arr);
}
}

141
tests/in/atomicOps.wgsl Normal file
View File

@@ -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<u32>,
atomic_arr: array<atomic<i32>, 2>,
}
@group(0) @binding(0)
var<storage, read_write> storage_atomic_scalar: atomic<u32>;
@group(0) @binding(1)
var<storage, read_write> storage_atomic_arr: array<atomic<i32>, 2>;
@group(0) @binding(2)
var<storage, read_write> storage_struct: Struct;
var<workgroup> workgroup_atomic_scalar: atomic<u32>;
var<workgroup> workgroup_atomic_arr: array<atomic<i32>, 2>;
var<workgroup> workgroup_struct: Struct;
@compute
@workgroup_size(2)
fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
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);
}

627
tests/out/analysis/access.info.ron
View File

@@ -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,

132
tests/out/glsl/atomicOps.cs_main.Compute.glsl Normal file
View File

@@ -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;
}

26
tests/out/hlsl/access.hlsl
View File

@@ -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)
{

2
tests/out/hlsl/access.hlsl.config
View File

@@ -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 )

112
tests/out/hlsl/atomicOps.hlsl Normal file
View File

@@ -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;
}

3
tests/out/hlsl/atomicOps.hlsl.config Normal file
View File

@@ -0,0 +1,3 @@
vertex=()
fragment=()
compute=(cs_main:cs_5_1 )

248
tests/out/ir/access.ron
View File

@@ -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,

27
tests/out/msl/access.msl
View File

@@ -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

126
tests/out/msl/atomicOps.msl Normal file
View File

@@ -0,0 +1,126 @@
// language: metal2.0
#include <metal_stdlib>
#include <simd/simd.h>
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;
}

106
tests/out/spv/access.spvasm
View File

@@ -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

240
tests/out/spv/atomicOps.spvasm Normal file
View File

@@ -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

25
tests/out/wgsl/access.wgsl
View File

@@ -165,31 +165,6 @@ fn foo_frag() -> @location(0) vec4<f32> {
return vec4<f32>(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<vec4<f32>,2>;

107
tests/out/wgsl/atomicOps.wgsl Normal file
View File

@@ -0,0 +1,107 @@
struct Struct {
atomic_scalar: atomic<u32>,
atomic_arr: array<atomic<i32>,2>,
}
@group(0) @binding(0)
var<storage, read_write> storage_atomic_scalar: atomic<u32>;
@group(0) @binding(1)
var<storage, read_write> storage_atomic_arr: array<atomic<i32>,2>;
@group(0) @binding(2)
var<storage, read_write> storage_struct: Struct;
var<workgroup> workgroup_atomic_scalar: atomic<u32>;
var<workgroup> workgroup_atomic_arr: array<atomic<i32>,2>;
var<workgroup> workgroup_struct: Struct;
@compute @workgroup_size(2, 1, 1)
fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
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;
}

4
tests/snapshots.rs
View File

@@ -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",